14.3 基于共享内存实现

在N-体计算中，最内层循环有足够的局部性和重用性，无须编程人员的介入即可让缓存有效运转，但是在CUDA构架上，使用共享内存来作为显式缓存使用能带来很大好处，[1]正如代码清单14-4展示的。内层的循环采用两层循环进行分块：第一层循环按照每次一个线程块方式遍历N个个体并加载到共享内存，第二层循环使用保存在共享内存的个体数据进行迭代计算。当同一线程束的线程访问共享内存中的同一个位置时，将得到广播机制的优化，因此这种使用模式非常适合GPU硬件的构架。

这种方法同样在Harris等人的报告中被提到。当N值很大时，该方法能够得到非常好的性能，几乎达到GPU的理论计算极限。

代码清单14-4 ComputeNBodyGravitation_shared

global void   
ComputeNBodyGravitation_Shared( float \*force, float \*posMass, float softeningSquared, size_t N) float4 \*posMass4  $=$  posMass; extern _shared float4 shPosMass[]; for ( int i  $=$ (blockIdx.x\*blockDim.x+threadIdx.x; i  $<  \mathbb{N}$  . i  $^ { \text{一} } =$  blockDim.x\*gridDim.x ) { float acc[3]  $= \{0\}$  ; float4 myPosMass  $=$  posMass4[i]; #pragma unroll 32 for (int j  $= 0$  ;j  $<  \mathbb{N}$  ;j  $^{+ + }$  blockDim.x){ shPosMass [threadIdx.x]  $=$  posMass4[j+threadIdx.x]; _syncthreads(); for (size_t k  $= 0$  ;k  $<$  blockDim.x;k++) { float fx,fy,fz; float4 bodyPosMass  $=$  shPosMass[k]; bodyBodyInteraction( &fx,&fy,&fz, myPosMass.x,myPosMass.y,myPosMass.z, bodyPosMass.x, bodyPosMass.y, bodyPosMass.z, bodyPosMass.w, softeningSquared); acc[0]  $+ =$  fx; acc[1]  $+ =$  fy; acc[2]  $+ =$  fz; } _syncthreads(); } force[3*i+0]  $=$  acc[0]; force[3*i+1]  $=$  acc[1]; force[3*i+2]  $=$  acc[2];

与之前用到的内核一样，循环展开得到了更好的性能。表14-2总结了在使用共享内存实现时循环展开的效果。此时最优展开因子为4，得到 $18\%$ 的性能提升。

表14-2 使用共享内存的内核的循环展开效果

[1] 共享内存在流处理器簇（SM）1.x构架的必备资源，它没有缓存。但是它被证明在所有的CUDA构架上都很有用，尽管在流处理器簇（SM）2.x和3.x上很轻微。