14.4_基于常量内存实现
14.4 基于常量内存实现
Stone等人描述了一种直接库伦求和（direct coulomb summation，DCS）的方法，使用共享内存来存储分子建模应用中的图像格点 [1]，因此必须使用常量内存储存个体的数据。代码清单14-5展示了一个CUDA内核，使用与我们的引力计算相同的方式进行处理。对于给定的一个内核，只有64KB的常量内存可以供开发人员使用，因此每一次内核调用只能处理大约4000个16字节长的个体数据。常量 g BodiesPerPass指定最内层循环可以考虑的个体数。
因为在最内层循环中，每个线程正在读取相同的个体数据，广播机制能够优化同一线程束的所有线程的读操作，使得常量内存能够获得很好的效果。
代码清单14-5 N-体内核（常量内存）
const int g BodiesPerPass = 4000;
constant __device __float4 g_constantBodies(g BodiesPerPass);
template<typename T> global _void
ComputeNBodyGravitation_GPU_AOS_const( T *force, T *posMass, T softeningSquared, size_t n, size_t N)
{ for (int i $=$ blockIdx.x\*blockDim.x + threadIdx.x; i $< \mathbb{N}$ . i $+ =$ blockDim.x\*gridDim.x ) { T acc[3] $= \{0\}$ ; float4 me $=$ ((float4 \*) posMass)[i]; T myX $=$ me.x; T myY $=$ me.y; T myZ $=$ me.z; for (int j $= 0$ ; j $< \mathbb{n}$ ; j++) { float4 body $=$ g_constantBodies[j]; float fx, fy, fz; bodyBodyInteraction( &fx, &fy, &fz, myX, myY, myZ, body.x, body.y, body.z, body.w, softeningSquared); acc[0] $+ =$ fx; acc[1] $+ =$ fy; acc[2] $+ =$ fz; } force[3*i+0] $+ =$ acc[0]; force[3*i+1] $+ =$ acc[1]; force[3*i+2] $+ =$ acc[2]; }正如代码清单14-6展示的，主机端的代码必须遍历所有的个体，在每一次调用内核前先调用CUDAMemcpyToSymbolAsync()函数复制数据到常量内存。
代码清单14-6 主机端代码（常量内存N-体）
float
ComputeNBodyGravitation_GPU_AOS_const( float \*force, float \*posMass, float softeningSquared, size_t N
1
{udaError_t status;udaEvent_t evStart $= 0$ ,evStop $= 0$ float ms $= 0.0$ size_t bodiesLeft $= \mathbf{N}$ .void \*p; CUDART_CHECK(cudaGetSymbolAddress(&p,g_constantBodies)); CUDART_CHECK(cudaEventCreate(&evStart)); CUDART_CHECK(cudaEventCreate(&evStop)); CUDART_CHECK(cudaEventRecord(evStart，NULL）); for（size_t i $= 0$ ;i<N；i+=g_bodiesPerPass）{ // bodiesThisPass $=$ max(bodiesLeft,g_bodiesPerPass); size_t bodiesThisPass $=$ bodiesLeft; if（ bodiesThisPass>g_bodiesPerPass）{ bodiesThisPass $=$ g_bodiesPerPass; } CUDART_CHECK(cudaMemcpyToSymbolAsync( g_constantBodies, ((float4\*)posMass)+i, bodiesThisPass*sizeof(float4), 0,udaMemcpyDeviceToDevice, NULL））; ComputeNBodyGravitation_GPU_AOS_const<<300,256>>>( force，posMass，softeningSquared,bodiesThisPass，N）; bodiesLeft-- bodiesThisPass; } CUDART_CHECK(cudaEventRecord(evStop，NULL）); CUDART_CHECK(cudaDeviceSynchronize()); CUDART_CHECK(cudaEventElapsedTime(&ms,evStart,evStop）); Error: CUDAEventDestroy(evStop); CUDAEventDestroy(evStart); return ms;