Autotuning CuTile

Ajoutez l'autotuning aux kernels CuTile en utilisant l'API exhaustive_search avec le pattern tune-once/cache/direct-launch.

Instructions

Suivez l'arbre de décision pour classifier le kernel, concevoir un espace de recherche, implémenter le pattern tune-once/cache/launch, et valider les performances.

1. Classifier — utilisez l'arbre de décision pour déterminer les dimensions de recherche (occupancy-only vs full tile search)
2. Concevoir l'espace de recherche — sélectionnez le template correspondant dans references/kernel-type-templates.md ; élaguer à ≤ 30 configs dans le code final via des filtres arch (les sondes d'exploration dirigée peuvent temporairement dépasser cette limite — voir Philosophie de conception)
3. Implémenter — ajouter exhaustive_search + cache + ct.launch en suivant le workflow pas à pas ; gérer les écritures sur place avec split-buffer si nécessaire
4. Tester — exécuter la correction avec l'autotuning activé et avec DISABLE_AUTOTUNE=1
5. Valider — benchmark A/B contre la meilleure config fixe connue ; voir references/search-strategies.md
6. Réduire — élaguer les configs morte-poids qui ne gagnent jamais, ciblant ≤ 8 configs par architecture pour minimiser le coût de compilation (étape 10)

Routeur de tâches — Accédez à ce dont vous avez besoin

Quick Reference — Autotuning Occupancy-Only (Tune-Once/Cache/Launch)

La plupart des kernels CuTile (elementwise, reduction, LayerNorm) ne nécessitent que le tuning d'occupancy. Copiez ce pattern :

from types import SimpleNamespace
from cuda.tile.tune import exhaustive_search
import cuda.tile as ct
import torch

def _my_autotune_configs():
    for occ in [1, 2, 4, 8]:
        yield SimpleNamespace(occupancy=occ)

# Cache au niveau du module : tune une fois, launch rapide à jamais après
_autotune_cache = {}

def my_op(x, output):
    stream = torch.cuda.current_stream()
    NUM_SM = torch.cuda.get_device_properties(x.device).multi_processor_count

    # Clé du cache : tout ce qui affecte la config optimale (utiliser str() pour device)
    cache_key = (x.shape, x.dtype, str(x.device))

    if cache_key not in _autotune_cache:
        configs = list(_my_autotune_configs())
        result = exhaustive_search(
            configs,
            stream,
            grid_fn=lambda cfg: (min(NUM_SM * cfg.occupancy, M), 1, 1),
            kernel=my_kernel,
            args_fn=lambda cfg: (x, output, ...),
            hints_fn=lambda cfg: {"occupancy": cfg.occupancy},
        )
        best_cfg = result.best.config
        tuned_kernel = my_kernel.replace_hints(occupancy=best_cfg.occupancy)
        _autotune_cache[cache_key] = (best_cfg, tuned_kernel)  # cache LES DEUX

    cfg, tuned_kernel = _autotune_cache[cache_key]
    grid = (min(NUM_SM * cfg.occupancy, M), 1, 1)
    ct.launch(stream, grid, tuned_kernel, (x, output, ...))

Règles clés :

• Tune une fois, cache, launch directement — exhaustive_search s'exécute seulement au premier appel par shape ; les appels suivants utilisent la config en cache + ct.launch sans surcharge
• Pour les kernels sur place, utilisez split-buffer pendant la recherche (tenseurs d'entrée/sortie séparés)
• Conserver ≤ 30 configs dans le code final (voir Philosophie de conception pour les sondes dirigées temporaires)
• exhaustive_search nécessite une Sequence (list/tuple) — convertir les générateurs avec list()
• L'espace de recherche doit inclure la config fixe originale — cela garantit que l'autotuning ne dégrade jamais les performances

Quand utiliser ce pattern : Le kernel a une taille de bloc fixe (non tunable par tile-size). Inclut : elementwise (SwiGLU, GeGLU), reduction (RMSNorm, LayerNorm), RoPE, et kernels persistants avec des tailles de bloc heuristiques (grouped GEMM).

Pour les kernels complexes (matmul avec tile sizes, FMHA, FP8 avec num_ctas), lisez le guide complet ci-dessous + kernel-type-templates.md.

⚠️ Trois pièges attrapent presque tout le monde — vérifiez avant de soumettre :

• replace_hints sur le chemin chaud ? → Cache BOTH config ET objet kernel depuis exhaustive_search. Appeler replace_hints() à chaque invocation recompile (100–500× plus lent) → Pitfall #7
• Kernel sur place (écrit dans le tenseur d'entrée) ? → DOIT utiliser le pattern split-buffer pendant la recherche → Pitfall #1
• Espace de recherche vide ? → Vérifiez les filtres arch et les contraintes num_ctas → Pitfall #5

Couverture minimale : Sur sm100+, les espaces de recherche FMHA/matmul/varlen doivent inclure à la fois num_ctas=1 et num_ctas=2. Pour les dimensions centrales (tile sizes, occupancy), conserver au moins 2 valeurs distinctes même si incertain lequel est meilleur — laisser exhaustive_search décider.

Quand arrêter le tuning : Une speedup moyenne dans [0.98, 1.02] signifie que votre espace de recherche actuel n'aide pas — mais ne signifie pas qu'aucune config n'aidera. Avant d'arrêter, vérifiez si vous avez couvert les dimensions clés pour ce type de kernel (consulter references/kernel-type-templates.md). Si l'espace de recherche couvre déjà les dimensions recommandées par le template et le meilleur résultat est toujours du bruit de fond, alors arrêtez — les micro-ajustements supplémentaires n'aideront pas. Si les dimensions clés manquent (par ex., jamais essayé num_ctas=2 pour un kernel dual-GEMM), étendez l'espace de recherche plutôt que d'abandonner.

Une fois que les tests de correction passent et le kernel autotuné montre une speedup par rapport à la baseline de config fixe, arrêtez — ne pas ré-exécuter pour "confirmer". Le timing de kernel GPU fluctue ±5–10 % entre les invocations en raison de la mise à l'échelle d'horloge et de la planification OS ; une baisse de timing ultérieure ne signifie pas que votre code est erroné.

Pour améliorer la speedup, modifiez seulement l'espace de recherche d'autotuning (configs, tile sizes, occupancy, num_ctas). Ne modifiez pas d'autre code (wrapper Python, gestion de stream, etc.) pour chasser la speedup — la performance de kernel est déterminée par la sélection de config, pas par le code côté hôte.

Guide de lecture

• Kernels occupancy-only (elementwise, reduction, persistent avec tailles de bloc fixes) : Quick Reference + Checklist des pièges est suffisant — ignorer les docs references/. Pour les kernels sur place, lisez aussi Pitfall #1.
• Kernels complexes (matmul avec tile sizes tunables, FMHA, FP8 avec num_ctas) : Quick Reference → Arbre de décision → Référence API → Workflow pas à pas → docs references/ pertinents.

Résumé en 5 étapes : Classifier kernel → Concevoir l'espace de recherche (parameter-space-design.md) → Implémenter en utilisant le template (kernel-type-templates.md) → Valider avec test A/B → Vérifier la checklist des pièges.

Lecture des références : Lire seulement la référence pertinente pour votre type de kernel — par ex., pour FMHA, lire la section Template 5 dans references/kernel-type-templates.md ; pour les contraintes matérielles, lire seulement la section de l'architecture cible. Éviter de lire toutes les références de bout en bout quand une recherche ciblée suffit.

Philosophie de conception

Construisez un petit espace de recherche précis de bas en haut — pas un grand espace taillé en bas. La compilation CuTile est beaucoup plus lourde que Triton (~0,5-1s par config), donc le code final devrait contenir ≤ 30 configs. L'approche est : classifier le type de kernel en premier, puis construire seulement les configs pertinentes pour ce type et cette architecture.

Exploration dirigée pendant le développement : Si les configs de template initiaux donnent une speedup < 1.0, vous pouvez exécuter une sonde temporaire plus grande (30–100 configs) via bash + python3 -c pour identifier les dimensions qui importent — mais cette sonde doit être directionnelle, pas un produit cartésien aveugle. Utilisez la classification de type de kernel pour décider quelles dimensions varier (par ex., pour dual-GEMM, sonder num_ctas × occupancy en fixant les tile sizes ; pour FMHA, sonder TILE_M × num_ctas en fixant TILE_N). Une fois que la sonde identifie la région gagnante, verrouiller l'espace de recherche du code final à ≤ 8 top candidats. Ne PAS écrire la grande sonde dans le fichier source — c'est un outil de diagnostic unique.

Arbre de décision : Quelles dimensions de recherche ce kernel a-t-il besoin ?

Tous les kernels doivent avoir l'autotuning ajouté. La question n'est pas si faire l'autotuning, mais quelles dimensions rechercher :

Quel type de kernel est-ce ?
├── Compute-bound (matmul, GEMM, FMHA) → A-t-il plusieurs dimensions tunables (tile sizes) ?
│   ├── OUI → Est-ce un kernel fused multi-GEMM (dual-GEMM, par ex. Linear+GLUAct) ?
│   │   ├── OUI → Template 9: low occupancy (1–2), conservative tiles (2× SHMEM/register pressure)
│   │   └── NON → Full search: TILE_M × TILE_N × (TILE_K) × occupancy × num_ctas
│   │              (voir templates matmul/FMHA dans kernel-type-templates.md)
│   └── NON → Occupancy-only search: [1, 2, 4, 8]
│             (voir Quick Reference ci-dessus)
├── Balanced (LayerNorm, reduction + compute) →
│   Occupancy-only search: [1, 2, 4, 8]
│   Expected benefit: 2-15%
└── Memory-bound (CE Loss, pure elementwise) →
    Occupancy-only search: [1, 2, 4, 8]
    Expected benefit: 0-15% (varies by kernel; zero-cost after tuning)

Pourquoi les kernels memory-bound ne recherchent que l'occupancy (pas num_ctas ou tile sizes) :

• num_ctas n'a aucun bénéfice : num_ctas > 1 active TMA multicast, où plusieurs CTAs partagent les données de tile en mémoire partagée (par ex., tiles A/B de matmul réutilisés entre CTAs). Les kernels memory-bound utilisent par-élément ct.gather/ct.scatter sans réutilisation de tile — la coopération multi-CTA ajoute une surcharge sans bénéfice de partage de données.
• Les tile sizes sont pré-déterminés : BLOCK_SIZE pour les kernels memory-bound est déterminé par un balayage hors-ligne (par ex., 1024 est globalement optimal sur B200 sur [256, 512, 1024, 2048, 4096, 8192]). C'est une constante, pas une tunable runtime.
• L'occupancy est le seul bouton efficace : Une occupancy plus haute permet à la GPU de masquer la latence mémoire en basculant vers un autre CTA tandis qu'un est bloqué sur une requête mémoire.

Preuve — expérience CE Loss : Une recherche de 12-config (occupancy × num_ctas) sur Cross-Entropy Loss a donné seulement 2,5% de gain (0,79x → 0,81x vs Triton). La dimension num_ctas n'a contribué à rien ; le résultat a été annulé car le coût de compilation dépassait le bénéfice marginal. Occupancy-only (4 configs) obtient le même résultat avec 3x moins de temps de compilation.

Note sur les kernels memory-bound : Ajouter l'autotuning occupancy-only vaut toujours la peine car :

• Le pattern tune-once/cache/launch n'a pas de surcharge runtime après le premier appel
• L'espace de recherche est minuscule (4 configs, ~2-4s de compilation)
• Même les petites améliorations ont de la valeur à grande échelle

Guide de sélection d'occupancy

L'occupancy contrôle combien de CTAs s'exécutent simultanément par SM. Utilisez ceci comme point de départ lors de la conception de l'espace de recherche d'occupancy :

Utilisez ces plages pour initialiser votre espace de recherche initial. Pour les kernels occupancy-only, [1, 2, 4, 8] couvre la plupart des cas — voir Quick Reference ci-dessus.

Référence API exhaustive_search

Voir references/api-reference.md pour la surface API complète exhaustive_search — signature actuelle, TuningResult, le pattern tune-once/cache/launch, replace_hints, hints de kernel, conception de search_space, et patterns grid_fn.

Workflow pas à pas

Voir references/workflow.md pour le workflow de bout en bout — ajouter l'autotuning à un nouveau kernel, gérer les configs multi-architecture existantes, intégration avec torch.autograd.Function, transfert de config cross-backend (Triton → CuTile), et optimiser une config existante.

Checklist des pièges

Voir references/pitfalls.md pour la liste complète des pièges courants — corruption de données sur place, timeout de compilation, skew de performance cache-froid, interférence de profilage NCU, épuisement du générateur search_space, perte de précision FP8, et recompilation replace_hints sur chemins chauds.

Portée et limites

Cette skill couvre seulement la configuration d'autotuning : conception de l'espace de recherche, invocation exhaustive_search, caching, et ct.launch avec hints tunés. Elle ne modifie pas le code du kernel.

Dans la portée (config d'autotuning) :

• Fonctions génératrices d'espace de recherche
• Appels exhaustive_search() et gestion de résultats
• kernel.replace_hints() pour appliquer les hints tunés
• Logique de cache (conception de clé, gestion de dict)
• ct.launch() avec kernel tuné
• Chemin de fallback DISABLE_AUTOTUNE

Hors de la portée (modifications de code de kernel — NE PAS faire ces changements) :

• Math flags (flush_to_zero, rounding_mode)
• Performance Hints (slice_hint, buffer_depth, copy_config)
• Patterns d'accès mémoire (conversion 2D→1D gather/scatter)
• Optimisations de codegen (safe_offs → padding_value)
• Changements d'algorithme (K-loop split, load balancing)

Suggestions d'optimisation supplémentaires

Après l'ajout de l'autotuning, les optimisations de kernel suivantes peuvent donner des gains additionnels. Celles-ci sont hors du domaine de cette skill — mentionnez-les à l'utilisateur comme prochaines étapes possibles, mais ne les implémentez pas dans le cadre de l'autotuning :

• Math flags : flush_to_zero=True + rounding_mode=APPROX peuvent fournir 34-72% d'amélioration pour les kernels de classe FMHA (défini via les variables d'environnement TILEIR_ENABLE_FTZ=1 TILEIR_ENABLE_APPROX=1 ou dans le code du kernel). Chaîne causale : les tiles plus grandes diminuent initialement la performance de 18-43% en raison de la surcharge de gestion des subnormaux ; l'activation de FTZ+APPROX sauve cela et inverse le résultat en +34-72%. Les math flags sont donc un préalable pour que les configs de tiles plus grandes soient efficaces sur les kernels de classe FMHA.
• Performance Hints : slice_hint, buffer_depth, copy_config — nécessite de modifier le code IR du kernel
• Patterns d'accès mémoire : Utiliser les charges TMA (ct.load) au lieu de ct.gather ; supprimer les vérifications de limites inutiles (check_bounds=False quand sûr)
• Qualité de codegen : Utiliser le paramètre padding_value au lieu du masquage ct.where manuel ; supprimer safe_offs
• Restructuration d'algorithme : K-loop split, load balancing, simplification algébrique

Différences avec l'autotuning Triton

Différences clés : Triton utilise le décorateur @triton.autotune avec les objets Config(...) ; CuTile utilise exhaustive_search() avec configs SimpleNamespace + cache séparé + ct.launch. CuTile n'a pas de num_warps/num_stages (le compilateur décide) — seulement tile sizes + occupancy + num_ctas. La compilation CuTile est plus lourde (conserver ≤30 configs dans le code final). Le cache CuTile est géré par l'utilisateur en mémoire (pas de persistance automatique). CuTile sépare args_fn (args du kernel) de hints_fn (hints du compilateur).

Documents de référence

Références de code source

Fichiers clés : ops/cutile/matmul.py (autotuning matmul), ops/cutile/attention.py (autotuning FMHA), suites/unsloth/cutile/ct_ops.py (shared autotune_configs() occupancy=[1,2,4,8]), suites/unsloth/cutile/swiglu.py (exemple elementwise), suites/unsloth/cutile/rope_embedding.py (pattern split-buffer), suites/unsloth/cutile/grouped_gemm.py (persistent GEMM, occupancy-only).

Exemples concrets

Chaque exemple montre le pattern avant → après : fixed_launch.py (ct.launch en dur) et autotuned_launch.py (refactorisé en tune-once/cache/launch).