283

CUDAの質問に対する回答とコメントを見ると、CUDAタグwikiで、すべてのAPI呼び出しの戻りステータスでエラーをチェックする必要があることがよく示されています。APIドキュメントには、、、、などcudaGetLastErrorの関数が含まれていますがcudaPeekAtLastError、cudaGetErrorStringこれらを組み合わせて、多くの追加コードを必要とせずにエラーを確実にキャッチして報告するための最良の方法は何ですか？

4

4 に答える 4

346

おそらく、ランタイムAPIコードのエラーをチェックする最良の方法は、次のようにアサートスタイルハンドラー関数とラッパーマクロを定義することです。

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

次に、各API呼び出しをgpuErrchkマクロでラップできます。マクロは、ラップするAPI呼び出しの戻りステータスを処理します。次に例を示します。

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

呼び出しでエラーが発生した場合、エラーを説明するテキストメッセージと、エラーが発生したコード内のファイルと行がに送信されstderr、アプリケーションが終了します。必要に応じて、より高度なアプリケーションをgpuAssert呼び出すのではなく、例外を発生させるように変更することもできます。exit()

2番目の関連する質問は、カーネル起動のエラーをチェックする方法です。これは、標準のランタイムAPI呼び出しのようにマクロ呼び出しで直接ラップすることはできません。カーネルの場合、次のようになります。

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

最初に無効な起動引数をチェックし、次にカーネルが停止して実行エラーをチェックするまでホストを強制的に待機させます。次のような後続のブロッキングAPI呼び出しがある場合は、同期を削除できます。

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

この場合、cudaMemcpy呼び出しは、カーネルの実行中に発生したエラーまたはメモリコピー自体からのエラーのいずれかを返す可能性があります。これは初心者にとって混乱を招く可能性があります。問題が発生している可能性がある場所を理解しやすくするために、デバッグ中にカーネルを起動した後に明示的な同期を使用することをお勧めします。

CUDA動的並列処理を使用する場合、非常に類似した方法を、デバイスカーネルでのCUDAランタイムAPIの使用、およびデバイスカーネルの起動後に適用できます。

#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);
   }
}
于 2012-12-26T09:35:49.810 に答える
73

上記のtalonmiesの答えは、assertスタイルの方法でアプリケーションを中止するための優れた方法です。

大規模なアプリケーションの一部として、C++コンテキストでエラー状態を報告して回復したい場合があります。

std::runtime_error以下を使用して派生したC++例外をスローすることにより、これを行うためのかなり簡潔な方法を次に示します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);
  }
}

これにより、ファイル名、行番号、およびの英語の説明がcudaError_t、スローされた例外の.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;
}

出力：

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

のクライアントはsome_function、必要に応じてCUDAエラーを他の種類のエラーと区別できます。

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

thrust::system_errorはであるためstd::runtime_error、前の例の精度が必要ない場合は、代わりに、幅広いクラスのエラーと同じ方法で処理できます。

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;
}
于 2012-12-27T07:31:21.220 に答える
30

C ++の標準的な方法：エラーをチェックしないでください...例外をスローするC++バインディングを使用してください。

私は以前、この問題に悩まされていました。TalonmiesとJaredの回答と同じように、以前はマクロ兼ラッパー関数ソリューションを使用していましたが、正直なところ？これにより、CUDAランタイムAPIの使用がさらに醜くCのようになります。

だから私はこれに別のより根本的な方法でアプローチしました。結果のサンプルとして、CUDAvectorAddサンプルの一部を次に示します。すべてのランタイムAPI呼び出しの完全なエラーチェックが含まれています。

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

繰り返しますが、すべての潜在的なエラーがチェックされ、エラーが発生した場合は例外になります（注意：カーネルが起動後にエラーを引き起こした場合、結果をコピーしようとした後ではなく、後でキャッチされます。カーネルが成功したことを確認するには、コマンドを使用して、起動とコピーの間のエラーをチェックする必要がありますcuda::outstanding_error::ensure_none()）。

上記のコードは私の

シンモダン-CUDAランタイムAPIライブラリ（Github）のC++ラッパー

例外には、呼び出しが失敗した後、文字列の説明とCUDAランタイムAPIステータスコードの両方が含まれていることに注意してください。

CUDAエラーがこれらのラッパーで自動的にチェックされる方法へのいくつかのリンク：

于 2013-12-09T18:58:16.583 に答える
13

ここで説明した解決策は私にとってうまくいきました。このソリューションは組み込みのcuda関数を使用しており、実装は非常に簡単です。

関連するコードを以下にコピーします。

#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;
}
于 2015-12-19T15:35:20.793 に答える