验证调用 cuda 内核的次数

Verfiy the number of times a cuda kernel is called

本文关键字：内核 调用 cuda 验证      更新时间：2023-10-16

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

__global__ void run2048Times(){ }

然后从主代码调用它：

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

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

我所做的是将指针传递给内核，并在每次内核运行时++d指针。

__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 之间的各种数字。

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

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

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

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

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

第一个解决方案的问题在于它被编译为以下 PTX 代码(有关 PTX 指令的说明，请参阅此处)：

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;

似乎该计数器的唯一要点是进行分析(即分析代码的运行方式)，而不是实际计算某些内容(即对程序没有功能优势)。

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

相关文章：