idiomatic kernel - Qual è il modo canonico per verificare la presenza di errori utilizzando l'API di runtime CUDA?





samples install (5)


Il C ++ - modo canonico: non controllare gli errori ... usa i collegamenti C ++ che generano eccezioni.

Ero irritato da questo problema; e avevo una soluzione con funzione macro-wrapper proprio come nelle risposte di Talonmies e Jared, ma, sinceramente? Rende l'uso dell'API di runtime CUDA ancora più brutto e simile a C.

Quindi mi sono avvicinato a questo in un modo diverso e più fondamentale. Per un esempio del risultato, ecco parte del campione vectorAdd CUDA - con controllo degli errori completo di ogni richiamo dell'API di runtime:

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

Ancora una volta - tutti i potenziali errori vengono controllati e segnalati tramite un'eccezione generata. Questo codice usa il mio

Wrapper Thin Modern-C ++ per la libreria API CUDA Runtime (Github)

Si noti che le eccezioni contengono sia una spiegazione di stringa sia il codice di stato dell'API di runtime CUDA dopo la chiamata fallita

Alcuni link su come gli errori CUDA vengono controllati automaticamente con questi wrapper:

Guardando attraverso le risposte e i commenti sulle domande CUDA, e nel tag wiki CUDA , vedo che spesso viene suggerito che lo stato di ritorno di ogni chiamata API dovrebbe essere controllato per gli errori. La documentazione dell'API contiene funzioni come cudaGetLastError , cudaPeekAtLastError e cudaGetErrorString , ma qual è il modo migliore per raggruppare questi elementi in modo affidabile per rilevare e segnalare errori senza richiedere molto codice aggiuntivo?




Probabilmente il modo migliore per verificare la presenza di errori nel codice API runtime è definire una funzione di gestione degli stili di assert e una macro wrapper come questa:

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

Puoi quindi racchiudere ogni chiamata API con la macro gpuErrchk , che elaborerà lo stato di ritorno dell'API chiamata wrapping, ad esempio:

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

Se c'è un errore in una chiamata, un messaggio testuale che descrive l'errore e il file e la linea nel codice in cui si è verificato l'errore saranno emessi su stderr e l'applicazione uscirà. Si potrebbe gpuAssert modificare gpuAssert per generare un'eccezione piuttosto che chiamare exit() in un'applicazione più sofisticata se fosse richiesta.

Una seconda domanda correlata riguarda il modo in cui verificare gli errori nei lanci del kernel, che non possono essere racchiusi direttamente in una chiamata macro come le chiamate API standard di runtime. Per i kernel, qualcosa del genere:

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

controllerà innanzitutto l'argomento di lancio non valido, quindi costringerà l'host ad attendere fino all'arresto del kernel ea verificare l'esistenza di un errore di esecuzione. La sincronizzazione può essere eliminata se hai una successiva chiamata API di blocco in questo modo:

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

nel qual caso la chiamata a cudaMemcpy può restituire gli errori che si sono verificati durante l'esecuzione del kernel o quelli dalla copia di memoria stessa. Questo può essere fonte di confusione per il principiante e raccomanderei l'uso della sincronizzazione esplicita dopo l'avvio del kernel durante il debug per rendere più facile capire dove potrebbero sorgere problemi.




La risposta di talonmie sopra è un ottimo modo per abortire un'applicazione in stile assert .

Occasionalmente potremmo voler segnalare e recuperare da una condizione di errore in un contesto C ++ come parte di un'applicazione più grande.

Ecco un modo ragionevolmente teso per farlo gettando un'eccezione C ++ derivata da std::runtime_error usando 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);
  }
}

Questo includerà il nome file, il numero di riga e una descrizione della lingua inglese di cudaError_t nel membro .what() dell'eccezione generata:

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

Il risultato:

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

Un client di some_function può distinguere gli errori CUDA da altri tipi di errori se lo si desidera:

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

Poiché thrust::system_error è uno std::runtime_error , possiamo alternativamente gestirlo nello stesso modo di un'ampia classe di errori se non richiediamo la precisione dell'esempio precedente:

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



La soluzione discussa here funzionato bene per me. Questa soluzione utilizza funzioni integrate di cuda ed è molto semplice da implementare.

Il codice rilevante è copiato di seguito:

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



Stilisticamente questo è molto bello. Nel mondo reale, mi aspetterei una probabilità del 60% di questa notazione invece di quella che hai dato:

C x c >>= f = C (value $ f x) (c + 1)

Ma è così piccolo che non vale la pena menzionarlo.

Su una nota più seria, non stilistica ma semantica: questa non è una monade. In realtà, viola tutte e tre le leggi della monade.

(1) return x >>= f  =  f x
(2) m >>= return    = m
(3) m >>= (f >=> g) = (m >>= f) >>= g

(Dove (>=>) è definito come f >=> g = \x -> fx >>= g . Se (>>=) è considerato un operatore "applicazione", quindi (>=>) è la composizione corrispondente operatore. Mi piace dichiarare la terza legge usando questo operatore perché fa emergere il significato della terza legge: associatività.)

Con questi calcoli:

(1):

return 0 >>= return 
  = C 0 0 >>= return
  = C (value $ return 0) 1
  = C 0 1
Not equal to return 0 = C 0 0

(2):

C 0 0 >>= return
  = C (value $ return 0) 1
  = C 0 1
Not equal to C 0 0

(3)

C 0 0 >>= (return >=> return)
  = C (value $ (return >=> return) 0) 1
  = C (value $ return 0 >>= return) 1
  = C (value $ C 0 1) 1
  = C 0 1

Is not equal to:

(C 0 0 >>= return) >>= return
  = C (value $ return 0) 1 >>= return
  = C 0 1 >>= return
  = C (value $ return 0) 2
  = C 0 2

Questo non è solo un errore nella tua implementazione - non c'è monade che "conta il numero di vincoli". Deve violare le leggi (1) e (2). Il fatto che il tuo violi la legge (3) è comunque un errore di implementazione.

Il guaio è che f nella definizione di (>>=) potrebbe restituire un'azione che ha più di un legame, e tu la stai ignorando. È necessario aggiungere il numero di binding dagli argomenti left e right:

C x c >>= f = C y (c+c'+1)
   where
   C y c' = f x

Questo conterà correttamente il numero di vincoli e soddisferà la terza legge, che è la legge di associatività. Non soddisferà gli altri due. Tuttavia, se si rilascia il +1 da questa definizione, si ottiene una vera monade, che è equivalente alla monade di Writer sul + monoid. Questo fondamentalmente somma insieme i risultati di tutte le subcomputazioni. Puoi usare questo per contare il numero di qualcosa , ma non per legare: il rilegare è troppo speciale per contare. Ma, ad esempio:

tick :: C ()
tick = C () 1

Quindi C conterà il numero di tick che si sono verificati nel calcolo.

In effetti, puoi sostituire Int con qualsiasi tipo e (+) con qualsiasi operatore associativo e ottenere una monade. Questo è ciò che una monade Writer è in generale. Se l'operatore non è associativo, fallirà la terza legge (capisci perché?).







cuda idiomatic error-checking