tilegym-cutile-autotuning

Adds autotuning to CuTile GPU kernels using exhaustive_search with tune-once, cache, and launch patterns.

3.2k|370|Updated Feb 25, 2026
One-click install
npx skills add https://github.com/NVIDIA/skills --skill tilegym-cutile-autotuning
Or copy as Structured Prompt for Agent
Please help me install this Agent Skill.
Skill: tilegym-cutile-autotuning
Source: https://github.com/NVIDIA/skills/tree/main/skills/tilegym-cutile-autotuning
Command: npx skills add https://github.com/NVIDIA/skills --skill tilegym-cutile-autotuning

SYSTEM DOCUMENTATION & REQUIREMENTS

💡 This Skill requires torch, cuda.tile, and includes references (resource) and assets (resource) components.

What problem does it solve?

CuTile GPU kernels ship with hardcoded launch configurations that leave performance on the table, and manually tuning tile sizes, occupancy, and num_ctas across GPU architectures (sm80 through sm120) is error-prone and slow. This Skill guides an agent through adding correct, cached autotuning to CuTile kernels while avoiding the pitfalls that cause data corruption, recompilation slowdowns, and compilation timeouts.

Core Features & Use Cases

  • Tune-Once/Cache/Launch Pattern: Implements exhaustive_search from cuda.tile.tune with module-level caching so tuning runs once per shape and subsequent launches have zero overhead.
  • Kernel-Type Templates: Provides copy-paste autotune templates for 9 kernel types including elementwise, in-place (split-buffer), matmul, persistent GEMM, FMHA, FP8, grouped GEMM, varlen attention, and dual-GEMM fusion.
  • Per-Architecture Configs: Supplies validated search spaces for sm80, sm90, sm100, sm103, and sm120 with hardware constraints on tile sizes, num_ctas, and TMA usage.
  • Pitfall Checklist: Documents 7 common failures such as in-place data corruption, replace_hints recompilation on hot paths, and empty search spaces, with fixes.
  • Use Case: A developer porting a Triton RMSNorm kernel to CuTile asks the agent to add autotuning; the agent classifies it as occupancy-only, applies the 4-config template, caches the tuned kernel, and validates with an A/B benchmark against the fixed-config baseline.

Quick Start

Ask your agent to add autotuning to a specific CuTile kernel, for example: add CuTile autotuning to my matmul kernel using exhaustive_search with per-architecture configs.

Frequently Asked Questions about tilegym-cutile-autotuning

High-intent search queries and answers about installing and using this skill.

FAQPage Schema
How do I add autotuning to a CuTile kernel?

Classify the kernel type with the decision tree, design a search space of SimpleNamespace configs, then call exhaustive_search from cuda.tile.tune once per shape and cache both the best config and the tuned kernel object. Subsequent calls use ct.launch with the cached kernel at zero overhead.

What is the difference between CuTile autotune and Triton autotune?

Triton uses a @triton.autotune decorator with Config objects, while CuTile uses exhaustive_search with SimpleNamespace configs plus a user-managed cache and explicit ct.launch. CuTile has no num_warps or num_stages knobs, only tile sizes, occupancy, and num_ctas, and compilation is heavier so search spaces should stay under 30 configs.

Why does my autotuned CuTile kernel get slower after the first run?

The usual cause is calling replace_hints on every invocation, which recompiles the kernel each time and runs 100 to 500 times slower. Cache both the best config and the tuned kernel object returned by replace_hints after exhaustive_search, then reuse them on the hot path.

How do I autotune an in-place CuTile kernel without data corruption?

Use the split-buffer pattern during the search phase: pass separate input and output tensors to exhaustive_search so repeated trial launches do not corrupt data. The final production ct.launch can then use the real in-place arguments with the cached tuned kernel.

Which num_ctas values are supported on each NVIDIA GPU architecture?

num_ctas greater than 1 requires CGA support, available on sm90 (H100), sm100 (B200), and sm103 (GB300) with values 1, 2, or 4. On sm120 (5090) and Ampere sm80/sm86, always use num_ctas=1 because multi-CTA cooperation gives no benefit or is unsupported.

Why is my CuTile autotune search space returning zero configs?

Empty search spaces usually come from architecture filters that yield nothing for the detected GPU capability, size guards that prune every config, or num_ctas constraints excluding all candidates. Check the conditional yield branches against torch.cuda.get_device_capability() and always include a num_ctas=1 fallback config.