为什么全球+共享比单独全球共享更快

Question

我需要一些帮助来理解Ron Farber代码的行为：http://www.drdobbs.com/parallel/cuda-supercomputing-for-the-masses-part/208801731?pgno=2

我不明白共享内存的使用如何比非共享内存版本提供更快的性能。也就是说，如果我添加更多的索引计算步骤，并使用添加另一个Rd/Wr周期来访问共享内存，这如何比仅使用全局内存更快？在任一情况下，相同数量的Rd/Wr周期访问全局mem。每个内核实例仍然只访问一次数据。数据仍然使用全局内存传入/传出。内核实例数相同。寄存器计数看起来是一样的。如何添加更多的处理步骤才能使其更快。(我们不会减去任何流程步骤。)从本质上讲，我们正在做更多的工作，而且完成得更快。

共享内存访问比全局访问快得多，但它不是零(或负)。我遗漏了什么？

“慢”代码：

__global__ void reverseArrayBlock(int *d_out, int *d_in) {
int inOffset  = blockDim.x * blockIdx.x;
int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
int in  = inOffset + threadIdx.x;
int out = outOffset + (blockDim.x - 1 - threadIdx.x);
d_out[out] = d_in[in];
}

“快速”代码：

__global__ void reverseArrayBlock(int *d_out, int *d_in) {
extern __shared__ int s_data[];

int inOffset  = blockDim.x * blockIdx.x;
int in  = inOffset + threadIdx.x;

// Load one element per thread from device memory and store it
// *in reversed order* into temporary shared memory
s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];

// Block until all threads in the block have written their data to shared mem
__syncthreads();

// write the data from shared memory in forward order,
// but to the reversed block offset as before
int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
int out = outOffset + threadIdx.x;
d_out[out] = s_data[threadIdx.x];
}

Answer 1

早期支持CUDA的设备(计算能力< 1.2)不会将“慢”版本中的d_outout写入视为合并写入。这些设备仅在“最好”的情况下合并存储器访问，其中半扭曲中的第i线程访问第i个字。结果，将发出16个存储器事务以服务于每半个warp的d_outout写入，而不是仅仅一个存储器事务。

从计算能力1.2开始，CUDA中的内存合并规则变得更加宽松。因此，“慢”版本中的d_outout写入也会被合并，并且不再需要使用共享内存作为便签簿。

您的代码示例的来源是2008年6月撰写的文章"CUDA，面向大众的超级计算:第5部分“。支持CUDA的计算能力1.2的设备是在2009年才上市的，所以这篇文章的作者明确谈到了计算能力小于1.2的设备。

有关更多详细信息，请参阅NVIDIA CUDA C Programming Guide中的F.3.2.1节。

Answer 2

这是因为共享内存更接近计算单元，因此延迟和峰值带宽不会成为此计算的瓶颈(至少在矩阵乘法的情况下)。

但最重要的是，最重要的原因是，磁贴中的许多数字正在被许多线程重用。因此，如果您从全局访问，您将多次检索这些数字。将它们写入共享内存一次将消除这种浪费的带宽使用

Answer 3

在查看全局内存访问时，速度较慢的代码向前读取，向后写入。快速的代码可以向前读和写。我认为更快的代码更快，因为缓存层次结构在某种程度上进行了优化，以便以降序(朝向更高的内存地址)访问全局内存。

CPU执行一些推测性获取，在程序访问数据之前，它们将从较高的内存地址填充高速缓存线。也许类似的事情也会发生在GPU上。