同一个线程两次访问同一个内存条会不会冲突？

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 相加，依此类推。

所有这些操作都需要 warp 中的每个线程从同一组中获取 2 个不同的值，但是，我没有遇到任何冲突：

我错过了什么？

Answer 1

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