Streams, événements et concurrence CPU-GPU

Pourquoi ce chapitre change la donne en production

Un kernel CUDAGPUIAProcesseur graphique parallélisant massivement les calculs matriciels ; indispensable à l'entraînement et à l'inférence des modèles de deep learning.Voir dans le glossaire bien optimisé qui tourne à 90 % d'occupation ne sert à rien si le GPU passe la moitié de son temps à attendre des donnéesdonnéesIAEnsemble d'informations structurées ou non utilisées pour entraîner, évaluer ou alimenter un modèle. La qualité, la quantité et la représentativité des données sont les facteurs décisifs pour les performances en apprentissage automatique.Voir dans le glossaire. Dans un pipeline d'inference réel — chargement d'un batch, copie vers le device, exécution du modèle, récupération des résultats — le code naïf enchaîne ces trois étapes en série : cudaMemcpy, kernel, cudaMemcpy retour, chacune bloquant jusqu'à la suivante. Le GPU reste inactif pendant les transferts, le CPU reste inactif pendant le calcul.

Les CUDA streams cassent cette dépendance artificielle. Un stream est une file d'opérations (copies, lancements de kernels, événements) exécutées dans l'ordre sur le device, mais deux streams différents peuvent progresser en parallèle sur des moteurs matériels distincts — le moteur de copie et le moteur de calcul. C'est ce mécanisme qui permet de traiter le bloc N+1 pendant que le bloc N est encore en train de calculer.

Le recouvrement H2D/calcul/D2H n'est possible qu'avec de la mémoire host verrouillée (pinned, allouée via cudaHostAlloc ou cudaMallocHost). Une copie asynchrone (cudaMemcpyAsync) depuis de la mémoire paginable classique est silencieusement rétrogradée en copie synchrone par le driver. Le code compile, s'exécute, ne recouvre rien — et rien dans les logs ne vous avertit.

Le flux par défaut : le piège numéro un

Toute opération CUDA lancée sans stream explicite atterrit sur le flux par défaut (stream 0). Ce comportement a une conséquence souvent ignorée : dans le mode legacy default stream (le comportement historique, toujours actif sauf configuration contraire), le flux par défaut se synchronise implicitement avec tous les autres streams du contextefenêtre de contexteIAQuantité de texte qu'un modèle peut prendre en compte simultanément : question, documents fournis et historique. Au-delà, les éléments les plus anciens sortent du champ.Voir dans le glossaire. Concrètement, si votre pipeline mélange des appels sur des streams explicites et un seul appel oublié sur le flux par défaut — un cudaMemset de debug, un kernel de warmup — cet appel agit comme une barrière globale et détruit tout le recouvrement que vous aviez patiemment construit.

Compilez avec --default-stream per-thread (ou définissez la macro CUDA_API_PER_THREAD_DEFAULT_STREAM) dès qu'un pipeline multi-stream est en jeu. Ce mode donne à chaque thread CPU son propre flux par défaut, qui ne se synchronise plus implicitement avec les streams explicites des autres threads. Sans ce flag, un seul appel mal placé sur stream 0 peut sérialiser silencieusement l'ensemble du pipeline en production.

Créer et organiser des streams

La création d'un stream non bloquant est directe :

cudaStream_t stream;
cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);

Le flag cudaStreamNonBlocking garantit que ce stream n'attendra pas le flux par défaut legacy, même si le code applicatif l'utilise ailleurs. Pour un pipeline d'inference, on crée généralement 2 à 4 streams dédiés à des étapes différentes (transfert, calcul, post-traitement) ou à des buffers différents dans un schéma de double/triple buffering.

Toutes les opérations qui composent une étape doivent être soumises au même stream pour préserver l'ordre relatif :

cudaMemcpyAsync(d_in, h_in_pinned, bytes, cudaMemcpyHostToDevice, stream);
inference_kernel<<<grid, block, 0, stream>>>(d_in, d_out);
cudaMemcpyAsync(h_out_pinned, d_out, bytes, cudaMemcpyDeviceToHost, stream);

L'appel host reprend la main immédiatement après chaque cudaMemcpyAsync et chaque lancement de kernel : ce sont des appels non bloquants du point de vue du CPU. La synchronisation, quand elle est nécessaire, se fait explicitement.

Les événements CUDA : synchroniser sans bloquer tout le monde

Un cudaEvent_t est un marqueur inséré dans un stream. Il sert à deux usages distincts qu'il ne faut pas confondre.

Mesure de performance : on enregistre un événement avant et après une séquence, puis on calcule l'écart avec cudaEventElapsedTime. C'est la méthode correcte pour chronométrer du code GPU — un std::chrono côté host ne mesure que le temps de soumission, pas le temps d'exécution réel, sauf si on force une synchronisation qui casse justement le recouvrement qu'on veut mesurer.

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, stream);
inference_kernel<<<grid, block, 0, stream>>>(d_in, d_out);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);

Synchronisation inter-stream : cudaStreamWaitEvent fait attendre un stream jusqu'à ce qu'un événement enregistré dans un autre stream soit atteint, sans bloquer le CPU et sans forcer les deux streams à se synchroniser entièrement. C'est l'outil correct quand une étape du pipeline B dépend réellement d'une étape du pipeline A, sans vouloir sacrifier l'indépendance du reste.

Pour valider qu'un recouvrement existe réellement — et pas seulement dans le code, mais dans l'exécution matérielle observée — utilisez Nsight Systems (nsys profile). La timeline affiche séparément le moteur de copie et le moteur de calcul : si les barres H2D et kernel apparaissent sur des lignes distinctes qui se chevauchent dans le temps, le recouvrement est réel. Ne faites jamais confiance uniquement au temps total mesuré côté host, qui peut masquer une sérialisation partielle.

Pipeliner l'inference batch en chunks

L'idée centrale du pipelining pour l'inference : au lieu de traiter tout le batch en un bloc (copie complète, calcul complet, copie retour complète), on découpe le batch en N chunks et on les fait progresser en décalage sur des streams distincts. Pendant que le chunk 2 est copié vers le device, le chunk 1 est calculé ; pendant que le chunk 1 est renvoyé vers le host, le chunk 2 est calculé et le chunk 3 est copié.

Pipeline à double buffering sur deux streams pour un batch découpé en n_chunks :

for (int i = 0; i < n_chunks; ++i) {
    cudaStream_t s = streams[i % 2];
    cudaMemcpyAsync(d_buf[i % 2], h_pinned[i], chunk_bytes,
                     cudaMemcpyHostToDevice, s);
    inference_kernel<<<grid, block, 0, s>>>(d_buf[i % 2], d_out[i % 2]);
    cudaMemcpyAsync(h_out_pinned[i], d_out[i % 2], out_bytes,
                     cudaMemcpyDeviceToHost, s);
}
for (auto& s : streams) cudaStreamSynchronize(s);

Avec deux buffers alternés, le chunk i peut être calculé pendant que le chunk i+1 est encore en cours de copie H2D, à condition que les buffers device ne soient pas partagés entre chunks consécutifs — sinon on introduit une dépendance fausse qui sérialise tout.

Le nombre optimal de chunks n'est pas arbitraire. Trop peu de chunks (2-3) laisse les étapes de bord (premier H2D, dernier D2H) peser lourd dans le total. Trop de chunks augmente l'overhead de lancement de kernel et de gestion de streams, qui n'est pas nul — chaque lancement coûte de l'ordre de quelques microsecondes de overhead driver. En pratique, on vise des chunks assez gros pour que le temps de calcul par chunk dépasse largement cet overhead, et on mesure avec Nsight Systems plutôt que de deviner.

Le nombre de buffers physiques (2 pour du double buffering, 3 pour du triple buffering) détermine combien d'étapes peuvent être en vol simultanément. Le triple buffering a du sens quand les trois étapes (H2D, calcul, D2H) ont des durées comparables ; avec un calcul très dominant, le double buffering suffit généralement à saturer le moteur de calcul.

Priorités de stream et cas multi-modèles

Tous les streams ne se valent pas forcément dans un service d'inference. Quand plusieurs modèles ou plusieurs clients partagent le même GPU, il est parfois nécessaire de garantir qu'une requête latence-critique passe devant un batch de traitement en arrière-plan. cudaStreamCreateWithPriority permet d'assigner une priorité relative à un stream :

int low, high;
cudaDeviceGetStreamPriorityRange(&low, &high);
cudaStream_t stream_prioritaire;
cudaStreamCreateWithPriority(&stream_prioritaire, cudaStreamNonBlocking, high);

Cette priorité influence l'ordonnancement des blocs sur le multiprocesseur streaming quand plusieurs kernels sont éligibles simultanément, mais elle ne préempte pas un kernel déjà en cours d'exécution : un kernel long à basse priorité qui a déjà démarré retardera quand même un kernel haute priorité soumis juste après. Pour de la latence garantie, la taille des kernels compte autant que la priorité elle-même — préférez des kernels courts et nombreux à un unique kernel monolithique si le partage de GPU entre charges hétérogènes est une contrainte réelle.

Dans un service d'inference qui sert plusieurs modèles, une pratique courante consiste à dédier un pool de streams par modèle plutôt que par requête individuelle, avec une file d'attente applicative en amont. Créer un stream par requête HTTP entrante sous forte charge génère un nombre de streams non borné, ce qui ajoute de la pression sur le driver sans bénéfice supplémentaire au-delà du nombre de moteurs matériels disponibles.

Streams et CUDA Graphs : quand aller plus loin

Lorsque le pipeline devient stable — même séquence d'opérations répétée à chaque chunk, seules les données changent — le coût de soumission CPU de chaque kernel et de chaque copie peut lui-même devenir un facteur limitant, en particulier pour des chunks petits et nombreux. Les CUDA Graphs permettent de capturer une séquence d'opérations sur un stream une seule fois (cudaStreamBeginCapture / cudaStreamEndCapture), puis de rejouer ce graphe à chaque itération avec un seul appel (cudaGraphLaunch), réduisant fortement l'overhead de soumission répétée. Ce mécanisme dépasse le cadre de ce chapitre, mais retenez qu'il s'appuie directement sur les streams et événements présentés ici : maîtriser leur usage manuel est un prérequis pour tirer parti des graphes efficacement, car un graphe capturé reproduit fidèlement — y compris ses défauts — le comportement de synchronisation du stream sur lequel il a été enregistré.

Tableau récapitulatif des primitives

Checklist avant mise en production

• Les buffers host utilisés dans des copies asynchrones sont-ils bien alloués en mémoire pinned ?
• Le binaire est-il compilé avec --default-stream per-thread, ou tous les appels sont-ils rigoureusement soumis à des streams explicites ?
• Le nombre de chunks a-t-il été choisi par mesure (Nsight Systems) plutôt que par intuition ?
• Les dépendances entre étapes sont-elles exprimées avec cudaStreamWaitEvent plutôt qu'avec des synchronisations globales ?
• Le recouvrement observé sur la timeline correspond-il au recouvrement attendu sur le papier — moteur copie et moteur calcul actifs simultanément ?
• Le nombre de streams créés reste-t-il raisonnable (généralement moins d'une dizaine) ? Chaque stream a un coût de gestion côté driver, et au-delà d'un certain nombre le GPU sature de toute façon ses moteurs matériels, qui sont en nombre fixe.

Limites structurelles à connaître

Le recouvrement ne fait pas disparaître le travail, il le chevauche. Si le calcul est déjà le facteur limitant (GPU-bound), ajouter des streams ne réduira pas le temps total au-delà du temps de calcul lui-même — le gain maximal théorique est borné par max(temps copie, temps calcul), pas par zéro. De même, le nombre de moteurs de copie matériels est limité (souvent un ou deux moteurs de copie selon l'architecture) : au-delà de ce nombre de copies simultanées, les transferts additionnels se mettent en file plutôt que de s'exécuter réellement en parallèle. Le pipelining déplace le goulot d'étranglement, il ne l'élimine pas — et c'est précisément pour cela qu'il faut mesurer, pas supposer.