Was ist die kanonische Methode, um mithilfe der CUDA-Laufzeit-API nach Fehlern zu suchen?

258

Wenn ich mir die Antworten und Kommentare zu CUDA-Fragen und im CUDA-Tag-Wiki ansehe , wird häufig empfohlen, den Rückgabestatus jedes API-Aufrufs auf Fehler zu überprüfen. Die API-Dokumentation enthält Funktionen wie cudaGetLastError, cudaPeekAtLastErrorund cudaGetErrorString, aber wie lassen sich diese am besten zusammenstellen, um Fehler zuverlässig abzufangen und zu melden, ohne dass viel zusätzlicher Code erforderlich ist?

Talonmien
quelle
13
Die CUDA- Beispiele von NVIDIA enthalten den Header helper_cuda.h, in dem Makros aufgerufen werden getLastCudaErrorund checkCudaErrorsder so ziemlich das tut, was in der akzeptierten Antwort beschrieben ist . Demonstrationen finden Sie in den Beispielen. Installieren Sie einfach die Beispiele zusammen mit dem Toolkit und Sie werden es haben.
Chappjc
@chappjc Ich denke nicht, dass diese Frage und Antwort vorgibt, originell zu sein, wenn dies das ist, was Sie meinen, aber es hat den Vorteil, dass die Leute über die CUDA-Fehlerprüfung informiert wurden.
JackOLantern
@ JackOLantern Nein, das habe ich nicht angedeutet. Diese Fragen und Antworten waren für mich sehr hilfreich und sicherlich einfacher zu finden als einige Header im SDK. Ich fand es wertvoll, darauf hinzuweisen, dass NVIDIA auch damit umgeht und wo man nach mehr suchen kann. Ich würde den Ton meines Kommentars mildern, wenn ich könnte. :)
chappjc
Debugging-Tools, mit denen Sie sich dem "Beginn" der Fehler nähern können, haben sich seit 2012 bei CUDA erheblich verbessert. Ich habe nicht mit GUI-basierten Debuggern gearbeitet, aber das CUDA-Tag-Wiki erwähnt die Befehlszeile cuda-gdb. Dies ist ein SEHR leistungsfähiges Tool, da es Ihnen ermöglicht, die tatsächlichen Warps und Threads auf der GPU selbst zu durchlaufen (erfordert jedoch meistens eine Architektur von 2.0+)
opetrenko
@bluefeet: Was war mit der Bearbeitung, die Sie zurückgesetzt haben? Es sah so aus, als hätte sich am Markdown nichts geändert, aber es wurde als Bearbeitung akzeptiert. Gab es etwas Schändliches bei der Arbeit?
Talonmies

Antworten:

304

Der wahrscheinlich beste Weg, um nach Fehlern im Laufzeit-API-Code zu suchen, besteht darin, eine Assert-Style-Handler-Funktion und ein Wrapper-Makro wie folgt zu definieren:

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

Anschließend können Sie jeden API-Aufruf mit dem gpuErrchkMakro umschließen, das den Rückgabestatus des umschlossenen API-Aufrufs verarbeitet. Beispiel:

gpuErrchk( cudaMalloc(&a_d, size*sizeof(int)) );

Wenn bei einem Aufruf ein Fehler auftritt, wird eine Textnachricht ausgegeben, die den Fehler sowie die Datei und Zeile in Ihrem Code beschreibt, in der der Fehler aufgetreten ist, stderrund die Anwendung wird beendet. Sie können möglicherweise Änderungen vornehmen gpuAssert, um eine Ausnahme auszulösen , anstatt exit()eine komplexere Anwendung aufzurufen, wenn dies erforderlich ist.

Eine zweite verwandte Frage ist, wie beim Starten von Kerneln nach Fehlern gesucht werden kann, die nicht direkt in einen Makroaufruf wie Standard-Laufzeit-API-Aufrufe eingeschlossen werden können. Für Kernel ungefähr so:

kernel<<<1,1>>>(a);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );

prüft zunächst, ob ein ungültiges Startargument vorliegt, und zwingt den Host dann, zu warten, bis der Kernel stoppt, und nach einem Ausführungsfehler zu suchen. Die Synchronisation kann beseitigt werden, wenn Sie einen nachfolgenden blockierenden API-Aufruf wie folgt haben:

kernel<<<1,1>>>(a_d);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaMemcpy(a_h, a_d, size * sizeof(int), cudaMemcpyDeviceToHost) );

In diesem Fall kann der cudaMemcpyAufruf entweder Fehler zurückgeben, die während der Kernelausführung aufgetreten sind, oder solche aus der Speicherkopie selbst. Dies kann für Anfänger verwirrend sein, und ich würde empfehlen, nach dem Start des Kernels während des Debuggens eine explizite Synchronisierung zu verwenden, um leichter zu verstehen, wo Probleme auftreten können.

Beachten Sie, dass bei Verwendung von CUDA Dynamic Parallelism eine sehr ähnliche Methode auf jede Verwendung der CUDA-Laufzeit-API in Gerätekerneln sowie nach dem Start eines Gerätekerns angewendet werden kann und sollte:

#include <assert.h>
#define cdpErrchk(ans) { cdpAssert((ans), __FILE__, __LINE__); }
__device__ void cdpAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      printf("GPU kernel assert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) assert(0);
   }
}
Talonmien
quelle
8
@ Harrism: Das glaube ich nicht. Das Community-Wiki ist für Fragen oder Antworten gedacht, die häufig bearbeitet werden. Dies ist nicht einer von denen
talonmies
1
sollten wir nicht hinzufügen, cudaDeviceReset()bevor wir auch beenden? Und eine Klausel zur Speicherfreigabe?
Aurelius
2
@talonmies: Muss für Async-CUDA-Laufzeitaufrufe wie cudaMemsetAsync und cudaMemcpyAsync auch das GPU-Gerät und der Host-Thread über den Aufruf von gpuErrchk (cudaDeviceSynchronize ()) synchronisiert werden?
Nurabha
2
Beachten Sie, dass die explizite Synchronisierung nach dem Start des Kernels nicht falsch ist, aber die Ausführungsleistung und die Verschachtelungssemantik erheblich beeinträchtigen kann. Wenn Sie Interleaving verwenden, kann eine explizite Synchronisierung zum Debuggen eine ganze Klasse von Fehlern verbergen, die im Release-Build möglicherweise schwer zu finden sind.
Masterxilo
Gibt es eine Möglichkeit, spezifischere Fehler für Kernelausführungen zu erhalten? Alle Fehler, die ich bekomme, geben mir nur die Zeilennummer aus dem Host-Code, nicht aus dem Kernel.
Azmisov
70

Die obige Antwort von talonmies ist eine gute Möglichkeit, eine Anwendung im assertStil abzubrechen .

Gelegentlich möchten wir möglicherweise einen Fehlerzustand in einem C ++ - Kontext als Teil einer größeren Anwendung melden und beheben.

Hier ist eine ziemlich knappe Möglichkeit, dies zu tun, indem Sie eine C ++ - Ausnahme auslösen, die von der std::runtime_errorVerwendung abgeleitet ist thrust::system_error:

#include <thrust/system_error.h>
#include <thrust/system/cuda/error.h>
#include <sstream>

void throw_on_cuda_error(cudaError_t code, const char *file, int line)
{
  if(code != cudaSuccess)
  {
    std::stringstream ss;
    ss << file << "(" << line << ")";
    std::string file_and_line;
    ss >> file_and_line;
    throw thrust::system_error(code, thrust::cuda_category(), file_and_line);
  }
}

Dies beinhaltet den Dateinamen, die Zeilennummer und eine Beschreibung der englischen Sprache cudaError_tin das .what()Mitglied der ausgelösten Ausnahme :

#include <iostream>

int main()
{
  try
  {
    // do something crazy
    throw_on_cuda_error(cudaSetDevice(-1), __FILE__, __LINE__);
  }
  catch(thrust::system_error &e)
  {
    std::cerr << "CUDA error after cudaSetDevice: " << e.what() << std::endl;

    // oops, recover
    cudaSetDevice(0);
  }

  return 0;
}

Die Ausgabe:

$ nvcc exception.cu -run
CUDA error after cudaSetDevice: exception.cu(23): invalid device ordinal

Ein Client von some_functionkann CUDA-Fehler auf Wunsch von anderen Arten von Fehlern unterscheiden:

try
{
  // call some_function which may throw something
  some_function();
}
catch(thrust::system_error &e)
{
  std::cerr << "CUDA error during some_function: " << e.what() << std::endl;
}
catch(std::bad_alloc &e)
{
  std::cerr << "Bad memory allocation during some_function: " << e.what() << std::endl;
}
catch(std::runtime_error &e)
{
  std::cerr << "Runtime error during some_function: " << e.what() << std::endl;
}
catch(...)
{
  std::cerr << "Some other kind of error during some_function" << std::endl;

  // no idea what to do, so just rethrow the exception
  throw;
}

Weil a thrust::system_errorist std::runtime_error, können wir es alternativ auf die gleiche Weise wie eine breite Klasse von Fehlern behandeln, wenn wir nicht die Genauigkeit des vorherigen Beispiels benötigen:

try
{
  // call some_function which may throw something
  some_function();
}
catch(std::runtime_error &e)
{
  std::cerr << "Runtime error during some_function: " << e.what() << std::endl;
}
Jared Hoberock
quelle
1
Die Schubköpfe scheinen neu angeordnet worden zu sein. <thrust/system/cuda_error.h>ist jetzt effektiv <thrust/system/cuda/error.h>.
Chappjc
Jared, ich denke, meine Wrapper-Bibliothek fasst Ihre vorgeschlagene Lösung zusammen - meistens - und ist leicht genug, um wahrscheinlich ersetzt zu werden. (Siehe meine Antwort)
einpoklum
27

Der C ++ - kanonische Weg: Nicht auf Fehler prüfen ... Verwenden Sie die C ++ - Bindungen, die Ausnahmen auslösen.

Früher ärgerte mich dieses Problem. und ich hatte früher eine Lösung mit Makro-Sperma-Wrapper-Funktion, genau wie in den Antworten von Talonmies und Jared, aber ehrlich? Dies macht die Verwendung der CUDA Runtime API noch hässlicher und C-ähnlicher.

Ich habe das also anders und grundlegender angegangen. Ein Beispiel für das Ergebnis finden Sie hier im CUDA- vectorAddBeispiel - mit vollständiger Fehlerprüfung für jeden Laufzeit-API-Aufruf:

// (... prepare host-side buffers here ...)

auto current_device = cuda::device::current::get();
auto d_A = cuda::memory::device::make_unique<float[]>(current_device, numElements);
auto d_B = cuda::memory::device::make_unique<float[]>(current_device, numElements);
auto d_C = cuda::memory::device::make_unique<float[]>(current_device, numElements);

cuda::memory::copy(d_A.get(), h_A.get(), size);
cuda::memory::copy(d_B.get(), h_B.get(), size);

// (... prepare a launch configuration here... )

cuda::launch(vectorAdd, launch_config,
    d_A.get(), d_B.get(), d_C.get(), numElements
);    
cuda::memory::copy(h_C.get(), d_C.get(), size);

// (... verify results here...)

Erneut - alle potenziellen Fehler werden überprüft, und eine Ausnahme, wenn ein Fehler aufgetreten ist (Einschränkung: Wenn der Kernel nach dem Start einen Fehler verursacht hat , wird er nach dem Versuch, das Ergebnis zu kopieren, nicht vorher abgefangen, um sicherzustellen, dass der Kernel erfolgreich war müssen zwischen dem Start und der Kopie mit einem cuda::outstanding_error::ensure_none()Befehl auf Fehler überprüft werden ).

Der obige Code verwendet my

Dünne Modern-C ++ - Wrapper für die CUDA Runtime API-Bibliothek (Github)

Beachten Sie, dass die Ausnahmen nach dem fehlgeschlagenen Aufruf sowohl eine Zeichenfolgenerklärung als auch den CUDA-Laufzeit-API-Statuscode enthalten.

Einige Links dazu, wie CUDA-Fehler mit diesen Wrappern automatisch überprüft werden:

einpoklum
quelle
10

Die hier diskutierte Lösung hat bei mir gut funktioniert. Diese Lösung verwendet integrierte Cuda-Funktionen und ist sehr einfach zu implementieren.

Der entsprechende Code wird unten kopiert:

#include <stdio.h>
#include <stdlib.h>

__global__ void foo(int *ptr)
{
  *ptr = 7;
}

int main(void)
{
  foo<<<1,1>>>(0);

  // make the host block until the device is finished with foo
  cudaDeviceSynchronize();

  // check for error
  cudaError_t error = cudaGetLastError();
  if(error != cudaSuccess)
  {
    // print the CUDA error message and exit
    printf("CUDA error: %s\n", cudaGetErrorString(error));
    exit(-1);
  }

  return 0;
}
jthomas
quelle