CUDA-Speicherzugriffe auf DGX Spark gezielt prüfen

Der CUDA-Kernel wurde bereits gestartet, doch der Fehler tritt erst beim Zurückkopieren des Ergebnisses auf. Mit kleineren Eingaben läuft dasselbe Programm dagegen durch. Bei solchen Symptomen sollten Sie zuerst die Gerätespeicherzugriffe prüfen, statt sofort den Treiber zu ändern. Auf DGX Spark lässt sich mit einem kleinen, eigenständigen Programm eine nachvollziehbare Diagnosekette aufbauen: einen Zugriff außerhalb der Speichergrenzen gezielt auslösen, die betroffene Quellzeile ermitteln, den Index korrigieren und anschließend prüfen, ob die Fehlererkennung weiterhin aktiv ist. Dafür sind weder Modellgewichte noch Eingriffe in produktive Dienste nötig.

1. Umgebung festhalten und Testgrenzen definieren

Der DGX Spark von TcpNode verwendet NVIDIA GB10 Grace Blackwell, die ARM64-Architektur und NVIDIA DGX OS und verfügt über 128GB LPDDR5x Unified Memory. CPU und GPU teilen sich diesen Speicher. Das bedeutet jedoch nicht, dass beliebige Zeiger über die Grenzen einer Speicherallokation hinaus zugreifen dürfen. Im folgenden Beispiel wird Gerätespeicher mit cudaMalloc reserviert, damit der gültige Adressbereich eindeutig feststeht.

Nach der Zahlung können Sie die Zugangsdaten in der Konsole einsehen und sich per SSH anmelden. Führen Sie den Test in einem eigenen Arbeitsverzeichnis mit einem normalen Benutzerkonto aus. Injizieren Sie keine Fehler in einen Inferenzprozess, der gerade Anfragen bearbeitet.

mkdir -p "$HOME/cuda-access-check"
cd "$HOME/cuda-access-check"
uname -m
nvidia-smi
command -v nvcc
command -v compute-sanitizer
nvcc --version
compute-sanitizer --version

Als Architektur sollte aarch64 erscheinen. Halten Sie die Werkzeugversionen und Treiberinformationen fest, damit spätere Kontrollläufe in derselben Umgebung stattfinden. Ein vorinstallierter CUDA-Stack bedeutet nicht, dass jedes Werkzeug bereits im aktuellen PATH verfügbar ist. Fehlt ein Befehl, prüfen Sie zunächst die installierte Toolchain und ihre Installationspfade, statt direkt den gesamten Treiber auszutauschen.

Führen Sie das Beispiel mit dem ungültigen Speicherzugriff ausschließlich in einem separaten Prozess aus, der anschließend verworfen werden kann. Nach einem unzulässigen Speicherzugriff müssen Sie den Prozess beenden und neu starten. Weitere Ergebnisse aus demselben CUDA-Kontext sind kein belastbarer Nachweis.

2. Einen Grenzübertritt mit nur einer veränderlichen Größe reproduzieren

Speichern Sie den folgenden Inhalt als repro.cu. Das Programm reserviert Speicher für genau eine Ganzzahl; der einzige gültige Index ist null. Über TEST_INDEX wird beim Kompilieren festgelegt, an welche Position geschrieben wird. Fehlerhafte und korrigierte Version verwenden damit denselben Quellcode, ohne gleichzeitig Eingabe, Allokationsgröße und Ausführungspfad zu verändern.

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

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

__global__ void write_value(int* data) {
    data[TEST_INDEX] = 7;
}

int main() {
    int* data = nullptr;
    int result = 0;
    checked(cudaMalloc(reinterpret_cast<void**>(&data), sizeof(int)));
    checked(cudaMemset(data, 0, sizeof(int)));
    write_value<<<1, 1>>>(data);
    checked(cudaGetLastError());
    checked(cudaDeviceSynchronize());
    checked(cudaMemcpy(&result, data, sizeof(int), cudaMemcpyDeviceToHost));
    checked(cudaFree(data));
    std::printf("result=%d\n", result);
    return result == 7 ? EXIT_SUCCESS : EXIT_FAILURE;
}

Kompilieren Sie die beiden Programme getrennt:

nvcc -arch=native -lineinfo -DTEST_INDEX=1 repro.cu -o repro-bad
nvcc -arch=native -lineinfo -DTEST_INDEX=0 repro.cu -o repro-good
sha256sum repro.cu repro-bad repro-good > build.sha256

-arch=native kompiliert für die aktuell sichtbare GPU. Das eignet sich zur lokalen Überprüfung, nicht zur Erstellung geräteübergreifend einsetzbarer Binärdateien. Unterstützt der Compiler das aktuelle Gerät nicht, prüfen Sie zuerst die Kompatibilität der Toolchain, statt auf Verdacht Parameter für eine andere Architektur zu übernehmen. -lineinfo erhält die Quellzeileninformationen des Gerätecodes. Für die erste Eingrenzung ist kein vollständiger Geräte-Debugmodus nötig, der das Ausführungsverhalten deutlich verändern würde.

Warum sowohl Kernel-Start als auch Synchronisierung geprüft werden müssen

cudaGetLastError prüft Fehler im Zusammenhang mit dem Kernel-Start. Der Kernel selbst wird jedoch in der Regel asynchron ausgeführt. cudaDeviceSynchronize wartet auf den Abschluss der zuvor eingereihten Gerätearbeit; Fehler während der Ausführung können an dieser Stelle zurückgemeldet werden. Beide Rückgabewerte müssen daher separat geprüft werden. Es reicht nicht, nur darauf zu achten, ob das Programm ein Ergebnis ausgibt.

3. Das Prüfergebnis über einen auswertbaren Exitcode erfassen

Führen Sie zuerst die fehlerhafte Version aus und speichern Sie Standardausgabe und Diagnoseprotokoll getrennt:

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-bad > bad.stdout 2> bad.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat bad.stdout bad.stderr

Zu erwarten ist ein Bericht über einen unzulässigen Schreibzugriff auf den globalen Speicher, der auf die Schreiboperation in write_value verweist. Der Bericht enthält üblicherweise auch die Zugriffsgröße, die Thread-Position und Angaben zur Lage der Adresse. Die genaue Darstellung hängt von der Werkzeugversion ab. Ein bestimmter Wortlaut sollte deshalb nicht das einzige Prüfkriterium sein.

Mit --error-exitcode 86 wirken sich erkannte Fehler auf das Ergebnis automatisierter Abläufe aus. Ist das getestete Programm selbst bereits fehlgeschlagen, kann stattdessen auch ein anderer von ihm zurückgegebener Wert ungleich null erscheinen. Das Abnahmekriterium lautet daher: „Exitcode ungleich null und Bericht zum erwarteten Zugriff außerhalb der Grenzen“ — nicht, dass jeder Fehler zwingend den Code 86 liefern muss.

Hängen Sie den Befehl nicht einfach an eine gewöhnliche Logging-Pipeline und werten anschließend nur den Status des letzten Protokollierungsbefehls aus. Die oben gezeigte Umleitung mit bedingter Verzweigung erhält das tatsächliche Ausführungsergebnis. Sie eignet sich auch für Skripte, die bei einem Fehler sofort abbrechen.

4. Anhand des Adressbereichs korrigieren, statt Symptome zu verdecken

Das Beispiel reserviert Platz für eine Ganzzahl, schreibt aber an Index eins und damit außerhalb des gültigen Bereichs. In einem realen Programm müssen Sie die im Bericht genannte Adresse der zugehörigen Speicherallokation zuordnen und anschließend Indexberechnung, Elementtyp, Schrittweite und Batch-Offset prüfen.

Typische Ursachen sind: Die Allokationsgröße wird in Bytes angegeben, der Index aber als Elementanzahl interpretiert; für den letzten Thread-Block reichen die Daten nicht für einen vollständigen Block; die Zeilenschrittweite eines zweidimensionalen Arrays weicht von seiner logischen Breite ab; oder ein anderer Task verwendet einen Zeiger weiter, nachdem der Speicher freigegeben wurde. Mehr Speicher zu reservieren kann den Fehler vorübergehend verdecken, beweist aber nicht die Korrektheit des Algorithmus.

Bei eindimensionalen Aufgaben sollte vor dem Zugriff üblicherweise index < count geprüft werden. Zusätzlich muss der für die Indexberechnung verwendete Ganzzahltyp das Produkt aufnehmen können. Bei Kernels, deren Threads über Shared Memory zusammenarbeiten, dürfen nicht beliebig einzelne Threads vorzeitig zurückkehren, während die übrigen anschließend eine blockweite Barriere erreichen. Das kann neue Synchronisierungsprobleme verursachen.

Tritt der Fehler nur in einem großen Workflow auf, können Sie vorübergehend nach verdächtigen Phasen Synchronisierungsprüfungen einfügen und den Bereich schrittweise eingrenzen. Ergänzen Sie jeweils nur einen Prüfpunkt und halten Sie den ersten Fehler fest. Reichen Sie nach dem Fehler keine weiteren Aufgaben ein, die zahlreiche Folgefehler erzeugen.

5. Die Korrektur mit Positiv- und Negativtests abnehmen

Führen Sie die korrigierte Version aus:

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-good > good.stdout 2> good.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat good.stdout good.stderr

In diesem Beispiel müssen drei Bedingungen gleichzeitig erfüllt sein: Der Exitcode ist null, der Diagnosebericht enthält keine Speicherzugriffsfehler, und die Ausgabe lautet result=7. „Keine Fehlermeldung“ allein genügt nicht. Auch ein nicht ausgeführter Kernel oder ein nicht tatsächlich geschriebenes Ergebnis kann einen Test oberflächlich erfolgreich erscheinen lassen.

Führen Sie danach die fehlerhafte Version erneut aus und prüfen Sie, ob sie weiterhin fehlschlägt. Dieser Negativtest deckt unter anderem auf, dass der Prüfbefehl entfernt wurde, der Test versehentlich eine alte Binärdatei verwendet oder die Pipeline den Fehlerstatus verschluckt. Reale Kernels sollten außerdem mit der kleinsten gültigen Eingabe, an Grenzen vollständiger Blöcke und mit einem Rest getestet werden, der keinen vollständigen Block füllt. Erlaubt die Schnittstelle leere Eingaben, muss auch deren Behandlung eindeutig festgelegt sein.

Archivieren Sie abschließend Quellcode, Kompilierbefehle, Versionsinformationen, Prüfsummen und die Protokolle beider Läufe. Das Prüfwerkzeug verändert den Ausführungsaufwand; die dabei gemessenen Laufzeiten lassen sich daher nicht unmittelbar als Leistungsbewertung verwenden. Auch ein bestandener memcheck schließt Datenrennen und Synchronisierungsprobleme nicht aus. Ein belastbarer Abschluss liefert separat überprüfbare Nachweise für gültige Speicherzugriffe, korrekte Rechenergebnisse und die Weitergabe von Fehlerzuständen — nicht nur einen Screenshot eines erfolgreichen Laufs.

Häufig gestellte Fragen