Jaki jest kanoniczny sposób sprawdzania błędów za pomocą środowiska wykonawczego CUDA API?

Przeglądając odpowiedzi i komentarze na pytania CUDA oraz na wiki tagu CUDA , często widzę, że często zaleca się sprawdzanie, czy w statusie zwrotu każdego wywołania API nie ma błędów. Dokumentacja API zawiera funkcje takie jak cudaGetLastError, cudaPeekAtLastErrori cudaGetErrorString, ale jaki jest najlepszy sposób, aby je połączyć, aby niezawodnie wychwytywać i zgłaszać błędy bez konieczności posiadania dodatkowego kodu?

Prawdopodobnie najlepszym sposobem sprawdzenia błędów w kodzie API środowiska wykonawczego jest zdefiniowanie funkcji obsługi stylu assert i makra opakowania w następujący sposób:

#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);
   }
}

Następnie można zawinąć każde wywołanie interfejsu API za pomocą gpuErrchkmakra, które przetworzy status powrotu zawiniętego wywołania interfejsu API, na przykład:

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

Jeśli wystąpi błąd w wywołaniu, zostanie wysłany komunikat tekstowy opisujący błąd oraz plik i wiersz w kodzie, w którym wystąpił błąd, stderri aplikacja zostanie zamknięta. Można sobie wyobrazić modyfikację, gpuAssertaby zgłosić wyjątek, zamiast wywoływać exit()bardziej wyrafinowane aplikacje, jeśli byłyby one wymagane.

Drugim powiązanym pytaniem jest, jak sprawdzić błędy w uruchomieniach jądra, których nie można bezpośrednio zawrzeć w wywołaniu makra, takim jak standardowe wywołania API środowiska wykonawczego. W przypadku jąder coś takiego:

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

najpierw sprawdzi, czy argument uruchamiania jest nieprawidłowy, a następnie zmusi host do czekania, aż jądro się zatrzyma i sprawdzi błąd wykonania. Synchronizację można wyeliminować, jeśli wystąpi kolejne blokujące wywołanie interfejsu API:

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

w takim przypadku cudaMemcpywywołanie może zwrócić albo błędy, które wystąpiły podczas wykonywania jądra, albo błędy z samej kopii pamięci. Może to być mylące dla początkujących i zaleciłbym użycie jawnej synchronizacji po uruchomieniu jądra podczas debugowania, aby łatwiej zrozumieć, gdzie mogą wystąpić problemy.

Należy pamiętać, że podczas korzystania z równoległego dynamicznego interfejsu CUDA bardzo podobna metodologia może i powinna być stosowana do każdego użycia interfejsu API środowiska wykonawczego CUDA w jądrach urządzeń, a także po uruchomieniu dowolnego jądra urządzenia:

#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);
   }
}

Powyższa odpowiedź talonmies to świetny sposób na przerwanie aplikacji w assertstylu.

Czasami możemy chcieć zgłosić błąd i usunąć go w kontekście C ++ jako część większej aplikacji.

Oto rozsądnie zwięzły sposób, aby to zrobić, zgłaszając wyjątek C ++ wynikający z std::runtime_errorużycia 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);
  }
}

Spowoduje to włączenie nazwy pliku, numeru wiersza i opisu w języku angielskim cudaError_telementu zgłoszonego wyjątku .what():

#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;
}

Wyjście:

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

Klient some_functionmoże w razie potrzeby odróżnić błędy CUDA od innych rodzajów błędów:

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;
}

Ponieważ thrust::system_errorjest to std::runtime_error, możemy alternatywnie obsłużyć to w ten sam sposób, co szeroką klasę błędów, jeśli nie wymagamy precyzji z poprzedniego przykładu:

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;
}

C ++ - kanoniczny sposób: nie sprawdzaj błędów ... użyj powiązań C ++, które generują wyjątki.

Byłem zirytowany tym problemem; i kiedyś miałem rozwiązanie z funkcją makro-cum-wrapper, tak jak w Talonmies i odpowiedziach Jareda, ale szczerze? To sprawia, że ​​korzystanie z interfejsu API środowiska wykonawczego CUDA jest jeszcze bardziej brzydkie i podobne do C.

Podszedłem więc do tego w inny i bardziej fundamentalny sposób. Oto przykładowy vectorAddprzykład CUDA - z pełnym sprawdzaniem błędów każdego wywołania API środowiska wykonawczego:

// (... 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...)

Ponownie - wszystkie potencjalne błędy są sprawdzane, a wyjątek występuje, jeśli wystąpił błąd (zastrzeżenie: jeśli jądro spowodowało jakiś błąd po uruchomieniu, zostanie przechwycone po próbie skopiowania wyniku, a nie wcześniej; aby upewnić się, że jądro zakończyło się powodzeniem, zrobiłbyś to trzeba sprawdzić pod kątem błędu między uruchomieniem a kopią za pomocą cuda::outstanding_error::ensure_none()polecenia).

Powyższy kod używa mojego

Cienkie opakowania Modern-C ++ dla biblioteki API środowiska wykonawczego CUDA (Github)

Należy zauważyć, że wyjątki zawierają zarówno wyjaśnienie ciągu, jak i kod stanu środowiska wykonawczego CUDA po nieudanym wywołaniu.

Kilka linków do automatycznego sprawdzania błędów CUDA za pomocą tych opakowań:

• Program testowy rzucający i wychwytujący wyjątki
• Dokumentacja funkcji związanych z błędami

Omówione tutaj rozwiązanie działało dla mnie dobrze. To rozwiązanie wykorzystuje wbudowane funkcje cuda i jest bardzo łatwe do wdrożenia.

Odpowiedni kod jest kopiowany poniżej:

#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;
}