Каков канонический способ проверки на ошибки с использованием API времени выполнения CUDA?

258

Просматривая ответы и комментарии на вопросы CUDA и вики-тег CUDA , я вижу, что часто предлагается проверять состояние возврата каждого вызова API на наличие ошибок. Документация по API содержит такие функции, как cudaGetLastError, cudaPeekAtLastErrorи cudaGetErrorString, но как лучше всего объединить их, чтобы надежно отлавливать и сообщать об ошибках, не требуя большого количества дополнительного кода?

talonmies
источник
13
Примеры NVIDIA CUDA содержат заголовок helper_cuda.h, в котором вызываются макросы getLastCudaErrorи checkCudaErrorsкоторый в значительной степени соответствует описанному в принятом ответе . Смотрите образцы для демонстраций. Просто установите образцы вместе с инструментарием, и он у вас будет.
chappjc
@chappjc Я не думаю, что этот вопрос и ответ претендуют на то, чтобы быть оригинальными, если это то, что вы имеете в виду, но имеет смысл, чтобы обучить людей, использующих проверку ошибок CUDA.
JackOLantern
@JackOLantern Нет, я не это имел в виду. Эти вопросы и ответы были очень полезны для меня, и их, безусловно, легче найти, чем заголовок в SDK. Я подумал, что было бы полезно отметить, что NVIDIA также справляется с этим и где искать больше. Я бы смягчил тон моего комментария, если бы мог. :)
chappjc
Инструменты отладки, позволяющие «приблизиться» к началу ошибок, значительно улучшились с 2012 года в CUDA. Я не работал с отладчиками на основе графического интерфейса, но вики-тег CUDA упоминает командную строку cuda-gdb. Это ОЧЕНЬ мощный инструмент, так как он позволяет вам просматривать реальные деформации и потоки на самом графическом процессоре (хотя в большинстве случаев требуется архитектура
2.0+
@bluefeet: какова была сделка с редактированием, которое вы откатили? Похоже, что на самом деле ничего не изменилось в уценке, но это было принято в качестве редактирования. Было ли что-то гнусное на работе?
talonmies

Ответы:

304

Вероятно, лучший способ проверить наличие ошибок в коде API времени выполнения - это определить функцию-обработчик стиля assert и макрос-обертку следующим образом:

#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()более сложное приложение, если бы это было необходимо.

Второй связанный с этим вопрос заключается в том, как проверять наличие ошибок при запуске ядра, которые нельзя напрямую обернуть в вызов макроса, как стандартные вызовы 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 очень похожая методология может и должна применяться к любому использованию API среды выполнения CUDA в ядрах устройств, а также после запуска любого ядра устройства:

#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
источник
8
@harrism: я так не думаю. Сообщество Wiki предназначено для вопросов или ответов, которые часто редактируются. Это не один из тех
talonmies
1
мы не должны добавить, cudaDeviceReset()прежде чем выйти также? И пункт для освобождения памяти?
Аврелий
2
@talonmies: для вызовов Async CUDA во время выполнения, таких как cudaMemsetAsync и cudaMemcpyAsync, требуется ли также синхронизация устройства gpu и потока хоста через вызов gpuErrchk (cudaDeviceSynchronize ())?
Нурабха
2
Обратите внимание, что явная синхронизация после запуска ядра не является неправильной, но может серьезно изменить производительность выполнения и семантику чередования. Если вы используете чередование, выполнение явной синхронизации для отладки может скрыть целый класс ошибок, которые может быть трудно отследить в сборке выпуска.
masterxilo
Есть ли способ получить более конкретные ошибки для выполнения ядра? Все ошибки, которые я получаю, дают мне номер строки из кода хоста, а не из ядра.
Азмисов
70

Ответ talonmies выше - отличный способ прервать приложение в assertстиле.

Иногда мы можем захотеть сообщить и исправить состояние ошибки в контексте C ++ как часть более крупного приложения.

Вот довольно краткий способ сделать это, создав исключение C ++, полученное std::runtime_errorпри использовании 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;
}
Джаред Хоберок
источник
1
Заголовки тяги, кажется, были переставлены. <thrust/system/cuda_error.h>сейчас эффективно <thrust/system/cuda/error.h>.
chappjc
Джаред, я думаю, моя библиотека-обертка включает в себя предложенное вами решение - в основном, и достаточно легка, чтобы ее можно было заменить. (См. Мой ответ)
einpoklum
27

C ++ - канонический способ: не проверять на ошибки ... используйте привязки C ++, которые генерируют исключения.

Раньше меня раздражала эта проблема; и я имел обыкновение иметь решение для функции макрокоманды, как в Talonmies и ответах Джареда, но, если честно? Это делает использование CUDA Runtime API еще более уродливым и похожим на C.

Так что я подошел к этому по-другому и более фундаментально. Для примера результата, вот часть vectorAddпримера CUDA - с полной проверкой ошибок каждого вызова 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()команды).

Код выше использует мой

Оболочки Thin Modern-C ++ для библиотеки API CUDA Runtime (Github)

Обратите внимание, что после неудачного вызова исключения содержат как строковое объяснение, так и код состояния API среды выполнения CUDA.

Несколько ссылок на то, как ошибки CUDA автоматически проверяются с помощью этих оболочек:

einpoklum
источник
10

Обсуждаемое здесь решение хорошо сработало для меня. Это решение использует встроенные функции 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;
}
jthomas
источник