我有一个这样的数组：
data[16] = {10,1,8,-1,0,-2,3,5,-2,-3,2,7,0,11,0,2}
我想使用 G80 GPU 上的共享内存计算此数组的减少量。
NVIDIA 文档中引用的内核是这样的：
__global__ void reduce1(int *g_idata, int *g_odata) {
extern __shared__ int sdata[];
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
sdata[tid] = g_idata[i];
__syncthreads();
// here the reduction :
for (unsigned int s=1; s < blockDim.x; s *= 2) {
int index = 2 * s * tid;
if (index < blockDim.x) {
sdata[index] += sdata[index + s];
}
__syncthreads();
}
论文作者表示，这种方法存在银行冲突的问题。我试图理解，但我不知道为什么？我知道银行冲突和广播访问的定义，但仍然无法理解这一点。
银行冲突
> G80 处理器是一个非常古老的支持 CUDA 的 GPU，在第一代 CUDA GPU 中，计算能力为 1.0。 最近的 CUDA 版本(6.5 之后)不再支持这些设备，因此在线文档不再包含了解这些设备中的库结构的必要信息。
因此，我将从 CUDA 1.x C 编程指南中摘录 cc 6.5 设备的必要信息：
G.3.3. 共享内存
共享内存有 16 个库，这些库的组织使得连续的 32 位字映射 到连续的银行。每个组的带宽为每两个时钟周期 32 位。
对 warp 的共享内存请求被拆分为两个内存请求，每个请求一个 半翘曲，独立发行。结果，不可能有银行 属于经线前半部分的线与属于 同一经线的后半部分。
在这些设备中，共享内存具有 16 个组结构，因此每个组的"宽度"为 32 位或 4 字节。 例如，每个银行的宽度与int或float数量相同。 因此，让我们设想可能存储在这种共享内存中的前 32 个 4 字节数量，以及它们相应的库(使用f而不是sdata作为数组名称)：
extern __shared__ int f[];
index: f[0] f[1] f[2] f[3] ... f[15] f[16] f[17] f[18] f[19] ... f[31]
bank: 0 1 2 3 ... 15 0 1 2 3 ... 15
共享内存中的前 16 个int数量属于组 0 到 15，共享内存中接下来的 16 个int数量也属于组 0 到 15(依此类推，如果我们的int数组中有更多数据)。
现在让我们看一下将触发银行冲突的代码行：
for (unsigned int s=1; s < blockDim.x; s *= 2) {
int index = 2 * s * tid;
if (index < blockDim.x) {
sdata[index] += sdata[index + s];
}
让我们考虑通过上述循环的第一次传递，其中s为 1。 这意味着index是2*1*tid，所以对于每个线程，index只是threadIdx.x值的两倍：
threadIdx.x: 0 1 2 3 4 5 6 7 8 9 10 11 ...
index: 0 2 4 6 8 10 12 14 16 18 20 22 ...
bank: 0 2 4 6 8 10 12 14 0 2 4 6 ...
因此，对于此读取操作：
+= sdata[index + s]
我们有：
threadIdx.x: 0 1 2 3 4 5 6 7 8 9 10 11 ...
index: 0 2 4 6 8 10 12 14 16 18 20 22 ...
index + s: 1 3 5 7 9 11 13 15 17 19 21 23 ...
bank: 1 3 5 7 9 11 13 15 1 3 5 7 ...
因此，在前 16 个线程中，我们有两个线程想要从库 1 读取，两个线程想要从库 3 读取，两个线程想要从库 5 读取，依此类推。 因此，此读取周期在前 16 个线程组中遇到 2 路库冲突。 请注意，同一行代码上的其他读取和写入操作同样存在银行冲突：
sdata[index] +=
因为这将读取，然后写入，每组 16 个线程两次到银行 0、2、4 等。
对于可能正在阅读此示例的其他人请注意：如前所述，它仅适用于cc 1.x 设备。 在 cc 2.x 和更新的设备上演示库冲突的方法可能相似，但细节不同，这是由于曲速执行差异以及这些较新设备具有 32 路库结构而不是 16 路库结构的事实。