Vérifier les dépendances entre streams CUDA sur DGX Spark

Lorsque deux streams CUDA partagent le même tampon sur le périphérique, un programme peut produire des résultats corrects plusieurs fois de suite, puis renvoyer des valeurs périmées après un changement de taille des données d’entrée. La cause n’est souvent pas la formule de calcul : le stream consommateur commence à lire avant que le stream producteur ait terminé ses écritures. Cet article présente un test minimal à deux streams sur DGX Spark pour vérifier séparément l’ordre de soumission, les dépendances d’exécution et le moment où les tampons peuvent être réutilisés. Une seule exécution réussie ne suffit pas à conclure.

1. Définir les dépendances de données avant de créer les streams

DGX Spark repose sur NVIDIA GB10 Grace Blackwell et dispose de 128 Go de mémoire unifiée LPDDR5x, partagée entre le CPU et le GPU. Cette unification physique de la mémoire n’établit toutefois pas automatiquement l’ordre d’exécution entre les streams de l’application. Les copies explicites, les lancements de noyaux et les accès depuis l’hôte restent soumis à leurs conditions de terminaison respectives.

Le test utilise une allocation explicite sur le périphérique et de la mémoire hôte verrouillée en pages pour rendre les limites plus faciles à observer. Le stream producteur copie les données d’entrée vers le périphérique, puis multiplie chaque élément par deux. Le stream consommateur ajoute un à chaque élément, puis recopie le résultat vers l’hôte. La chaîne complète de dépendances est la suivante :

Entrée prête → copie terminée → multiplication par deux terminée → début de l’addition de un → copie de retour terminée → vérification par l’hôte.

Les opérations d’un même stream s’exécutent dans l’ordre du stream ; entre deux streams, la dépendance doit être établie explicitement. L’événement doit être enregistré dans le stream producteur après le noyau de multiplication. Le stream consommateur doit attendre cet événement avant de lancer le noyau d’addition.

Le fait que l’hôte soumette d’abord les opérations du stream producteur, puis celles du stream consommateur, ne définit que l’ordre de soumission. Les dépendances de données entre streams exigent un mécanisme de synchronisation explicite. Elles ne doivent pas reposer sur un ordre d’exécution qui ne se produit par hasard que lorsque la charge est faible.

2. Construire un programme minimal indépendant du stream par défaut

Après le paiement sur TcpNode, consultez les identifiants de connexion dans la console, puis connectez-vous à NVIDIA DGX OS via SSH. Exécutez d’abord uname -m, nvidia-smi et nvcc --version, puis conservez respectivement les informations sur l’architecture de l’hôte, le périphérique et le pilote, ainsi que la chaîne de compilation. L’architecture attendue est aarch64, qui correspond à ARM64. La version du compilateur ne doit pas être déduite de la sortie du pilote.

Enregistrez le contenu suivant dans stream_contract.cu. Les deux streams sont créés avec l’indicateur non bloquant afin que la synchronisation implicite du stream par défaut historique ne masque pas une dépendance manquante. L’événement sert uniquement à ordonner les opérations, sans mesurer leur durée.

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>

static void check(cudaError_t result) {
    if (result != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(result));
        std::exit(1);
    }
}

__global__ void scale(float* data, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count) data[i] *= 2.0f;
}

__global__ void add(float* data, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count) data[i] += 1.0f;
}

int main() {
    constexpr int count = 4096;
    const size_t bytes = count * sizeof(float);
    float *input = nullptr, *output = nullptr, *device = nullptr;
    cudaStream_t producer, consumer;
    cudaEvent_t ready;

    check(cudaMallocHost(reinterpret_cast<void**>(&input), bytes));
    check(cudaMallocHost(reinterpret_cast<void**>(&output), bytes));
    check(cudaMalloc(reinterpret_cast<void**>(&device), bytes));
    check(cudaStreamCreateWithFlags(&producer, cudaStreamNonBlocking));
    check(cudaStreamCreateWithFlags(&consumer, cudaStreamNonBlocking));
    check(cudaEventCreateWithFlags(&ready, cudaEventDisableTiming));

    for (int i = 0; i < count; ++i) {
        input[i] = static_cast<float>(i);
        output[i] = -1.0f;
    }

    check(cudaMemcpyAsync(device, input, bytes,
                          cudaMemcpyHostToDevice, producer));
    scale<<<(count + 255) / 256, 256, 0, producer>>>(device, count);
    check(cudaGetLastError());
    check(cudaEventRecord(ready, producer));

    check(cudaStreamWaitEvent(consumer, ready, 0));
    add<<<(count + 255) / 256, 256, 0, consumer>>>(device, count);
    check(cudaGetLastError());
    check(cudaMemcpyAsync(output, device, bytes,
                          cudaMemcpyDeviceToHost, consumer));
    check(cudaStreamSynchronize(consumer));

    int failures = 0;
    for (int i = 0; i < count; ++i) {
        float expected = 2.0f * static_cast<float>(i) + 1.0f;
        if (output[i] != expected) ++failures;
    }
    std::printf("elements=%d mismatches=%d\n", count, failures);

    check(cudaEventDestroy(ready));
    check(cudaStreamDestroy(consumer));
    check(cudaStreamDestroy(producer));
    check(cudaFree(device));
    check(cudaFreeHost(output));
    check(cudaFreeHost(input));
    return failures == 0 ? 0 : 2;
}

3. Valider à la fois la sortie et le code de retour

Compilez et exécutez le programme dans le même répertoire :

nvcc -O2 stream_contract.cu -o stream_contract
./stream_contract
printf 'exit=%s\n' "$?"

Le résultat attendu est elements=4096 mismatches=0, avec un code de retour nul. Une erreur d’appel à l’API ou de synchronisation renvoie le code 1 ; une différence numérique renvoie le code 2. Vérifier uniquement que le processus ne plante pas laisse passer des calculs terminés dont le résultat est faux.

Toutes les valeurs de cet exemple sont des entiers exactement représentables en simple précision, ce qui autorise une comparaison d’égalité directe. Ce critère ne doit pas être repris tel quel pour un modèle de calcul en virgule flottante réel : il faut définir, selon l’algorithme, les tolérances d’erreur absolue et relative ainsi que les règles de traitement des valeurs non finies.

Remplacez ensuite le nombre d’éléments successivement par 1, 257 et 4097, puis recompilez et relancez le programme à chaque fois. Ces tailles couvrent le cas d’un seul élément, la fin d’un bloc de threads et les limites entre plusieurs blocs. Ce sont des entrées de test, pas des mesures de performances. Pour chaque essai, conservez le fichier source, la commande de compilation, les informations sur la chaîne de compilation et la sortie d’exécution, plutôt qu’une simple capture d’écran montrant un succès.

Vérifier que la détection des erreurs fonctionne réellement

Remplacez temporairement le 1.0f à la fin de l’expression du résultat attendu par 2.0f. Le programme doit signaler des différences et renvoyer le code 2. Rétablissez ensuite l’expression d’origine. Ce test permet de vérifier que le contrôle des résultats fonctionne. En revanche, ne considérez pas que « supprimer l’attente doit nécessairement provoquer une erreur » constitue un test négatif valable : l’apparition d’une condition de concurrence dépend de l’ordonnancement, et un programme auquel il manque une dépendance peut tout de même réussir provisoirement.

4. Relier les événements à la durée de vie des tampons

Le retour de cudaMemcpyAsync indique généralement seulement que l’opération a été soumise ; il n’autorise pas à modifier les données d’entrée. Dans cet exemple, l’entrée reste inchangée depuis la soumission de la copie jusqu’à la fin de la synchronisation du stream consommateur. De même, la sortie n’est lue qu’après la synchronisation du stream chargé de la copie de retour. Cette stratégie est prudente, mais ses limites sont claires : elle constitue une bonne référence initiale pour établir la correction du programme.

La mémoire verrouillée en pages contribue à remplir les conditions nécessaires aux transferts asynchrones, mais ne constitue pas une notification de fin. Le tampon sur le périphérique ne doit pas non plus être libéré tant que d’autres streams y accèdent. Disposer de 128 Go de mémoire unifiée ne signifie pas davantage qu’il faut verrouiller en pages de grandes quantités de mémoire hôte. La taille des lots doit être limitée en fonction du jeu de données réellement utilisé.

L’événement du producteur ne couvre que les opérations qui le précèdent, et non l’addition ni la copie de retour effectuées par le stream consommateur. Lorsque cet événement est terminé, la phase de production est donc achevée, mais rien ne garantit encore que le tampon sur le périphérique puisse être écrasé par le lot suivant. Pour passer à une exécution en pipeline, enregistrez un événement de fin supplémentaire après la copie de retour dans le stream consommateur et utilisez-le pour contrôler la remise à disposition de l’ensemble des tampons.

Chaque emplacement réservé à un lot traité en parallèle doit disposer de ses propres tampons d’entrée, de sortie et de périphérique, ainsi que de son propre état de terminaison. Validez d’abord le fonctionnement avec un seul emplacement avant d’en ajouter d’autres. N’augmentez pas simplement le nombre de streams tout en continuant à partager le même pointeur sans protection.

5. Réduire progressivement la portée de la synchronisation pendant le diagnostic

Si les résultats sont parfois incorrects, vérifiez d’abord que l’attente précède bien le noyau consommateur, que l’événement est enregistré après le noyau producteur et que l’enregistrement et l’attente portent sur le même événement. Déplacer l’enregistrement de l’événement d’une ligne vers le haut peut suffire à remplacer, à tort, la garantie « les données ont été entièrement écrites » par la seule garantie « la copie est terminée ».

Contrôlez ensuite les erreurs après chaque lancement de noyau ainsi que la valeur de retour de la synchronisation finale. Le premier contrôle détecte certains problèmes de lancement ; le second récupère les erreurs survenues pendant l’exécution asynchrone. Aucun ne remplace l’autre. Si vous vérifiez uniquement le lancement sans attendre la fin de l’exécution, une erreur peut ne se manifester que lors d’un appel ultérieur à l’API.

Dans une copie du programme destinée au diagnostic, vous pouvez ajouter temporairement une synchronisation du stream pendant la phase de production pour déterminer si le problème vient de l’ordre d’exécution entre streams. Si l’erreur disparaît, poursuivez l’examen des dépendances et de la propriété des tampons au lieu de vous contenter d’une synchronisation de tout le périphérique. Une synchronisation globale modifie l’ordonnancement et peut masquer le problème initial de durée de vie des tampons.

La version finale doit conserver un événement explicite de fin de production, une attente côté consommateur, un événement autorisant la réutilisation des tampons, ainsi qu’une vérification reproductible des résultats. Ce n’est que lorsque chaque dépendance et la propriété de chaque zone mémoire sont clairement établies qu’un programme multistream dispose d’une base solide pour accroître encore la concurrence.

Questions fréquentes