Apa cara kanonik untuk memeriksa kesalahan menggunakan API runtime CUDA?

258

Melihat melalui jawaban dan komentar pada pertanyaan CUDA, dan di wiki tag CUDA , saya melihat sering disarankan bahwa status pengembalian setiap panggilan API harus diperiksa untuk kesalahan. Dokumentasi API berisi fungsi-fungsi seperti cudaGetLastError,, cudaPeekAtLastErrordan cudaGetErrorString, tetapi apa cara terbaik untuk menyatukannya untuk menangkap dan melaporkan kesalahan secara andal tanpa memerlukan banyak kode tambahan?

talonmies
sumber
13
Sampel CUDA dari NVIDIA berisi header, helper_cuda.h, yang memiliki makro yang dipanggil getLastCudaErrordan checkCudaErrors, yang melakukan cukup banyak seperti yang dijelaskan dalam jawaban yang diterima . Lihat contoh untuk demonstrasi. Cukup pilih untuk menginstal sampel bersama dengan toolkit dan Anda akan memilikinya.
chappjc
@ chappjc Saya tidak berpikir pertanyaan dan jawaban ini berpura-pura menjadi orisinal, jika ini yang Anda maksudkan, tetapi memiliki manfaat untuk mendidik orang menggunakan pemeriksaan kesalahan CUDA.
JackOLantern
@JackOLantern Tidak, bukan itu yang saya maksudkan. T&J ini sangat membantu saya dan tentu saja lebih mudah ditemukan daripada beberapa header di SDK. Saya pikir itu berharga untuk menunjukkan ini juga bagaimana NVIDIA menanganinya dan di mana mencari lebih banyak. Saya akan melunakkan nada komentar saya jika saya bisa. :)
chappjc
Alat debugging memungkinkan Anda untuk "mendekati" di mana kesalahan mulai telah meningkat banyak sejak 2012 di CUDA. Saya belum pernah bekerja dengan debugger berbasis GUI tetapi tag CUDA wiki menyebutkan baris perintah cuda-gdb. Ini adalah alat yang SANGAT ampuh karena memungkinkan Anda untuk melangkah melewati lungsin dan utas yang sebenarnya pada GPU itu sendiri (sebagian besar membutuhkan arsitektur 2.0+)
opetrenko
@bluefeet: apa kesepakatan dengan suntingan yang Anda putar kembali? Tampaknya tidak ada yang benar-benar berubah pada penurunan harga, tetapi diterima sebagai hasil edit. Apakah ada sesuatu yang jahat di tempat kerja?
talonmies

Jawaban:

304

Mungkin cara terbaik untuk memeriksa kesalahan dalam kode API runtime adalah dengan mendefinisikan fungsi penangan gaya tegas dan pembungkus makro seperti ini:

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

Anda kemudian dapat membungkus setiap panggilan API dengan gpuErrchkmakro, yang akan memproses status pengembalian dari panggilan API yang dibungkus, misalnya:

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

Jika ada kesalahan dalam panggilan, pesan teks yang menggambarkan kesalahan dan file dan baris dalam kode Anda di mana kesalahan terjadi akan dipancarkan ke stderrdan aplikasi akan keluar. Anda dapat memodifikasi gpuAssertuntuk meningkatkan pengecualian daripada memanggil exit()aplikasi yang lebih canggih jika diperlukan.

Pertanyaan kedua yang terkait adalah bagaimana memeriksa kesalahan dalam peluncuran kernel, yang tidak dapat langsung dibungkus dengan panggilan makro seperti panggilan API runtime standar. Untuk kernel, kira-kira seperti ini:

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

pertama-tama akan memeriksa argumen peluncuran yang tidak valid, kemudian memaksa host untuk menunggu sampai kernel berhenti dan memeriksa kesalahan eksekusi. Sinkronisasi dapat dihilangkan jika Anda memiliki panggilan API pemblokiran berikutnya seperti ini:

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

dalam hal ini cudaMemcpy panggilan dapat mengembalikan kesalahan yang terjadi selama eksekusi kernel atau dari salinan memori itu sendiri. Ini bisa membingungkan bagi pemula, dan saya akan merekomendasikan menggunakan sinkronisasi eksplisit setelah peluncuran kernel selama debugging untuk membuatnya lebih mudah untuk memahami di mana masalah mungkin timbul.

Perhatikan bahwa ketika menggunakan CUDA Dynamic Parallelism , metodologi yang sangat mirip dapat dan harus diterapkan pada penggunaan API runtime CUDA di kernel perangkat, serta setelah kernel perangkat diluncurkan:

#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);
   }
}
talonmies
sumber
8
@harrism: Saya rasa tidak. Wiki Komunitas dimaksudkan untuk pertanyaan atau jawaban yang sering diedit. Ini bukan salah satunya
talonmies
1
tidakkah kita menambahkan cudaDeviceReset()sebelum keluar juga? Dan klausa untuk alokasi memori?
Aurelius
2
@talonmies: Untuk panggilan runtime CUDA Async, seperti cudaMemsetAsync dan cudaMemcpyAsync, apakah itu juga memerlukan sinkronisasi perangkat gpu dan host thread melalui panggilan ke gpuErrchk (cudaDeviceSynchronize ())?
nurabha
2
Perhatikan bahwa sinkronisasi eksplisit setelah peluncuran kernel tidak salah tetapi sangat dapat mengubah kinerja eksekusi dan semantik interleaving. Jika Anda menggunakan interleaving, melakukan sinkronisasi eksplisit untuk debugging mungkin menyembunyikan seluruh kelas bug yang mungkin sulit dilacak dalam rilis Release.
masterxilo
Apakah ada cara untuk mendapatkan kesalahan yang lebih spesifik untuk eksekusi kernel? Semua kesalahan yang saya dapatkan hanya memberi saya nomor baris dari kode host, bukan dari kernel.
Azmisov
70

Jawaban talonmies di atas adalah cara yang baik untuk membatalkan aplikasi dengan assertcara-gaya.

Terkadang kami mungkin ingin melaporkan dan memulihkan dari kondisi kesalahan dalam konteks C ++ sebagai bagian dari aplikasi yang lebih besar.

Berikut cara yang cukup singkat untuk melakukan itu dengan melempar pengecualian C ++ yang berasal dari std::runtime_errorpenggunaan 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);
  }
}

Ini akan memasukkan nama file, nomor baris, dan deskripsi bahasa Inggris dari anggota cudaError_tpengecualian yang dilemparkan .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;
}

Hasil:

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

Klien some_functiondapat membedakan kesalahan CUDA dari jenis kesalahan lainnya jika diinginkan:

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

Karena thrust::system_errora std::runtime_error, kita dapat menanganinya dengan cara yang sama dari kelas kesalahan yang luas jika kita tidak memerlukan ketepatan contoh sebelumnya:

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
sumber
1
Header dorong tampaknya telah diatur ulang. <thrust/system/cuda_error.h>sekarang efektif <thrust/system/cuda/error.h>.
chappjc
Jared, saya pikir perpustakaan pembungkus saya menggunakan solusi yang disarankan - kebanyakan, dan cukup ringan untuk menjadi pengganti. (Lihat jawaban saya)
einpoklum
27

C ++ - cara kanonik: Jangan periksa kesalahan ... gunakan binding C ++ yang melempar pengecualian.

Dulu saya kesal dengan masalah ini; dan saya dulu punya solusi fungsi makro-cum-wrapper seperti di Talonmies dan jawaban Jared, tapi, jujur? Itu membuat menggunakan API CUDA Runtime lebih jelek dan seperti-C.

Jadi saya sudah mendekati ini dengan cara yang berbeda dan lebih mendasar. Untuk sampel hasil, inilah bagian dari vectorAddsampel CUDA - dengan pengecekan kesalahan lengkap dari setiap panggilan API 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...)

Lagi - semua kesalahan potensial diperiksa, dan pengecualian jika kesalahan terjadi (peringatan: Jika kernel menyebabkan beberapa kesalahan setelah peluncuran, itu akan ditangkap setelah upaya untuk menyalin hasil, bukan sebelumnya; untuk memastikan kernel berhasil Anda akan perlu memeriksa kesalahan antara peluncuran dan salin dengan acuda::outstanding_error::ensure_none() perintah).

Kode di atas menggunakan my

Pembungkus Thin Modern-C ++ untuk perpustakaan CUDA Runtime API (Github)

Perhatikan bahwa pengecualian membawa penjelasan string dan kode status API CUDA runtime setelah panggilan gagal.

Beberapa tautan ke bagaimana kesalahan CUDA diperiksa secara otomatis dengan pembungkus ini:

einpoklum
sumber
10

Solusi yang dibahas di sini bekerja dengan baik untuk saya. Solusi ini menggunakan fungsi cuda bawaan dan sangat mudah diterapkan.

Kode yang relevan disalin di bawah ini:

#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
sumber