---

name: nvidia-tilegym-cutile-autotuning description: 'Use when adding, modifying, optimizing, or debugging CuTile autotuning code. Trigger signals: exhaustive_search / replace_hints / hints_fn / cuda.tile.tune in code, autotune in filenames, or correctness/performance issues in autotuned CuTile kernels. Covers: tune-once/cache/launch pattern, per-architecture configs (sm80–sm120), parameter space design (tile sizes, occupancy, num_ctas), and 7 common pitfalls with solutions.' license: CC-BY-4.0 AND Apache-2.0 metadata: author: NVIDIA

CuTile Autotuning

Add autotuning to CuTile kernels using the exhaustive_search API with tune-once/cache/direct-launch pattern.

Instructions

Follow the decision tree to classify the kernel, design a search space, implement the tune-once/cache/launch pattern, and validate performance.

1. Classify — use the Decision Tree to determine search dimensions (occupancy-only vs full tile search)
2. Design search space — select the matching template from references/kernel-type-templates.md; prune to ≤ 30 configs via arch filters
3. Implement — add exhaustive_search + cache + ct.launch following the Step-by-Step Workflow; handle in-place writes with split-buffer if needed
4. Test — run correctness with autotune enabled and with DISABLE_AUTOTUNE=1
5. Validate — A/B benchmark against fixed best-known config; see references/search-strategies.md

Task Router — Jump to What You Need

What are you trying to do?	Go to
Add autotune to a new kernel (most common)	Quick Reference below → Workflow: Adding Autotune → references/kernel-type-templates.md (pick by kernel type: T1=elementwise, T2=in-place, T3=matmul, T4=persistent, T5=FMHA, T6=FP8, T7=grouped GEMM, T8=varlen attention, T9=dual-GEMM fusion)
Debug: data corruption / wrong results after first run	Pitfall #1 (In-Place Kernel)
Debug: autotune taking 5+ minutes	Pitfall #2 (Compilation Timeout)
Debug: search space generator returning zero configs	Pitfall #5 first; also check arch filters, size guards, and num_ctas constraints
Optimize an existing autotune config	Workflow: Optimizing an Existing Config

Quick Reference — Occupancy-Only Autotune (Tune-Once/Cache/Launch)

Most CuTile kernels (elementwise, reduction, LayerNorm) need only occupancy tuning. Copy this pattern:

from types import SimpleNamespace
from cuda.tile.tune import exhaustive_search
import cuda.tile as ct
import torch

def _my_autotune_configs():
    for occ in [1, 2, 4, 8]:
        yield SimpleNamespace(occupancy=occ)

# Module-level cache: tune once, launch fast forever after
_autotune_cache = {}

def my_op(x, output):
    stream = torch.cuda.current_stream()
    NUM_SM = torch.cuda.get_device_properties(x.device).multi_processor_count

    # Cache key: anything that affects optimal config (use str() for device)
    cache_key = (x.shape, x.dtype, str(x.device))

    if cache_key not in _autotune_cache:
        configs = list(_my_autotune_configs())
        result = exhaustive_search(
            configs,
            stream,
            grid_fn=lambda cfg: (min(NUM_SM * cfg.occupancy, M), 1, 1),
            kernel=my_kernel,
            args_fn=lambda cfg: (x, output, ...),
            hints_fn=lambda cfg: {"occupancy": cfg.occupancy},
        )
        best_cfg = result.best.config
        tuned_kernel = my_kernel.replace_hints(occupancy=best_cfg.occupancy)
        _autotune_cache[cache_key] = (best_cfg, tuned_kernel)  # cache BOTH

    cfg, tuned_kernel = _autotune_cache[cache_key]
    grid = (min(NUM_SM * cfg.occupancy, M), 1, 1)
    ct.launch(stream, grid, tuned_kernel, (x, output, ...))

Key rules:

• Tune once, cache, launch directly — exhaustive_search runs only on first call per shape; subsequent calls use cached config + ct.launch with zero overhead
• For in-place kernels use split-buffer during search (separate input/output tensors)
• Keep ≤ 30 configs total
• exhaustive_search requires a Sequence (list/tuple) — convert generators with list()
• Search space must include the original fixed config — this guarantees autotuning never makes performance worse

When to use this pattern: Kernel has fixed block size (not tile-size tunable). Includes: elementwise (SwiGLU, GeGLU), reduction (RMSNorm, LayerNorm), RoPE, and persistent kernels with heuristic block sizes (grouped GEMM).

For complex kernels (matmul with tile sizes, FMHA, FP8 with num_ctas), read the full guide below + kernel-type-templates.md.

⚠️ Two pitfalls catch almost everyone — check before submitting:

• In-place kernel (writes back to input tensor)? → MUST use split-buffer pattern during search → Pitfall #1
• Search space empty? → Check arch filters and num_ctas constraints → Pitfall #5

Reading Guide

• Occupancy-only kernels (elementwise, reduction, persistent with fixed block sizes): Quick Reference + Pitfall Checklist is sufficient — skip references/ docs. For in-place kernels, also read Pitfall #1.
• Complex kernels (matmul with tunable tile sizes, FMHA, FP8 with num_ctas): Quick Reference → Decision Tree → API Reference → Step-by-Step Workflow → relevant references/ docs.

5-step summary: Classify kernel → Design search space (parameter-space-design.md) → Implement using template (kernel-type-templates.md) → Validate with A/B test → Check Pitfall Checklist.

Design Philosophy

Build a small, precise search space bottom-up — not a large space trimmed down. CuTile compilation is much heavier than Triton (~0.5-1s per config), so 30 configs is the hard upper limit. The approach is: classify the kernel type first, then construct only the relevant configs for that type and architecture. Never start with a large cartesian product and prune — start with the minimum viable space and expand only if data shows it's needed.

Decision Tree: What Search Dimensions Does This Kernel Need?

All kernels should have autotuning added. The question is not whether to autotune, but what dimensions to search:

What type of kernel is this?
├── Compute-bound (matmul, GEMM, FMHA) → Does it have multiple tunable dimensions (tile sizes)?
│   ├── YES → Is it a fused multi-GEMM kernel (dual-GEMM, e.g. Linear+GLUAct)?
│   │   ├── YES → Template 9: low occupancy (1–2), conservative tiles (2× SHMEM/register pressure)
│   │   └── NO  → Full search: TILE_M × TILE_N × (TILE_K) × occupancy × num_ctas
│   │             (see matmul/FMHA templates in kernel-type-templates.md)
│   └── NO  → Occupancy-only search: [1, 2, 4, 8]
│             (see Quick Reference above)
├── Balanced (LayerNorm, reduction + compute) →
│   Occupancy-only search: [1, 2, 4, 8]
│   Expected benefit: 2-15%
└── Memory-bound (CE Loss, pure elementwise) →
    Occupancy-only search: [1, 2, 4, 8]
    Expected benefit: 0-15% (varies by kernel; zero-cost after tuning)

Why memory-bound kernels only search occupancy (not num_ctas or tile sizes):

• num_ctas has zero benefit: num_ctas > 1 enables TMA multicast, where multiple CTAs share tile data in shared memory (e.g., matmul A/B tiles reused across CTAs). Memory-bound kernels use per-element ct.gather/ct.scatter with no tile reuse — multi-CTA cooperation adds overhead with no data sharing benefit.
• Tile sizes are pre-determined: BLOCK_SIZE for memory-bound kernels is determined by offline sweep (e.g., 1024 is globally optimal on B200 across [256, 512, 1024, 2048, 4096, 8192]). This is a constant, not a runtime tunable.
• Occupancy is the only effective knob: Higher occupancy lets the GPU hide memory latency by switching to another CTA while one is stalled on a memory request.

Evidence — CE Loss experiment: A 12-config search (occupancy × num_ctas) on Cross-Entropy Loss yielded only 2.5% gain (0.79x → 0.81x vs Triton). The num_ctas dimension contributed nothing; the result was reverted because compilation cost outweighed the marginal benefit. Occupancy-only (4 configs) achieves the same result at 3x less compilation time.

Note on memory-bound kernels: Adding occupancy-only autotune is always worthwhile because:

• The tune-once/cache/launch pattern has zero runtime overhead after the first call
• The search space is tiny (4 configs, ~2-4s compilation)
• Even small improvements have value at scale

Occupancy Selection Guide

Occupancy controls how many CTAs run concurrently per SM. Use this as a starting point when designing the occupancy search space:

Occupancy Range	Best For	Example Kernels
1–4	Compute-bound (heavy math)	Complex transforms, matmul
4–8	Balanced (GEMM, TMA)	Matrix multiply, FMHA
8–16	Memory-bound (reductions)	Softmax, LayerNorm
16–32	Very light (copies, casts)	Type conversions, elementwise

Use these ranges to seed your initial search space. For occupancy-only kernels, [1, 2, 4, 8] covers most cases — see Quick Reference above.

exhaustive_search API Reference

⚠️ Deprecated API: cuda.tile_experimental.autotune_launch() (aka ct_experimental.autotune_launch) is deprecated and should NOT be used. It combines search + launch in one call with random sampling, which produces less reproducible results and worse config selection compared to exhaustive_search. Always use cuda.tile.tune.exhaustive_search (the current API below) with explicit caching and ct.launch.

Current API (cuda.tile.tune)

from cuda.tile.tune import exhaustive_search, TuningResult

result: TuningResult = exhaustive_search(
    search_space,   # Sequence[T] — list or tuple of configs (NOT a generator)
    stream,         # torch.cuda.current_stream()
    grid_fn,        # callable(cfg) → tuple[int, ...]
    kernel,         # @ct.kernel decorated function
    args_fn,        # callable(cfg) → tuple of kernel args
    hints_fn=None,  # callable(cfg) → {"occupancy": int, "num_ctas": int}
    *,
    quiet=False     # suppress output
)

TuningResult

@dataclass
class TuningResult[T]:
    best: Measurement       # best config + timing (mean_us, error_margin_us, num_samples)
    successes: Sequence[Measurement]   # all successful configs (sorted by performance)
    failures: Sequence[tuple[T, str, str]]  # (config, exception_type, message)

Key properties:

• Exhaustive: evaluates ALL configs in order — no random sampling, no skipped configs
• Search only: does not perform the final production launch — it executes trial runs internally for benchmarking, but you call ct.launch separately for the actual production invocation
• No built-in cache: you manage caching explicitly (see tune-once/cache/launch pattern)
• Deterministic: same search space always produces the same evaluation order

Tune-Once / Cache / Launch Pattern

This is the recommended pattern for all autotuned kernels. It ensures:

• First call: runs exhaustive_search to find the best config (~2-30s depending on space size)
• Subsequent calls: uses cached config with ct.launch — zero overhead (identical to a fixed ct.launch)

_cache = {}

def run_kernel_autotuned(x, ...):
    stream = torch.cuda.current_stream()
    cache_key = (x.shape, x.dtype, str(x.device))

    if cache_key not in _cache:
        configs = list(_my_autotune_configs())
        result = exhaustive_search(
            configs, stream,
            grid_fn=lambda cfg: ...,
            kernel=my_kernel,
            args_fn=lambda cfg: ...,
            hints_fn=lambda cfg: {"occupancy": cfg.occupancy},
        )
        best_cfg = result.best.config
        tuned_kernel = my_kernel.replace_hints(occupancy=best_cfg.occupancy)
        _cache[cache_key] = (best_cfg, tuned_kernel)  # cache BOTH config and compiled kernel

    cfg, tuned_kernel = _cache[cache_key]
    grid = compute_grid(cfg)
    ct.launch(stream, grid, tuned_kernel, (x, ...))

Why this pattern matters: The ct.launch call in the fast path is identical to what you'd write for a fixed-config kernel. There is zero per-call overhead — no lock, no hash lookup, no lambda invocation. The only cost is the Python dict lookup for _cache[cache_key].

⚠️ Critical: always cache the tuned kernel object, not just the config. replace_hints() returns a new kernel object with its own independent JIT cache. Calling it on every invocation triggers recompilation each time, degrading performance by 100–500×. Call replace_hints() once after exhaustive_search, store the returned kernel in the cache alongside the config, and reuse it directly on the fast path. See Pitfall #7.

replace_hints

After finding the best config, use kernel.replace_hints() to create a kernel variant with the optimal hints:

# For occupancy-only:
tuned_kernel = my_kernel.replace_hints(occupancy=cfg.occupancy)

# For occupancy + num_ctas:
tuned_kernel = my_kernel.replace_hints(occupancy=cfg.occupancy, num_ctas=cfg.num_ctas)

replace_hints accepts only occupancy and num_ctas — these are the only compiler hints controllable via the autotune API.

ByTarget wrapping for cross-architecture portability: When creating tuned kernel variants via ct.kernel(), prefer wrapping hint values in ct.ByTarget for portability across GPU architectures:

# Preferred: explicit architecture targeting (portable)
tuned_kernel = ct.kernel(
    my_kernel._pyfunc,
    occupancy=ct.ByTarget(sm_100=best_cfg.occupancy),
    num_ctas=ct.ByTarget(sm_100=best_cfg.num_ctas, default=1),
)

# Also acceptable: plain integers (when targeting a single architecture)
tuned_kernel = ct.kernel(my_kernel._pyfunc, occupancy=best_cfg.occupancy)

When targeting only the current GPU (the common case in autotuning), plain integers work fine. Use ByTarget when the code may run on multiple architectures or when following production conventions (TileGym production code consistently uses ByTarget).

Kernel Hints

CuTile kernel performance is controlled by two compile-time hints:

• occupancy: Number of CTAs per SM. Higher occupancy = more parallelism but less shared memory per CTA.
• num_ctas: Number of CTAs in a CGA (Cooperative Group Array). Used for multi-CTA cooperation (e.g., TMA multicast). Only supported on sm90+.

Three ways to set hints:

# 1. Fixed value in decorator (no autotune needed)
@ct.kernel(occupancy=2, num_ctas=1)
def my_kernel(...): ...

# 2. Architecture-specific fixed value (no autotune needed)
@ct.kernel(num_ctas=ct.ByTarget(sm_100=2, sm_120=1, default=1))
def my_kernel(...): ...

# 3. Runtime autotune via exhaustive_search + replace_hints
# IMPORTANT: Remove fixed hints from decorator first!
@ct.kernel
def my_kernel(...): ...

# Then in the host wrapper:
tuned_kernel = my_kernel.replace_hints(occupancy=best_occ, num_ctas=best_ctas)
ct.launch(stream, grid, tuned_kernel, args)

Important: replace_hints correctly overrides decorator hints (it uses dataclasses.replace() internally). However, if you forget to call replace_hints, the decorator's fixed values are used instead of the autotuned values. To avoid this confusion, always remove fixed hints from the @ct.kernel(...) decorator before adding autotuning — this makes it explicit that hints come only from the autotune path.

search_space Design

The search space is a list of SimpleNamespace objects. Each namespace holds config fields that grid_fn, args_fn, and hints_fn can read.

from types import SimpleNamespace

# Occupancy-only (elementwise kernels)
def autotune_configs():
    for occ in [1, 2, 4, 8]:
        yield SimpleNamespace(occupancy=occ)

# Full matmul search space — see parameter-space-design.md for complete per-architecture configs
# Pattern: yield SimpleNamespace(TILE_SIZE_M=..., TILE_SIZE_N=..., TILE_SIZE_K=..., num_ctas=..., occupancy=...)

Note: exhaustive_search requires a Sequence (list/tuple), not a generator. Always convert with list():

configs = list(autotune_configs())
result = exhaustive_search(configs, ...)

grid_fn Patterns

from math import ceil

# Pattern A: Simple tile coverage (matmul, elementwise)
grid_fn=lambda cfg: (ceil(M / cfg.TILE_SIZE_M) * ceil(N / cfg.TILE_SIZE_N), 1, 1)

# Pattern B: Persistent matmul (static_persistent_matmul_kernel)
NUM_SMS = torch.cuda.get_device_properties("cuda").multi_processor_count
grid_fn=lambda cfg: (
    min(NUM_SMS // cfg.num_ctas, ceil(M / cfg.TILE_M) * ceil(N / cfg.TILE_N)) * cfg.occupancy,
    1, 1,
)

# Pattern C: 2D grid (FMHA — one dim for seq tiles, one for batch*heads)
grid_fn=lambda cfg: (ceil(q_len / cfg.TILE_M), batch_size * num_heads, 1)

# Pattern D: 1D elementwise (cdiv = math.ceil(a/b), from ct_ops.py)
grid_fn=lambda cfg: (cdiv(n_elements, BLOCK_SIZE),)

# Pattern E: Grouped GEMM persistent (grid fixed at NUM_SMS, occupancy via hints_fn only)
grid_fn=lambda cfg: (NUM_SMS, 1, 1)

Step-by-Step Workflow

Adding Autotune to a New Kernel

1. Classify the kernel using the decision tree above.

   • VERIFY: You know whether this is occupancy-only or requires tile-size tuning.

2. Remove hardcoded hints from decorator (strongly recommended): If the kernel currently has hardcoded hints in its decorator (e.g. @ct.kernel(occupancy=2, num_ctas=1)), remove those fixed hints and change to bare @ct.kernel before adding autotuning. While replace_hints does correctly override decorator values at runtime, leaving them creates a silent fallback trap: if any code path (e.g., DISABLE_AUTOTUNE, error handling, or a future refactor) skips replace_hints, the decorator's fixed hints are used instead of the autotuned values — and this produces no error, just silently worse performance. Removing them makes the failure mode explicit (missing hints → compiler defaults) rather than silent (wrong fixed hints used).

   • VERIFY: The @ct.kernel decorator has no occupancy= or num_ctas= arguments before proceeding. Use bare @ct.kernel instead.

3. Check for in-place writes: If the kernel modifies input tensors in-place, you MUST use the split-buffer pattern during exhaustive_search — see Pitfall #1.

   • VERIFY: Either the kernel is not in-place, or you have added a split-buffer scratch tensor for the search phase.

4. Select the template from kernel-type-templates.md based on kernel type.

5. Design the search space following parameter-space-design.md:

   • Start from reference configs, not from scratch. Clone configs from existing production kernels of the same type (e.g., ops/cutile/matmul.py for GEMM) and adapt. For GEMM-class kernels, nvMatmulHeuristics can suggest 8-16 high-quality candidates that reach 96-99% peak performance — see parameter-space-design.md for details.
   • Detect the current GPU architecture with torch.cuda.get_device_capability().
   • Target one architecture at a time. Generate configs only for the detected arch. Do NOT add branches for other architectures — they cannot be tested on this machine and untested code paths are unreliable. If multi-arch support is needed later, add it in a separate pass on the appropriate hardware.
   • Identify tunable parameters (tile sizes, occupancy, num_ctas)
   • Ensure the search space includes the original fixed config (or an equivalent). This guarantees that the autotuned result is at least as good as the original — no performance regression is possible.
   • If the generated set exceeds 30, apply tile size filters and pruning rules to reduce it to ≤ 30
   • VERIFY: Total configs ≤ 30 (hard limit: CuTile compilation is heavy, >30 configs will timeout).

6. Implement the tune-once/cache/launch pattern:

   • Define a _cache dict at module level
   • Define a cache key that captures all parameters affecting optimal config (shapes, dtypes, device, any flags like is_causal). ⚠️ Use str(x.device) not x.device in the cache key — torch.device objects are not reliably hashable and can cause TypeError: unhashable type at runtime. Always convert to string: cache_key = (..., x.dtype, str(x.device)). Tip: For GEMM-class kernels, round dimensions to the next power of 2 in the cache key (e.g., cache_key = (next_pow2(M), next_pow2(N), next_pow2(K), dtype, str(device))) to reduce unique key count and avoid re-tuning for similar shapes.
   • Call exhaustive_search(list(configs), ...) only when cache misses
   • Store result.best.config in cache
   • Use kernel.replace_hints(...) to create the tuned kernel variant
   • Use ct.launch() for the actual kernel invocation
   • grid_fn correctly computes grid from config
   • args_fn passes all kernel arguments including tile sizes as ct.Constant[int]
   • hints_fn passes occupancy and/or num_ctas from config
   • VERIFY: exhaustive_search receives a list() of configs, not a raw generator.

7. (MANDATORY) Add DISABLE_AUTOTUNE support for CI and profiling: check os.environ.get("DISABLE_AUTOTUNE", "0") == "1" — when set, skip exhaustive_search entirely and fall back to ct.launch with the first valid config. This is required for:

   • CI determinism (autotune adds variable wall time)
   • NCU profiling (prevents autotune trial runs from cluttering the trace — see Pitfall #4)
   • Debugging (isolates kernel correctness from autotune behavior) Place the check before the cache lookup so that DISABLE_AUTOTUNE=1 bypasses all autotune logic. Provide a hardcoded fallback config in case the generator yields zero configs.
   • VERIFY: Running with DISABLE_AUTOTUNE=1 produces correct results and does not call exhaustive_search.

8. Test: Run correctness tests first (pytest -k "test_op and cutile"), then benchmark.

   • VERIFY: Correctness passes with autotune enabled AND with DISABLE_AUTOTUNE=1.

9. Validate with A/B test: Compare autotune version vs fixed best-known config. See search-strategies.md for methodology.

   • VERIFY: Autotune version ≥ baseline (or within noise). If worse, check that the search space includes the original fixed config, and that replace_hints is being used correctly.

10. (MANDATORY) Run the test and verify performance before submitting.

    Execute the provided test script (e.g. ENABLE_TILE=1 python3 test.py) and check:

    • correctness: PASS
    • speedup_over_fixed >= 1.0 (autotuned must not be slower than fixed baseline)

    If speedup_over_fixed < 1.0:

    • Check that the search space includes the original fixed config (this guarantees no regression)
    • Check if replace_hints is being called on every code path — revisit Step 2 (if any path skips replace_hints, the decorator's fixed hints are used instead of autotuned values)
    • Expand search space if all configs perform similarly (see references/parameter-space-design.md → "Adapting Search Space")

    ⚠️ DO NOT submit without running the test at least once. Writing correct-looking code is not sufficient — autotuning bugs (silent hint override, split-buffer omission) are only caught at runtime.

Integration with torch.autograd.Function

When the kernel is used inside a torch.autograd.Function:

• Place the tune-once/cache/launch logic in forward() only. The cached config is reused across calls.
• In backward(), using ct.launch with a fixed or cached config is often sufficient. However, if backward has its own independent search space (e.g. grouped GEMM dX and dW have separate optimal configs), autotuning is appropriate there too.
• Example: rope_embedding.py — forward uses exhaustive_search + cache with split-buffer, backward uses ct.launch with same-buffer (Q_in=Q_out).

Cross-Backend Config Transfer (Triton → CuTile)

Use src/tilegym/autotune.py: maps BLOCK_SIZE_M/N/K → TILE_SIZE_M/N/K; num_warps/num_stages have no CuTile equivalent.

Optimizing an Existing Autotune Config

1. Profile first: Use NCU (set DISABLE_AUTOTUNE=1).
2. Expand (too narrow): add tile sizes, num_ctas (sm90+), swap_ab.
3. Prune (too slow): remove suboptimal configs, use arch-conditional yield, add size filters.
4. Re-validate: A/B test to confirm improvement.

Pitfall Checklist

Before submitting code with autotune, verify these:

Pitfall #1: In-Place Kernel Data Corruption

Problem: exhaustive_search runs the kernel multiple times to benchmark. If the kernel modifies input tensors in-place, the data is corrupted after the first trial run.

Solution: Split-buffer pattern — use separate read-only input and write-only output during search:

# During exhaustive_search: use separate output buffer
Q_scratch = torch.empty_like(Q)
configs = list(_rope_autotune_configs())
result = exhaustive_search(
    configs, stream,
    grid_fn=...,
    kernel=rope_kernel,
    args_fn=lambda cfg: (Q, Q_scratch, ...),  # Q_in != Q_out
    hints_fn=...,
)

# After search: launch with in-place args using tuned config
cfg = result.best.config
tuned_kernel = rope_kernel.replace_hints(occupancy=cfg.occupancy)
ct.launch(stream, grid, tuned_kernel, (Q, Q, ...))  # Q_in == Q_out (in-place)

Real example: rope_embedding.py — Search uses split-buffer, final launch uses same-buffer.

Also wrong: Using Q.clone() in args_fn — this adds ~4us per clone, which is fatal for small kernels (~5us). The clone+copy pattern caused 0.48x performance in RoPE.

Tip — isolating output buffers in args_fn: For kernels that write to a dedicated output tensor (not in-place), use c.clone() inside args_fn to prevent trial runs from overwriting the final output buffer:

# Output tensor c will be overwritten by each trial — clone it so trials don't
# corrupt the buffer the caller expects to use after exhaustive_search returns.
result = exhaustive_search(
    configs, stream,
    grid_fn=...,
    kernel=my_kernel,
    args_fn=lambda cfg: (a, b, c.clone()),  # each trial gets a fresh output
    hints_fn=...,
)

This is safe because the clone cost (~4us) is negligible relative to compute-bound kernel execution time (~50us+). Only avoid clone() for very small, memory-bound kernels where 4us is a significant fraction of runtime — in that case, pre-allocate a single scratch buffer outside args_fn (as in the split-buffer pattern above).

Pitfall #2: Compilation Timeout

Problem: >30 configs causes compilation to exceed 5 minutes. CuTile compilation is heavier than Triton.

Solution:

• Keep total search space ≤ 30 configs — apply arch filters, tile size filters, and pruning rules until you're under the limit
• Use architecture-conditional yield to only generate relevant configs
• Prune the search space using architecture-conditional yield and size filters until total configs ≤ 30

Real example: Grouped GEMM expanded from 4 to 32 configs → all backward tests timed out. Reverted to occupancy-only (4 configs) with no performance loss.

Pitfall #3: Cold-Cache Performance Skew

Problem: First process run is slower due to driver/JIT caches. Can cause wrong config selection.

Solution: Always warm up before measuring. exhaustive_search has built-in warmup, but first-process cold start is unavoidable. Re-run if you suspect the initial result was affected.

Pitfall #4: NCU Profiling Interference

Problem: NCU profiles autotune trial runs, cluttering the trace.

Solution: Set DISABLE_AUTOTUNE=1 before profiling, or use ncu --launch-skip N.

Pitfall #5: search_space as Generator (Exhaustion)

Problem: exhaustive_search requires a Sequence (list/tuple), not a generator. Passing a generator directly will fail or produce unexpected results.

Solution: Always convert to list:

# CORRECT: convert generator to list
configs = list(_matmul_autotune_configs())
result = exhaustive_search(configs, ...)

# WRONG: passing generator directly
result = exhaustive_search(_matmul_autotune_configs(), ...)

Pitfall #6: FP8 Precision Loss

Problem: Hardware / breaks FP8 quantization bucket boundaries.

Solution: Use ct.truediv(x, y, rounding_mode=RoundingMode.FULL) for IEEE-compliant division in FP8 kernels. Never use / operator for FP8 scale computation.

Pitfall #7: replace_hints on Hot Path (Recompilation)

Problem: replace_hints() returns a new kernel object with its own JIT cache (internally uses dataclasses.replace() which creates a fresh instance). Calling it on every kernel invocation — even with the same arguments — triggers recompilation every time. This is the most common autotune performance bug: cutile_ms jumps from ~0.04ms to 16–39ms (100–500× slower).

Incorrect (recompiles on every call):

_cache[key] = result.best.config  # only stores config

cfg = _cache[key]
tuned = my_kernel.replace_hints(occupancy=cfg.occupancy)  # NEW kernel each time!
ct.launch(stream, grid, tuned, ...)

Correct (compile once, reuse forever):

best_cfg = result.best.config
tuned = my_kernel.replace_hints(occupancy=best_cfg.occupancy)  # compile ONCE
_cache[key] = (best_cfg, tuned)  # cache both

cfg, tuned = _cache[key]
ct.launch(stream, grid, tuned, ...)  # reuse compiled kernel

Rule: Call replace_hints exactly once per config (immediately after exhaustive_search), cache the returned kernel object, and never call replace_hints again on the fast path.

Scope and Boundaries

This skill covers only autotune configuration: search space design, exhaustive_search invocation, caching, and ct.launch with tuned hints. It does not modify kernel code.

In scope (autotune config):

• Search space generator functions
• exhaustive_search() calls and result handling
• kernel.replace_hints() for applying tuned hints
• Cache logic (key design, dict management)
• ct.launch() with tuned kernel
• DISABLE_AUTOTUNE fallback path

Out of scope (kernel code modifications — do NOT make these changes):

• Math flags (flush_to_zero, rounding_mode)
• Performance Hints (slice_hint, buffer_depth, copy_config)
• Memory access patterns (2D→1D gather/scatter conversion)
• Codegen optimizations (safe_offs → padding_value)
• Algorithm changes (K-loop split, load balancing)

Further Optimization Suggestions

After adding autotuning, the following kernel-level optimizations may yield additional gains. These are outside the scope of this skill — mention them to the user as potential next steps, but do not implement them as part of autotuning:

• Math flags: flush_to_zero=True + rounding_mode=APPROX can provide 34-72% improvement for FMHA-class kernels (set via environment variables TILEIR_ENABLE_FTZ=1 TILEIR_ENABLE_APPROX=1 or in kernel code). Causal chain: larger tiles initially decrease performance by 18-43% due to subnormal handling overhead; enabling FTZ+APPROX rescues this and flips the result to +34-72%. Math flags are therefore a prerequisite for large-tile configs to be effective on FMHA-class kernels.
• Performance Hints: slice_hint, buffer_depth, copy_config — requires modifying kernel IR code
• Memory access patterns: Using TMA loads (ct.load) instead of ct.gather; removing unnecessary bounds checks (check_bounds=False when safe)
• Codegen quality: Using padding_value parameter instead of manual ct.where masking; removing safe_offs
• Algorithm restructuring: K-loop split, load balancing, algebraic simplification

Differences from Triton Autotune

Key differences: Triton uses @triton.autotune decorator with Config(...) objects; CuTile uses exhaustive_search() with SimpleNamespace configs + separate cache + ct.launch. CuTile has no num_warps/num_stages (compiler decides) — only tile sizes + occupancy + num_ctas. CuTile compilation is heavier (keep ≤30 configs). CuTile cache is user-managed in-memory (no automatic persistence). CuTile separates args_fn (kernel args) from hints_fn (compiler hints).

Reference Documents

Category	Document	Content
Parameter Design	parameter-space-design.md	Per-kernel-type parameter spaces, cross-arch patterns, grid_fn patterns, pruning rules
Search Strategies	search-strategies.md	Exhaustive search, A/B test methodology, DISABLE_AUTOTUNE pattern
Templates	kernel-type-templates.md	Copy-paste autotune templates for 8 kernel types
Hardware	hardware-constraints.md	Per-architecture constraints, tile size ranges, num_ctas rules, TMA requirements

Source Code References

Key files: ops/cutile/matmul.py (matmul autotune), ops/cutile/attention.py (FMHA autotune), suites/unsloth/cutile/ct_ops.py (shared autotune_configs() occupancy=[1,2,4,8]), suites/unsloth/cutile/swiglu.py (elementwise example), suites/unsloth/cutile/rope_embedding.py (split-buffer pattern), suites/unsloth/cutile/grouped_gemm.py (persistent GEMM, occupancy-only).

Worked Examples

Each example shows the before → after pattern: fixed_launch.py (hardcoded ct.launch) and autotuned_launch.py (refactored to tune-once/cache/launch).

Directory	Kernel	Autotune Pattern	Complexity	Key Teaching Point
assets/examples/01_rmsnorm_occupancy_only/	RMSNorm (reduction)	Occupancy-only [1,2,4,8]	Low	Most common pattern — no tile tuning, just find best occupancy. Grid = NUM_SM * cfg.occupancy. Not in-place.
assets/examples/02_matmul_full_search/	GEMM C=A@B	Full: TILE_M/N/K + occupancy + num_ctas (sm90+)	High	Compute-bound kernel with multiple tunable dimensions. args_fn passes tile sizes as ct.Constant[int]. grid_fn depends on cfg. ≤30 configs.
assets/examples/03_rope_inplace_splitbuffer/	RoPE embedding (in-place)	Occupancy-only, with split-buffer	Medium	In-place kernel MUST use split-buffer during search to avoid corruption. Search writes to scratch; final ct.launch uses real in-place args.