1 · Premier noyau CUDA¶

Le programme minimal, la chaîne d'outils, et le modèle mental hôte/périphérique.

1.1 Le programme complet¶

#include <cstdio>

// Le noyau : s'exécute sur le GPU, un exemplaire par thread
__global__ void saxpy(int n, float a, const float* x, float* y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        y[i] = a * x[i] + y[i];
    }
}

int main() {
    const int n = 1 << 20;               // 1 048 576 éléments
    const size_t octets = n * sizeof(float);

    // 1. Allouer et remplir côté hôte
    float *h_x = (float*)malloc(octets);
    float *h_y = (float*)malloc(octets);
    for (int i = 0; i < n; ++i) { h_x[i] = 1.0f; h_y[i] = 2.0f; }

    // 2. Allouer côté périphérique
    float *d_x, *d_y;
    cudaMalloc(&d_x, octets);
    cudaMalloc(&d_y, octets);

    // 3. Copier hôte → périphérique
    cudaMemcpy(d_x, h_x, octets, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y, octets, cudaMemcpyHostToDevice);

    // 4. Lancer
    const int threads_par_bloc = 256;
    const int blocs = (n + threads_par_bloc - 1) / threads_par_bloc;
    saxpy<<<blocs, threads_par_bloc>>>(n, 2.0f, d_x, d_y);

    // 5. Copier périphérique → hôte (synchronise implicitement)
    cudaMemcpy(h_y, d_y, octets, cudaMemcpyDeviceToHost);

    printf("y[0] = %f  (attendu 4.0)\n", h_y[0]);

    cudaFree(d_x); cudaFree(d_y);
    free(h_x); free(h_y);
    return 0;
}

Compilation :

nvcc -O3 -arch=sm_90 saxpy.cu -o saxpy   # sm_90 = Hopper
./saxpy

Ce programme contient tous les concepts du modèle CUDA. Le reste du cours ne fait que les raffiner.

1.2 Les cinq étapes, décortiquées¶

Étape 0 : deux mémoires, deux mondes¶

h_x vit dans la RAM du CPU. d_x vit dans la HBM du GPU. Ce sont des espaces d'adressage différents. Déréférencer d_x depuis le CPU est un segfault ; déréférencer h_x depuis un noyau en est un aussi (avec un message moins clair).

La convention h_ / d_ est universelle dans le monde CUDA. Utilisez-la : elle évite des heures de débogage.

Les exceptions

Sur les plateformes à mémoire unifiée (cudaMallocManaged) ou cohérente (Grace Hopper, MI300A, Apple Silicon), il n'y a qu'un espace d'adressage. C'est plus simple, parfois plus lent, et cela masque des problèmes de performance. À réserver au prototypage tant qu'on apprend.

Étape 4 : la syntaxe `<<<>>>`¶

noyau<<<grille, bloc, mem_partagee, flux>>>(arguments);

Paramètre	Type	Rôle
`grille`	`dim3`	nombre de blocs (jusqu'à 3D)
`bloc`	`dim3`	threads par bloc (jusqu'à 3D, max 1 024 au total)
`mem_partagee`	`size_t`	mémoire partagée dynamique, en octets (défaut 0)
`flux`	`cudaStream_t`	flux d'exécution (défaut : le flux nul)

Les deux derniers sont optionnels. On les retrouvera au chapitre 6.

La formule du nombre de blocs¶

int blocs = (n + threads_par_bloc - 1) / threads_par_bloc;

C'est une division entière arrondie vers le haut. Avec n = 1000 et threads = 256 : \((1000 + 255)/256 = 4\) blocs, soit 1 024 threads pour 1 000 éléments. Les 24 threads en trop sont neutralisés par la garde if (i < n).

Ne jamais oublier la garde de bord

Sans if (i < n), les 24 threads excédentaires écrivent hors du tableau. Sur GPU, cela ne provoque pas toujours un plantage : cela corrompt silencieusement de la mémoire voisine. C'est le bug le plus coûteux du débutant, et compute-sanitizer est fait pour le détecter :

compute-sanitizer --tool memcheck ./saxpy

Le lancement est asynchrone¶

saxpy<<<blocs, threads>>>(...);   // retourne immédiatement
// le CPU continue ici pendant que le GPU travaille
cudaMemcpy(...);                  // ← attend implicitement la fin du noyau

C'est une source de confusion majeure lors des mesures de temps :

// FAUX : mesure le temps de mise en file, pas d'exécution
auto t0 = now();
noyau<<<g, b>>>(...);
auto t1 = now();

// CORRECT
auto t0 = now();
noyau<<<g, b>>>(...);
cudaDeviceSynchronize();
auto t1 = now();

Voir Performance · Profilage pour la méthode de mesure correcte (échauffement, répétitions, événements CUDA).

1.3 Les qualificateurs de fonction¶

Qualificateur	Appelé depuis	S'exécute sur
`__global__`	hôte (ou périphérique en dynamic parallelism)	périphérique
`__device__`	périphérique	périphérique
`__host__`	hôte	hôte (défaut)
`__host__ __device__`	les deux	les deux (compilé deux fois)

__global__ désigne un point d'entrée de noyau. Sa signature doit retourner void — la valeur de retour n'aurait aucun sens pour des milliers de threads.

__host__ __device__ est très utile pour du code partagé (fonctions mathématiques, structures) :

__host__ __device__ inline float sigmoid(float x) {
    return 1.0f / (1.0f + expf(-x));
}

1.4 Les variables intégrées¶

Accessibles dans tout noyau, elles constituent l'unique moyen pour un thread de savoir qui il est :

Variable	Type	Contenu
`threadIdx.{x,y,z}`	`uint3`	position du thread dans son bloc
`blockIdx.{x,y,z}`	`uint3`	position du bloc dans la grille
`blockDim.{x,y,z}`	`dim3`	taille du bloc
`gridDim.{x,y,z}`	`dim3`	taille de la grille
`warpSize`	`int`	32 sur NVIDIA (voir la mise en garde du chapitre SIMT)

D'où la formule universelle en 1D :

int i = blockIdx.x * blockDim.x + threadIdx.x;

En 2D :

int col = blockIdx.x * blockDim.x + threadIdx.x;
int lig = blockIdx.y * blockDim.y + threadIdx.y;
int idx = lig * largeur + col;    // ← attention à l'ordre

L'erreur d'ordre en 2D

En CUDA, x est la dimension la plus rapide : les threads de x consécutifs sont dans le même warp. Pour un tableau stocké par lignes (row-major, la convention C et PyTorch), il faut donc que x indexe les colonnes. Inverser produit un noyau correct mais 8 à 32× plus lent, parce que les accès ne sont plus coalescés.

C'est probablement l'erreur de performance la plus fréquente chez les débutants, et elle est totalement invisible : le résultat est juste.

1.5 La chaîne de compilation¶

mon_noyau.cu
    │
    ├─→ code hôte (C++)  ──→ gcc/clang ──→ objet hôte
    │
    └─→ code périphérique ──→ nvcc frontend
                                  │
                                  ├─→ PTX (assembleur virtuel, portable)
                                  │      │
                                  │      └─→ ptxas ──→ SASS (binaire, par architecture)
                                  │
                                  └─→ fatbinary (PTX + SASS, embarqué dans l'exécutable)

Deux niveaux d'assembleur, et c'est important :

PTX est un assembleur virtuel, indépendant de l'architecture réelle. Il est stable, documenté, et peut être compilé à l'exécution par le pilote (JIT) pour une architecture plus récente que celle connue à la compilation.
SASS est le binaire réel de la puce. Non documenté officiellement, différent à chaque génération, et c'est ce que le matériel exécute.

D'où la distinction dans les options de compilation :

# Génère du SASS pour sm_90 ET du PTX pour compatibilité future
nvcc -gencode arch=compute_90,code=sm_90 \
     -gencode arch=compute_90,code=compute_90 fichier.cu

Les raccourcis utiles :

Option	Effet
`-arch=sm_90`	SASS pour Hopper uniquement
`-arch=sm_90a`	Hopper + fonctionnalités « accélérées » (`wgmma`, TMA)
`-arch=native`	détecte la carte présente
`-Xptxas -v`	affiche registres, mémoire partagée, spills
`-lineinfo`	métadonnées pour le profileur (indispensable)
`-G`	débogage complet (désactive les optimisations : ne jamais profiler avec)

Le suffixe a de sm_90a

Certaines fonctionnalités (TMA, wgmma, tcgen05) ne sont accessibles qu'avec le suffixe a (architecture-specific). Le code produit n'est pas compatible en avant : un binaire sm_90a ne tournera pas sur Blackwell. C'est un compromis assumé : ces instructions sont trop liées au matériel pour être virtualisées en PTX portable.

Pour inspecter le SASS généré :

cuobjdump -sass ./mon_binaire | less
nvdisasm -c mon_noyau.cubin

Lire du SASS est un savoir-faire avancé mais parfois indispensable : c'est le seul moyen de vérifier qu'une instruction MMA a bien été émise, ou de compter les instructions réellement exécutées dans une boucle.

1.6 La gestion des erreurs¶

Presque toutes les fonctions CUDA renvoient un cudaError_t. Presque personne ne le vérifie, et c'est une mauvaise idée.

La macro standard, à copier dans tous vos projets :

#include <cstdio>
#include <cstdlib>

#define CUDA_CHECK(appel)                                                \
    do {                                                                 \
        cudaError_t err_ = (appel);                                      \
        if (err_ != cudaSuccess) {                                       \
            fprintf(stderr, "Erreur CUDA %s:%d : %s\n",                  \
                    __FILE__, __LINE__, cudaGetErrorString(err_));       \
            exit(EXIT_FAILURE);                                          \
        }                                                                \
    } while (0)

Usage :

CUDA_CHECK(cudaMalloc(&d_x, octets));
CUDA_CHECK(cudaMemcpy(d_x, h_x, octets, cudaMemcpyHostToDevice));

Le lancement de noyau est un cas spécial : <<<>>> ne renvoie rien. Il faut interroger l'erreur après :

saxpy<<<blocs, threads>>>(n, 2.0f, d_x, d_y);
CUDA_CHECK(cudaGetLastError());        // erreurs de lancement (config invalide)
CUDA_CHECK(cudaDeviceSynchronize());   // erreurs d'exécution (accès illégal)

Les erreurs collantes

Certaines erreurs CUDA — notamment cudaErrorIllegalAddress — sont collantes (sticky) : une fois survenues, tout le contexte CUDA est corrompu et tous les appels suivants échouent. Il faut détruire le contexte et recommencer.

Conséquence pratique : si vous voyez une cascade d'erreurs, seule la première a du sens. Cherchez-la.

1.7 Ce que fait vraiment `cudaMalloc`¶

cudaMalloc est un appel synchrone et coûteux (plusieurs dizaines de microsecondes) : il synchronise le périphérique et peut prendre un verrou global dans le pilote. Ne jamais l'appeler dans une boucle chaude.

Les alternatives modernes :

// Allocation asynchrone sur un flux, avec pool de mémoire (CUDA 11.2+)
cudaMallocAsync(&d_ptr, octets, flux);
cudaFreeAsync(d_ptr, flux);

C'est ce que fait PyTorch en interne avec son caching allocator : il alloue de gros blocs une fois et les redistribue lui-même. C'est aussi pourquoi nvidia-smi montre PyTorch consommant toute la VRAM même si votre modèle est petit.

Résumé du chapitre¶

À retenir

Cinq étapes : allouer hôte, allouer périphérique, copier, lancer, copier en retour.
Le lancement est asynchrone. Toute mesure de temps sans synchronisation est fausse.
i = blockIdx.x * blockDim.x + threadIdx.x, toujours suivi de if (i < n).
En 2D, x doit indexer la dimension contiguë en mémoire, sinon les accès ne sont pas coalescés.
PTX est virtuel et portable, SASS est réel et spécifique. sm_90a donne accès aux instructions avancées au prix de la compatibilité en avant.
Vérifier les erreurs, y compris après <<<>>> avec cudaGetLastError().

Vérifiez que vous avez compris¶

Ce programme affiche 2.0 au lieu de 4.0. Pourquoi ?

saxpy<<<blocs, threads>>>(n, 2.0f, d_x, d_y);
cudaMemcpy(h_y, d_y, octets, cudaMemcpyDeviceToHost);

Si le résultat est la valeur initiale, c'est que le noyau n'a pas tourné : erreur de configuration de lancement (par exemple threads > 1 024, ou blocs = 0), ou noyau compilé pour une autre architecture.

cudaMemcpy synchronise mais ne signale pas l'échec du noyau de manière évidente. D'où l'obligation d'appeler cudaGetLastError() juste après le lancement : il aurait renvoyé cudaErrorInvalidConfiguration.

Pourquoi -G ne doit-il jamais être utilisé pour mesurer les performances ?

-G désactive toutes les optimisations du compilateur de périphérique et ajoute des métadonnées de débogage volumineuses. Un noyau compilé avec -G peut être 5 à 20× plus lent. Pour profiler, utilisez -lineinfo, qui donne la correspondance ligne source ↔ instruction sans toucher aux optimisations.

Vous voulez traiter 10 milliards d'éléments. blocs = n/256 donne 39 millions de blocs. Est-ce acceptable ?

La limite de gridDim.x est \(2^{31}-1\), donc techniquement oui. Mais c'est une mauvaise idée : chaque bloc a un coût d'ordonnancement fixe, et 39 millions de blocs qui font chacun un travail minuscule sont dominés par ce coût.

La solution idiomatique est la boucle à parcours de grille (grid-stride loop), traitée au chapitre suivant : on lance juste assez de blocs pour remplir le GPU, et chaque thread traite plusieurs éléments dans une boucle.

Chapitre suivant : 2 · Grille, blocs, warps