[C++知识库]CUDA C编程（十三）使用共享内存作为可编程管理缓存减少全局内存访问

??使用共享内存的主要原因之一是要缓存片上的数据，从而减少核函数中全局内存访问的次数，接下来将重新使用并行归约核函数，以共享内存作为可编程管理缓存来减少全局内存的访问。

使 用 共 享 内 存 的 归 约
??下面的reduceGmem核函数将被作为基准性能的起点。实现并行归约只使用全局内存，输入元素的内循环是完全展开的。核函数如下：

__global__ void reduceGmem(int *g_idata,int *g_odata,unsigned int n)
{
   unsigned int tid = threadIdx.x;
   //计算数据块的偏移量
   int *idata = g_idata + blokcIdx.x * blockDim.x;
   
   unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if(idx >= n)
      return;
   
   //核函数执行一个使用全局内存的原地归约。并将其归约到32个元素
   if(blockDim.x >= 1024 && tid < 512)
      idata[tid] += idata[tid + 512];
   __synthreads();
   if(blockDim.x > 512 && tid < 256)
      idata[tid] += idata[tid + 256];
   __synthreads();
   if(blockDim.x > 256 && tid < 128)
      idata[tid] += idata[tid + 128];
   __synthreads();
   if(blockDim.x > 128 && tid < 64)
      idata[tid] += idata[tid + 64];
   __synthreads();
   
   //核函数执行原地归约，这个过程仅使用每个线程块的第一个线程束
   //在循环展开的部分，volatile修饰符用来确保当线程束在锁步中执行，只有最新数值能被读取
   if(tid < 32)
   {
      volatile int *vsmem = idata;
      vsmem[tid] += vsmem[tid + 32];
      vsmem[tid] += vsmem[tid + 16];
      vsmem[tid] += vsmem[tid + 8];
      vsmem[tid] += vsmem[tid + 4];
      vsmem[tid] += vsmem[tid + 2];
      vsmem[tid] += vsmem[tid + 1];
   }
   
   //最后，分配给该线程块的输入数据块总数被写回到全局内存中
   if(tid == 0)
      g_odata[blockIdx.x] = idata[0];
}

??接下来测试下面的原地归约核函数reduceSmem，它增加了带有共享内存的全局内存操作，这个核函数和原来的reduceGmem核函数几乎相同，然而，reduceSmem函数没有使用全局内存中的输入数组子集来执行原地归约，而是使用共享数组smem。

//使用下述的宏，将块大小设置为恒定的128个线程
#define DIM 128
__global__ void reduceSmem(int *g_idata,int *g_odata,unsigned int n)
{
   //smem被声明为与每个线程块具体有相同的维度。
   __shared__ int smem[DIM];
   unsigned int tid = threadIdx.x;
   //计算数据块的偏移量
   int *idata = g_idata + blokcIdx.x * blockDim.x;
   
   unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if(idx >= n)
      return;

   smem[tid] = idata[tid];
   __synthreads();
   
   //核函数执行一个使用全局内存的原地归约。并将其归约到32个元素
   if(blockDim.x >= 1024 && tid < 512)
       smem[tid] += idata[tid + 512];
   __synthreads();
   if(blockDim.x > 512 && tid < 256)
       smem[tid] += idata[tid + 256];
   __synthreads();
   if(blockDim.x > 256 && tid < 128)
       smem[tid] += idata[tid + 128];
   __synthreads();
   if(blockDim.x > 128 && tid < 64)
       smem[tid] += idata[tid + 64];
   __synthreads();
   
   //核函数执行原地归约，这个过程仅使用每个线程块的第一个线程束
   //在循环展开的部分，volatile修饰符用来确保当线程束在锁步中执行，只有最新数值能被读取
   if(tid < 32)
   {
      volatile int *vsmem = smem;
      vsmem[tid] += vsmem[tid + 32];
      vsmem[tid] += vsmem[tid + 16];
      vsmem[tid] += vsmem[tid + 8];
      vsmem[tid] += vsmem[tid + 4];
      vsmem[tid] += vsmem[tid + 2];
      vsmem[tid] += vsmem[tid + 1];
   }
   
   //最后，分配给该线程块的输入数据块总数被写回到全局内存中
   if(tid == 0)
      g_odata[blockIdx.x] = idata[0];
}

??经过测试可以得到，使用共享内存明显减少了全局内存访问。

使 用 展 开 的 并 行 归 约
??在前面的核函数中，每个线程块处理一个数据块，接下来通过展开线程块来提升内核性能。以下内核展开了4个线程块，即每个线程处理来自于4个数据块的数据元素。通过展开，可以获得如下优势：1.通过在每个线程中提供更多的并行I/O,增加全局内存的吞吐量；2.全局内存存储事务减少了1/4；3.整体内核性能提升。具体的核函数如下：

#define DIM 128
__global__ void reduceSmemUnroll(int *g_idata,int *g_odata,unsigned int n)
{
   //smem被声明为与每个线程块具体有相同的维度。
   __shared__ int smem[DIM];
   unsigned int tid = threadIdx.x;
   //计算数据块的偏移量
   int *idata = g_idata + blokcIdx.x * blockDim.x;
   
   unsigned int idx = blockIdx.x * blockDim.x * 4 + threadIdx.x;
   
   int tmpSum = 0;
   if(idx + 3 * blokcDim.x <= n)
   {
      int a1 = g_idata[idx];
      int a2 = g_idata[idx + blockDim.x];
      int a3 = g_idata[idx + blockDim.x * 2];
      int a4 = g_idata[idx + blockDim.x * 3];
      tmpSum = a1 + a2 + a3 + a4;
   }

   smem[tid] = tmpSum;
   __synthreads();
   
   //核函数执行一个使用全局内存的原地归约。并将其归约到32个元素
   if(blockDim.x >= 1024 && tid < 512)
       smem[tid] += idata[tid + 512];
   __synthreads();
   if(blockDim.x > 512 && tid < 256)
       smem[tid] += idata[tid + 256];
   __synthreads();
   if(blockDim.x > 256 && tid < 128)
       smem[tid] += idata[tid + 128];
   __synthreads();
   if(blockDim.x > 128 && tid < 64)
       smem[tid] += idata[tid + 64];
   __synthreads();
   
   //核函数执行原地归约，这个过程仅使用每个线程块的第一个线程束
   //在循环展开的部分，volatile修饰符用来确保当线程束在锁步中执行，只有最新数值能被读取
   if(tid < 32)
   {
      volatile int *vsmem = smem;
      vsmem[tid] += vsmem[tid + 32];
      vsmem[tid] += vsmem[tid + 16];
      vsmem[tid] += vsmem[tid + 8];
      vsmem[tid] += vsmem[tid + 4];
      vsmem[tid] += vsmem[tid + 2];
      vsmem[tid] += vsmem[tid + 1];
   }
   
   //最后，分配给该线程块的输入数据块总数被写回到全局内存中
   if(tid == 0)
      g_odata[blockIdx.x] = idata[0];
}

使 用 动 态 共 享 内 存 的 并 行 归 约
??并行归约核函数还可以使用动态共享内存来执行，通过以下声明，在reduceSmemUnroll中使用动态内存取代静态共享内存：extern __shared__ int smem[];

有 效 带 宽
??由于归约核函数是受内存带宽约束的，所以评估它们时所使用的适当的性能指标是有效带宽。有效带宽是在核函数的完整执行时间内I/O的数量(以字节为单位)。对于内存约束的应用程序，有效带宽是一个估算实际带宽利用率的很好的指标。它可以表示为:
??????有效带宽 = (读字节 + 写字节) ÷ (运行时间 × 10⁹) GB/s
??下表总结了每个核函数已取得的有效带宽，显然，可以通过展开块来获得有效带宽的显著改进。这样做使每个线程在运行时间中同时有多个请求，这会导致内存总线高饱和。