Tile Language

https://tilelang.com
Admin
Tile Language is a concise, Pythonic domain-specific language for developing high-performance...
Tokens:317,794
Snippets:3,061
Trust Score:5.7
Update:2 months ago
Show doc for...
Tokens:
Raw
### Define Kernel Setup

Source: https://tilelang.com/compiler_internals/tensor_checks.html

Initial setup for the kernel environment used in subsequent examples.

```python
# Convention:
# A: float16 [M, K]
# B: float16 [K, N]
# C: float16 [M, N]
# Target: CUDA (device_type=2)
fn = matmul_relu_kernel  # your compiled function
M = N = K = 1024
```

--------------------------------

### Verify environment setup

Source: https://tilelang.com/_sources/tools/iket.md.txt

Command to verify that TileLang, IKET, and PyTorch are correctly installed in the current Python environment.

```bash
python -c "import tilelang, iket, torch"
```

--------------------------------

### Access and Verify Installation

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Enter the running container and verify the installation by checking the package version.

```bash
docker exec -it tilelang_b200 /bin/zsh
# Inside the container:
python -c "import tilelang; print(tilelang.__version__)"
```

--------------------------------

### Install tilelang from GitHub

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Installs the latest development version directly from the source repository.

```bash
pip install git+https://github.com/tile-ai/tilelang.git
```

--------------------------------

### Install OS-level prerequisites

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Install required system packages on Ubuntu or Debian-based Linux distributions.

```bash
apt-get update
apt-get install -y python3 python3-dev python3-setuptools gcc zlib1g-dev build-essential cmake libedit-dev
```

--------------------------------

### Install Visualization Dependencies

Source: https://tilelang.com/_sources/tools/layout_visualization.md.txt

Install the necessary dependencies for PNG, PDF, and SVG output.

```bash
pip install "tilelang[vis]"
```

--------------------------------

### Install tilelang from a local wheel

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Installs a specific prebuilt wheel file downloaded from the project's release page.

```bash
pip install tilelang-0.0.0.dev0+ubuntu.20.4.cu120-py3-none-any.whl
```

--------------------------------

### GemmWarpPolicy Examples

Source: https://tilelang.com/_sources/autoapi/tilelang/tileop/base/index.rst.txt

Demonstrates how to use the `from_warp_partition` method to infer the warp policy. Examples show cases for full row partitioning, full column partitioning, and balanced square partitioning.

```python
GemmWarpPolicy.from_warp_partition(4, 1)  # All warps in rows
GemmWarpPolicy.FullRow
GemmWarpPolicy.from_warp_partition(1, 4)  # All warps in columns
GemmWarpPolicy.FullCol
GemmWarpPolicy.from_warp_partition(2, 2)  # Balanced distribution
GemmWarpPolicy.Square
```

--------------------------------

### Install Tilelang Nightly Version

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Use pip to install the latest nightly build from the specified repository URL.

```bash
pip install tilelang -f https://tile-ai.github.io/whl/nightly
# or pip install tilelang --find-links https://tile-ai.github.io/whl/nightly
```

--------------------------------

### Install tilelang via PyPI

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Installs the latest stable version of tilelang from the Python Package Index.

```bash
pip install tilelang
```

--------------------------------

### Access Pointer Creation Examples

Source: https://tilelang.com/_sources/autoapi/tilelang/language/builtin/index.rst.txt

Examples of using T.access_ptr to compute buffer pointers with varying extents.

```python
T.access_ptr(A[i], "r")
```

```python
T.access_ptr(A[i], "r", 16)
```

```python
T.access_ptr(A[i, j], "r", m, n)
```

--------------------------------

### Install with host CUDA toolchain

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Clone the repository and install using pip when a host CUDA environment is available.

```bash
git clone --recursive https://github.com/tile-ai/tilelang.git
cd tilelang
pip install . -v
```

--------------------------------

### Verify tilelang installation

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Checks that the library is correctly installed by printing the current version.

```bash
python -c "import tilelang; print(tilelang.__version__)"
```

--------------------------------

### Matmul ReLU Kernel Example

Source: https://tilelang.com/_sources/compiler_internals/tensor_checks.md.txt

A reference example demonstrating a matmul and ReLU kernel, including initialization of kernel context, shared memory allocation, and computation.

```python
@T.prim_func
def matmul_relu_kernel(
    A: T.Tensor((M, K), dtype),
    B: T.Tensor((K, N), dtype),
    C: T.Tensor((M, N), dtype),
):
    # Initialize Kernel Context
    with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):
        A_shared = T.alloc_shared((block_M, block_K), dtype)
        B_shared = T.alloc_shared((block_K, block_N), dtype)
        C_local = T.alloc_fragment((block_M, block_N), accum_dtype)
        T.clear(C_local)
        for ko in T.Pipelined(T.ceildiv(K, block_K), num_stages=0):
            T.copy(A[by * block_M, ko * block_K], A_shared)
            T.copy(B[ko * block_K, bx * block_N], B_shared)
            T.gemm(A_shared, B_shared, C_local)
        T.copy(C_local, C[by * block_M, bx * block_N])
```

--------------------------------

### Install in development mode

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Install the package in editable mode to reflect Python file changes immediately.

```bash
pip install -e . -v
```

--------------------------------

### StorageRewrite Inplace Detection Examples

Source: https://tilelang.com/_sources/autoapi/tilelang/transform/pass_config/index.rst.txt

Examples showing the difference in generated IR when StorageRewrite inplace detection is disabled versus enabled.

```text
read_buf = T.alloc_buffer((1,), T.int32, scope="local.var")
write_buf = T.alloc_buffer((1,), T.int32, scope="local.var")
write_buf[0] = read_buf[0] * 2
f(write_buf[0])
```

```text
read_buf = T.alloc_buffer((1,), T.int32, scope="local.var")
read_buf[0] = read_buf[0] * 2
f(read_buf[0])
```

--------------------------------

### Run the Docker Container

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Start a container with GPU access and shared memory configuration.

```bash
docker run -itd \
  --shm-size 32g \
  --gpus all \
  -v /home/tilelang:/home/tilelang \
  --name tilelang_b200 \
  tilelang-cu120 \
  /bin/zsh
```

--------------------------------

### Verify Developer Mode Installation

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Command to verify that TileLang is loading from the local development build directory.

```console
$ python -c 'import tilelang'
2025-10-14 11:11:29  [TileLang:tilelang.env:WARNING]: Loading tilelang libs from dev root: /Users/yyc/repo/tilelang/build
```

--------------------------------

### Launch Standard CUDA Kernels

Source: https://tilelang.com/_sources/autoapi/tilelang/language/kernel/index.rst.txt

Examples for launching 1-D and 2-D grid kernels using T.Kernel.

```python
with T.Kernel(T.ceildiv(N, 128), threads=128) as bx:
    # bx is the blockIdx.x binding (also iterable as (bx,))
    ...
```

```python
with T.Kernel(grid_x, grid_y, threads=(64, 2)) as (bx, by):
    tx, ty = T.get_thread_bindings()
    ...
```

--------------------------------

### Build with existing TVM installation

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Specify a custom TVM repository path using the TVM_ROOT environment variable during installation.

```bash
TVM_ROOT=<your-tvm-repo> pip install . -v
```

--------------------------------

### List Profiler Output Files

Source: https://tilelang.com/tools/iket.html

Example of the file types generated in the output directory.

```text
iket_pid_0x....pftrace
iket_pid_0x....pftrace.gz
iket_pid_0x....trace.json
iket_pid_0x....html
```

--------------------------------

### Serve Trace Files

Source: https://tilelang.com/_sources/tools/iket.md.txt

Starts a local HTTP server to host the generated trace files for browser-based inspection.

```bash
cd /tmp/tilelang_iket_all_features_profile
python3 -m http.server 8080
```

--------------------------------

### Runtime Log Level Configuration

Source: https://tilelang.com/tutorials/logging.html

Example of setting the environment variable to control log output at runtime.

```bash
TVM_LOG_DEBUG=1 python3 code.py
```

--------------------------------

### Compile instrumented kernel

Source: https://tilelang.com/_sources/tools/iket.md.txt

Example of defining an instrumented kernel using iket.range and iket.mark, then compiling it within an iket.session.

```python
import tilelang
import tilelang.language as T
from tilelang.tools.cuda import iket


def instrumented_add(n: int, threads: int = 128):
    @T.prim_func
    def main(
        A: T.Tensor((n,), T.float32),
        B: T.Tensor((n,), T.float32),
        C: T.Tensor((n,), T.float32),
    ):
        with T.Kernel(T.ceildiv(n, threads), threads=threads) as bx:
            with iket.range("block_total"):
                for tx in T.Parallel(threads):
                    i = bx * threads + tx
                    if i < n:
                        iket.mark("before_store")
                        C[i] = A[i] + B[i]
                        iket.mark("after_store")

    return main


with iket.session(output_dir="/tmp/tilelang_iket"):
    program = instrumented_add(1024)
    kernel = tilelang.compile(
        program,
        out_idx=-1,
        target="cuda",
        execution_backend="cython",
    )
```

--------------------------------

### Launch Cluster Kernels

Source: https://tilelang.com/_sources/autoapi/tilelang/language/kernel/index.rst.txt

Example for launching a kernel with a CUDA thread block cluster for SM90+ hardware.

```python
with T.ClusterKernel(grid_x, grid_y, cluster_dims=2, threads=128) as (bx, by):
    ...
```

--------------------------------

### Compile and Benchmark Configurations

Source: https://tilelang.com/tutorials/auto_tuning.html

Initialize the AutoTuner with the kernel and configurations, then execute the benchmark.

```python
autotuner = AutoTuner.from_kernel(
    kernel=kernel, configs=get_configs(M, N, K, with_roller)).set_compile_args(
        out_idx=[-1],
        supply_type=tl.TensorSupplyType.Integer,
        ref_prog=ref_program,
        skip_check=False,
        target="auto",
    )
result = autotuner.run(warmup=3, rep=20)
out_c = result.kernel(a, b)
```

--------------------------------

### Device ID Mismatch Example (Multi-GPU)

Source: https://tilelang.com/_sources/compiler_internals/tensor_checks.md.txt

Demonstrates a 'device_id' mismatch error in a multi-GPU setup where tensors reside on different GPUs. Place all tensors on the same GPU for correct execution.

```python
import torch

A = torch.empty((M, K), device='cuda:0', dtype=torch.float16)
B = torch.empty((K, N), device='cuda:1', dtype=torch.float16)
C = torch.empty((M, N), device='cuda:0', dtype=torch.float16)
fn(A, B, C)
```

--------------------------------

### Install with pip-provided CUDA toolchain

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Methods for installing without a host CUDA installation using pip-provided packages.

```bash
git clone --recursive https://github.com/tile-ai/tilelang.git
cd tilelang
pip install -r requirements-dev.txt
pip install "nvidia-cuda-nvcc>=13" "nvidia-cuda-cccl>=13" "nvidia-cuda-nvrtc>=13"
pip install . -v --no-build-isolation
```

```bash
# Point to the cu<ver> directory inside another venv's site-packages
export WITH_PIP_CUDA_TOOLCHAIN=/path/to/venv/lib/python3.x/site-packages/nvidia/cu13
pip install . -v
```

--------------------------------

### recommend_hints

Source: https://tilelang.com/_sources/autoapi/tilelang/carver/template/base/index.rst.txt

Provides a list of recommended hardware-aware configurations.

```APIDOC
## recommend_hints

### Description
Provides a list of recommended hardware-aware configurations.

### Method
`recommend_hints(topk = 10)`

### Parameters
#### Path Parameters
None

#### Query Parameters
- **topk** (int, optional) - Number of top configurations to return. Defaults to 10.

### Returns
- **List[Hint]** - A list of recommended configurations.
```

--------------------------------

### Build Package with Version Labeling

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Example of the build output showing the embedded SDK and version information in the generated wheel filename.

```bash
$ python -mbuild -w
...
Successfully built tilelang-0.1.6.post1+cu116.git0d4a74be-cp38-abi3-linux_x86_64.whl
```

--------------------------------

### main

Source: https://tilelang.com/_sources/autoapi/tilelang/testing/index.rst.txt

Entry point for running tilelang tests.

```APIDOC
## main

### Description
Entry point for running tilelang tests.

### Parameters
None.
```

--------------------------------

### Configure IKET Output Helpers

Source: https://tilelang.com/_sources/tools/iket.md.txt

Use these functions to set the output directory, retrieve paths, and manage profiling commands.

```python
iket.set_output_dir("/tmp/tilelang_iket")
iket.output_dir()
iket.output_path("kernel.cu")
iket.trace_files()
iket.profile_command([...], directory="/tmp/tilelang_iket")
```

--------------------------------

### Implement vector addition with @jit

Source: https://tilelang.com/programming_guides/language_basics.html

A complete end-to-end example showing kernel definition and host-side execution using PyTorch tensors.

```python
import tilelang
import tilelang.language as T
from tilelang import jit

@jit  # infers target from tensors at first call
def add(N: int, block: int = 256, dtype: str = 'float32'):

    @T.prim_func
    def add_kernel(
        A: T.Tensor((N,), dtype),
        B: T.Tensor((N,), dtype),
        C: T.Tensor((N,), dtype),
    ):
        with T.Kernel(T.ceildiv(N, block), threads=block) as bx:
            for i in T.Parallel(block):
                gi = bx * block + i
                # Optional — LegalizeSafeMemoryAccess inserts a guard when an access may be OOB
                C[gi] = A[gi] + B[gi]

    return add_kernel

# Host side (PyTorch shown; NumPy/DLPack also supported)
import torch
N = 1 << 20
A = torch.randn(N, device='cuda', dtype=torch.float32)
B = torch.randn(N, device='cuda', dtype=torch.float32)
C = torch.empty(N, device='cuda', dtype=torch.float32)

kernel = add(N)
kernel(A, B, C)  # runs on GPU
torch.testing.assert_close(C, A + B)
```

--------------------------------

### profile_command(command, *, directory, postprocess='all', clobber=True)

Source: https://tilelang.com/_sources/autoapi/tilelang/tools/cuda/iket/cli/index.rst.txt

Builds an IKET CLI command for exporting traces.

```APIDOC
## profile_command(command, *, directory, postprocess='all', clobber=True)

### Description
Build an IKET CLI command that exports traces to the specified directory.

### Parameters
- **command** (str) - Required - The command to profile.
- **directory** (str) - Required - The target directory for exported traces.
- **postprocess** (str) - Optional - Post-processing mode, defaults to 'all'.
- **clobber** (bool) - Optional - Whether to overwrite existing files, defaults to True.
```

--------------------------------

### Launch a 2-D grid with thread bindings

Source: https://tilelang.com/autoapi/tilelang/language/kernel/index.html

Shows how to launch a 2-D grid and retrieve thread bindings within the kernel context.

```python
with T.Kernel(grid_x, grid_y, threads=(64, 2)) as (bx, by):
    tx, ty = T.get_thread_bindings()
    ...
```

--------------------------------

### start_workers()

Source: https://tilelang.com/_sources/autoapi/tilelang/autodd/index.rst.txt

Starts the asynchronous worker processes.

```APIDOC
### start_workers()

Starts the asynchronous worker processes.
```

--------------------------------

### get_cuda_version

Source: https://tilelang.com/_sources/autoapi/tilelang/contrib/nvcc/index.rst.txt

Retrieves the version of the installed CUDA toolkit.

```APIDOC
## get_cuda_version(cuda_path=None)

### Description
Utility function to detect and return the CUDA version.

### Parameters
- **cuda_path** (Optional[str]) - Optional - Path to the CUDA root directory. If None, uses find_cuda_path().

### Returns
- **float** - The detected CUDA version.
```

--------------------------------

### JITKernel.__init__

Source: https://tilelang.com/_sources/autoapi/tilelang/jit/kernel/index.rst.txt

Initializes a new JITKernel instance to compile and wrap a TileLang function.

```APIDOC
## JITKernel(func=None, out_idx=None, execution_backend='tvm_ffi', target='auto', target_host=None, verbose=False, pass_configs=None, from_database=False, compile_flags=None)

### Description
Creates a wrapper class for compiling and invoking TileLang (TVM TIR) functions as PyTorch-compatible functions.
```

--------------------------------

### iket.range_push / iket.range_pop

Source: https://tilelang.com/_sources/tools/iket.md.txt

Explicitly starts and ends a range.

```APIDOC
## iket.range_push(name) / iket.range_pop(name)

### Description
Explicitly marks the start and end of a range. `range_start` and `range_end` are aliases for these methods.

### Parameters
- **name** (string) - Required - The name of the range.

### Example
```python
iket.range_push("compute")
# TileLang statements
iket.range_pop("compute")
```
```

--------------------------------

### arch property

Source: https://tilelang.com/_sources/autoapi/tilelang/carver/template/base/index.rst.txt

Gets the current architecture of the template.

```APIDOC
## arch

### Description
Returns the current architecture.

### Property
`arch`

### Returns
- **TileDevice** - The architecture of this template.
```

--------------------------------

### Create a 1-D CUDA kernel launch

Source: https://tilelang.com/autoapi/tilelang/language/kernel/index.html

Demonstrates initializing a 1-D kernel launch frame and unpacking the block index.

```python
with T.Kernel(T.ceildiv(N, 128), threads=128) as bx:
    # bx is the blockIdx.x binding (also iterable as (bx,))
    ...
```

--------------------------------

### find_rocm_path

Source: https://tilelang.com/_sources/autoapi/tilelang/contrib/rocm/index.rst.txt

A utility function to find the ROCm installation path.

```APIDOC
## find_rocm_path()

### Description
Utility function to find ROCm path.

### Returns
- **path** (str) - Path to ROCm root.
```

--------------------------------

### Configure and Apply Compilation Targets

Source: https://tilelang.com/_sources/get_started/targets.md.txt

Demonstrates passing a target configuration dictionary to the compiler or JIT decorator to specify device-specific options.

```python
target = {"kind": "cuda", "arch": "sm_90"}
kernel = tilelang.compile(func, target=target, execution_backend="cython")
# or
@tilelang.jit(target=target)
def compiled_kernel(*args):
    return func(*args)
```

--------------------------------

### parallel(start, stop=None, *, annotations=None)

Source: https://tilelang.com/_sources/autoapi/tilelang/language/tir/ir/index.rst.txt

Creates a parallel For statement.

```APIDOC
## parallel(start, stop=None, *, annotations=None)

### Description
The parallel For statement.

### Parameters
- **start** (PrimExpr) - Required - The minimum value of iteration.
- **stop** (PrimExpr) - Optional - The maximum value of iteration.
- **annotations** (Dict[str, Any]) - Optional - The optional annotations of the For statement.

### Returns
- **res** (frame.ForFrame) - The ForFrame.
```

--------------------------------

### serial(start, stop=None, *, annotations=None)

Source: https://tilelang.com/_sources/autoapi/tilelang/language/tir/ir/index.rst.txt

Creates a serial For statement.

```APIDOC
## serial(start, stop=None, *, annotations=None)

### Description
The serial For statement.

### Parameters
- **start** (PrimExpr) - Required - The minimum value of iteration.
- **stop** (PrimExpr) - Optional - The maximum value of iteration.
- **annotations** (Dict[str, Any]) - Optional - The optional annotations of the For statement.

### Returns
- **res** (frame.ForFrame) - The ForFrame.
```

--------------------------------

### output_nodes property

Source: https://tilelang.com/_sources/autoapi/tilelang/carver/template/base/index.rst.txt

Gets the output nodes associated with the template.

```APIDOC
## output_nodes

### Description
Returns the output nodes associated with this template.

### Property
`output_nodes`

### Returns
- **List[OutputNode]** - The output nodes.
```

--------------------------------

### thread_binding(start, stop=None, thread=None, *, annotations=None)

Source: https://tilelang.com/_sources/autoapi/tilelang/language/tir/ir/index.rst.txt

Creates a thread-binding For statement.

```APIDOC
## thread_binding(start, stop=None, thread=None, *, annotations=None)

### Description
The thread-binding For statement.

### Parameters
- **start** (PrimExpr) - Required - The minimum value of iteration.
- **stop** (PrimExpr) - Optional - The maximum value of iteration.
- **thread** (str) - Optional - The thread for loop variable to bind.
- **annotations** (Dict[str, Any]) - Optional - The optional annotations of the For statement.

### Returns
- **res** (frame.ForFrame) - The ForFrame.
```

--------------------------------

### main(args)

Source: https://tilelang.com/_sources/autoapi/tilelang/autodd/index.rst.txt

Main entry point for the asynchronous execution.

```APIDOC
### main(args)

Main entry point.

#### Parameters
- **args** (Args) - The arguments object.
```

--------------------------------

### CuTe Elementwise Addition Setup

Source: https://tilelang.com/_sources/deeplearning_operators/elementwise.md.txt

Initializes tensors and buffers for elementwise addition using CuTe. This code sets up local tiles for registers and prepares for computation.

```c++
template<int NUM_ELE_PER_THREAD=8>
__global__ void elementwise_add(nv_bfloat16* C,
                                 const nv_bfloat16* A,
                                 const nv_bfloat16* B,
                                 int N) {
  using namespace cute;

  const int idx = threadIdx.x + blockIdx.x * blockDim.x;

  Tensor t_C = make_tensor(make_gmem_ptr(C), make_shape(N));
  Tensor t_A = make_tensor(make_gmem_ptr(A), make_shape(N));
  Tensor t_B = make_tensor(make_gmem_ptr(B), make_shape(N));

  Tensor t_C_tile = local_tile(t_C, make_shape(Int<NUM_ELE_PER_THREAD>{}), make_coord(idx));
  Tensor t_A_tile = local_tile(t_A, make_shape(Int<NUM_ELE_PER_THREAD>{}), make_coord(idx));
  Tensor t_B_tile = local_tile(t_B, make_shape(Int<NUM_ELE_PER_THREAD>{}), make_coord(idx));

  Tensor reg_buffer_A = make_tensor_like(t_A_tile);
  Tensor reg_buffer_B = make_tensor_like(t_B_tile);

```

--------------------------------

### get_num_sms

Source: https://tilelang.com/_sources/autoapi/tilelang/carver/arch/driver/cuda_driver/index.rst.txt

Get the number of streaming multiprocessors (SMs) on the CUDA device.

```APIDOC
## get_num_sms(device_id = 0)

### Description
Get the number of streaming multiprocessors (SMs) on the CUDA device.

### Parameters
#### Path Parameters
- **device_id** (int) - Optional - The CUDA device ID. Defaults to 0.

### Returns
- **int**: The number of SMs on the device.

### Raises
- **RuntimeError**: If unable to get the device properties.
```

--------------------------------

### mark

Source: https://tilelang.com/_sources/autoapi/tilelang/tools/cuda/iket/frontend/index.rst.txt

Emit an IKET instant marker at the current program point.

```APIDOC
## function mark(name, payload=None)

### Description
Emits an instant marker event at the current execution point.

### Parameters
- **name** (str) - The name of the marker.
- **payload** (Any, optional) - Optional payload to attach to the marker.
```

--------------------------------

### Enforce Non-Nullable Tensor Input

Source: https://tilelang.com/compiler_internals/tensor_checks.html

Example where a tensor is accessed, requiring it to be non-NULL.

```python
@T.prim_func
def main(A: T.Tensor((M, K), dtype)):
    A[0] = 1
```

--------------------------------

### Macro Decorator Usage

Source: https://tilelang.com/autoapi/tilelang/language/eager/builder/index.html

Example of defining a macro and using it within a prim_func.

```python
@macro
def my_macro(x: T.int32) -> T.int32:
    return x ** 2
@prim_func
def my_func(A: T.Tensor((10,), T.int32), B: T.Tensor((10,), T.int32)):
    with T.Kernel(1) as _:
        for i in T.serial(10):
            B[i] = my_macro(A[i])
```

--------------------------------

### Clone the Tilelang Repository

Source: https://tilelang.com/_sources/get_started/Installation.md.txt

Download the source code and initialize submodules.

```bash
git clone --recursive https://github.com/tile-ai/tilelang
cd tilelang
```

--------------------------------

### Define a Primitive Function

Source: https://tilelang.com/_sources/tutorials/debug_tools_for_tilelang.md.txt

A simplified example of a primitive function definition in TileLang.

```python
@T.prim_func
def main(Q: T.Tensor(shape_q, dtype)):
    # ...existing code...
```

--------------------------------

### Example Clang Remark Output

Source: https://tilelang.com/_sources/autoapi/tilelang/contrib/hip_resource_info/index.rst.txt

The format of the clang remarks parsed by this module.

```text
remark: src.cc:9:0: Function Name: main_kernel [-Rpass-analysis=kernel-resource-usage]
remark: src.cc:9:0:     TotalSGPRs: 16 [-Rpass-analysis=kernel-resource-usage]
remark: src.cc:9:0:     VGPRs: 5 [-Rpass-analysis=kernel-resource-usage]
remark: src.cc:9:0:     ScratchSize [bytes/lane]: 0 [-Rpass-analysis=kernel-resource-usage]
remark: src.cc:9:0:     SGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
remark: src.cc:9:0:     VGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
...
```

--------------------------------

### tvm_access_ptr

Source: https://tilelang.com/_sources/autoapi/tilelang/language/tir/op/index.rst.txt

Gets the head access address with memory access pattern information.

```APIDOC
## tvm_access_ptr(ptype, data, offset, extent, rw_mask)

### Description
Get head access address with memory access pattern info.

### Parameters
- **ptype** (Expr) - The data type of pointer.
- **data** (DType*) - The data of pointer.
- **offset** (int) - The offset of pointer.
- **extent** (int) - The extent of pointer.
- **rw_mask** (int) - The read write mask.

### Returns
- **call** (PrimExpr) - The call expression.
```

--------------------------------

### Generate Configurations with Carver

Source: https://tilelang.com/tutorials/auto_tuning.html

Use the Carver framework to recommend tiling strategies based on architecture templates.

```python
# Configure Matmul template
arch = CUDA("cuda")
carve_template = MatmulTemplate(
    M=M,
    N=N,
    K=K,
    in_dtype="float16",
    out_dtype="float16",
    accum_dtype="float",
).with_arch(arch)

# Generate top-k optimization hints (topk=10 recommended)
roller_hints = carve_template.recommend_hints(topk=10)

# Configure candidate parameters
for hint in roller_hints:

    # ...existing code...

    config["block_M"] = block_m
    config["block_N"] = block_n
    config["block_K"] = hint.rstep[0]
    config["num_stages"] = hint.pipeline_stage
    config["thread_num"] = block_rows * block_cols * 32
    config["enable_rasteration"] = hint.rasterization_plan is not NoRasterization
```

--------------------------------

### get_registers_per_block

Source: https://tilelang.com/_sources/autoapi/tilelang/carver/arch/driver/cuda_driver/index.rst.txt

Get the maximum number of 32-bit registers available per block.

```APIDOC
## get_registers_per_block(device_id = 0)

### Description
Get the maximum number of 32-bit registers available per block.

### Parameters
#### Path Parameters
- **device_id** (int) - Optional - The ID of the CUDA device. Defaults to 0.
```

--------------------------------

### TVMFFIKernelAdapter.from_database

Source: https://tilelang.com/autoapi/tilelang/jit/adapter/tvm_ffi/index.html

Creates an adapter instance from a pre-compiled database entry.

```APIDOC
## from_database(params, result_idx, target, func_or_mod, host_kernel_source, device_kernel_source, kernel_lib_path, verbose=False, pass_configs=None, compile_flags=None)

### Description
Class method to instantiate a TVMFFIKernelAdapter using information retrieved from a kernel database.

### Parameters
- **params** (list[tvm.relax.TensorType]) - Required - List of tensor types.
- **result_idx** (list[int]) - Required - List of result indices.
- **target** (str) - Required - The target device.
- **func_or_mod** (tvm.tirx.PrimFunc | tilelang.tvm.IRModule) - Required - The function or module.
- **host_kernel_source** (CachedTextSource) - Required - Cached host source.
- **device_kernel_source** (CachedTextSource) - Required - Cached device source.
- **kernel_lib_path** (str) - Required - Path to the kernel library.
- **verbose** (bool) - Optional - Enable verbose logging.
```

--------------------------------

### Avoided C++ naming conventions

Source: https://tilelang.com/_sources/developer_guide/cpp_style.md.txt

Examples of naming patterns to avoid, such as lowerCamelCase for functions.

```cpp
Layout makeLinearLayout(Array<PrimExpr> shape);
bool isSharedBuffer(const Buffer& buffer);
Stmt Lower(const LowerArgs& T, arith::Analyzer* analyzer) const;
```

--------------------------------

### Analyze a Tiled GEMM Kernel

Source: https://tilelang.com/_sources/tools/analyzer.md.txt

Demonstrates the full workflow of defining a tiled GEMM kernel, generating its TIR, and running the analyzer to obtain performance metrics.

```python
import tilelang
import tilelang.language as T
from tilelang.carver.arch import CUDA
from tilelang.tools import Analyzer

M = N = K = 1024


@tilelang.jit
def matmul(A, B, block_M, block_N, block_K):
    A: T.Tensor((M, K), T.float16)
    B: T.Tensor((N, K), T.float16)
    C = T.empty((M, N), T.float16)

    with T.Kernel(
        T.ceildiv(N, block_N),
        T.ceildiv(M, block_M),
        threads=128,
    ) as (bx, by):
        A_shared = T.alloc_shared((block_M, block_K), T.float16)
        B_shared = T.alloc_shared((block_N, block_K), T.float16)
        C_local = T.alloc_fragment((block_M, block_N), T.float32)

        T.clear(C_local)
        for k in T.serial(T.ceildiv(K, block_K)):
            T.copy(A[by * block_M, k * block_K], A_shared)
            T.copy(B[bx * block_N, k * block_K], B_shared)
            T.gemm(A_shared, B_shared, C_local, transpose_B=True)
        T.copy(C_local, C[by * block_M, bx * block_N])

    return C


tir = matmul.get_tir(block_M=128, block_N=128, block_K=32)
device = CUDA("cuda")
result = Analyzer.analysis(tir, device)

print(f"FLOPs: {result.total_flops}")
print(f"Global bytes: {result.total_global_bytes}")
print(f"Estimated seconds: {result.estimated_time}")
print(f"Model peak TFLOPS: {result.expected_tflops}")
print(f"Model bandwidth (GB/s): {result.expected_bandwidth_GBps}")
```

--------------------------------

### Configuring StorageRewrite Inplace Detection

Source: https://tilelang.com/_sources/autoapi/tilelang/transform/pass_config/index.rst.txt

Example of how to enable the TL_STORAGE_REWRITE_DETECT_INPLACE configuration using PassContext.

```python
from tilelang.transform import PassContext, PassConfigKey

with PassContext(
    config={PassConfigKey.TL_STORAGE_REWRITE_DETECT_INPLACE.value: True}
):
    mod = tilelang.transform.StorageRewrite()(mod)
```

--------------------------------

### session(*, reset_events = True, override = True, disable_on_exit = True, output_dir = None, runtime_payloads = None, disable_cache = True)

Source: https://tilelang.com/_sources/autoapi/tilelang/tools/cuda/iket/session/index.rst.txt

Configures the IKET environment around kernel construction and compilation.

```APIDOC
## session(*, reset_events = True, override = True, disable_on_exit = True, output_dir = None, runtime_payloads = None, disable_cache = True)

### Description
Configure IKET around kernel construction and compilation.

### Parameters
- **reset_events** (bool) - Optional - Defaults to True.
- **override** (bool) - Optional - Defaults to True.
- **disable_on_exit** (bool) - Optional - Defaults to True.
- **output_dir** (str) - Optional - Defaults to None.
- **runtime_payloads** (bool) - Optional - Defaults to None.
- **disable_cache** (bool) - Optional - Defaults to True.
```

--------------------------------

### LDS Transpose Reads

Source: https://tilelang.com/_sources/autoapi/tilelang/language/builtin/index.rst.txt

Examples of performing LDS transpose reads for gfx950 hardware.

```python
>>> val = T.ds_read_tr16_b64(smem[i])
```

```python
>>> val = T.ds_read_tr8_b64(smem[i])
```

--------------------------------

### Vectorized(start, stop=None, *, annotations=None)

Source: https://tilelang.com/autoapi/tilelang/language/loop/index.html

Creates a vectorized loop frame. This is an alias for T.vectorized.

```APIDOC
## Vectorized(start, stop=None, *, annotations=None)

### Description
Creates a vectorized loop frame, allowing for SIMD-style iteration over a range. This function is an alias for T.vectorized.

### Parameters
- **start** (tvm.tirx.PrimExpr) - Required - The minimum value of iteration.
- **stop** (tvm.tirx.PrimExpr | None) - Optional - The maximum value of iteration.
- **annotations** (dict[str, Any] | None) - Optional - The optional annotations of the For statement.

### Returns
- **res** (tvm.tirx.script.builder.frame.ForFrame) - The resulting ForFrame.
```

--------------------------------

### range_push

Source: https://tilelang.com/_sources/autoapi/tilelang/tools/cuda/iket/frontend/index.rst.txt

Emit an IKET range-start event at the current program point.

```APIDOC
## function range_push(name, payload=None)

### Description
Starts a new range event at the current program point.

### Parameters
- **name** (str) - The name of the range.
- **payload** (Any, optional) - Optional payload to attach to the range start.
```

--------------------------------

### Serial Loop

Source: https://tilelang.com/_sources/programming_guides/language_basics.md.txt

Implement a standard for-loop using `T.serial(start, stop[, step])`.

```python
for i in T.serial(N):
    ...
```

--------------------------------

### tilelang.language.kernel.Kernel

Source: https://tilelang.com/autoapi/tilelang/language/kernel/index.html

Constructs a target-neutral kernel launch frame.

```APIDOC
## tilelang.language.kernel.Kernel(*blocks, threads=None, prelude=None)

### Description
Constructs a kernel launch frame in a target-neutral form. The launch nest is emitted as thread_binding loops, allowing the same kernel to be compiled for different backends.

### Parameters
- **blocks** (int) - Required - A list of extents (1-3 dimensions) representing gridDim.(x|y|z).
- **threads** (int) - Optional - An integer representing blockDim.x or a list of integers representing blockDim.(x|y|z). If -1, threadIdx.x binding is skipped.
- **prelude** (str) - Optional - Import C code injected before the generated kernel code.
```

--------------------------------

### retrieve_ptr

Source: https://tilelang.com/_sources/autoapi/tilelang/utils/language/index.rst.txt

Retrieves a pointer to the start of a buffer region, supporting Buffers, BufferRegions, and BufferLoads.

```APIDOC
## retrieve_ptr(obj, access_type='r', ignore_last_ndim=0)

### Description
Retrieve a pointer to the start of a (possibly sliced) buffer region. It handles different buffer-like objects by calculating offsets based on region minima or indices.

### Parameters
- **obj** (Buffer-like) - Required - The buffer, BufferRegion, or BufferLoad object.
- **access_type** (str) - Optional - TVM Buffer access mask (e.g., "r", "w", "rw"). Defaults to 'r'.
- **ignore_last_ndim** (int) - Optional - Number of dimensions to ignore when calculating the offset.
```

--------------------------------

### Generate Candidate Configurations with Combinatorial Traversal

Source: https://tilelang.com/_sources/tutorials/auto_tuning.md.txt

Use `itertools.product` to generate all combinations of parameters for auto-tuning. This is useful for exploring a wide range of configurations.

```python
import itertools

block_M = [64, 128, 256]
block_N = [64, 128, 256]
block_K = [32, 64]
num_stages = [0, 1, 2, 3]
thread_num = [128, 256]
enable_rasterization = [True, False]
_configs = list(
    itertools.product(
        block_M,
        block_N,
        block_K,
        num_stages,
        thread_num,
        enable_rasterization,
    ))

configs = [
    {
        "block_M": c[0],
        "block_N": c[1],
        "block_K": c[2],
        "num_stages": c[3],
        "thread_num": c[4],
        "enable_rasteration": c[5]
    } for c in _configs
]

```

--------------------------------

### get_max_dynamic_shared_size_bytes

Source: https://tilelang.com/_sources/autoapi/tilelang/carver/arch/driver/cuda_driver/index.rst.txt

Get the maximum dynamic shared memory size in bytes, kilobytes, or megabytes.

```APIDOC
## get_max_dynamic_shared_size_bytes(device_id = 0, format = 'bytes')

### Description
Get the maximum dynamic shared memory size in bytes, kilobytes, or megabytes.

### Parameters
#### Path Parameters
- **device_id** (int) - Optional - The ID of the CUDA device. Defaults to 0.
- **format** (str) - Optional - The unit of the returned value (e.g., 'bytes', 'KB', 'MB'). Defaults to 'bytes'.
```

--------------------------------

### Create Bank Swizzled Layouts

Source: https://tilelang.com/autoapi/tilelang/layout/swizzle/index.html

Examples of creating full, half, and quarter bank swizzled layouts for a given buffer.

```python
make_full_bank_swizzled_layout(buffer)
```

```python
make_half_bank_swizzled_layout(buffer)
```

```python
make_quarter_bank_swizzled_layout(buffer)
```

--------------------------------

### get_rocm_arch

Source: https://tilelang.com/_sources/autoapi/tilelang/contrib/rocm/index.rst.txt

Retrieves the AMD GPU architecture. It takes an optional `rocm_path` to the ROCm installation directory.

```APIDOC
## get_rocm_arch(rocm_path='/opt/rocm')

### Description
Utility function to get the AMD GPU architecture.

### Parameters
#### Arguments
- **rocm_path** (str) - The path to rocm installation directory.

### Returns
- **gpu_arch** (str) - The AMD GPU architecture.
```

--------------------------------

### cli_main(argv=None)

Source: https://tilelang.com/_sources/autoapi/tilelang/autodd/index.rst.txt

Main entry point for the command-line interface.

```APIDOC
### cli_main(argv=None)

Main entry point for CLI.

#### Parameters
- **argv** (list, optional) - Command line arguments.
```

--------------------------------

### Access Trace URL

Source: https://tilelang.com/tools/iket.html

Example URL format for accessing the generated HTML trace viewer.

```text
http://localhost:8080/iket_pid_0x....html
```

--------------------------------

### Preferred C++ naming conventions

Source: https://tilelang.com/_sources/developer_guide/cpp_style.md.txt

Examples of preferred PascalCase for functions and methods, and lower_snake_case for parameters.

```cpp
Layout MakeLinearLayout(Array<PrimExpr> shape);
bool IsSharedBuffer(const Buffer& buffer);

Stmt Lower(const LowerArgs& args, arith::Analyzer* analyzer) const;
LayoutMap InferLayout(const LayoutInferArgs& args, InferLevel level) const;
```

--------------------------------

### tilelang.autodd.cli_main

Source: https://tilelang.com/autoapi/tilelang/autodd/index.html

Entry point for the autodd command line interface.

```APIDOC
## tilelang.autodd.cli_main(argv=None)

### Description
Executes the main CLI logic for the autodd module.

### Parameters
- **argv** (Sequence[str] | None) - Optional - List of command line arguments.
```

--------------------------------

### Run the Pass Visualizer

Source: https://tilelang.com/_sources/autoapi/tilelang/tools/pass_visualizer/viewer/index.rst.txt

Execute the visualizer from the command line to generate an HTML report for a specific TileLang kernel script.

```bash
python -m tilelang.tools.pass_visualizer.viewer \
    tilelang/tools/pass_visualizer/examples/gemm_relu.py \
    --set M=1024 --set N=1024 --set K=1024 \
    --set block_M=128 --set block_N=128 --set block_K=32 \
    --out gemm_relu_passes.html
```

--------------------------------

### Define Replayable Bind Dependencies

Source: https://tilelang.com/_sources/programming_guides/software_pipeline.md.txt

Example of defining scalar dependencies that are replayed before consumer statements.

```python
base: T.int32 = ko * BK
offset: T.int32 = base + tx
T.copy(A[offset], A_shared[tx])
```

--------------------------------

### Handle Residual Tiles

Source: https://tilelang.com/_sources/programming_guides/control_flow.md.txt

Example of using T.Parallel with explicit boundary guards for 2D kernels.

```python
for i, j in T.Parallel(M, N):
    gi = by * BM + i
    gj = bx * BN + j
    if T.all_of(gi < M, gj < N):     # optional in many cases
        C[gi, gj] = A[gi, gj] + B[gi, gj]
```

--------------------------------

### Create an instant event marker

Source: https://tilelang.com/_sources/tools/iket.md.txt

Use this to record a single point in time during execution.

```python
iket.mark("load_inputs")
```

--------------------------------

### PTX MMA SM70 Operation

Source: https://tilelang.com/_sources/autoapi/tilelang/language/builtin/index.rst.txt

Example of invoking a PTX MMA operation for SM70 architecture.

```python
>>> T.ptx_mma_sm70(
...     "float16",
...     "m16n16k4",
...     "row",
...     "col",
...     "fp16",
...     "fp16",
...     "fp16",
...     A_local.data,
...     0,
...     B_local.data,
...     0,
...     C_local.data,
...     0,
... )
```

--------------------------------

### Use a Generic Mapping Class

Source: https://tilelang.com/_sources/autoapi/tilelang/autotuner/tuner/index.rst.txt

Example of using a generic class within a function signature.

```python
def lookup_name(mapping: Mapping[KT, VT], key: KT, default: VT) -> VT:
    try:
        return mapping[key]
    except KeyError:
        return default
```

--------------------------------

### tilelang.jit.compile

Source: https://tilelang.com/_sources/autoapi/tilelang/jit/index.rst.txt

Compiles a given TileLang PrimFunc with TVM and builds a JITKernel.

```APIDOC
## tilelang.jit.compile(func=None, out_idx=None, execution_backend=None, target=None, target_host=None, verbose=None, pass_configs=None, compile_flags=None)

### Description
Compile the given TileLang PrimFunc with TVM and build a JITKernel.

### Parameters
- **func** (tvm.tirx.PrimFunc, optional) - The TileLang TIR function to compile and wrap.
- **out_idx** (Union[List[int], int], optional) - Index(es) of the output tensors to return.
- **execution_backend** (Literal["auto", "tvm_ffi", "cython", "nvrtc", "torch", "cutedsl"], optional) - Execution backend to use for kernel execution.
- **target** (str, dict, or tvm.target.Target, optional) - Compilation target.
- **target_host** (str, dict, or tvm.target.Target, optional) - Target host for cross-compilation.
- **verbose** (bool, optional) - Whether to enable verbose output.
- **pass_configs** (dict, optional) - Additional keyword arguments to pass to the Compiler PassContext.
```
Tokens:
Raw