### Install OS-Level Prerequisites on Ubuntu/Debian

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

Installs necessary system packages for building tilelang from source on Ubuntu/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 Tilelang from GitHub Repository

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

Installs the latest development version of tilelang directly from its GitHub repository.

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

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

### 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 Development Requirements and Install Editable

Source: https://tilelang.com/get_started/Installation.html

Install development dependencies and then install TileLang in editable mode. This is useful for developers who need to recompile frequently.

```bash
pip install -r requirements-dev.txt

# For first time compilation
pip install -e . -v --no-build-isolation
```

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

### Install Tilelang from Prebuilt Wheel

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

Installs a specific prebuilt version of tilelang from a local wheel file. Ensure the filename matches the desired version and CUDA/OS compatibility.

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

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

### 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])
```

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

### Verify Tilelang Installation

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

Checks if tilelang has been installed correctly by importing it and printing its version.

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

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

### Clone and Install TileLang from Source

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

Clone the repository, update submodules, and install TileLang using pip with no build isolation. This method is suitable for development and allows for custom build configurations via environment variables.

```bash
git clone --recursive https://github.com/tile-ai/tilelang.git /opt/tilelang
cd /opt/tilelang
git submodule update --init --recursive
export CMAKE_ARGS="-DUSE_CUDA=OFF -DUSE_ROCM=ON -DROCM_PATH=/opt/rocm -DLLVM_CONFIG=${LLVM_CONFIG}"

pip install -e . -v --no-build-isolation --no-deps

pip install "apache-tvm-ffi>=0.1.6" "z3-solver>=4.13.0"
# If you already installed torch-c-dlpack-ext and hit `libtorch_cuda.so` errors:
# pip uninstall -y torch-c-dlpack-ext

# If you hit Cython compile errors like `PyLong_SHIFT`/`digit` not declared,
# disable the stable ABI (abi3) for editable builds:
# export CMAKE_ARGS="-DUSE_CUDA=OFF -DUSE_ROCM=ON -DROCM_PATH=/opt/rocm -DLLVM_CONFIG=${LLVM_CONFIG} -DSKBUILD_SABI_VERSION="
# pip install -e . -v --no-build-isolation --no-deps

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

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

### Access Docker Container and Verify Installation

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

Execute commands inside the running tilelang Docker container to verify the installation. This includes accessing the container and printing the installed tilelang version.

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

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

### Clone and Install TileLang from Source (ROCm)

Source: https://tilelang.com/get_started/Installation.html

Clone the TileLang repository and install it using pip with ROCm support enabled. This command assumes ROCm is installed at /opt/rocm and skips dependency resolution, expecting PyTorch to be pre-installed.

```bash
git clone --recursive https://github.com/tile-ai/tilelang.git /opt/tilelang
cd /opt/tilelang
git submodule update --init --recursive
export CMAKE_ARGS="-DUSE_CUDA=OFF -DUSE_ROCM=ON -DROCM_PATH=/opt/rocm -DLLVM_CONFIG=${LLVM_CONFIG}"

# Avoid pulling CUDA wheels / reinstalling torch by skipping dependency resolution.
# Assume torch is already installed in the container.
pip install -e . -v --no-build-isolation --no-deps

# Manually install required runtime deps when using --no-deps.
# Note: skip torch-c-dlpack-ext on ROCm (its wheel expects CUDA libs).
pip install "apache-tvm-ffi>=0.1.6" "z3-solver>=4.13.0"
# If you already installed torch-c-dlpack-ext and hit `libtorch_cuda.so` errors:
# pip uninstall -y torch-c-dlpack-ext
```

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

### Set PYTHONPATH and Verify Installation

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

Adds the tilelang repository root to the PYTHONPATH environment variable and verifies the installation by printing the version. This is for developers working directly from the source tree.

```bash
export PYTHONPATH=/path/to/tilelang:$PYTHONPATH
python -c "import tilelang; print(tilelang.__version__)"
```

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

### Install TileLang with Nightly Builds

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

Install the latest nightly build of TileLang using pip with a specified find-links URL. This provides access to the most recent features and bug fixes before official releases.

```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 in Development Mode

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

Installs tilelang in editable mode using pip. Changes to Python files are reflected immediately without reinstallation. C++ changes still require rebuilding.

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

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

### End-to-End Example: After InjectFenceProxy

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

This code demonstrates the same TileLang prim_func after the InjectFenceProxy pass. A `T.fence_proxy_async()` instruction has been inserted.

```python
@T.prim_func
def kernel():
    with T.Kernel(1):
        desc = T.decl_buffer((1,), "uint64", scope="local.descriptor")
        smem = T.decl_buffer((128,), "float16", scope="shared")
        T.initialize_wgmma_descriptor(desc, T.uint64(0), 2, 1, 32)
        smem[0] = T.float16(0)
        T.fence_proxy_async()
        T.ptx_wgmma_ss(
            "float16",
            "m64n64k16",
            T.bool(True),
            T.bool(True),
            "fp16",
            "fp16",
            "fp16",
            desc.data,
            T.int32(0),
            desc.data,
            T.int32(0),
            smem.data,
            T.int32(0),
            T.bool(True),
            1,
            1,
        )
```

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

### Install with Custom TVM Path

Source: https://tilelang.com/get_started/Installation.html

Use this command to install tilelang when you have an existing TVM codebase. Set the TVM_ROOT environment variable to your TVM repository location.

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

```

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

### TileLang Layout Inference Text Output Example

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

Example of textual output from TileLang's layout inference, showing the mapping between buffer shapes, thread indices, and data access patterns.

```text
C_local inferenced layout:
  Shape: [32, 32] -> [8]
  Thread: _j // 16 * 64 + _i // 16 * 32 + _i % 8 * 4 + _j % 8 // 2
  Index:  [_j % 16 // 8 * 4 + _i % 16 // 8 * 2 + _j % 2]
```

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

### 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.
```

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

### 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)
```

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

### Start Profile Intrinsic

Source: https://tilelang.com/autoapi/tilelang/language/tir/op/index.html

Starts profiling for a given intrinsic ID.

```APIDOC
## start_profile_intrinsic(id: int)

### Description
Start profile intrinsic.

### Parameters
* **id**: The intrinsic id.

### Returns
**call** – The call expression.

### Return type
PrimExpr
```

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

### 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.
```

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

### main

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

Entry point for running tests.

```APIDOC
## main()

### Description
Main function to execute tests.
```

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

### End-to-End Example: Before InjectFenceProxy

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

This code shows a TileLang prim_func before the InjectFenceProxy pass is applied. It includes descriptor initialization and a `ptx_wgmma_ss` call.

```python
@T.prim_func
def kernel():
    with T.Kernel(1):
        desc = T.decl_buffer((1,), "uint64", scope="local.descriptor")
        smem = T.decl_buffer((128,), "float16", scope="shared")
        T.initialize_wgmma_descriptor(desc, T.uint64(0), 2, 1, 32)
        smem[0] = T.float16(0)
        T.ptx_wgmma_ss(
            "float16",
            "m64n64k16",
            T.bool(True),
            T.bool(True),
            "fp16",
            "fp16",
            "fp16",
            desc.data,
            T.int32(0),
            desc.data,
            T.int32(0),
            smem.data,
            T.int32(0),
            T.bool(True),
            1,
            1,
        )
```

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

### RewriteName Example

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

Example of a NodeTransformer that rewrites all occurrences of name lookups (e.g., 'foo') to 'data['foo']'.

```APIDOC
## RewriteName(NodeTransformer)

### Description
An example transformer that rewrites all occurrences of name lookups (``foo``) to ``data['foo']``.

### Methods
#### visit_Name(node)

### Code Example
```python
class RewriteName(NodeTransformer):

    def visit_Name(self, node):
        return Subscript(
            value=Name(id='data', ctx=Load()),
            slice=Constant(value=node.id),
            ctx=node.ctx
        )
```
```

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

### compile

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

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

```APIDOC
## compile

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

### Parameters
- `func` (PrimFunc) - The TileLang PrimFunc to compile.
- `out_idx` (int, optional) - The output index.
- `execution_backend` (ExecutionBackend, optional) - The execution backend to use.
- `target` (str, optional) - The compilation target.
```

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

### Build Tilelang from Source with Pip-Provided CUDA

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

Installs tilelang from source using pip-provided CUDA toolchain packages. This option avoids the need for a host CUDA installation. Requires installing development requirements and specific CUDA 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
```

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

### tilelang.language.kernel.Kernel

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

Tools to quickly construct a GPU kernel launch frame.

```APIDOC
## tilelang.language.kernel.Kernel

Tools to quickly construct a GPU kernel launch frame.

### Parameters

- **blocks** (_int_) – A list of extent, can be 1-3 dimension, representing gridDim.(x|y|z)
- **threads** (_int_) – A integer representing blockDim.x Or a list of integers representing blockDim.(x|y|z) if the value is -1, we skip the threadIdx.x binding.
- **cluster_dims** (_int_ _|__tuple_ _[__int_ _,__int_ _,__int_ _]__|__list_ _[__int_ _]__|__None_) – The cluster dimensions for SM90+ cluster launch. For example, use 2 or (2, 1, 1) to create 2-CTA clusters. When specified, the kernel will be launched using cudaLaunchKernelEx with cudaLaunchAttributeClusterDimension.
- **prelude** (_str_) – The import c code of the kernel, will be injected before the generated kernel code.
- **is_cpu** (_bool_)

### Returns

- **res** – The result LaunchThreadFrame.
  Return type: Tuple[frame.LaunchThreadFrame]

### Examples

Create a 1-D CUDA kernel launch and unpack the single block index:
```python
with T.Kernel(T.ceildiv(N, 128), threads=128) as bx:
    # bx is the blockIdx.x binding (also iterable as (bx,))
    ...
```

Launch a 2-D grid while requesting two thread dimensions:
```python
with T.Kernel(grid_x, grid_y, threads=(64, 2)) as (bx, by):
    tx, ty = T.get_thread_bindings()
    ...
```

Emit a CPU kernel where thread bindings are skipped:
```python
with T.Kernel(loop_extent, is_cpu=True) as (i,):
    ...
```
```

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

### Install Latest Tilelang via Pip

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

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

```bash
pip install tilelang
```

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

### Main Execution Functions

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

Entry points for running tilelang tasks and CLI.

```APIDOC
## main

### Description
Main function for executing tilelang tasks.

## cli_main

### Description
Main entry point for the tilelang command-line interface.
```

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

### Example of Allowed Pipelined Loop Outside Parallel

Source: https://tilelang.com/autoapi/tilelang/analysis/nested_loop_checker/index.html

This example shows an allowed loop nesting where a T.Pipelined loop contains a T.Parallel loop.

```python
for i in T.Pipelined(K):
    for j in T.Parallel(N): # allowed, ok
        …
```

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

### DSLMutator Transformer Example

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

Illustrates a DSLMutator transformer that rewrites name lookups to data['foo']. This is a specific example of AST transformation for data access.

```python
Here is an example transformer that rewrites all occurrences of name lookups
   (``foo``) to ``data['foo']``::

      class RewriteName(NodeTransformer):

          def visit_Name(self, node):
              return Subscript(
                  value=Name(id='data', ctx=Load()),
                  slice=Constant(value=node.id),
                  ctx=node.ctx
              )
```

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

### Run TileLang Docker Container

Source: https://tilelang.com/get_started/Installation.html

Start a tilelang Docker container with GPU access and volume mounting. Adjust the host path and container name as needed.

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

```

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

### Run TileLang in Developer Mode

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

Execute a Python command to import TileLang and observe developer mode logs, which indicate the loading path from the development root.

```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
```

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

### Example of Forbidden Pipelined Loop Inside Parallel

Source: https://tilelang.com/autoapi/tilelang/analysis/nested_loop_checker/index.html

This example illustrates a forbidden nested loop pattern where a T.Pipelined loop is nested inside a T.Parallel loop.

```python
for i in T.Parallel(M):
    for j in T.Pipelined(K): # forbidden!
        …
```

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

### tilelang.autodd.cli_main

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

Parses command-line arguments and runs the main TileLang process.

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

### Description
Parses command-line arguments and runs the main TileLang process.

### Parameters
#### Path Parameters
- **argv** (Sequence[str] | None) - A list of command-line arguments. If None, sys.argv is used.

### Return type
None
```

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

### tilelang.transform.SplitHostDevice()

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

Splits functions into host and device components, even for empty kernels.

```APIDOC
## tilelang.transform.SplitHostDevice()

### Description
Split host/device functions even for empty kernels.

### Returns
- **fpass** (tvm.transform.Pass) - The result pass.
```

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

### Generic Type Usage Example

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

Illustrates how to use a generic type, such as the 'Mapping' example, in function signatures to provide type hints for generic collections. This enhances code readability and maintainability.

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

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

### Build Tilelang from Source with Host CUDA

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

Clones the tilelang repository (including submodules) and installs it using pip with verbose output. Requires a host CUDA toolchain.

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

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

### Initialize TileLang Environment

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

Sets up the TileLang environment, including defining export macros and initializing an error buffer. This function should be called before other TileLang operations.

```c++
"""
      #ifdef _WIN32
      #define TL_EXPORT __declspec(dllexport)
      #else
      #define TL_EXPORT
      #endif
      
      #define ERROR_BUF_SIZE 1024
      static char error_buf[ERROR_BUF_SIZE];
      
      extern "C" TL_EXPORT const char* get_last_error() {
          return error_buf;
      }
      
      extern "C" TL_EXPORT int init() {
          error_buf[0] = '\0';
          {0}
          return 0;
      }
      """
```

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

### Example of Forbidden Parallel Nested Loops

Source: https://tilelang.com/autoapi/tilelang/analysis/nested_loop_checker/index.html

This example demonstrates a forbidden nested loop pattern where a T.Parallel loop is directly followed by a T.copy operation, which is considered a TileOp with 'parallel' behaviors.

```python
for i in T.Parallel(M):
    T.copy(A, B) # forbidden!
```

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

### Run in Editable Mode

Source: https://tilelang.com/get_started/Installation.html

Execute a Python command to import tilelang and observe development logs, indicating the loading of libraries from the development root.

```python
python -c 'import tilelang'
```

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

### Hint Class

Source: https://tilelang.com/autoapi/tilelang/carver/roller/hint/index.html

Configuration for tiling and computation hints.

```APIDOC
## Hint

### Description
Configuration for tiling and computation hints, specifying how operations should be scheduled and optimized.

### Attributes
- `arch`: Target architecture.
- `use_tc`: Whether to use tensor cores.
- `block`: Block dimensions.
- `thread`: Thread dimensions.
- `warp`: Warp dimensions.
- `rstep`: Reduction step size.
- `reduce_thread`: Thread reduction configuration.
- `rasterization_plan`: Plan for rasterization.
- `cached_tensors`: List of tensors to cache.
- `output_strides`: Strides for the output tensor.
- `schedule_stages`: Number of scheduling stages.
- `block_reduction_depth`: Depth of block reduction.
- `split_k_factor`: Factor for splitting the K dimension.
- `vectorize`: Vectorization factor.
- `pipeline_stage`: Pipeline stage configuration.
- `use_async`: Whether to use asynchronous operations.
- `opt_shapes`: Optimized shapes for operations.
- `intrin_info`: Intrinsic information object.
- `shared_scope`: Scope for shared memory.
- `pass_context`: Context for passing information.
- `raxis_order`: Order of reduction axes.
- `step`: Step size for operations.

### Methods
- `to_dict()`: Converts the Hint object to a dictionary.
- `from_dict(data)`: Creates a Hint object from a dictionary.
- `tensorcore_legalization()`: Legalizes the hint for tensor core usage.
- `complete_config()`: Completes the hint configuration.
- `__repr__()`: Returns a string representation of the Hint object.
```

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

### Build TileLang using python-build

Source: https://tilelang.com/get_started/Installation.html

Build TileLang using the `python-build` tool. This command will generate a wheel file, potentially including SDK and version information in its name.

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

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

### create_staticlib()

Source: https://tilelang.com/autoapi/tilelang/contrib/cc/index.html

Create a static library from input files.

```APIDOC
## create_staticlib(_output_, _inputs_, _ar_=None)

### Description
Create static library.

### Parameters
- **output** (str) - The target shared library.
- **inputs** (List[str]) - List of inputs files. Each input file can be a tarball of objects or an object file.
- **ar** (Optional[str]) - Path to the ar command to be used.
```

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

### Example of Allowed Strict Continuous Parallel Loops

Source: https://tilelang.com/autoapi/tilelang/analysis/nested_loop_checker/index.html

This example shows an allowed nested loop pattern where two T.Parallel loops are strictly continuous, indicating they can be fused into a single T.Parallel loop.

```python
for i in T.Parallel(M):
    for j in T.Parallel(N):
        …
```

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

### Install System Libraries for ROCm Container Build

Source: https://tilelang.com/get_started/Installation.html

Install necessary system libraries, build tools, and Python packages within a ROCm container. This is required for building tilelang from source in a ROCm environment.

```bash
# Inside the container (as root)
apt-get update && apt-get install -y --no-install-recommends \
  build-essential git wget curl ca-certificates gnupg \
  libgtest-dev libgmock-dev \
  libprotobuf-dev protobuf-compiler libgflags-dev libsqlite3-dev \
  python3 python3-dev python3-setuptools python3-pip \
  gcc libtinfo-dev zlib1g-dev libedit-dev libxml2-dev \
  cmake ninja-build pkg-config libstdc++6 \
  && rm -rf /var/lib/apt/lists/*

# Prefer the container venv (avoid system pip)
export PATH="/opt/venv/bin:${PATH}"

# Build GoogleTest static libs (Ubuntu package ships sources only)
cmake -S /usr/src/googletest -B /tmp/build-gtest -DBUILD_GTEST=ON -DBUILD_GMOCK=ON -DCMAKE_BUILD_TYPE=Release
cmake --build /tmp/build-gtest -j"$(nproc)"
cp -v /tmp/build-gtest/lib/*.a /usr/lib/x86_64-linux-gnu/
rm -rf /tmp/build-gtest

# Keep setuptools < 80 (compat with some base images)
pip install --upgrade "setuptools>=77.0.3,<80" wheel cmake ninja scikit-build-core

# Locate ROCm llvm-config (install LLVM 18 if missing)
LLVM_CONFIG_PATH=""
for p in /opt/rocm/llvm/bin/llvm-config /opt/rocm/llvm-*/bin/llvm-config /opt/rocm-*/llvm*/bin/llvm-config; do
  if [ -x "$p" ]; then LLVM_CONFIG_PATH="$p"; break; fi
done
if [ -z "$LLVM_CONFIG_PATH" ]; then
  echo "ROCm llvm-config not found; installing LLVM 18..."
  curl -fsSL https://apt.llvm.org/llvm.sh -o /tmp/llvm.sh
  chmod +x /tmp/llvm.sh
  /tmp/llvm.sh 18
  LLVM_CONFIG_PATH="$(command -v llvm-config-18)"
  if [ -z "$LLVM_CONFIG_PATH" ]; then
    echo "ERROR: llvm-config-18 not found after install"
    exit 1
  fi
fi
export LLVM_CONFIG="$LLVM_CONFIG_PATH"
export PATH="$(dirname "$LLVM_CONFIG"):/usr/local/bin:${PATH}"

# Optional shim for tools that expect llvm-config-16
mkdir -p /usr/local/bin
printf "#!/usr/bin/env bash\nexec \"%s\" \"\$@\"\n" "$LLVM_CONFIG_PATH" > /usr/local/bin/llvm-config-16
chmod +x /usr/local/bin/llvm-config-16

# TVM Python bits need Cython (for system Python used by the build)
pip install --no-cache-dir "cython>=0.29.36,<3.0"

# Clone + build TileLang (ROCm)

```

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

### Provide Inputs via Context Manager

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

Use a context manager to supply fixed inputs for autotuning across different configurations. This ensures consistent data for benchmarking.

```python
with set_autotune_inputs(A, B, C):
    tuned = matmul(M, N, K)
```

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

### 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
]

```

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

### WarpSpecialize Function Examples

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

Demonstrates how to use the WarpSpecialize function to construct warp group frames based on thread indices. The examples show conditional thread assignments based on warp group indices.

```python
T.ws(0)
# if tx < 128
```

```python
T.ws(1)
# if tx >= 128 and tx < 256
```

```python
T.ws(0, 1)
# if tx < 128 or (tx >= 128 and tx < 256)
```

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

### JITKernel Instance Methods

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

Methods for invoking the compiled function, profiling, and retrieving source code.

```APIDOC
## JITKernel Instance Methods

### __call__

#### Description
Invokes the compiled function with the given arguments.

#### Parameters
- ***args** (Any) – Positional arguments for the function.
- ****kwds** (Any) – Keyword arguments for the function.

#### Returns
The result of the function execution.

#### Return type
Any

### get_profiler

#### Description
Creates a profiler to benchmark the compiled runtime module.

#### Parameters
- **tensor_supply_type** (TensorSupplyType, optional) - The type of input tensors to supply for profiling (default: TensorSupplyType.Auto).

#### Returns
A Profiler instance for benchmarking the runtime module.

#### Return type
Profiler

### get_kernel_source

#### Description
Returns the source code of the compiled kernel function.

#### Returns
The source code of the compiled kernel function.

#### Return type
str

#### Parameters
- **kernel_only** (bool) - If True, returns only the kernel source.

### get_host_source

#### Description
Returns the source code of the host function.

#### Return type
str

### run_once

#### Parameters
- **func** (Callable | None) - The function to run.

#### Return type
None

### show_source

#### Description
Print generated source code to stdout.

#### Parameters
- **which** (Literal["kernel", "host", "both"], optional) – Select which source to print. Defaults to “kernel”.

#### Return type
None

#### Examples
```python
jit_kernel.show_source()            # print kernel source
jit_kernel.show_source("host")      # print host source
jit_kernel.show_source("both")      # print both sources
```

### export_sources

#### Description
Export generated source code to files.

#### Parameters
- **kernel_path** (Optional[str]) – Destination file path to write the kernel source. If None, skips writing kernel code.
- **host_path** (Optional[str]) – Destination file path to write the host source. If None, skips writing host code.
```

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

### Install with Custom TVM Path

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

Use this command to install tilelang when you have an existing TVM codebase. Set the TVM_ROOT environment variable to your TVM repository location. Note that this may require additional path configurations.

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

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

### Launch Work with T.Kernel Context

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

Declare a launch context and create block/thread bindings using `with T.Kernel(...)`. For GPU backends, specify grid and threads per block. Use structured loops like `T.serial`, `T.unroll`, `T.Parallel`, or `T.Pipelined` inside the kernel.

```python
with T.Kernel(grid_x, grid_y, threads=128) as (bx, by):
    ...  # bx/by are blockIdx.x/y

```

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

### Compile and Benchmark with AutoTuner

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

Configure JIT compilation and benchmarking settings using `AutoTuner`. The `run` method executes the tuning process, and the result contains the optimized kernel.

```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)


```

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

### Install System Libraries for ROCm Build

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

Install necessary system libraries, build tools, and Python packages within a ROCm container for building tilelang from source. This command is intended to be run as root inside the container.

```bash
# Inside the container (as root)
apt-get update && apt-get install -y --no-install-recommends \
  build-essential git wget curl ca-certificates gnupg \
  libgtest-dev libgmock-dev \
  libprotobuf-dev protobuf-compiler libgflags-dev libsqlite3-dev \
  python3 python3-dev python3-setuptools python3-pip \
  gcc libtinfo-dev zlib1g-dev libedit-dev libxml2-dev \
  cmake ninja-build pkg-config libstdc++6 \
  && rm -rf /var/lib/apt/lists/*

# Prefer the container venv (avoid system pip)
export PATH="/opt/venv/bin:${PATH}"

# Build GoogleTest static libs (Ubuntu package ships sources only)
cmake -S /usr/src/googletest -B /tmp/build-gtest -DBUILD_GTEST=ON -DBUILD_GMOCK=ON -DCMAKE_BUILD_TYPE=Release
cmake --build /tmp/build-gtest -j"$(nproc)"
cp -v /tmp/build-gtest/lib/*.a /usr/lib/x86_64-linux-gnu/
rm -rf /tmp/build-gtest

# Keep setuptools < 80 (compat with some base images)
pip install --upgrade "setuptools>=77.0.3,<80" wheel cmake scikit-build-core

# Locate ROCm llvm-config (install LLVM 18 if missing)
LLVM_CONFIG_PATH=""
for p in /opt/rocm/llvm/bin/llvm-config /opt/rocm/llvm-*/bin/llvm-config /opt/rocm-*/llvm*/bin/llvm-config; do
  if [ -x "$p" ]; then LLVM_CONFIG_PATH="$p"; break; fi
done
if [ -z "$LLVM_CONFIG_PATH" ]; then
  echo "ROCm llvm-config not found; installing LLVM 18..."
  curl -fsSL https://apt.llvm.org/llvm.sh -o /tmp/llvm.sh
  chmod +x /tmp/llvm.sh
  /tmp/llvm.sh 18
  LLVM_CONFIG_PATH="$(command -v llvm-config-18)"
  if [ -z "$LLVM_CONFIG_PATH" ]; then
    echo "ERROR: llvm-config-18 not found after install"
    exit 1
  fi
fi
export LLVM_CONFIG="$LLVM_CONFIG_PATH"
export PATH="$(dirname "$LLVM_CONFIG"):/usr/local/bin:${PATH}"

# Optional shim for tools that expect llvm-config-16
mkdir -p /usr/local/bin
printf "#!/usr/bin/env bash\nexec \"%s\" \"\$@\"\n" "$LLVM_CONFIG_PATH" > /usr/local/bin/llvm-config-16
chmod +x /usr/local/bin/llvm-config-16

# TVM Python bits need Cython (for system Python used by the build)
pip install --no-cache-dir "cython>=0.29.36,<3.0"

# Clone + build TileLang (ROCm)

```

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

### Initialize Tilelang JIT Environment

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

Initializes the tilelang JIT environment, setting up error reporting and potentially other configurations. This function is designed to be called once at the start of the application.

```cpp
"""
#ifdef _WIN32
#define TL_EXPORT __declspec(dllexport)
#else
#define TL_EXPORT
#endif

#define ERROR_BUF_SIZE 1024
static char error_buf[ERROR_BUF_SIZE];

extern "C" TL_EXPORT const char* get_last_error() {
    return error_buf;
}

extern "C" TL_EXPORT int init() {
    error_buf[0] = '\0';
    {0}
    return 0;
}
"""
```

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

### AutoDD Example: Minimize Buggy GEMM

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

Demonstrates using AutoDD to reduce a complex TileLang program containing a GEMM shape mismatch bug to its minimal reproducible form. This example specifies multiple parallel jobs.

```bash
python -m tilelang.autodd buggy_matmul.py --err-msg "Dimension mismatch" -o minimized.py -j 4

```

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

### get_block

Source: https://tilelang.com/autoapi/tilelang/carver/common_schedules/index.html

Get the target block from a schedule.

```APIDOC
## get_block(sch, blocks, name)

### Description
Get the target block from a schedule.

### Parameters
#### Path Parameters
- **sch** (tir.Schedule) - Required - The TIR schedule used to get target block.
- **name** (str) - Required - The name of the target block.
- **blocks** (list[tilelang.carver.analysis.BlockInfo]) - Required

### Returns
- **target_block** (BlockRV) - The target block.
```

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

### Build Native Extension via CMake and Make

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

Builds the tilelang native extension (`libtilelang.so`) using CMake and Make. This is a prerequisite for working from the source tree via PYTHONPATH.

```bash
mkdir -p build
cd build
cmake .. -DUSE_CUDA=ON
make -j
```

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

### get_pass_context

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

Get the current pass context.

```APIDOC
## get_pass_context

### Description
Get the current pass context.

### Function Signature
`get_pass_context()`

### Returns
- **pass_context** - The current pass context.
```

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

### GemmWarpPolicy Examples

Source: https://tilelang.com/autoapi/tilelang/tileop/base/index.html

Demonstrates how to use GemmWarpPolicy to specify warp partitioning strategies for GEMM operations. Use from_block_row_cols to create policies based on row/column distribution of warps.

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

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

### create_executable()

Source: https://tilelang.com/autoapi/tilelang/contrib/cc/index.html

Create an executable binary from object files.

```APIDOC
## create_executable(_output_, _objects_, _options_=None, _cc_=None, _cwd_=None, _ccache_env_=None)

### Description
Create executable binary.

### Parameters
- **output** (str) - The target executable.
- **objects** (List[str]) - List of object files.
- **options** (Optional[List[str]]) - The list of additional options string.
- **cc** (Optional[str]) - The compiler command.
- **cwd** (Optional[str]) - The urrent working directory.
- **ccache_env** (Optional[Dict[str, str]]) - The environment variable for ccache. Set None to disable ccache by default.
```

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

### Run TileLang Docker Container

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

Start a tilelang Docker container with GPU access and shared memory configured. Adjust the volume mount path to your local tilelang directory. This command starts the container in detached mode with zsh as the shell.

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

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

### start_profile_intrinsic

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

Starts profiling for a given intrinsic ID.

```APIDOC
## start_profile_intrinsic(id)

### Description
Start profile intrinsic.

### Parameters
#### Positional Arguments
- **id** (int) - The intrinsic id.

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

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

### WMMAIntrinEmitter.get_store_index_map

Source: https://tilelang.com/autoapi/tilelang/rocm/intrinsics/wmma_macro_generator/index.html

Gets the index map for store operations.

```APIDOC
## WMMAIntrinEmitter.get_store_index_map()

### Description
Retrieves the index mapping used for store operations within WMMA.

### Method
N/A (Method Call)

### Endpoint
N/A

### Parameters
N/A

### Request Example
N/A

### Response
N/A
```

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

### WMMAIntrinEmitter.get_ldmatrix_index_map

Source: https://tilelang.com/autoapi/tilelang/rocm/intrinsics/wmma_macro_generator/index.html

Gets the index map for `ldmatrix` operations.

```APIDOC
## WMMAIntrinEmitter.get_ldmatrix_index_map()

### Description
Retrieves the index mapping used for `ldmatrix` (load matrix) operations within WMMA.

### Method
N/A (Method Call)

### Endpoint
N/A

### Parameters
N/A

### Request Example
N/A

### Response
N/A
```

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

### Naive TileLang GEMV Kernel Implementation

Source: https://tilelang.com/deeplearning_operators/gemv.html

A naive GEMV kernel implemented in TileLang, adapting a GEMM tiling strategy. This provides a foundation for more optimized kernels.

```python
def naive_gemv(
    N: int,
    K: int,
    BLOCK_N: int,
    BLOCK_K: int,
    dtype: str = "float16",
    accum_dtype: str = "float",
):

    @T.prim_func
    def main(
            A: T.Buffer((K,), dtype),
            B: T.Buffer((N, K), dtype),
            C: T.Buffer((N,), dtype),
    ):
        with T.Kernel(T.ceildiv(N, BLOCK_N)) as bn:
            tn = T.get_thread_binding(0)  # tn = threadIdx.x
            A_shared = T.alloc_shared((BLOCK_K,), dtype)
            B_shared = T.alloc_shared((BLOCK_N, BLOCK_K), dtype)
            C_reg = T.alloc_local((1,), accum_dtype)
            T.clear(C_reg)
            for bk in T.serial(T.ceildiv(K, BLOCK_K)):
                for tk in T.serial(BLOCK_K):
                    A_shared[tk] = A[bk * BLOCK_K + tk]
                    B_shared[tn, tk] = B[bn * BLOCK_N + tn, bk * BLOCK_K + tk]
                for tk in T.serial(BLOCK_K):
                    C_reg[0] += A_shared[tk].astype(accum_dtype) * B_shared[tn,
                                                                            tk].astype(accum_dtype)
            C[bn * BLOCK_N + tn] = C_reg[0]

    return main

```

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

### get_index_map

Source: https://tilelang.com/autoapi/tilelang/carver/matmul_analysis/index.html

Gets the index maps for a given block.

```APIDOC
## get_index_map(block, layout)

### Description
Gets the index maps for a given block.

### Parameters
#### Path Parameters
- **block** (any) - Description not available
- **layout** (any) - Optional - Description not available
```

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

### Build TileLang with Custom CMake Arguments

Source: https://tilelang.com/get_started/Installation.html

Example of building TileLang with specific CMake arguments, such as disabling CUDA and enabling ROCm support, and disabling the stable ABI for editable builds to avoid Cython compile errors.

```bash
export CMAKE_ARGS="-DUSE_CUDA=OFF -DUSE_ROCM=ON -DROCM_PATH=/opt/rocm -DLLVM_CONFIG=${LLVM_CONFIG} -DSKBUILD_SABI_VERSION="
pip install -e . -v --no-build-isolation --no-deps
```

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

### 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.
```

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

### 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 Programmatic Stream Serialization

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

Sets up launch configuration attributes to enable programmatic stream serialization. This involves defining the attribute ID and value for serialization.

```python
num_attrs = 1
attrs = [CUlaunchAttribute()]
attrs[0].id = CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_STREAM_SERIALIZATION
attrs[0].value.programmaticStreamSerializationAllowed = 1

config.numAttrs = num_attrs
config.attrs = attrs
```

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

### create_staticlib

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

Creates a static library from a list of input files.

```APIDOC
## create_staticlib

### Description
Create static library.

### Parameters
- **output** (str) - The target shared library.
- **inputs** (List[str]) - List of inputs files. Each input file can be a tarball of objects or an object file.
- **ar** (Optional[str]) - Path to the ar command to be used.
```

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

### infinity

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

Get the infinity value for a given data type.

```APIDOC
## infinity(dtype, span = None)

### Description
infinity value of dtype

### Parameters
#### Path Parameters
- **dtype** (str) - The data type.
- **span** (Optional[Span]) - The location of this operator in the source code.

### Returns
- **value** (tvm.Expr) - The infinity value of dtype.
```

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

### max_value

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

Get the maximum value for a given data type.

```APIDOC
## max_value(dtype, span = None)

### Description
maximum value of dtype

### Parameters
#### Path Parameters
- **dtype** (str) - The data type.
- **span** (Optional[Span]) - The location of this operator in the source code.

### Returns
- **value** (tvm.Expr) - The maximum value of dtype.
```

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

### min_value

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

Get the minimum value for a given data type.

```APIDOC
## min_value(dtype, span=None)

### Description
minimum value of dtype

### Parameters
#### Path Parameters
- **dtype** (str) - The data type.
- **span** (Optional[Span]) - The location of this operator in the source code.

### Returns
- **value** (tvm.Expr) - The minimum value of dtype.
```