CUDA 共享内存寻址答案

【问题标题】：CUDA shared memory addressingCUDA 共享内存寻址
【发布时间】：2012-07-01 23:37:19
【问题描述】：

我了解当我在内核中声明一个共享内存数组时，所有线程都会声明相同大小的数组。类似的代码

__shared__ int s[5];

将在每个线程中创建一个 20 字节的数组。我理解解决共享内存的方式是它在所有线程中都是通用的。所以，如果我按如下方式处理下标 10

s[10] = 1900;

它是所有线程中完全相同的内存位置。不同的线程不会为下标10访问不同的共享内存地址。这是正确的吗？编译器当然会抛出下标超出范围的警告。

【问题讨论】：

顾名思义，共享内存数组是在共享内存中，并不是在每个线程中创建的。也就是说，共享数组s只有一个，这就是为什么不同线程访问共享内存中的同一个地址时会发生bank冲突的原因。
共享内存由块中的线程共享，而不是内核的所有线程。它可以被视为可编程缓存。

【解决方案1】：

实际上它将创建一个 20 字节的数组每个块，而不是每个线程。

块中的每个线程都可以访问这 20 个字节。因此，如果您需要每个线程有 N 个字节，以及一个有 M 个线程的块，则需要为每个块创建一个 N*M 缓冲区。

在你的情况下，如果有 128 个线程，你会有

__shared__ int array[5*128];

并且array[10] 将是块内任何线程的有效地址。

【讨论】：