CUDA 内核和内存访问（一个内核不完全执行，下一个不启动）

Posted 2023-04-15

【中文标题】CUDA 内核和内存访问（一个内核不完全执行，下一个不启动）

【英文标题】：CUDA kernels and memory access (one kernel doesn't execute entirely and the next doesn't get launched)

【发布时间】：2012-09-05 23:27:45

【问题描述】：

我在这里遇到了麻烦。我启动了两个内核，检查某个值是否是预期的值（memcpy 到主机），如果是我停止，如果不是我再次启动这两个内核。

第一个内核：

__global__  void aco_step(const KPDeviceData* data)

int obj = threadIdx.x;
int ant = blockIdx.x;
int id = threadIdx.x + blockIdx.x * blockDim.x;

*(data->added) = 1;

while(*(data->added) == 1)

    *(data->added) = 0;

    //check if obj fits
    int fits = (data->obj_weights[obj] + data->weight[ant] <= data->max_weight);
    fits = fits * !(getElement(data->selections, data->selections_pitch, ant, obj));

    if(obj == 0)
        printf("ant %d going..\n", ant);
    __syncthreads();

...

代码在此之后继续。但是那个 printf 永远不会被打印出来，那个 syncthreads 只是为了调试目的。

“添加”变量是共享的，但由于共享内存是 PITA 并且通常会在代码中引发错误，所以我暂时将其删除。这个“添加”变量不是最聪明的做法，但它比替代方法更快，后者检查数组中的任何变量是否是主机上的某个值并决定是否继续迭代。

getElement，只是简单地用音高做矩阵内存计算以访问正确的位置并返回那里的元素：

int* el = (int*) ((char*)mat + row * pitch) + col;
return *el;

obj_weights 数组具有正确的大小，n*sizeof(int)。权重数组 ants*sizeof(float) 也是如此。所以它们并没有越界。

这个之后的内核在开始时有一个 printf，它也不会被打印，并且在 printf 之后它在设备内存上设置一个变量，并且在内核完成后将该内存复制到 CPU，当我在 CPU 代码中打印它时，它不是正确的值。所以我认为这个内核正在做一些非法的事情，第二个甚至没有启动。

我正在测试一些实例，当我启动 8 个块和 512 个线程时，它运行正常。 32 个块，512 个线程，OK。但是8块1024线程，这样一来，内核就不行了，32块1024线程也不行。

我做错了吗？内存访问？我是否启动了太多线程？

编辑：尝试删除“添加”变量和 while 循环，因此它应该只执行一次。仍然不起作用，没有任何内容被打印，即使 printf 就在三个初始行之后并且下一个内核也没有打印任何内容。

编辑：另一件事，我使用的是 GTX 570，因此根据http://en.wikipedia.org/wiki/CUDA，“每个块的最大线程数”为 1024。也许我会坚持最大 512 或检查我能把这个值放多高。

【参考方案1】：

__syncthreads() 仅当条件在块的所有线程上的计算结果相同时才允许使用内部条件代码。

在您的情况下，条件会遇到竞争条件并且是不确定的，因此它很可能会针对不同的线程计算出不同的结果。

printf() 输出仅在内核成功完成后显示。在这种情况下，它不是由于上面提到的问题，所以输出永远不会出现。您可以通过测试所有 CUDA 函数调用的错误返回码来解决这个问题。

【讨论】：

好吧，我明白了syncthreads点，完全忘记了。但是关于内存写入，如果多个线程向同一个地址写入同一个值，这个值会被更新，不知道会被更新多少次。取自这里：***.com/questions/5953955/…

赞成，因为我不知道内核启动返回的错误，所以我结束了发现问题。在此处获取错误捕获代码：code.google.com/p/stanford-cs193g-sp2010/wiki/…

我不完全确定 *(data-> added) 测试的目的是什么。但至少要消除竞争条件，请在 while() 之前插入一个 __syncthreads()。由于 __syncthreads() 同步每个块，这还需要将标志移回共享内存（不要害怕 - 如果编程正确，共享内存没有任何问题）。

我使用的是动态分配的共享数组，他们所做的只是修改了我的代码。我将添加的变量移回共享，它现在可以工作了。我还稍微更改了代码。每个块启动 512 个线程很好，但显然启动 1024 不是，即使我的 GPU 支持每个块 1024。

其他一些资源（即寄存器或共享内存）的短缺可能会阻止您启动每个块支持的最大线程数。您可以使用 Nvidia 的 Occupancy Calculator 电子表格来检查特定案例的限制。