CUDA：使用 atomicAdd 遇到非法内存访问答案

【问题标题】：CUDA:illegal memory access was encountered with atomicAddCUDA：使用 atomicAdd 遇到非法内存访问
【发布时间】：2021-08-12 06:00:23
【问题描述】：

只要我不增加网格大小或再次调用内核，我的程序中的原子操作就可以正常工作。怎么会这样？也许共享内存不会自动释放？

  __global__ void DevTest() {
     __shared__  int* k1;
    k1 = new int(0);
    atomicAdd( k1, 1);
  }
int main()
{
for (int i = 0; i < 100; i++) DevTest << < 50, 50 >> > (); 
}

【问题讨论】：

标签： c cuda

【解决方案1】：

这个：

 __shared__  int* k1;

在共享内存中为指针创建存储。指针在那里未初始化；它没有指向任何东西。

这个：

k1 = new int(0);

将该指针设置为指向设备堆上的某个位置，而不是共享内存中。 The device heap is limited by default to 8MB。此外，还有一个分配粒度，因此单个int 分配可能会占用超过 4 个字节的设备堆空间（它会）。

在 C++ 中，每个 new 操作都有一个对应的 delete 通常是一种很好的做法。你的内核代码没有这个，所以当你增加网格大小时，你会用掉越来越多的设备堆内存。您最终会遇到 8MB 的限制。

所以fix这个至少有2个选项：

delete 在线程代码末尾使用new 创建的分配
增加设备堆的限制，说明链接在上面的文档中

顺便说一句，共享内存由线程块中的所有线程共享。因此，对于 <<<50,50>>> 的内核启动，每个线程块中有 50 个线程。这 50 个线程中的每一个都将看到 same k1 指针，并且每个线程都会尝试将其设置为单独的位置/分配，因为每个线程都执行 new 操作。这没有任何意义，它会阻止上面的第 1 项正常工作（分配的 50 个指针值中的 49 个将丢失）。

所以你的代码没有任何意义。你在这里展示的东西不是一件明智的事情，也没有简单的方法来解决它。你可以这样做：

  __global__ void DevTest() {
    __shared__  int* k1;
    if (threadIdx.x == 0) k1 = new int(0);
    __syncthreads();
    atomicAdd( k1, 1);
    __syncthreads();
    if (threadIdx.x == 0) delete k1;
  }

如果您启动了足够多的块（即，如果有足够的常驻块，并考虑到未定义的分配粒度），即使这样的安排在理论上（也许在某些未来的大型 GPU 上）最终也会达到 8MB 的限制。所以正确的做法大概是1和2都做。

【讨论】：

谢谢，它有效。之前尝试过使用delete，但觉得只在一个线程中分配内存不值得。