c++ - 验证调用 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. thread 1 将其本地拷贝增加 1
6. thread 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;