CUDA-Stream-Abhängigkeiten auf DGX Spark prüfen

Wenn zwei CUDA-Streams denselben Gerätepuffer verwenden, kann ein Programm mehrfach korrekte Ergebnisse liefern und nach einer Änderung der Eingabegröße plötzlich veraltete Werte lesen. Die Ursache liegt häufig nicht in der Rechenvorschrift: Der Verbraucher-Stream beginnt bereits zu lesen, während der Produzenten-Stream noch schreibt. Dieser Artikel entwickelt auf DGX Spark ein minimales Zwei-Stream-Beispiel und prüft Einreichungsreihenfolge, Ausführungsabhängigkeiten und den Zeitpunkt der Pufferwiederverwendung getrennt. Ein einzelner erfolgreicher Durchlauf reicht dabei nicht als Nachweis.

1. Erst die Datenabhängigkeiten festlegen, dann die Streams anlegen

DGX Spark verwendet NVIDIA GB10 Grace Blackwell und bietet 128GB LPDDR5x Unified Memory. CPU und GPU teilen sich denselben Speicherpool. Physisch gemeinsamer Speicher stellt jedoch nicht automatisch die streamübergreifende Ausführungsreihenfolge für eine Anwendung her. Explizite Kopiervorgänge, Kernel-Starts und Host-Zugriffe müssen weiterhin ihre jeweiligen Abschlussbedingungen einhalten.

Das Experiment verwendet explizit zugewiesenen Gerätespeicher und seitenfixierten Host-Speicher, damit die Grenzen der einzelnen Schritte klar erkennbar sind. Der Produzenten-Stream kopiert die Eingabe auf das Gerät und multipliziert jedes Element mit zwei. Der Verbraucher-Stream addiert zu jedem Element eins und kopiert das Ergebnis zurück auf den Host. Die vollständige Abhängigkeitskette lautet:

Eingabe bereit → Kopieren abgeschlossen → Multiplikation mit zwei abgeschlossen → Addition von eins beginnt → Rückübertragung abgeschlossen → Prüfung auf dem Host.

Innerhalb eines Streams werden Operationen in Stream-Reihenfolge ausgeführt; zwischen zwei Streams muss die Abhängigkeit dagegen explizit hergestellt werden. Das Event muss im Produzenten-Stream nach dem Multiplikationskernel aufgezeichnet werden. Im Verbraucher-Stream muss vor dem Einreichen des Additionskernels das Warten auf dieses Event eingereiht werden.

Dass der Host zuerst Aufrufe für den Produzenten-Stream und danach für den Verbraucher-Stream ausführt, legt lediglich die Einreichungsreihenfolge fest. Streamübergreifende Datenabhängigkeiten brauchen eine explizite Synchronisationsgrundlage. Eine Ausführungsreihenfolge, die bei geringer Last zufällig passt, ist keine solche Grundlage.

2. Ein Minimalprogramm ohne Abhängigkeit vom Default-Stream erstellen

Nach der Zahlung bei TcpNode rufen Sie die Zugangsdaten in der Konsole ab und melden sich per SSH bei NVIDIA DGX OS an. Führen Sie zunächst uname -m, nvidia-smi und nvcc --version aus. Speichern Sie damit getrennt die Host-Architektur, die Geräte- und Treiberinformationen sowie die Angaben zur Compiler-Toolchain. Als Architektur sollte aarch64 für ARM64 erscheinen. Die Compilerversion darf nicht aus der Treiberausgabe abgeleitet werden.

Speichern Sie den folgenden Inhalt als stream_contract.cu. Beide Streams werden mit dem Non-Blocking-Flag angelegt, damit die implizite Synchronisation des herkömmlichen Default-Streams keine fehlenden Abhängigkeiten verdeckt. Das Event dient ausschließlich der Reihenfolge, nicht der Zeitmessung.

#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. Ausgabe und Exit-Code gemeinsam prüfen

Kompilieren und starten Sie das Programm im selben Verzeichnis:

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

Der Test ist bestanden, wenn die Ausgabe elements=4096 mismatches=0 lautet und der Exit-Code null ist. Bei einem fehlgeschlagenen API-Aufruf oder einer fehlgeschlagenen Synchronisation wird eins zurückgegeben, bei abweichenden Zahlenwerten zwei. Wer nur prüft, ob der Prozess abstürzt, übersieht abgeschlossene Berechnungen mit falschen Ergebnissen.

Alle Beispielwerte liegen im Bereich der Ganzzahlen, die sich mit einfacher Gleitkommagenauigkeit exakt darstellen lassen. Deshalb wird hier direkt auf Gleichheit geprüft. Für reale Gleitkommamodelle darf dieses Kriterium nicht unverändert übernommen werden. Absolute und relative Fehlertoleranzen sowie der Umgang mit nicht endlichen Werten müssen sich nach dem jeweiligen Algorithmus richten.

Ändern Sie anschließend die Elementanzahl nacheinander auf 1, 257 und 4097 und kompilieren und starten Sie das Programm jeweils erneut. Diese Größen decken ein einzelnes Element, den Rand eines Threadblocks und Grenzen über mehrere Blöcke hinweg ab. Es handelt sich um Testeingaben, nicht um Leistungsdaten. Speichern Sie für jeden Durchlauf die Quelldatei, den Compileraufruf, die Toolchain-Informationen und die Programmausgabe, statt nur einen Screenshot eines erfolgreichen Laufs aufzubewahren.

Prüfen, ob der Fehlerpfad tatsächlich funktioniert

Ersetzen Sie vorübergehend das abschließende 1.0f im Ausdruck für den Erwartungswert durch 2.0f. Das Programm muss Abweichungen melden und zwei zurückgeben. Stellen Sie danach den ursprünglichen Ausdruck wieder her. So prüfen Sie, ob die Ergebniskontrolle tatsächlich greift. Verwenden Sie dagegen nicht die Forderung „Nach dem Entfernen des Wartens muss ein Fehler auftreten“ als Negativtest: Ob eine Race Condition sichtbar wird, hängt vom Scheduling ab. Auch ein Programm mit fehlender Abhängigkeit kann vorübergehend bestehen.

4. Event-Grenzen auf die Speicherlebensdauer übertragen

Wenn cudaMemcpyAsync zurückkehrt, bedeutet das in der Regel nur, dass der Aufruf eingereicht wurde. Daraus folgt nicht, dass die Eingabe bereits überschrieben werden darf. Im Beispiel bleibt sie vom Einreichen des Kopiervorgangs bis zum Abschluss der Synchronisation des Verbraucher-Streams unverändert. Auch die Ausgabe wird erst gelesen, nachdem der Stream mit der Rückübertragung synchronisiert wurde. Diese Strategie ist eher konservativ, setzt aber klare Grenzen und eignet sich als Ausgangsbasis für die Korrektheitsprüfung.

Seitenfixierter Speicher hilft, die Voraussetzungen für asynchrone Übertragungen zu erfüllen, ist aber keine Abschlussbenachrichtigung. Ebenso darf der Gerätepuffer nicht vorzeitig freigegeben werden, solange andere Streams noch darauf zugreifen. Auch 128GB Unified Memory sind kein Grund, große Mengen des Host-Speichers vollständig zu fixieren. Die Batch-Größe sollte sich am tatsächlichen Arbeitssatz orientieren.

Das Produzenten-Event erfasst nur die davor eingereihten Operationen, nicht die Addition und Rückübertragung im Verbraucher-Stream. Nach Abschluss dieses Events steht daher fest, dass die Produktionsphase beendet ist. Es steht aber noch nicht fest, dass der Gerätepuffer für den nächsten Batch überschrieben werden darf. Beim Ausbau zur Pipeline muss nach der Rückübertragung im Verbraucher-Stream ein weiteres Abschluss-Event aufgezeichnet werden. Dieses steuert die Freigabe der gesamten Puffergruppe zur Wiederverwendung.

Jeder Slot für einen parallel verarbeiteten Batch benötigt eigene Eingabe-, Ausgabe- und Gerätepuffer sowie einen eigenen Abschlussstatus. Prüfen Sie zunächst einen einzelnen Slot und erhöhen Sie erst danach die Anzahl. Nur mehr Streams anzulegen, während weiterhin ungeschützt derselbe Zeiger gemeinsam verwendet wird, reicht nicht aus.

5. Bei der Fehlersuche den Synchronisationsumfang schrittweise eingrenzen

Treten sporadisch falsche Ergebnisse auf, prüfen Sie zuerst drei Punkte: Wird das Warten vor dem Verbraucher-Kernel eingereiht? Wird das Event nach dem Produzenten-Kernel aufgezeichnet? Beziehen sich Aufzeichnung und Warten auf dasselbe Event? Bereits eine um eine Zeile vorgezogene Event-Aufzeichnung kann aus „Daten vollständig geschrieben“ fälschlich nur noch „Kopiervorgang abgeschlossen“ machen.

Prüfen Sie danach die Fehler nach jedem Kernel-Start und den Rückgabewert der abschließenden Synchronisation. Die erste Prüfung erkennt bestimmte Startprobleme, die zweite erfasst Fehler während der asynchronen Ausführung. Keine der beiden ersetzt die andere. Wer nur den Status nach dem Start prüft, aber nicht auf den Abschluss wartet, sieht den Fehler möglicherweise erst bei einem späteren API-Aufruf.

In einer Diagnosekopie des Programms können Sie vorübergehend eine Stream-Synchronisation in die Produzentenphase einfügen. So lässt sich eingrenzen, ob das Problem mit der streamübergreifenden Reihenfolge zusammenhängt. Verschwindet der Fehler, prüfen Sie weiter die Abhängigkeiten und die Zuständigkeit für die Puffer, statt die Lösung mit einer geräteweiten Synchronisation abzuschließen. Globale Synchronisation verändert das Scheduling und kann bestehende Probleme mit der Speicherlebensdauer verdecken.

In der endgültigen Fassung sollten ein explizites Produzenten-Event, das Warten im Verbraucher-Stream, ein Event zur Freigabe der Puffer und eine wiederholbare Ausgabeprüfung verbleiben. Erst wenn sich sowohl die Abhängigkeitsbeziehungen als auch die Zuständigkeiten für den Speicher eindeutig erklären lassen, hat ein Multi-Stream-Programm eine belastbare Grundlage für mehr Parallelität.

Häufig gestellte Fragen