8 · AMD, HIP et ROCm¶

L'état réel de l'alternative à CUDA en 2026 : ce qui fonctionne, ce qui ne fonctionne pas, et ce que coûte un portage.

8.1 La pile¶

Application
    │
    ├─ PyTorch / JAX / TensorFlow    ← support amont
    ├─ Triton (back-end ROCm)        ← fonctionnel
    ├─ Bibliothèques : rocBLAS, MIOpen, rocFFT, RCCL, hipSPARSE
    │
    ├─ HIP  (Heterogeneous Interface for Portability)
    │       ← API quasi identique à CUDA, compile pour AMD ET NVIDIA
    │
    ├─ ROCm runtime + pilote (amdgpu)
    └─ Matériel : CDNA (Instinct) ou RDNA (Radeon)

HIP est le cœur du dispositif : une API dont les noms sont ceux de CUDA avec le préfixe cuda remplacé par hip.

CUDA	HIP
`cudaMalloc`	`hipMalloc`
`cudaMemcpy`	`hipMemcpy`
`__syncthreads()`	`__syncthreads()`
`<<<grille, bloc>>>`	`<<<grille, bloc>>>`
`cublasSgemm`	`hipblasSgemm`
`nvcc`	`hipcc`

Un outil de conversion automatique est fourni :

hipify-perl mon_noyau.cu > mon_noyau.hip
hipify-clang --inplace mon_projet/     # version basée sur clang, plus fiable

AMD annonce que l'outillage HIP-NVCC couvre « environ 92 % des API device de CUDA 12.5 » sans intervention.

8.2 Le matériel¶

Génération	Cible	Cartes	Caractéristiques
CDNA 2	`gfx90a`	MI250X	2 dies, FP64 fort
CDNA 3	`gfx942`	MI300X, MI300A	8 XCD, 192 Go HBM3, 5,3 To/s
CDNA 4	`gfx950`	MI350X, MI355X	256 CU, 288 Go HBM3E, 8,0 To/s, MXFP8/6/4 natifs
RDNA	`gfx11xx`, `gfx12xx`	Radeon	grand public, wavefront 32

Le MI355X (CDNA 4) apporte :

256 unités de calcul, 160 Ko de LDS par CU (contre 64 Ko sur CDNA 3) ;
le doublement du débit des Matrix Cores pour les types ≤ 16 bits ;
le support natif de MXFP8, MXFP6 et MXFP4 avec échelle par blocs d'exposant.

Les deux arguments techniques solides d'AMD

1. La capacité mémoire. 288 Go par carte contre 192 sur B200. Pour l'inférence de très grands modèles, cela change le nombre de GPU nécessaires.

2. Le FP64. Le MI300X délivre 81,7 TFLOPS en double précision, contre ~34 sur H100 — un rapport de 2,4×. En HPC traditionnel, c'est décisif, et c'est ce qui explique le choix d'AMD pour les supercalculateurs Frontier et El Capitan.

8.3 Les différences qui comptent pour un portage¶

Le wavefront de 64¶

C'est la différence structurelle. Sur CDNA, un wavefront fait 64 threads (32 sur RDNA).

Tout code qui suppose 32 doit être audité :

// CUDA
for (int off = 16; off > 0; off >>= 1)
    x += __shfl_down_sync(0xffffffff, x, off);

// HIP sur CDNA : une étape de plus, et pas de masque _sync
for (int off = 32; off > 0; off >>= 1)
    x += __shfl_down(x, off, 64);

Les conséquences en cascade : taille de bloc optimale différente, nombre de partiels dans une réduction hiérarchique, calcul d'occupancy, seuils de divergence.

Pas de TMA¶

Aucun équivalent au Tensor Memory Accelerator. Le calcul d'adresses reste à la charge des threads, ce qui consomme registres et instructions.

Conséquence : le motif producteur/consommateur, si rentable sur Hopper, l'est beaucoup moins ici — son gain principal venait de ce qu'un seul thread suffisait à alimenter le bloc.

MFMA est synchrone¶

Pas d'équivalent à wgmma ni à tcgen05. Les instructions MFMA :

sont synchrones ;
prennent leurs opérandes dans les registres vectoriels ;
opèrent à l'échelle du wavefront de 64.

Le pipelining doit donc s'obtenir par ordonnancement d'instructions et préchargement explicite dans le LDS, pas par des mécanismes asynchrones matériels.

Pas de clusters¶

Aucun équivalent aux thread block clusters ni à la mémoire partagée distribuée. La coopération entre CU passe par le L2 ou la HBM.

C'est précisément le vide que le papier Fleet cherche à combler, en introduisant une notion de chiplet-task qui lie travail et données à un XCD donné — puisque « les modèles de programmation actuels (CUDA/HIP) exposent une hiérarchie d'exécution plate qui ne peut pas exprimer la localité ni la synchronisation au niveau du chiplet ».

Les XCD¶

Un MI300X est composé de 8 Accelerator Compute Dies, chacun avec son propre L2. Deux blocs sur des XCD différents ne partagent pas de cache.

C'est à la fois un problème et une opportunité :

problème : la localité de données au niveau du die n'est pas exprimable dans HIP ;
opportunité : en dupliquant les données par die, on obtient des gains substantiels. Kog l'a fait pour son monokernel ; Fleet mesure un passage du taux de réussite L2 de 12 % à 54 % (lot 32) et de 39 % à 61 % (lot 64), avec jusqu'à 37 % de trafic HBM en moins.

8.4 L'état de l'écosystème logiciel¶

Domaine	État en 2026
PyTorch	support amont, mûr
JAX	fonctionnel
Triton	back-end ROCm supporté par AMD, utilisé en production
Helion	documenté par AMD dans son AI Developer Hub
vLLM, SGLang	supportés, FP8 disponible sur MI300X
FlashAttention	portages disponibles, souvent en retard d'une version
CUTLASS	pas d'équivalent direct ; `composable_kernel` joue ce rôle
Débogage	`rocgdb`, `rocprof`, `Omniperf`
ONNX Runtime	intégré

composable_kernel (CK) est la réponse d'AMD à CUTLASS : une bibliothèque de gabarits C++ pour construire des GEMM et des convolutions fusionnées. Elle est moins documentée et moins connue que CUTLASS, mais c'est le bon outil pour les noyaux de pointe sur CDNA.

8.5 Comment porter, concrètement¶

Étape 1 — la conversion mécanique.

hipify-clang --inplace -- -I/chemin/includes src/*.cu

Cela traite les appels d'API et la plupart des intrinsèques. Selon AMD, ~92 % des API device sont couvertes.

Étape 2 — l'audit du wavefront.

grep -rn "32\|0xffffffff\|warpSize\|shfl" src/

Chaque occurrence est un point de vérification. C'est le travail manuel principal.

Étape 3 — les fonctionnalités sans équivalent.

TMA, clusters, DSMEM, wgmma, setmaxnreg : à réécrire selon un autre modèle, pas à traduire.

Étape 4 — le réglage.

Les tailles de blocs, de tuiles et le nombre d'étages optimaux sont différents. Rejouez l'autotuning complet.

Étape 5 — la validation numérique.

Les ordres d'accumulation diffèrent, les bibliothèques mathématiques aussi. Testez avec des tolérances, pas avec l'égalité.

L'estimation honnête du coût

Pour un code CUDA classique (pas d'utilisation des fonctionnalités Hopper) : quelques jours à quelques semaines, dominés par l'audit du wavefront et le réglage.

Pour un code optimisé Hopper (TMA, wgmma, spécialisation des warps, clusters) : c'est une réécriture, pas un portage. Le modèle de performance est différent.

8.6 Les megakernels côté AMD¶

Deux travaux méritent d'être connus, tous deux traités en partie 8.

Kog a construit un monokernel d'inférence sur MI300X, atteignant plus de 3 000 jetons/s par requête sur un modèle de 2 milliards de paramètres en FP16, sur un nœud à 8 GPU. Leurs techniques marquantes :

synchronisation par sentinelle NaN au lieu d'atomiques : 7,59-7,88 µs → 0,80-0,93 µs ;
duplication des tenseurs par XCD pour éviter le trafic inter-chiplet ;
streaming continu des poids vers le LDS avec des hints non-temporels.

Fleet propose une abstraction de tâches hiérarchique liée aux chiplets, évaluée sur MI350 avec Qwen3-8B : latence de décodage 1,3-1,5× inférieure à vLLM aux lots 1-8, et 1,27-1,30× de mieux qu'un megakernel non conscient des chiplets aux lots plus grands.

Ce que cela dit du portage

Les meilleurs résultats AMD ne viennent pas de portages de techniques NVIDIA, mais de techniques conçues pour les particularités d'AMD : les XCD, le LDS de 160 Ko, l'absence de TMA.

C'est la bonne façon d'aborder AMD : non pas « comment reproduire mon noyau Hopper », mais « quelle est la meilleure structure pour cette machine ».

Résumé du chapitre¶

À retenir

HIP est CUDA renommé, avec hipify pour la conversion automatique (~92 % des API device couvertes).
Différences structurelles : wavefront de 64, pas de TMA, MFMA synchrone, pas de clusters ni de DSMEM.
CDNA 4 (MI355X) : 256 CU, 288 Go à 8 To/s, MXFP8/6/4 natifs, débit Matrix Core doublé ≤ 16 bits.
Arguments solides d'AMD : capacité mémoire (288 Go) et FP64 (81,7 TFLOPS sur MI300X, 2,4× un H100).
Un MI300X est 8 XCD avec des L2 séparés : la localité de die est un levier majeur et non exprimable dans HIP.
Portage : quelques jours pour du CUDA classique, réécriture pour du code optimisé Hopper.
Les meilleurs résultats AMD viennent de techniques conçues pour AMD, pas de portages.

Vérifiez que vous avez compris¶

Pourquoi le wavefront de 64 change-t-il plus que le nombre d'étapes d'une réduction ?

Parce qu'il modifie toutes les granularités :

une garde de bord qui laisse 16 threads inactifs gaspille 25 % d'un warp NVIDIA mais 25 % de deux fois plus de threads ;
le seuil de rentabilité de la divergence change ;
le nombre de wavefronts par bloc pour une taille donnée est divisé par deux, ce qui modifie l'occupancy et le calcul des registres ;
une réduction hiérarchique a 6 étapes au lieu de 5, et le nombre de partiels change ;
les accès mémoire sont fusionnés sur 64 threads : les motifs de coalescence optimaux diffèrent.

C'est pour cela que le réglage doit être entièrement refait, et pas seulement ajusté.

Pourquoi la duplication de tenseurs par XCD améliore-t-elle les performances, alors qu'elle augmente le volume mémoire ?

Parce que le goulot n'est pas la capacité mais la bande passante et la localité de cache.

Sans duplication, un tenseur partagé réside dans un seul emplacement. Les blocs des 8 XCD y accèdent, et 7 XCD sur 8 subissent un défaut de L2 suivi d'un accès HBM.

Avec une copie par XCD, chaque bloc trouve ses données dans son L2 local. Fleet mesure le taux de réussite L2 passant de 12 % à 54 % et jusqu'à 37 % de trafic HBM en moins.

C'est le compromis classique espace/temps, appliqué à un niveau de hiérarchie que le modèle de programmation n'expose pas.

Votre code CUDA utilise cooperative_groups::grid_group::sync(). Que devient-il en HIP ?

HIP a un équivalent : hipLaunchCooperativeKernel et les cooperative groups HIP. La sémantique est la même, avec la même contrainte de résidence de tous les blocs.

Mais le coût diffère, et pas qu'un peu. Kog mesure sur MI300X une synchronisation de grille naïve à 7,59-7,88 µs, et attribue à ce mécanisme « environ 35 % du temps total de génération d'un jeton » dans leur implémentation initiale.

Autrement dit : ce qui était coûteux sur NVIDIA est prohibitif sur AMD, à cause de la topologie multi-die. D'où leur remplacement par une synchronisation par sentinelle, 9× plus rapide.

Chapitre suivant : 9 · Portable : SYCL, OpenCL, Vulkan, WebGPU, Metal