Запуск ошибки времени выполнения в ядре CUDA - PullRequest
Новая
       53

Запуск ошибки времени выполнения в ядре CUDA

0 голосов
/ 08 июня 2018

В CUDA мы не можем исключать;но - мы можем и действительно иногда достигаем исключительных ситуаций, в которых мы не можем продолжить работу, и на хосте мы бы выдавали исключение.

Так что, как второй лучший вариант, мы можем по крайней мере вызвать ошибку времени выполненияпрекратить выполнять необоснованную работу и указать, что что-то пошло не так.

Что хорошо делать в ядре CUDA, которое:

  1. Не вызывает неопределенного поведения
  2. Остановит выполнение ядра после достижения
  3. Не вызовет предупреждение / ошибку компилятора

?

1 Ответ

0 голосов
/ 08 июня 2018

Варианты 1 Утверждения:

Все поддерживаемые в настоящее время графические процессоры включают механизм утверждений в ядре, описанный здесь .

Непосредственно из документации: 

#include <assert.h>

__global__ void testAssert(void)
{
    int is_one = 1;
    int should_be_one = 0;

    // This will have no effect
    assert(is_one);

    // This will halt kernel execution
    assert(should_be_one);
}

int main(int argc, char* argv[])
{
    testAssert<<<1,1>>>();
    cudaDeviceSynchronize();

    return 0;
}

Существует специальная ошибка времени выполнения CUDA cudaErrorAssert, о которой будет сообщать любое ядро, которое запускает вызов подтверждения во время выполнения.Как и во всех других ошибках времени выполнения на стороне устройства, контекст ошибки будет уничтожен, и потребуется создать новый контекст (путем вызова cudaDeviceReset()).

Обратите внимание, что (к сожалению) не поддерживается в MacOSиз-за ограничений водителя.

Вариант 2 Недопустимая инструкция

Вы можете использовать встроенные ptx и asm ("trap;"), чтобы вызвать недопустимую инструкцию .

Вот некоторыекод, демонстрирующий: 

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <cstdio>
#include <cstdlib>

__global__ void kernel(int i) {
    if(i > 0) {
        asm("trap;");
    }

    ::printf("%d\n", i);
}

inline void error_check(cudaError_t err, const char* file, int line) {
    if(err != cudaSuccess) {
        ::fprintf(stderr, "CUDA ERROR at %s[%d] : %s\n", file, line, cudaGetErrorString(err));
        abort();
    }
}
#define CUDA_CHECK(err) do { error_check(err, __FILE__, __LINE__); } while(0)


int main() {
    kernel<<<1, 1>>>(0);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());


    kernel<<<1, 1>>>(1);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());

}

, который выводит:

0

CUDA ERROR в ... kernel.cu [31]: обнаружена недопустимая инструкция

...