skills.rest
NVIDIA
NVIDIAOFFICIAL
Explore all skills by this author →
WebsiteGitHub

tilegym-converting-cutile-to-triton

Converts cuTile GPU kernels to Triton with TMA optimization and debugging guidance.

3.2k|370|Updated Feb 25, 2026
Search Tags:#conversion#CUDA#kernel optimization#cutile#GPU kernel#Triton#TMA
One-click install
npx skills add https://github.com/NVIDIA/skills --skill tilegym-converting-cutile-to-triton
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: tilegym-converting-cutile-to-triton
Source: https://github.com/NVIDIA/skills/tree/main/skills/tilegym-converting-cutile-to-triton
Command: npx skills add https://github.com/NVIDIA/skills --skill tilegym-converting-cutile-to-triton

SYSTEM DOCUMENTATION & REQUIREMENTS

💡 This Skill includes references (resource) components.

What problem does it solve?

Porting NVIDIA cuTile kernels (@ct.kernel) to Triton (@triton.jit) is error-prone: wrong API mappings, missing TMA usage, and layout-flag mistakes cause crashes or 5-50x performance regressions. This Skill provides a phase-gated conversion workflow with verified API mappings, checklists, and debugging guides.

Core Features & Use Cases

  • Guided Conversion Workflow: A gated analyze → convert → validate → test → benchmark pipeline with a mandatory completion checklist covering correctness, TMA usage, and performance gates.
  • API Mapping & Gotchas: Complete cuTile-to-Triton mapping tables (ct.load/ct.store, ct.mma, ct.launch, ct.Constant) plus documented translation errors like hardcoded pointer types and int32 stride overflow.
  • Performance & Debugging References: TMA (tl.make_tensor_descriptor) enforcement for 2D+ loads, dual-kernel layout flag patterns, autotune grid via META, and fixes for cudaErrorIllegalAddress and numerical mismatches.
  • Use Case: An engineer porting a TileGym softmax or attention kernel from cuTile to Triton follows the checklist, applies the API mapping, runs pytest gates, and benchmarks Triton against the cuTile baseline.

Quick Start

Ask your agent to convert the cuTile kernel in my repository to Triton and validate it with the TileGym pytest suite.

Frequently Asked Questions about tilegym-converting-cutile-to-triton

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

FAQPage Schema
How do I convert a cuTile kernel to Triton?▼

Follow the phase-gated workflow: run pre-flight grep analysis on the cuTile source, apply the cuTile-to-Triton API mapping, convert the kernel and host wrapper, then validate with pytest and benchmark against the cuTile baseline. A mandatory completion checklist gates correctness, TMA usage, and performance.

Why is my converted Triton kernel 5-20x slower than cuTile?▼

The most common cause is using raw tl.load with pointer offsets and masks for 2D+ block-shaped tile loads instead of TMA via tl.make_tensor_descriptor. Other causes include missing autotune, broadcast_to with tl.dot for batched matmul, and extract_slice chains.

How do I fix cudaErrorIllegalAddress after converting to Triton?▼

Check for hardcoded pointer types in tl.load(ptr_table + idx).to(tl.pointer_type(...)) calls, which must use a DTYPE constexpr matching the actual tensor dtype. Also verify stride tensors use int64 instead of int32 to avoid overflow on large matrices.

Does the conversion handle transpose or dual layout flags?▼

Yes, ops exposing transpose or transpose_v flags require two separate @triton.jit kernels with different math and TMA layouts per mode, not one kernel with tl.trans. Autotuned launches must compute the grid from META parameters via a lambda.

What tests validate a converted Triton kernel?▼

Run the TileGym pytest suite filtered for the Triton backend, for example pytest tests/ops/test_<op>.py -k "triton" -vs, requiring zero failures. Performance gates require Triton to reach at least 80% of cuTile TFLOPS on benchmark tests.

Related Skills

Explore More Software Engineering Skills→
yo-stevenyo-steven

adding-cutile-kernel

Adds a cuTile GPU kernel operator to TileGm with full integration.

Community
Advanced
moirai-internalmoirai-internal

diffusion-kernel

Integrate optimized CUDA and Triton kernels for diffusion model inference.

Official
Advanced
igamenovoerigamenovoer

krnopt-hw-aware-optimization

Identify GPU compute capability and workload shape for CUDA kernel optimization.

Community
Advanced
spacecat2002spacecat2002

add-sgl-kernel

Add AOT CUDA or C++ kernels to sgl-kernel with CMake and pytest.

Community
Advanced