超过 1 块的 CUDA 共享内存停止工作答案

【问题标题】：CUDA shared memory more than 1 block stops working超过 1 块的 CUDA 共享内存停止工作
【发布时间】：2016-04-30 22:52:01
【问题描述】：

我正在使用共享内存为学校编写程序。我使该功能正常工作，但它仅适用于每个网格大小为 1 的块。如果每个网格有超过 1 个块，则该功能不再起作用。

例如，如果我发送一个大小为 10、具有 1 个块和 10 个线程的数组，则此函数可以正常工作。如果我在同一个数组中发送 2 个块和每个块 5 个线程，那么它不再起作用。当我说不再工作时，我得到的是我发送的相同数组，然后又返回，就像什么都没发生一样。

__global__ void rankSortShared(int* a, int n)
{
    int threadId = threadIdx.x + blockDim.x * blockIdx.x;
    int x = 0;

    // Make dynamic sized shared array
    // size will be determined from the 3rd parameter in the 
    // kernal call 'ranksortShared<<<blocksPerGrid, threadsPerBlock, *(size of shared mem)*>>>()'
    extern __shared__ int b[];

    // copy data from incomming array to shared array
    // then sync the threads so all threads have put their 
    // values into the shared array
    b[threadId] = a[threadId];
    __syncthreads();

    // now use shared array for faster lookups
    for (int j = 0; j < n; j++)
    {
        // handle duplicate values
        if (b[j] < b[threadId] || (b[threadId] == b[j] && j < threadId))
        {
            x++;
        }   
    }
    // put data back into array to be transferred back to CPU
    a[x] = b[threadId];
}

对于我的生活，我无法弄清楚为什么。有没有一种特殊的方法可以跨多个块使用共享内存？

只是为了尝试调试它，我用a[threadId] = threadId 替换了所有代码，并且数组看似正确（打印数字 0 到 9）。所以我不明白为什么如果 threadId 似乎是正确的并且它适用于 1 个块，这将不起作用。

【问题讨论】：

标签： cuda shared-memory

【解决方案1】：

这里有问题：

int threadId = threadIdx.x + blockDim.x * blockIdx.x;
b[threadId] = ...;

您的threadId 变量是一个全局唯一的线程索引。这意味着您在网格中包含的线程越多（例如，通过更多块），该索引就会越高。

但是对于共享内存，每个块中的索引从零开始。所以最终，随着您添加更多块，您的 threadId 变量将大于块中的共享内存量。

通常的解决方案是这样做：

b[threadIdx.x] = ...;

因为threadIdx.x 变量在每个块中从零开始。（您可以使用此策略将每次出现的b[threadId] 替换为b[threadIdx.x]，而不仅仅是一次出现。）

您的代码中可能还存在其他问题。由于您没有展示完整的代码，因此无法诊断其他人。

基于这样的陈述：

当我说不再工作时，我得到的是我发送的相同数组，然后又返回，就像什么都没发生一样。

我猜你没有做proper cuda error checking。在向他人寻求帮助之前，你真的应该这样做。即使你不理解错误输出，它也会帮助其他试图帮助你的人。

【讨论】：