Tilelang Skill

TileLang development

Use the local TileLang checkout as the primary source for APIs, implementation patterns, and repository-specific build instructions. Prefer current source and examples over remembered signatures because the DSL, compiler, and backend integrations evolve quickly.

Locate the checkout

Resolve the directory containing this SKILL.md, then use its repos/tilelang/ child. The installer links that path to agent-gpu-skills/third_party/tilelang/ or to the checkout supplied through TILELANG_REPO.

In commands below, replace TILELANG_REPO with the resolved absolute path:

TILELANG_REPO=/absolute/path/to/tilelang-skill/repos/tilelang

If the checkout is missing, run this from the agent-gpu-skills repository and reinstall the Skill:

bash update-repos.sh tilelang
bash install.sh --skill tilelang-skill

The default checkout is a source-reference snapshot. It does not initialize TileLang's nested third-party submodules. Use a complete development checkout through TILELANG_REPO when building TileLang or changing its compiler.

Choose the source surface

Read quick-reference.md when mapping an operation, backend, or compiler question to a concrete path in the validated checkout.

Query workflow

1. Classify the task as kernel DSL, JIT/runtime, autotuning/profiling, or compiler/backend work.
2. Find the closest current example for the operation, target backend, architecture, dtype, and shape regime.
3. Verify each DSL operation against its Python definition, then trace lowering or backend code only as far as the question requires.
4. Preserve the workload's shapes, dtypes, layouts, memory scopes, target, launch configuration, and numerical contract.
5. Establish correctness and confirm the actual generated path before tuning.

Discover current documentation and examples before relying on a remembered filename:

find "$TILELANG_REPO/docs" -type f -name '*.md' | sort
find "$TILELANG_REPO/examples" -type f -name '*.py' | sort

Search kernel patterns and their definitions together:

rg -n '@tilelang\.jit|@T\.prim_func|T\.Kernel|T\.Pipelined' \
  "$TILELANG_REPO/examples"

rg -n 'T\.copy|T\.gemm|T\.alloc_shared|T\.alloc_fragment' \
  "$TILELANG_REPO/examples"

rg -n '^def (copy|gemm|alloc_shared|alloc_fragment)|class Kernel' \
  "$TILELANG_REPO/tilelang/language"

Trace JIT, lowering, and target selection:

rg -n 'def (compile|lower)|class .*Kernel|class .*Adapter' \
  "$TILELANG_REPO/tilelang/jit" \
  "$TILELANG_REPO/tilelang/engine"

rg -n 'target|backend|codegen' \
  "$TILELANG_REPO/tilelang/backend" \
  "$TILELANG_REPO/tilelang/cuda" \
  "$TILELANG_REPO/tilelang/rocm" \
  "$TILELANG_REPO/tilelang/metal" \
  "$TILELANG_REPO/tilelang/cpu"

Trace compiler transformations and generated source:

rg -n 'Pass|Lower|Legalize|Layout' \
  "$TILELANG_REPO/src/transform" \
  "$TILELANG_REPO/tilelang/engine"

rg -n 'CodeGen|codegen|Build' \
  "$TILELANG_REPO/src/cuda" \
  "$TILELANG_REPO/src/rocm" \
  "$TILELANG_REPO/src/metal" \
  "$TILELANG_REPO/src/cpu"

Search tuning, profiling, and nearby tests:

rg -n 'autotune|Profiler|do_bench|get_profiler' \
  "$TILELANG_REPO/examples" \
  "$TILELANG_REPO/tilelang/autotuner" \
  "$TILELANG_REPO/tilelang/profiler"

rg -n 'copy|gemm|pipelined|layout' \
  "$TILELANG_REPO/testing/python/language" \
  "$TILELANG_REPO/testing/python"

Implementation discipline

Keep these layers separate during diagnosis:

TileLang Python kernel and compile-time parameters
  → TileLang IR, semantic checks, and compiler transforms
  → backend source generation and native compilation
  → runtime dispatch on the selected target

A TileLang source construct does not prove which hardware instruction or memory path the generated kernel uses. Inspect get_kernel_source(), compiler traces, or profile data when that distinction matters. Add cuda-skill for PTX semantics, compute capability, Nsight, or Compute Sanitizer details.

TileLang contains a CuTeDSL integration under tilelang/contrib/cutedsl/. Use this Skill when tracing how TileLang selects or calls that integration. Use cutlass-skill when the task is about CuTeDSL APIs or implementations themselves.

For correctness work:

• compare against an independent reference over representative and boundary shapes;
• test tails, dynamic dimensions, layouts, dtype conversions, and backend-specific paths;
• separate DSL parsing, lowering, native compilation, runtime, and numerical failures;
• reproduce the original target and compile-time parameters before minimizing the case.

For performance work:

• freeze the benchmark shape set, warmup, repetition count, and synchronization method;
• confirm the measured call dispatches to the intended compiled kernel;
• inspect generated source before attributing a result to TMA, tensor cores, tcgen05, or a CuTeDSL path;
• change one tile, thread count, stage count, layout, or pass configuration at a time.

Build and test routing

Before building, installing, or testing TileLang itself, read the upstream repository instructions:

.agents/skills/tilelang-build/SKILL.md

Those instructions are versioned with the checkout and are authoritative for the current build commands. Most compiler and kernel tests require an appropriate device. Verify the active import path, build directory, target backend, and GPU before interpreting a result.

Updating the source

From the agent-gpu-skills repository:

bash update-repos.sh tilelang
python3 scripts/validate_repo.py --require-sources

The checkout follows TileLang main, while third_party/UPSTREAMS.toml records the commit last accepted by this Skill. Review source-map drift before updating that record.