Conversion cuTile → Triton

Convertir les kernels @ct.kernel en @triton.jit. Mappage API : references/api-mapping.md (cuTile → Triton).

Dans le Markdown de cette skill, la syntaxe de lancement Triton kernel［grid］(…) utilise des crochets Unicode pour que les vérificateurs de liens ne parsent pas [grid](…) comme un hyperlien ; utilisez des crochets ASCII normaux dans le vrai code Triton.

Instructions

Suivez le workflow par phases dans translations/workflow.md. Chaque conversion doit passer par analyze → convert → validate → test → benchmark, avec des gates explicites avant de progresser. Utilisez les documents dans Workflow Selection quand la tâche correspond à un cas particulier (erreurs, flags de layout, perf).

0. Stratégie d'optimisation (perf-sensible / attention) — Si l'op est attention, FMHA, sliding window, soft cap, ou GQA (ex. Gemma gemma_attention), lisez references/optimization-strategy.md avant de convertir la boucle interne, puis appliquez §4 Gemma FMHA checklist. Pour les autres kernels GEMM/BMM/attention-adjacents, survolez quand même §2–§3 de ce fichier après que TMA soit fait.

1. Sélectionner le chemin — Op TileGym existant : mode standard dans translations/workflow.md. Si la source cuTile utilise transpose / transpose_v, layouts duaux, ou chemins de style MLA, lisez translations/advanced-patterns.md avant d'écrire Triton (deux kernels + grille META, pas un kernel + tl.trans).

2. Pré-vol — Exécutez les commandes grep Pre-flight Analysis sur la source cuTile. Comptez les définitions @ct.kernel ; notez les ct.load/ct.store, ct.launch, Constant et flags de layout pertinents pour TMA.

3. Lire le mappage — Gardez references/api-mapping.md ouvert pour les paires API cuTile → Triton. Pour les défaillances à l'exécution (adresse illégale, dtype, strides), utilisez references/debugging.md.

4. Convertir — Copiez la Conversion Checklist dans une liste de tâches et exécutez dans l'ordre. Structure et placement de fichiers : translations/file-structure.md. Obligatoire : tout load/store de tuile 2D+ de forme bloc utilise tl.make_tensor_descriptor (TMA), pas du tl.load(ptr+offs, mask=…) brut pour tuiles complètes—sauter cela est la source la plus courante de grandes régressions. Côté host : lancement Triton entre crochets <code>kernel［grid］(args)</code> avec tuple ou lambda META: (…) pour autotune ; pas de ct.launch.

5. Valider — Vérification syntaxique du nouveau module Triton ; exécutez les targets pytest TileGym pertinents pour l'op : pytest tests/ops/test_<op>.py -k "triton" -vs. Corrigez les défaillances avant le benchmark.

6. Benchmark — Comparez Triton vs cuTile sur les tests de perf. Si Triton est clairement plus lent, suivez PERFORMANCE ANALYSIS (Phase c2t-5) dans translations/workflow.md et references/optimizing-reference.md pour GEMM/BMM/attention ; utilisez references/optimization-strategy.md comme checklist ordonnée. Si vous voyez des ralentissements de 10–50×, lisez d'abord CRITICAL PERFORMANCE PATTERNS dans le même fichier de workflow.

Règles d'exécution (OBLIGATOIRE) :

• Créez et suivez la conversion checklist (ex. TodoWrite) avant de modifier le code du kernel ; complétez les étapes dans l'ordre—ne sautez pas le pré-vol ou les décisions TMA.
• Pour attention / FMHA / Gemma / GQA / soft cap / sliding window : lisez references/optimization-strategy.md et appliquez §4 avant de traiter la conversion comme optimisée.
• Ne mettez pas en production des loads de tuile 2D+ brut pointeur+mask où TMA s'applique ; documentez toute exception intentionnelle.
• Si les tests ou benchmarks échouent à un gate, arrêtez et corrigez avant de déclarer la conversion terminée—ne mettez pas en pile de changements non vérifiés.

Workflow Selection

• Op TileGym existant → Mode Standard : translations/workflow.md
• Erreurs (cudaErrorIllegalAddress, mismatch de forme, mismatch numérique) → references/debugging.md
• Patterns avancés (TMA, flags de layout dual transpose, autotune + grille META, Array.slice, ct.gather().item()) → translations/advanced-patterns.md (deux kernels de style MLA, évitez une régression de 3–15× sur transpose=False).
• Performance (kernel Triton plus lent que cuTile, autotuning, profiling) → translations/workflow.md (section PERFORMANCE ANALYSIS (Phase c2t-5))
• Hub de stratégie d'optimisation (checklist ordonnée : advanced-patterns + optimizing-reference) → references/optimization-strategy.md — lisez d'abord pour attention/FMHA/Gemma ; puis creusez les deux docs source au besoin
• Optimiser GEMM/BMM/attention (après TMA, ou Triton 10–20% plus lent) → references/optimizing-reference.md — EVEN_K fast path, transpose via arithmétique de pointeur, layout de grille, largeur d'autotune, epilogue subtile ; utilisez ces patterns pendant la conversion et avant la validation de perf (résumé dans optimization-strategy §2–§3)
• Conversion Gemma attention / GQA FMHA → references/optimization-strategy.md §4
• Optimisation Blackwell (kernels complexes avec algorithmes itératifs, register pressure, loop unrolling) → references/optimizing-reference.md §9 — descripteurs TMA, loop_unroll_factor, autotuning d'occupancy, block sizes compatible TMEM, slab allocator, conception dual-path kernel
• ⚠️ Régression 10-50× (ralentissement catastrophique après conversion) → translations/workflow.md — section CRITICAL PERFORMANCE PATTERNS (AVOID 10-50x REGRESSION)
• ⚠️ Bonne perf sur transpose=True uniquement, effondrement sur transpose=False (ou inverse) → translations/advanced-patterns.md — §1 Flag de layout dual ; deux kernels @triton.jit + grid = lambda META: (... META["BLOCK_H"] ...)

Pre-flight Analysis (Exécuter AVANT de convertir)

# Compter les kernels (seul le kernel principal reçoit @triton.jit, les helpers restent def classique)
grep "@ct\.kernel" source.py | wc -l

# Vérifier les patterns nécessitant un traitement spécial
grep "ct\.transpose\|ct\.permute" source.py   # → utiliser tl.trans/tl.permute
grep "ct\.astype" source.py                    # → utiliser .to(dtype)
grep "ct\.load\|ct\.store" source.py          # → TMA pour 2D+ (tl.make_tensor_descriptor), PAS du tl.load(ptr+offs) brut
grep "ct\.launch" source.py                    # → lancement entre crochets : kernel puis [grid] puis (args)
grep "ct\.Constant\|ct\.ConstInt" source.py    # → tl.constexpr
grep "ct\.cdiv" source.py                      # → triton.cdiv (host) ou 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 si nécessaire
grep "transpose\|transpose_v" source.py       # → si hit, lire translations/advanced-patterns.md (kernels duaux + grille META)

Conversion Checklist

Copiez cette checklist et suivez la progression :

Conversion Progress:
 [ ] Step 0 (attention / Gemma FMHA / GQA / soft cap / sliding window): Lire [references/optimization-strategy.md](./references/optimization-strategy.md) et appliquer la checklist §4 avant la boucle interne Triton
 [ ] Step 1: Pré-vol — exécuter les commandes grep ci-dessus, noter les patterns spéciaux et les loads 2D+ (→ TMA)
 [ ] Step 2: Analyser le kernel cuTile source (identifier les patterns, shapes, dtypes)
 [ ] Step 3: Créer un fichier Triton avec la structure correcte (voir translations/file-structure.md)
 [ ] Step 4: Convertir la signature du kernel (args tensor → args pointer, Constant → constexpr)
 [ ] Step 4b: TMA (OBLIGATOIRE pour les loads 2D+) — utiliser tl.make_tensor_descriptor pour chaque load/store de tuile 2D+ ; ne PAS mettre en production du tl.load(ptr+offs,mask) brut pour l'accès de forme bloc (voir workflow.md § TMA OPTIMIZATION)
 [ ] Step 5: Convertir le corps du kernel (appliquer la table des gotchas ci-dessous + mappage API)
 [ ] Step 6: Convertir le wrapper host (grille tuple/lambda, lancement entre crochets : kernel, grid, puis arguments ; pas de ct.launch) ; appeler triton.set_allocator(alloc_fn) si utilisation de TMA
 [ ] Step 7: Valider — vérification syntaxique ou run pytest sur le fichier Triton
 [ ] Step 8: Tester — run pytest, vérifier X passed 0 failed
 [ ] Step 9: Si test échoue → corriger → re-valider → re-tester (boucle jusqu'à vert)
 [ ] Step 10: Benchmark — run test de perf, comparer vs cuTile (voir workflow.md § PERFORMANCE ANALYSIS)
 [ ] Step 10b: Si GEMM/BMM/attention et Triton > 20% plus lent → parcourir [references/optimization-strategy.md](./references/optimization-strategy.md) §2–§3 puis [references/optimizing-reference.md](./references/optimizing-reference.md) (EVEN_K, transpose, grid, autotune, epilogue subtile), puis re-benchmark
 [ ] Step 10c: Si op a `transpose` / flag de layout → lire [translations/advanced-patterns.md](./translations/advanced-patterns.md) ; vérifier **kernels séparés** par layout (pas transpose-kernel + `tl.trans`) ; les lancers **autotuned** utilisent `lambda META: (triton.cdiv(..., META["BLOCK_H"]), ...)` — pas de `BLOCK_H`/`BLOCK_N` fixe via `apply()` sauf si autotune est désactivé

Post-conversion Verification (TMA est obligatoire pour les loads 2D+):
 [ ] TMA: Tous les loads de tuile 2D+ utilisent tl.make_tensor_descriptor(...).load([...]) ; pas de ptr+mask brut pour l'accès 2D+ de forme bloc (sinon régression 5-20×)
 [ ] Grid utilise tuple ou lambda (pas 3-tuple obligatoire comme cuTile)
 [ ] Autotune Triton ajouté si l'op cuTile utilise kernel_configs/autotune (voir workflow § PERFORMANCE ANALYSIS)
 [ ] Grid host utilise triton.cdiv où approprié (pas seulement (a+b-1)//b)
 [ ] Indexage pointeur/offset : Triton utilise les offsets d'élément (ptr + offs), pas l'index de bloc dans tl.load (ou utiliser le descripteur TMA)
 [ ] ct.astype(x, dtype) → x.to(dtype) dans Triton
 [ ] ct.mma(a, b, acc=acc) → tl.dot(a, b, acc) (pas de keyword dans Triton)
 [ ] Args optionnels/None : Triton permet None dans les args du kernel si souhaité (cuTile nécessitait dummy+flag)
 [ ] Masquage appliqué quand BLOCK_SIZE > dimension réelle (idem cuTile) ; avec TMA, les masks peuvent souvent être supprimés pour les tuiles complètes
 [ ] Diviseur de réduction utilise actual_size, PAS BLOCK_SIZE
 [ ] fp32/tf32 : Triton permet par défaut allow_tf32=True ; matcher le comportement cuTile si vous aviez un cast tf32 explicite
 [ ] Si un load 2D+ quelconque utilise du ptr+mask brut (exception uniquement) : documenter POURQUOI TMA n'a pas été utilisé
 [ ] tl.assume() hints d'alignement ajoutés pour strides et pointeurs

Gotchas (Erreurs de traduction les plus courantes) {#gotchas-most-common-translation-errors}

Table exhaustive de patterns qui cassent fréquemment ou régressent lors du portage de @ct.kernel vers @triton.jit — accumulateur mma, cast de type, grille, usage TMA, gestion dtype, flags de layout, matmul batched, etc.

Voir : references/gotchas.md — lisez ceci AVANT d'écrire le kernel Triton.

Performance Gotchas (Risque de régression 10-50×) {#performance-gotchas-10-50x-regression-risk}

⚠️ Ces patterns causent des ralentissements CATASTROPHIQUES. Vérifiez AVANT le benchmark.

Patterns et leur impact : TMA vs ptr+mask brut (5-20×), autotune vs sizes de tuile fixe (2-3×), broadcast_to + tl.dot (10-50×), chaînes extract_slice (2-5×), et plus.

Voir : references/performance-gotchas.md — table complète de risque de régression.

Détails complets : translations/workflow.md — section CRITICAL PERFORMANCE PATTERNS (AVOID 10-50x REGRESSION).

Mappage API complet : references/api-mapping.md.

Dtype math Triton (erf/erfc/exp/log/sqrt) et le pattern « ne remplacez pas erf par tanh » : references/debugging.md — section Triton Math Function Dtype Requirements (CRITICAL).

Stratégie d'optimisation (hub)

Fichier : references/optimization-strategy.md

Résume translations/advanced-patterns.md (flags de layout, kernels duaux, autotune+META, lancer batched, pointeurs Blackwell) et references/optimizing-reference.md (micro-opts post-TMA, §9) en §1–§3 plus une checklist obligatoire §4 Gemma FMHA.

Règle : Pour les conversions attention / FMHA / style Gemma, ouvrez optimization-strategy dans la même session que workflow — ne vous fiez pas à TMA seul pour la validation de perf.

Reference Documents {#reference-documents}

Lisez depuis la perspective cuTile → Triton. Les fichiers principaux vivent dans cette skill sous ``.

Worked Examples

Utilisez cutile_kernel.py comme source et triton_kernel.py comme cible :

Lisez d'abord cutile_kernel.py, puis triton_kernel.py, pour voir le mappage inverse.

⚠️ MANDATORY COMPLETION CHECKLIST (NE PAS SAUTER)

Une conversion N'EST PAS COMPLÈTE tant que TOUS les éléments ne sont pas cochés. Copiez et complétez :

MANDATORY COMPLETION GATES:
 [ ] 1. CORRECTNESS: pytest passe avec 0 défaillances
     Commande : python -m pytest {test_path} -k "test_op and triton" -vs --tb=short
     Gate: "X passed, 0 failed"

 [ ] 2. TMA OPTIMIZATION: Tous les loads de tuile 2D+ utilisent tl.make_tensor_descriptor
     Vérifier : grep -n "tl.load.*mask" triton_file.py | wc -l  # Devrait être 0 pour ops 2D+
     Sauter = régression de performance de 5-20×

 [ ] 3. PERFORMANCE TEST: Triton à moins de 20% de la baseline cuTile
     Commande : python -m pytest {test_path} -k "test_perf" --print-record -v
     OU : Exécuter le script benchmark : cd tests/benchmark && python bench_{op}.py
     Gate: Triton TFLOPS >= 0,8 * CuTile TFLOPS

 [ ] 4. PERFORMANCE COMPARISON RECORDED:
     Documenter les résultats :
     | Config | Triton (TFLOPS) | CuTile (TFLOPS) | Ratio |
     |--------|-----------------|-----------------|-------|
     | [fill] | [fill]          | [fill]          | [fill]|

CONVERSION COMPLETE: Tous les 4 gates sont passés ? → OUI / NON

Pourquoi c'est important :

• Gate 1 attrape les bugs fonctionnels
• Gate 2 prévient les régressions catastrophiques de 5-20× (erreur la plus courante)
• Gate 3 valide que l'optimisation était efficace
• Gate 4 crée un enregistrement de responsabilité

Si un gate échoue : Corrigez et re-vérifiez avant de déclarer la conversion complète.