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

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_out [out]写入合并写入.这些设备只会在"最好"的情况下合并内存访问,其中半经线中的第i个线程访问第i个字.结果,将发出16个内存事务来为每半个warp而不是仅仅一个内存事务服务d_out [out]写入.

从计算能力1.2开始,CUDA中内存合并的规则变得更加轻松.因此,"慢"版本中的d_out [out]写入也将合并,并且不再需要使用共享内存作为暂存区.

您的代码示例的来源是文章"CUDA,面向大众的超级计算:第5部分",该文章于2008年6月编写.具有计算能力1.2的CUDA设备仅在2009年上市,因此该文章的作者清楚地谈到了关于计算能力<1.2的设备.

有关更多详细信息,请参阅NVIDIA CUDA C编程指南中的F.3.2.1部分.