cuda - 同一个线程两次访问同一个内存库会导致冲突吗？

Question

我正在研究一个可以减少向量的内核。它基本上将向量中的所有位置相加并将结果存储在位置 0 中。

我正在遵循这个方案，包含 512 个浮点元素块：

编码：

//scratch[] is a vector located in shared memory with all 512 elements
NUM_ELEMENTS = 512;
for( stride=NUM_ELEMENTS/2; stride>=1; stride = stride/2 ) {
  if (threadIdx.x < stride){
    scratch[threadIdx.x] += scratch[threadIdx.x + stride];
  }
  __syncthreads();
}

奇怪的是，我预计会出现共享的银行冲突，但我没有。在第一次迭代中，线程 0 将位置 0 和位置 256 相加，它们位于同一组中。线程 1 将位置 1 和位置 257 相加，依此类推。

所有这些操作都需要扭曲中的每个线程从同一银行获得 2 个不同的值，但是，我没有任何冲突：

我错过了什么？

Answer 1

银行冲突的计算是基于每个请求的每个内存指令。共享加载（右侧）和共享存储（左侧）作为单独的指令执行，相隔许多时钟周期。