验证调用 cuda 内核的次数

Question

假设你有一个想要运行 2048 次的 cuda 内核，那么你可以这样定义你的内核：

__global__ void run2048Times(){ }

然后你从你的主代码中调用它：

run2048Times<<<2,1024>>>();

到目前为止一切似乎都很好。但是现在说，当您调用内核数百万次时，出于调试目的，您想要验证您实际调用内核的次数是多少次。

我所做的是将一个指针传递给内核，并在每次内核运行时对指针进行 ++ 处理。

__global__ void run2048Times(int *kernelCount){ 
    kernelCount[0]++; // Add to the pointer
}

但是，当我将该指针复制回主函数时，我得到“2”。

起初它让我感到困惑，然后在 5 分钟的咖啡和来回踱步之后，我意识到这可能是有道理的，因为 cuda 内核同时运行 1024 个自身实例，这意味着内核覆盖了“kernelCount” [0]" 而不是真正添加它。

所以我决定这样做：

__global__ void run2048Times(int *kernelCount){

   // Get the id of the kernel
   int id = blockIdx.x * blockDim.x + threadIdx.x;

   // If the id is bigger than the pointer overwrite it
   if(id > kernelCount[0]){
        kernelCount[0] = id; 
   }
}

天才！！我认为这可以保证工作。直到我运行它并得到了 0 到 2000 之间的各种数字。

这告诉我上面提到的问题在这里仍然发生。

有没有办法做到这一点，即使它涉及强制内核暂停并等待彼此运行？

Answer 1

假设这是一个简化的示例，并且您实际上并没有像其他人已经建议的那样尝试进行分析，而是想在更复杂的场景中使用它，您可以使用atomicAdd 来实现您想要的结果，即将确保增量操作作为单个原子操作执行：

__global__ void run2048Times(int *kernelCount){ 
    atomicAdd(kernelCount, 1); // Add to the pointer
}

为什么您的解决方案不起作用：

您的第一个解决方案的问题在于它被编译成以下 PTX 代码（有关 PTX 指令的描述，请参阅 here）：

ld.global.u32   %r1, [%rd2];
add.s32     %r2, %r1, 1;
st.global.u32   [%rd2], %r2;

您可以通过使用--ptx 选项调用nvcc 来验证这一点，以仅生成中间表示。

这里可能发生的是以下时间线，假设您启动 2 个线程（注意：这是一个简化的示例，并不完全是 GPU 的工作原理，但足以说明问题）：

1. 线程0从kernelCount读取0
2. 线程1从kernelCount读取0
3. 线程0 将其本地副本增加1
4. 线程0存储1返回kernelCount
5. 线程1 将其本地副本增加1
6. 线程1存储1返回kernelCount

即使启动了 2 个线程，你也会得到 1。

即使线程按顺序启动，您的第二个解决方案也是错误的，因为线程索引是从 0 开始的。所以我假设你想这样做：

__global__ void run2048Times(int *kernelCount){

   // Get the id of the kernel
   int id = blockIdx.x * blockDim.x + threadIdx.x;

   // If the id is bigger than the pointer overwrite it
   if(id + 1 > kernelCount[0]){
        kernelCount[0] = id + 1; 
   }
}

这将编译成：

    ld.global.u32   %r5, [%rd1];
    setp.lt.s32 %p1, %r1, %r5;
    @%p1 bra    BB0_2;

    add.s32     %r6, %r1, 1;
    st.global.u32   [%rd1], %r6;

BB0_2:
    ret;

这里可能发生的是以下时间表：

1. 线程0从kernelCount读取0
2. 线程1从kernelCount读取0
3. 线程1 比较0 和1 + 1 并将2 存储到kernelCount
4. 线程0 比较0 和0 + 1 并将1 存储到kernelCount

你最终得到错误的结果 1。

如果您想更好地理解同步和非原子操作的问题，我建议您选择一本好的并行编程/CUDA 书籍。

编辑：

为了完整起见，使用atomicAdd的版本编译成：

    atom.global.add.u32     %r1, [%rd2], 1;

Answer 2

似乎该计数器的唯一目的是进行分析（即分析代码如何运行）而不是实际计算某些东西（即对程序没有功能优势）。

有专为此任务设计的分析工具。例如，nvprof 给出了调用次数，以及代码库中每个内核的一些时间指标。