cudaGetLastError。哪个内核执行引发了它？

Question

我已经实现了一个管道，其中许多内核在特定流中启动。内核被排入流中，并在调度程序决定最佳时执行。

在我的代码中，在每次内核入队之后，我通过调用 cudaGetLastError 检查是否有任何错误，根据文档，“它从运行时调用返回最后一个错误。这个调用也可能从以前的异步返回错误代码发射”。因此，如果内核只是入队而不执行，我理解返回的错误仅指内核正确入队（参数检查、网格和块大小、共享内存等...）。

我的问题是：我将许多不同的内核排入队列，而无需等待每个内核的执行完成。现在想象一下，我的一个内核（我们称之为 Kernel1）中有一个错误，它会导致非法内存访问（例如）。如果我在入队后立即检查 cudaGetLastError，则返回值为成功，因为它已正确入队。所以我的 CPU 线程继续前进并继续将内核排队到流中。在某些时候 Kernel1 被执行并引发了非法内存访问。因此，下次我检查 cudaGetLastError 时，我会得到 cuda 错误，但到那时，CPU 线程是代码中的另一个点。因此，我知道有一个错误，但我不知道是哪个内核引发了它。

一个选项是同步（阻塞 CPU 线程）直到每个内核的执行完成，然后检查错误代码，但出于性能原因，这不是一个选项。

问题是，有什么方法可以查询哪个内核引发了 cudaGetLastError 返回的给定错误代码？如果不是，您认为最好的处理方法是什么？

Answer 1

有一个environment variableCUDA_​LAUNCH_​BLOCKING，您可以使用它来序列化内核启动的其他异步序列的内核执行。这应该允许您通过主机代码中的内部错误检查或通过像cuda-memcheck 这样的外部工具来隔离导致错误的内核实例。

Answer 2

我测试了 3 个不同的选项：

1. 将 CUDA_​LAUNCH_​BLOCKING 环境变量设置为 1。这会强制阻塞 CPU 线程，直到内核执行完成。我们可以在每次执行后检查是否有错误捕获确切的故障点。虽然这对性能有明显影响，但这可能有助于在生产环境中限制错误，而无需在客户端执行任何更改。
2. 分发使用标志 -lineinfo 编译的生产代码并使用 cuda-memncheck 再次运行代码。这对性能没有影响，我们也不需要在客户端中执行任何更改。虽然，我们必须在稍微不同的环境中执行二进制文件，并且在某些情况下，例如运行 GPU 任务的服务，可能很难实现。

3. 在每次内核调用后插入一个回调。在 userData 参数中，包括内核调用的唯一 id，以及可能使用的参数的一些信息。这可以直接分布在生产环境中，并始终为我们提供准确的故障点，我们不需要在客户端执行任何更改。虽然，这种方法的性能影响是巨大的。显然，回调函数由驱动程序线程处理并导致性能影响。我写了一个代码来测试它

   #include <cuda_runtime.h>
   
   #include <vector>
   #include <chrono>
   #include <iostream>
   
   #define BLOC_SIZE       1024
   #define NUM_ELEMENTS    BLOC_SIZE * 32
   #define NUM_ITERATIONS  500
   
   __global__ void KernelCopy(const unsigned int *input, unsigned int *result) {
     unsigned int pos = blockIdx.x * BLOC_SIZE + threadIdx.x;
     result[pos] = input[pos];
   }
   
   void CUDART_CB myStreamCallback(cudaStream_t stream, cudaError_t status, void *data) {
     if (status) {
       std::cout << "Error: " << cudaGetErrorString(status) << "-->";
     }
   }
   
   #define CUDA_CHECK_LAST_ERROR   cudaStreamAddCallback(stream, myStreamCallback, nullptr, 0)
   
   int main() {
     cudaError_t c_ret;
     c_ret = cudaSetDevice(0);
     if (c_ret != cudaSuccess) {
       return -1;
     }
   
     unsigned int *input;
     c_ret = cudaMalloc((void **)&input, NUM_ELEMENTS * sizeof(unsigned int));
     if (c_ret != cudaSuccess) {
       return -1;
     }
   
     std::vector<unsigned int> h_input(NUM_ELEMENTS);
     for (unsigned int i = 0; i < NUM_ELEMENTS; i++) {
       h_input[i] = i;
     }
   
     c_ret = cudaMemcpy(input, h_input.data(), NUM_ELEMENTS * sizeof(unsigned int), cudaMemcpyKind::cudaMemcpyHostToDevice);
     if (c_ret != cudaSuccess) {
       return -1;
     }
   
     unsigned int *result;
     c_ret = cudaMalloc((void **)&result, NUM_ELEMENTS * sizeof(unsigned int));
     if (c_ret != cudaSuccess) {
       return -1;
     }
   
     cudaStream_t stream;
     c_ret = cudaStreamCreate(&stream);
     if (c_ret != cudaSuccess) {
       return -1;
     }
   
     std::chrono::steady_clock::time_point start;
     std::chrono::steady_clock::time_point end;
   
     start = std::chrono::steady_clock::now();
     for (unsigned int i = 0; i < 500; i++) {
       dim3 grid(NUM_ELEMENTS / BLOC_SIZE);
       KernelCopy <<< grid, BLOC_SIZE, 0, stream >>> (input, result);
       CUDA_CHECK_LAST_ERROR;
     }
     cudaStreamSynchronize(stream);
     end = std::chrono::steady_clock::now();
     std::cout << "With callback took (ms): " << std::chrono::duration<float, std::milli>(end - start).count() << '\n';
   
     start = std::chrono::steady_clock::now();
     for (unsigned int i = 0; i < 500; i++) {
       dim3 grid(NUM_ELEMENTS / BLOC_SIZE);
       KernelCopy <<< grid, BLOC_SIZE, 0, stream >>> (input, result);
       c_ret = cudaGetLastError();
       if (c_ret) {
         std::cout << "Error: " << cudaGetErrorString(c_ret) << "-->";
       }
     }
     cudaStreamSynchronize(stream);
     end = std::chrono::steady_clock::now();
     std::cout << "Without callback took (ms): " << std::chrono::duration<float, std::milli>(end - start).count() << '\n';
   
     c_ret = cudaStreamDestroy(stream);
     if (c_ret != cudaSuccess) {
       return -1;
     }
     c_ret = cudaFree(result);
     if (c_ret != cudaSuccess) {
       return -1;
     }
     c_ret = cudaFree(input);
     if (c_ret != cudaSuccess) {
       return -1;
     }
   
     return 0;
   }
   

输出：

需要回调（毫秒）：47.8729

没有回调（毫秒）：1.9317

（CUDA 9.2、Windows 10、Visual Studio 2015、Nvidia Tesla P4）

对我来说，在生产环境中，唯一有效的方法是 2。