mcpField
T
tilegym-converting-cutile-to-triton

tilegym-converting-cutile-to-triton

github.com/nvidia开发3.8k1.5k 次安装今天更新
代码仓库

关于 tilegym-converting-cutile-to-triton

tilegym-converting-cutile-to-triton 是将 cuTile 内核转换为 Triton 的技能,面向维护或移植 GPU 内核的开发者。它采用分阶段流程,并强调测试、基于 TMA 的内存访问及性能检查。

主要功能

  1. 01按阶段执行分析、转换和验证
  2. 02提供 cuTile 与 Triton API 对照
  3. 03指导 2D 及以上访问使用 TMA
  4. 04涵盖布局和注意力等特殊模式
  5. 05通过测试与基准验证结果

适用场景

  1. 01将 TileGym 的 cuTile 算子移植到 Triton
  2. 02排查转换后的数值或运行时错误
  3. 03调查转换后性能退化
  4. 04处理注意力或双布局内核转换

安装

复制下面的命令,在终端中运行。

$npx skills add https://github.com/nvidia/skills --skill tilegym-converting-cutile-to-triton

安装后在 Claude Code 中输入/tilegym-converting-cutile-to-triton确认已加载。

包含的文件

除 SKILL.md 外,安装时一起下载的参考资料和脚本。

BENCHMARK.md
skill-card.md
skill.oms.sig
evals/
evals.json
examples/01_vector_add/
cutile_kernel.py
triton_kernel.py
examples/02_softmax/
cutile_kernel.py
triton_kernel.py
examples/03_layernorm/
cutile_kernel.py
triton_kernel.py
examples/04_matmul/
cutile_kernel.py
triton_kernel.py
examples/05_attention/
cutile_kernel.py
triton_kernel.py
references/
api-mapping.md
debugging.md
gotchas.md
harness-integration.md
optimization-strategy.md
optimizing-reference.md
performance-gotchas.md
translations/
advanced-patterns.md
file-structure.md
workflow.md

详细介绍

cuTile → Triton 转换

将 @ct.kernel 内核转换为 @triton.jit。API 映射:references/api-mapping.md(cuTile → Triton)。

在此技能的 Markdown 中,Triton 启动语法 kernel[grid](…) 使用 Unicode 方括号,避免链接检查器将 [grid](…) 解析为超链接;在实际 Triton 代码中请使用普通 ASCII 方括号。

说明

请遵循 translations/workflow.md 中按阶段设门的工作流。每次转换都应经过 分析 → 转换 → 验证 → 测试 → 基准测试,并在进入下一阶段前明确通过关卡。任务符合特殊情况(错误、布局标志、性能)时,请参阅工作流选择中的文档。

  1. 优化策略(性能敏感 / attention) — 如果算子涉及 attention、FMHA、滑动窗口、soft cap 或 GQA(例如 Gemma gemma_attention),请在转换内层循环之前阅读 references/optimization-strategy.md,然后应用 §4 Gemma FMHA 检查清单。对于其他 GEMM/BMM/attention 相邻内核,在完成 TMA 后,也请快速浏览该文件的 §2–§3。

  2. 选择路径 — 对于已有的 TileGym 算子:使用 translations/workflow.md 中的标准模式。如果 cuTile 源码使用 transpose / transpose_v、双布局或 MLA 风格路径,请在编写 Triton 代码之前阅读 translations/advanced-patterns.md(使用两个内核 + META 网格,而非一个内核 + tl.trans)。

  3. 预检 — 在 cuTile 源码上运行转换前预检分析中的 grep 命令。统计 @ct.kernel 定义数量;记录与 TMA 相关的 ct.load/ct.store、ct.launch、Constant 和布局标志。

  4. 阅读映射文档 — 打开 references/api-mapping.md,查阅 cuTile → Triton API 对照。对于运行时故障(非法地址、dtype、步长),请使用 references/debugging.md。

  5. 转换 — 将转换检查清单复制到待办列表,并按顺序执行。代码结构和文件位置: translations/file-structure.md。**必须遵守:**任何 2D 及以上的块状 tile 加载/存储都必须使用 tl.make_tensor_descriptor(TMA),不能对完整 tile 使用原始 tl.load(ptr+offs, mask=…)——跳过此步骤是造成严重性能回退的最常见原因。主机端:使用元组或用于自动调优的 lambda META: (…) 进行 Triton 方括号启动 kernel[grid](args);不要使用 ct.launch。

  6. 验证 — 对新的 Triton 模块进行语法检查;运行该算子对应的 TileGym pytest 目标:pytest tests/ops/test_<op>.py -k "triton" -vs。基准测试前修复所有失败。

  7. 基准测试 — 在性能测试中比较 Triton 与 cuTile。如果 Triton 明显较慢,请遵循 translations/workflow.md 中的 性能分析(阶段 c2t-5),并在 GEMM/BMM/attention 场景下参阅 references/optimizing-reference.md;使用 references/optimization-strategy.md 中按顺序排列的检查清单。如果发现 10–50× 的性能下降,请先阅读同一工作流文件中的 关键性能模式。

执行规则(必须遵守):

  • 在编辑内核代码之前创建并跟踪转换检查清单(例如使用 TodoWrite);按顺序完成各步骤,不得跳过预检或 TMA 决策。
  • 对于 attention / FMHA / Gemma / GQA / soft cap / 滑动窗口:请阅读 references/optimization-strategy.md 并应用 §4,之后才能将转换视为已优化。
  • 在适用 TMA 的情况下,不得提交使用原始指针 + 掩码的 2D 及以上 tile 加载;任何有意例外都必须记录说明。
  • 如果测试或基准测试未通过关卡,请停止并修复,之后再宣布转换完成;不要叠加未经验证的更改。

工作流选择

  • 已有 TileGym op → 标准模式:translations/workflow.md
  • 错误(cudaErrorIllegalAddress、形状不匹配、数值不匹配)→ references/debugging.md
  • 高级模式(TMA、双布局标志 transpose、自动调优 + META 网格、Array.slice、ct.gather().item())→ translations/advanced-patterns.md(MLA 风格的双内核,避免在 transpose=False 时出现 3–15× 的性能回退)。
  • 性能(Triton 内核比 cuTile 慢、自动调优、性能分析)→ translations/workflow.md(**性能分析(阶段 c2t-5)**部分)
  • 优化策略中心(有序检查清单:advanced-patterns + optimizing-reference)→ references/optimization-strategy.md — 对 attention/FMHA/Gemma,请先阅读;然后按需深入阅读两份源文档
  • 优化 GEMM/BMM/attention(使用 TMA 后,或 Triton 慢 10–20%)→ references/optimizing-reference.md — EVEN_K 快速路径、通过指针运算实现转置、网格布局、自动调优范围、尾处理子块;转换期间和性能验收前使用这些模式(总结见 optimization-strategy §2–§3)
  • Gemma attention / GQA FMHA 转换 → references/optimization-strategy.md §4
  • Blackwell 优化(包含迭代算法、寄存器压力、循环展开的复杂内核)→ references/optimizing-reference.md §9 — TMA 描述符、loop_unroll_factor、占用率自动调优、适配 TMEM 的块大小、slab 分配器、双路径内核设计
  • ⚠️ 性能回退 10–50 倍(转换后出现灾难性降速)→ translations/workflow.md — **关键性能模式(避免性能回退 10–50 倍)**部分
  • ⚠️ 仅在 transpose=True 时性能良好,在 transpose=False 时骤降(或相反)→ translations/advanced-patterns.md — §1 双布局标志;两个 @triton.jit 内核 + grid = lambda META: (... META["BLOCK_H"] ...)

转换前分析(开始转换前运行)

# Count kernels (only main kernel gets @triton.jit, helpers stay plain def)
grep "@ct\.kernel" source.py | wc -l

# Check for patterns needing special handling
grep "ct\.transpose\|ct\.permute" source.py   # → use tl.trans/tl.permute
grep "ct\.astype" source.py                    # → use .to(dtype)
grep "ct\.load\|ct\.store" source.py          # → TMA for 2D+ (tl.make_tensor_descriptor), NOT raw tl.load(ptr+offs)
grep "ct\.launch" source.py                    # → bracket launch: kernel then [grid] then (args)
grep "ct\.Constant\|ct\.ConstInt" source.py    # → tl.constexpr
grep "ct\.cdiv" source.py                      # → triton.cdiv (host) or Python (a+b-1)//b
grep "ct\.bid\|ct\.num_blocks" source.py       # → tl.program_id/tl.num_programs
grep "1 << .*\.bit_length" source.py           # → triton.next_power_of_2 if needed
grep "transpose\|transpose_v" source.py       # → if hit, read translations/advanced-patterns.md (dual kernels + META grid)

转换检查清单

复制此检查清单并跟踪进度:

Conversion Progress:
 [ ] Step 0 (attention / Gemma FMHA / GQA / soft cap / sliding window): Read [references/optimization-strategy.md](./references/optimization-strategy.md) and apply §4 checklist before inner-loop Triton
 [ ] Step 1: Pre-flight — run grep commands above, note special patterns and 2D+ loads (→ TMA)
 [ ] Step 2: Analyze source cuTile kernel (identify patterns, shapes, dtypes)
 [ ] Step 3: Create Triton file with correct structure (see translations/file-structure.md)
 [ ] Step 4: Convert kernel signature (tensor args → pointer args, Constant → constexpr)
 [ ] Step 4b: TMA (MANDATORY for 2D+ loads) — use tl.make_tensor_descriptor for every 2D+ tile load/store; do NOT ship raw tl.load(ptr+offs,mask) for block-shaped access (see workflow.md § TMA OPTIMIZATION)
 [ ] Step 5: Convert kernel body (apply gotchas table below + API mapping)
 [ ] Step 6: Convert host wrapper (grid tuple/lambda, bracket-style launch: kernel, grid, then arguments; no ct.launch); call triton.set_allocator(alloc_fn) if using TMA
 [ ] Step 7: Validate — run pytest or syntax check on Triton file
 [ ] Step 8: Test — run pytest, verify X passed 0 failed
 [ ] Step 9: If test fails → fix → re-validate → re-test (loop until green)
 [ ] Step 10: Benchmark — run perf test, compare vs cuTile (see workflow.md § PERFORMANCE ANALYSIS)
 [ ] Step 10b: If GEMM/BMM/attention and Triton &gt;20% slower → walk [references/optimization-strategy.md](./references/optimization-strategy.md) §2–§3 then [references/optimizing-reference.md](./references/optimizing-reference.md) (EVEN_K, transpose, grid, autotune, epilogue subtile), then re-benchmark
 [ ] Step 10c: If op has `transpose` / layout flag → read [translations/advanced-patterns.md](./translations/advanced-patterns.md); verify **separate kernels** per layout (not transpose-kernel + `tl.trans`); **autotuned** launches use `lambda META: (triton.cdiv(..., META["BLOCK_H"]), ...)` — no fixed `BLOCK_H`/`BLOCK_N` through `apply()` unless autotune is disabled

Post-conversion Verification (TMA is mandatory for 2D+ loads):
 [ ] TMA: All 2D+ tile loads use tl.make_tensor_descriptor(...).load([...]); no raw ptr+mask for block-shaped 2D+ access (else 5x-20x regression)
 [ ] Grid uses tuple or lambda (not 3-tuple required like cuTile)
 [ ] Triton autotune added if cuTile op used kernel_configs/autotune (see workflow § PERFORMANCE ANALYSIS)
 [ ] Host grid uses triton.cdiv where appropriate (not (a+b-1)//b only)
 [ ] Pointer/offset indexing: Triton uses element offsets (ptr + offs), not block index in tl.load (or use TMA descriptor)
 [ ] ct.astype(x, dtype) → x.to(dtype) in Triton
 [ ] ct.mma(a, b, acc=acc) → tl.dot(a, b, acc) (no keyword in Triton)
 [ ] Optional/None args: Triton allows None in kernel args if desired (cuTile required dummy+flag)
 [ ] Masking applied when BLOCK_SIZE > actual dimension (same as cuTile); with TMA, masks can often be removed for full tiles
 [ ] Reduction divisor uses actual_size, NOT BLOCK_SIZE
 [ ] fp32/tf32: Triton defaults allow_tf32=True; match cuTile behavior if you had explicit tf32 cast
 [ ] If any 2D+ load uses raw ptr+mask (exception only): document WHY TMA was not used
 [ ] tl.assume() alignment hints added for strides and pointers

易错点(最常见的转换错误) {#gotchas-most-common-translation-errors}

列出将 @ct.kernel 移植到 @triton.jit 时常见的出错或性能回退模式的综合表格 — mma 累加器、类型转换、网格、TMA 用法、dtype 处理、布局标志、批量矩阵乘法等。

参见: references/gotchas.md — 编写 Triton 内核前请先阅读此文档。

性能易错点(存在 10–50 倍回退风险) {#performance-gotchas-10-50x-regression-risk}

⚠️ 这些问题会导致灾难性降速。基准测试前请先检查。

模式及其影响:TMA 与原始指针 + 掩码(5–20×)、自动调优与固定分块大小(2–3×)、broadcast_to + tl.dot(10–50×)、extract_slice 链(2–5×)等。

参见: references/performance-gotchas.md — 完整的性能回退风险表。

完整详情: translations/workflow.md — **关键性能模式(避免性能回退 10–50 倍)**部分。

完整 API 映射:references/api-mapping.md。

Triton 数学函数 dtype(erf/erfc/exp/log/sqrt)以及“不要用 tanh 替代 erf”模式:references/debugging.md — **Triton 数学函数的 dtype 要求(关键)**部分。

优化策略(总览)

文件: references/optimization-strategy.md

将 translations/advanced-patterns.md(布局标志、双内核、自动调优 + META、批量启动、Blackwell 指针)和 references/optimizing-reference.md(TMA 后微优化、第 9 节)汇总为 §1–§3,并附上必需的 §4 Gemma FMHA 检查清单。

规则: 对于 attention / FMHA / Gemma 风格的转换,请在同一会话中同时打开 optimization-strategy 和 workflow;性能验收不能只依赖 TMA。

参考文档 {#reference-documents}

从 cuTile → Triton 的视角阅读。核心文件位于此技能的 `` 下。

类别文档内容
策略optimization-strategy.md有序总览: advanced-patterns + optimizing-reference;§4 Gemma FMHA 必需检查清单
工作流translations/workflow.md标准 c2t 转换(阶段 + 检查清单)
translations/file-structure.md从 cuTile 转换时 Triton 文件的放置位置
translations/advanced-patterns.md双布局标志(转置)、autotune + META 网格、MLA 风格的双内核
APIapi-mapping.mdcuTile → Triton 映射
optimizing-reference.mdGEMM/BMM/attention 优化(EVEN_K、转置、网格、autotune、epilogue 子块)
注意事项gotchas.md常见的 cuTile→Triton 翻译错误(mma、dtype、网格、TMA、布局标志)
performance-gotchas.md10-50× 性能回退风险表(TMA 与 ptr+mask、broadcast_to、extract_slice 链、autotune)
测试与错误references/debugging.mdTriton 运行时错误(cudaErrorIllegalAddress、指针类型、步长溢出)

完整示例

将 cutile_kernel.py 用作源文件,将 triton_kernel.py 用作目标文件:

示例目录复杂度
向量加法examples/01_vector_add/基础
Softmaxexamples/02_softmax/中级
LayerNormexamples/03_layernorm/中级
MatMulexamples/04_matmul/高级
Attentionexamples/05_attention/高级

先阅读 cutile_kernel.py,再阅读 triton_kernel.py,以了解反向映射。

⚠️ 必须完成的检查清单(请勿跳过)

所有项目均已勾选后,转换才算完成。请复制并填写:

MANDATORY COMPLETION GATES:
 [ ] 1. CORRECTNESS: pytest passes with 0 failures
     Command: python -m pytest {test_path} -k "test_op and triton" -vs --tb=short
     Gate: "X passed, 0 failed"

 [ ] 2. TMA OPTIMIZATION: All 2D+ tile loads use tl.make_tensor_descriptor
     Verify: grep -n "tl.load.*mask" triton_file.py | wc -l  # Should be 0 for 2D+ ops
     Skip = 5-20x performance regression

 [ ] 3. PERFORMANCE TEST: Triton within 20% of cuTile baseline
     Command: python -m pytest {test_path} -k "test_perf" --print-record -v
     OR: Run benchmark script: cd tests/benchmark && python bench_{op}.py
     Gate: Triton TFLOPS >= 0.8 * CuTile TFLOPS

 [ ] 4. PERFORMANCE COMPARISON RECORDED:
     Document results:
     | Config | Triton (TFLOPS) | CuTile (TFLOPS) | Ratio |
     |--------|-----------------|-----------------|-------|
     | [fill] | [fill]          | [fill]          | [fill]|

CONVERSION COMPLETE: All 4 gates passed? → YES / NO

这很重要,因为:

  • Gate 1 可发现功能缺陷
  • Gate 2 可防止灾难性的 5-20x 性能回退(最常见的错误)
  • Gate 3 可验证优化是否有效
  • Gate 4 可生成责任记录

如果任何关卡未通过: 请先修复并重新验证,再宣布完成。

评论

支持 Markdown