--- name: cutile-autotuning-cab309aea5c6 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 --- # 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: ```python 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`](references/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`](references/parameter-space-design.md)) → Implement using template ([`kernel-type-templates.md`](references/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`) ```python 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 ```python @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`) ```python _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: ```python # 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: ```python # 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: ```python # 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. ```python 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()`: ```python configs = list(autotune_configs()) result = exhaustive_search(configs, ...) ``` ### grid_fn Patterns ```python 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`](references/kernel-type-templates.md) based on kernel type. 5. **Design the search space** following [`parameter-space-design.md`](references/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`](references/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`](references/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: ```python # 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: ```python # 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: ```python # 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): ```python _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): ```python 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`](references/parameter-space-design.md) | Per-kernel-type parameter spaces, cross-arch patterns, grid_fn patterns, pruning rules | | **Search Strategies** | [`search-strategies.md`](references/search-strategies.md) | Exhaustive search, A/B test methodology, DISABLE_AUTOTUNE pattern | | **Templates** | [`kernel-type-templates.md`](references/kernel-type-templates.md) | Copy-paste autotune templates for 8 kernel types | | **Hardware** | [`hardware-constraints.md`](references/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/`](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/`](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/`](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. | --- Source attribution: yo-steven Original project: https://github.com/yo-steven/skills-exploration-20260522/blob/6e1127ffb7690aa8038b172d0ae7cc720754870d/skills/TileGym/cutile-autotuning/SKILL.md License: Apache-2.0 License information: https://www.apache.org/licenses/LICENSE-2.0 Original upstream instructions; downloaded from Mega AI Hub. ## Upstream license notice SPDX-License-Identifier: Apache-2.0 AND CC-BY-4.0 Copyright (c) 2026 NVIDIA Corporation This project is licensed under a dual license: - Apache License 2.0 (for code and configuration) - Creative Commons Attribution 4.0 International (CC BY 4.0) (for documentation) You may use the contents of this repository under the terms of either license, or both, as applicable to the type of content. ================================================================================ Apache License Version 2.0, January 2004 http://www.apache.org/licenses/ TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION 1. Definitions. "License" shall mean the terms and conditions for use, reproduction, and distribution as defined by Sections 1 through 9 of this document. "Licensor" shall mean the copyright owner or entity authorized by the copyright owner that is granting the License. "Legal Entity" shall mean the union of the acting entity and all other entities that control, are controlled by, or are under common control with that entity. For the purposes of this definition, "control" means (i) the power, direct or indirect, to cause the direction or management of such entity, whether by contract or otherwise, or (ii) ownership of fifty percent (50%) or more of the outstanding shares, or (iii) beneficial ownership of such entity. "You" (or "Your") shall mean an individual or Legal Entity exercising permissions granted by this License. "Source" form shall mean the preferred form for making modifications, including but not limited to software source code, documentation source, and configuration files. "Object" form shall mean any form resulting from mechanical transformation or translation of a Source form, including but not limited to compiled object code, generated documentation, and conversions to other media types. "Work" shall mean the work of authorship, whether in Source or Object form, made available under the License, as indicated by a copyright notice that is included in or attached to the work (an example is provided in the Appendix below). "Derivative Works" shall mean any work, whether in Source or Object form, that is based on (or derived from) the Work and for which the editorial revisions, annotations, elaborations, or other modifications represent, as a whole, an original work of authorship. For the purposes of this License, Derivative Works shall not include works that remain separable from, or merely link (or bind by name) to the interfaces of, the Work and Derivative Works thereof. "Contribution" shall mean any work of authorship, including the original version of the Work and any modifications or additions to that Work or Derivative Works thereof, that is intentionally submitted to the Licensor for inclusion in the Work by the copyright owner or by an individual or Legal Entity authorized to submit on behalf of the copyright owner. For the purposes of this definition, "submitted" means any form of electronic, verbal, or written communication sent to the Licensor or its representatives, including but not limited to communication on electronic mailing lists, source code control systems, and issue tracking systems that are managed by, or on behalf of, the Licensor for the purpose of discussing and improving the Work, but excluding communication that is conspicuously marked or otherwise designated in writing by the copyright owner as "Not a Contribution." "Contributor" shall mean Licensor and any individual or Legal Entity on behalf of whom a Contribution has been received by the Licensor and subsequently incorporated within the Work. 2. Grant of Copyright License. Subject to the terms and conditions of this License, each Contributor hereby grants to You a perpetual, worldwide, non-exclusive, no-charge, royalty-free, irrevocable copyright license to reproduce, prepare Derivative Works of, publicly display, publicly perform, sublicense, and distribute the Work and such Derivative Works in Source or Object form. 3. Grant of Patent License. Subject to the terms and conditions of this License, each Contributor hereby grants to You a perpetual, worldwide, non-exclusive, no-charge, royalty-free, irrevocable (except as stated in this section) patent license to make, have made, use, offer to sell, sell, import, and otherwise transfer the Work, where such license applies only to those patent claims licensable by such Contributor that are necessarily infringed by their Contribution(s) alone or by combination of their Contribution(s) with the Work to which such Contribution(s) was submitted. If You institute patent litigation against any entity (including a cross-claim or counterclaim in a lawsuit) alleging that the Work or a Contribution incorporated within the Work constitutes direct or contributory patent infringement, then any patent licenses granted to You under this License for that Work shall terminate as of the date such litigation is filed. 4. Redistribution. You may reproduce and distribute copies of the Work or Derivative Works thereof in any medium, with or without modifications, and in Source or Object form, provided that You meet the following conditions: (a) You must give any other recipients of the Work or Derivative Works a copy of this License; and (b) You must cause any modified files to carry prominent notices stating that You changed the files; and (c) You must retain, in the Source form of any Derivative Works that You distribute, all copyright, patent, trademark, and attribution notices from the Source form of the Work, excluding those notices that do not pertain to any part of the Derivative Works; and (d) If the Work includes a "NOTICE" text file as part of its distribution, then any Derivative Works that You distribute must include a readable copy of the attribution notices contained within such NOTICE file, excluding any notices that do not pertain to any part of the Derivative Works, in at least one of the following places: within a NOTICE text file distributed as part of the Derivative Works; within the Source form or documentation, if provided along with the Derivative Works; or, within a display generated by the Derivative Works, if and wherever such third-party notices normally appear. The contents of the NOTICE file are for informational purposes only and do not modify the License. You may add Your own attribution notices within Derivative Works that You distribute, alongside or as an addendum to the NOTICE text from the Work, provided that such additional attribution notices cannot be construed as modifying the License. You may add Your own copyright statement to Your modifications and may provide additional or different license terms and conditions for use, reproduction, or distribution of Your modifications, or for any such Derivative Works as a whole, provided Your use, reproduction, and distribution of the Work otherwise complies with the conditions stated in this License. 5. Submission of Contributions. Unless You explicitly state otherwise, any Contribution intentionally submitted for inclusion in the Work by You to the Licensor shall be under the terms and conditions of this License, without any additional terms or conditions. Notwithstanding the above, nothing herein shall supersede or modify the terms of any separate license agreement you may have executed with Licensor regarding such Contributions. 6. Trademarks. This License does not grant permission to use the trade names, trademarks, service marks, or product names of the Licensor, except as required for reasonable and customary use in describing the origin of the Work and reproducing the content of the NOTICE file. 7. Disclaimer of Warranty. Unless required by applicable law or agreed to in writing, Licensor provides the Work (and each Contributor provides its Contributions) on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied, including, without limitation, any warranties or conditions of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A PARTICULAR PURPOSE. You are solely responsible for determining the appropriateness of using or redistributing the Work and assume any risks associated with Your exercise of permissions under this License. 8. Limitation of Liability. In no event and under no legal theory, whether in tort (including negligence), contract, or otherwise, unless required by applicable law (such as deliberate and grossly negligent acts) or agreed to in writing, shall any Contributor be liable to You for damages, including any direct, indirect, special, incidental, or consequential damages of any character arising as a result of this License or out of the use or inability to use the Work (including but not limited to damages for loss of goodwill, work stoppage, computer failure or malfunction, or any and all other commercial damages or losses), even if such Contributor has been advised of the possibility of such damages. 9. Accepting Warranty or Additional Liability. While redistributing the Work or Derivative Works thereof, You may choose to offer, and charge a fee for, acceptance of support, warranty, indemnity, or other liability obligations and/or rights consistent with this License. However, in accepting such obligations, You may act only on Your own behalf and on Your sole responsibility, not on behalf of any other Contributor, and only if You agree to indemnify, defend, and hold each Contributor harmless for any liability incurred by, or claims asserted against, such Contributor by reason of your accepting any such warranty or additional liability. END OF TERMS AND CONDITIONS ================================================================================ Creative Commons Attribution 4.0 International License (CC BY 4.0) https://creativecommons.org/licenses/by/4.0/ You are free to: Share — copy and redistribute the material in any medium or format for any purpose, even commercially. Adapt — remix, transform, and build upon the material for any purpose, even commercially. The licensor cannot revoke these freedoms as long as you follow the license terms. Under the following terms: Attribution — You must give appropriate credit, provide a link to the license, and indicate if changes were made. You may do so in any reasonable manner, but not in any way that suggests the licensor endorses you or your use. No additional restrictions — You may not apply legal terms or technological measures that legally restrict others from doing anything the license permits. Notices: You do not have to comply with the license for elements of the material in the public domain or where your use is permitted by an applicable exception or limitation. No warranties are given. The license may not give you all of the permissions necessary for your intended use. For example, other rights such as publicity, privacy, or moral rights may limit how you use the material. For the full legal text, see: https://creativecommons.org/licenses/by/4.0/legalcode # Mega AI Hub resource metadata Creator / publisher: yo-steven Original project: https://github.com/yo-steven/skills-exploration-20260522/blob/6e1127ffb7690aa8038b172d0ae7cc720754870d/skills/TileGym/cutile-autotuning/SKILL.md Attribution basis: Source repository record Fields of work: Development & Engineering Roles: Backend Developer, Data Engineer, ML / AI Engineer Specialties: Backend Tools: Cuda Country applicability: Country not specified Classification basis: Source catalog labels Instruction SHA-256: 7f2cb01b99fc59d04ae979b415bf6fea96dd781901a8da82f6dfc53eafb04bbc