在 CUDA 内核中触发运行时错误

Question

在CUDA中，我们不能抛出异常；但是 - 我们可以并且确实偶尔会遇到无法继续的异常情况，并且在主机上我们会抛出异常。

因此，作为第二好的，我们至少可以触发运行时错误以停止做不合理的工作并表明出现问题。

在 CUDA 内核中这样做有什么好处，它：

1. 不会导致未定义的行为
2. 一旦达到将停止内核执行
3. 不会触发编译器警告/错误

?

Answer 1

选项 1 断言：

当前支持的所有 GPU 都包含内核断言机制，此处描述。

直接从文档：

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

cudaErrorAssert任何在执行期间触发断言调用的内核都会报告一个专用的 CUDA 运行时错误。与所有其他设备端运行时错误一样，上下文将在出现错误时被销毁，并且需要创建一个新的上下文（通过调用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

...kernel.cu[31] 处的 CUDA 错误：遇到非法指令