Loading...
Loading...
Compare original and translation side by side
@ct.kernel@triton.jitkernel[grid](…)[grid](…)@ct.kernel@triton.jitkernel[grid](…)[grid](…)gemma_attentiontranslations/workflow.mdtransposetranspose_vMETAtl.trans@ct.kernelct.loadct.storect.launchConstanttl.make_tensor_descriptortl.load(ptr+offs, mask=…)lambda META: (…)ct.launchpytest tests/ops/test_<op>.py -k "triton" -vsgemma_attentiontranslations/workflow.mdtransposetranspose_vMETAtl.transtl.make_tensor_descriptortl.load(ptr+offs, mask=…)lambda META: (…)ct.launchpytest tests/ops/test_<op>.py -k "triton" -vscudaErrorIllegalAddresstransposeMETAtranspose=Falseloop_unroll_factortranspose=Truetranspose=False@triton.jitgrid = lambda META: (... META["BLOCK_H"] ...)cudaErrorIllegalAddresstransposeMETAtranspose=Falseloop_unroll_factortranspose=Truetranspose=False@triton.jitgrid = lambda META: (... META["BLOCK_H"] ...)undefinedundefinedundefinedundefinedConversion 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 >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转换进度:
[ ] 步骤0(注意力/Gemma FMHA/GQA/软上限/滑动窗口):阅读[references/optimization-strategy.md](./references/optimization-strategy.md)并应用§4检查清单,再编写Triton内循环
[ ] 步骤1:预检查 — 运行上述grep命令,记录特殊模式和二维及以上加载(→ TMA)
[ ] 步骤2:分析源cuTile内核(识别模式、形状、数据类型)
[ ] 步骤3:创建结构正确的Triton文件(参考translations/file-structure.md)
[ ] 步骤4:转换内核签名(张量参数→指针参数,Constant→constexpr)
[ ] 步骤4b:TMA(二维及以上加载强制要求)—— 所有二维及以上分片加载/存储使用tl.make_tensor_descriptor;不得对块形状访问使用原始tl.load(ptr+offs,mask)(参考workflow.md § TMA优化)
[ ] 步骤5:转换内核主体(应用下方陷阱表+API映射)
[ ] 步骤6:转换主机包装器(网格元组/ lambda,方括号式启动:kernel、网格、然后参数;无需ct.launch);若使用TMA需调用triton.set_allocator(alloc_fn)
[ ] 步骤7:验证 — 对Triton文件运行pytest或语法检查
[ ] 步骤8:测试 — 运行pytest,验证全部通过、无失败
[ ] 步骤9:若测试失败→修复→重新验证→重新测试(循环至全部通过)
[ ] 步骤10:基准测试 — 运行性能测试,与cuTile对比(参考workflow.md § 性能分析)
[ ] 步骤10b:若为GEMM/BMM/注意力且Triton慢于cuTile20%以上→浏览[references/optimization-strategy.md](./references/optimization-strategy.md) §2–§3,再参考[references/optimizing-reference.md](./references/optimizing-reference.md)(EVEN_K、转置、网格、自动调优、结尾子分片),然后重新基准测试
[ ] 步骤10c:若算子包含`transpose`/布局标志→阅读[translations/advanced-patterns.md](./translations/advanced-patterns.md);验证每个布局使用**独立内核**(而非转置内核+`tl.trans`);**自动调优**启动使用`lambda META: (triton.cdiv(..., META["BLOCK_H"]), ...)` — 除非禁用自动调优,否则不得通过`apply()`固定`BLOCK_H`/`BLOCK_N`
转换后验证(二维及以上加载必须使用TMA):
[ ] TMA:所有二维及以上分片加载使用tl.make_tensor_descriptor(...).load([...]);块形状二维及以上访问不得使用原始ptr+mask(否则会出现5-20倍性能退化)
[ ] 网格使用元组或lambda(无需像cuTile那样必须使用三元组)
[ ] 若cuTile算子使用kernel_configs/autotune,需添加Triton自动调优(参考工作流§ 性能分析)
[ ] 主机网格在合适场景下使用triton.cdiv(不得仅使用(a+b-1)//b)
[ ] 指针/偏移索引:Triton使用元素偏移(ptr + offs),而非tl.load中的块索引(或使用TMA描述符)
[ ] ct.astype(x, dtype) → Triton中使用x.to(dtype)
[ ] ct.mma(a, b, acc=acc) → Triton中使用tl.dot(a, b, acc)(无关键字参数)
[ ] 可选/None参数:TMA允许内核参数为None(cuTile需要占位符+标志)
[ ] 当BLOCK_SIZE > 实际维度时应用掩码(与cuTile相同);使用TMA时,完整分片通常可移除掩码
[ ] 归约除数使用实际大小,而非BLOCK_SIZE
[ ] fp32/tf32:Triton默认allow_tf32=True;若cuTile有显式tf32转换需匹配其行为
[ ] 若任何二维及以上加载使用原始ptr+mask(仅例外情况):记录未使用TMA的原因
[ ] 为步长和指针添加tl.assume()对齐提示@ct.kernel@triton.jit@ct.kernel@triton.jitbroadcast_to + tl.dotextract_slicebroadcast_to + tl.dotextract_sliceMETAMETA| Category | Document | Content |
|---|---|---|
| Strategy | optimization-strategy.md | Ordered hub: advanced-patterns + optimizing-reference; §4 Gemma FMHA mandatory checklist |
| Workflows | translations/workflow.md | Standard c2t conversion (phases + checklist) |
| translations/file-structure.md | Where to place Triton files when converting from cuTile | |
| translations/advanced-patterns.md | Dual layout flags (transpose), autotune + | |
| API | api-mapping.md | cuTile → Triton mapping |
| optimizing-reference.md | GEMM/BMM/attention optimizations (EVEN_K, transpose, grid, autotune, epilogue subtile) | |
| Gotchas | gotchas.md | Common cuTile→Triton translation errors (mma, dtype, grid, TMA, layout flags) |
| performance-gotchas.md | 10-50× regression-risk table (TMA vs ptr+mask, broadcast_to, extract_slice chains, autotune) | |
| Testing & errors | references/debugging.md | Triton runtime errors (cudaErrorIllegalAddress, pointer type, stride overflow) |
| 分类 | 文档 | 内容 |
|---|---|---|
| 策略 | optimization-strategy.md | 有序中心文档: 高级模式+优化参考;§4 Gemma FMHA强制检查清单 |
| 工作流 | translations/workflow.md | 标准c2t转换(阶段+检查清单) |
| translations/file-structure.md | 从cuTile转换时Triton文件的放置规则 | |
| translations/advanced-patterns.md | 双布局标志(transpose)、自动调优+ | |
| API | api-mapping.md | cuTile → Triton映射表 |
| optimizing-reference.md | GEMM/BMM/注意力优化(EVEN_K、转置、网格、自动调优、结尾子分片) | |
| 陷阱 | gotchas.md | cuTile→Triton转换常见错误(mma、数据类型、网格、TMA、布局标志) |
| performance-gotchas.md | 10-50倍退化风险表(TMA vs ptr+mask、broadcast_to、extract_slice链、自动调优) | |
| 测试与错误 | references/debugging.md | Triton运行时错误(cudaErrorIllegalAddress、指针类型、步长溢出) |
| Example | Directory | Complexity |
|---|---|---|
| Vector Add | examples/01_vector_add/ | Basic |
| Softmax | examples/02_softmax/ | Intermediate |
| LayerNorm | examples/03_layernorm/ | Intermediate |
| MatMul | examples/04_matmul/ | Advanced |
| Attention | examples/05_attention/ | Advanced |
cutile_kernel.pytriton_kernel.py| 示例 | 目录 | 复杂度 |
|---|---|---|
| 向量加法 | examples/01_vector_add/ | 基础 |
| Softmax | examples/02_softmax/ | 中级 |
| LayerNorm | examples/03_layernorm/ | 中级 |
| 矩阵乘法 | examples/04_matmul/ | 高级 |
| 注意力 | examples/05_attention/ | 高级 |
cutile_kernel.pytriton_kernel.pyMANDATORY 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强制完成检查点:
[ ] 1. 正确性:pytest全部通过,无失败
命令:python -m pytest {test_path} -k "test_op and triton" -vs --tb=short
检查点:“X passed, 0 failed”
[ ] 2. TMA优化:所有二维及以上分片加载使用tl.make_tensor_descriptor
验证:grep -n "tl.load.*mask" triton_file.py | wc -l # 二维及以上算子结果应为0
跳过此步骤会导致5-20倍性能退化
[ ] 3. 性能测试:Triton性能达到cuTile基准的80%以上
命令:python -m pytest {test_path} -k "test_perf" --print-record -v
或:运行基准脚本:cd tests/benchmark && python bench_{op}.py
检查点:Triton TFLOPS >= 0.8 * CuTile TFLOPS
[ ] 4. 性能对比已记录:
记录结果:
| 配置 | Triton (TFLOPS) | CuTile (TFLOPS) | 比值 |
|--------|-----------------|-----------------|-------|
| [填写] | [填写] | [填写] | [填写]|
转换完成:所有4个检查点均通过?→ 是 / 否