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

Question

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

因此，作为次优选择，我们至少可以触发运行时错误以停止执行不合理的工作并指示出现问题。

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

不会导致未定义的行为
一旦达到将停止内核执行
不会触发编译器warning/error

?

Answer 1

选项 1 断言：

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

直接来自文档：

#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;") 来触发 illegal instruction.

这里有一些代码证明了这一点：

#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 at ...kernel.cu[31] : an illegal instruction was encountered

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

Triggering a runtime error within a CUDA kernel

exception

error-handling

cuda

custom-error-handling

选项 1 断言：

选项 2 非法指令