为什么全球+共享的速度比全球独自

我需要一些帮助理解的罗恩·法伯的代码的行为：http://www.drdobbs.com/parallel/cuda-supercomputing-for-the-masses-part/208801731?pgno=2 为什么全球+共享的速度比全球独自

我不理解如何使用共享MEM的是给在非共享内存更快的性能版。即，如果我再添加一些索引计算步骤并使用另一个Rd/Wr循环来访问共享内存，那么如何比单独使用全局内存更快？在这两种情况下，相同的数字或Rd/Wr循环访问全局内存。每个内核实例只能访问一次数据。数据仍然使用全局内存进/出。内核实例的数量是相同的。寄存器数量看起来是一样的。如何添加更多的处理步骤使其更快。（我们没有减去任何流程步骤。）基本上我们正在做更多的工作，并且它正在更快地完成。

共享内存访问速度比全球快得多，但它不是零（或负值）。我错过了什么？

的“慢”的代码：

__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]; 
}

来源

2012-08-13 Doug

你在使用什么卡？它可以产生显着的差异。 – 2012-08-13 18:13:08

由于显着不同，较旧的卡仅支持在“正向顺序”模式下有效读取全局内存。较新的卡片不应该受此影响。（我相信这在2.x发布时有所变化） – 2012-08-13 19:32:16

该文章写于2008年。硬件是费米前。我认为这是'老'卡。 – Doug 2012-08-13 19:45:40

早期支持CUDA的设备（计算能力< 1.2）不会将您的“慢”版本中的d_out [out]写入视为合并写入。这些设备只会在“最好”的情况下合并内存访问，在这种情况下，第二个半字节中的第i个线程访问第i个字。结果，将发出16个内存事务来服务每半个warp的d_out [out]写入，而不是仅仅一次内存事务。

从计算能力1.2开始，CUDA中的内存合并规则变得更放松了。因此，“慢”版本中的d_out [out]写入也会合并，并且不再需要使用共享内存作为便笺。

代码示例的源代码是2008年6月编写的文章“CUDA，大众超级计算：第5部分”。具有计算能力1.2的CUDA支持的设备仅在2009年上市，所以作者该文章清楚地讨论了具有计算能力的设备< 1.2。

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

来源

2012-08-13 19:37:46

这是因为共享的存储器更接近计算单元，因此等待时间和峰值带宽将不是这个计算的瓶颈（至少在矩阵乘法的情况下）

但最重要的是，最重要的原因是瓷砖中的很多数字被重复使用很多线程。所以如果你从全球进行访问，你会多次检索这些数字。将它们写入共享内存将消除浪费的带宽使用

来源

2012-08-13 18:21:20 lurscher

这里没有数据重用。它是从一个独特的全球读取的，它被写入另一个独特的全球一次。在这两个代码块中都会发生相同数量的全局访问。即相同的全局带宽。 – Doug 2012-08-13 18:29:35

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

CPU执行一些推测性提取，在数据被程序触及之前，它们将从更高的存储器地址填充缓存行。也许在GPU上发生类似的事情。

来源

2012-08-13 19:23:03

CC 2.0和3.0数据高速缓存层次结构不预取。预取通过软件说明得到支持（参见PTX手册）。 SM到L1的访问都是使用128B的事务。 L1到L2访问使用32B事务。未缓存的L2（未缓存在L1中）可以使用32B事务完成。 – 2012-08-14 02:48:36

为什么全球+共享的速度比全球独自

回答

相关问题