当前位置：首页 > news >正文

CUDA编程 | 5.3减少全局内存访问

news 2024/12/13 21:35:27

5.3减少全局内存访问

使用共享内存的并行归约

使用展开的并行归约

使用动态共享内存的并行归约

有效带宽

前言：从今天开始将会陆续更新CUDA编程的内容。（后面是一些废话，可以直接跳到正文部分）其实学习CUDA编程已经有一段时间，但因为各种拖延一直没有更新博客，由于昨天好友的push终于将这件事情提上日程~很开心开启一段新的学习之旅。之前没有更完的C++系列后续也会陆续缓慢更新，大家敬请期待~

注：此博客是对谭生的博客的一些学习感悟，详细内容请移步谭生的博客进行学习

谭升的博客

5.3减少全局内存访问

使用共享内存的主要原因就是减少对全局内存的访问，来减少不必要的延迟。

需要集中解决以下两个问题：

如何重新安排数据访问模式以避免线程束分化
如何展开循环以保证有足够的操作使指令和内存带宽饱和

使用共享内存的并行归约

下面分别是使用全局内存和使用共享内存进行并行归约的代码

__global__ void reduceGmem(int * g_idata,int * g_odata,unsigned int n)
{//set thread IDunsigned int tid = threadIdx.x;unsigned int idx = blockDim.x*blockIdx.x+threadIdx.x;//boundary checkif (tid >= n) return;//convert global data pointer to theint *idata = g_idata + blockIdx.x*blockDim.x;//in-place reduction in global memoryif(blockDim.x>=1024 && tid <512)idata[tid]+=idata[tid+512];__syncthreads();if(blockDim.x>=512 && tid <256)idata[tid]+=idata[tid+256];__syncthreads();if(blockDim.x>=256 && tid <128)idata[tid]+=idata[tid+128];__syncthreads();if(blockDim.x>=128 && tid <64)idata[tid]+=idata[tid+64];__syncthreads();//write result for this block to global memif(tid<32){volatile int *vmem = idata;vmem[tid]+=vmem[tid+32];vmem[tid]+=vmem[tid+16];vmem[tid]+=vmem[tid+8];vmem[tid]+=vmem[tid+4];vmem[tid]+=vmem[tid+2];vmem[tid]+=vmem[tid+1];}if (tid == 0)g_odata[blockIdx.x] = idata[0];}

__global__ void reduceSmem(int * g_idata,int * g_odata,unsigned int n)
{//set thread ID__shared__ int smem[DIM];unsigned int tid = threadIdx.x;//unsigned int idx = blockDim.x*blockIdx.x+threadIdx.x;//boundary checkif (tid >= n) return;//convert global data pointer to theint *idata = g_idata + blockIdx.x*blockDim.x;smem[tid]=idata[tid];__syncthreads();//in-place reduction in global memoryif(blockDim.x>=1024 && tid <512)smem[tid]+=smem[tid+512];__syncthreads();if(blockDim.x>=512 && tid <256)smem[tid]+=smem[tid+256];__syncthreads();if(blockDim.x>=256 && tid <128)smem[tid]+=smem[tid+128];__syncthreads();if(blockDim.x>=128 && tid <64)smem[tid]+=smem[tid+64];__syncthreads();//write result for this block to global memif(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] = smem[0];}

区别：

reduceGmem 使用全局内存（global memory）进行就地归约，这意味着它直接在全局内存中修改数据。
reduceSmem 则使用了共享内存（shared memory）smem 来存储每个线程从其负责的全局内存位置读取的数据，并在共享内存中进行归约操作

代码几乎相同，只是多了一个共享内存的声明，各线程将全局写入共享内存，以及后面的同步指令。

由这张图可以看出，使用共享内存减少全局内存的访问可以明显降低代码运行的时间。

接下来看一些指标

gld_transactions 和 gst_transactions 是CUDA编程中用于性能分析的指标，它们分别代表了全局内存加载（Global Memory Load）事务数和全局内存存储（Global Memory Store）事务数。

gld_transactions（全局内存加载事务数）：
- 此指标衡量了在执行CUDA程序期间，从全局内存加载数据的总事务数。全局内存是GPU上的一种内存类型，其访问速度相对较慢，但容量较大，用于存储程序中的大部分数据。
- 当CUDA程序需要从全局内存读取数据时，会产生全局内存加载事务。这些事务的数量会影响程序的性能，因为全局内存的访问延迟较高。
gst_transactions（全局内存存储事务数）：
- 与gld_transactions相对应，此指标衡量了在执行CUDA程序期间，向全局内存存储数据的总事务数。
- 当CUDA程序需要将数据写入全局内存时，会产生全局内存存储事务。同样地，这些事务的数量也会影响程序的性能。

在CUDA程序的性能优化中，通常会关注这些内存访问事务的数量，因为它们能够反映程序对全局内存的访问模式，从而帮助开发者识别潜在的性能瓶颈。例如，过多的全局内存访问事务可能意味着程序没有有效地利用GPU上的其他内存类型（如共享内存或纹理内存），或者程序的内存访问模式不够高效。通过优化内存访问模式，可以减少全局内存访问事务的数量，从而提高程序的性能。

针对上面两段代码，两个指标数值如下

可以看出使用共享内存的，比使用全局内存的高很多

使用展开的并行归约

代码如下

__global__ void reduceUnroll4Smem(int * g_idata,int * g_odata,unsigned int n)
{//set thread ID__shared__ int smem[DIM];unsigned int tid = threadIdx.x;unsigned int idx = blockDim.x*blockIdx.x*4+threadIdx.x;//boundary checkif (tid >= n) return;//convert global data pointer to theint tempSum=0;if(idx+3 * blockDim.x<=n){int a1=g_idata[idx];int a2=g_idata[idx+blockDim.x];int a3=g_idata[idx+2*blockDim.x];int a4=g_idata[idx+3*blockDim.x];tempSum=a1+a2+a3+a4;}smem[tid]=tempSum;__syncthreads();//in-place reduction in global memoryif(blockDim.x>=1024 && tid <512)smem[tid]+=smem[tid+512];__syncthreads();if(blockDim.x>=512 && tid <256)smem[tid]+=smem[tid+256];__syncthreads();if(blockDim.x>=256 && tid <128)smem[tid]+=smem[tid+128];__syncthreads();if(blockDim.x>=128 && tid <64)smem[tid]+=smem[tid+64];__syncthreads();//write result for this block to global memif(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] = smem[0];}

多了其他三块的求和

unsigned int idx = blockDim.x*blockIdx.x*4+threadIdx.x;
//boundary check
if (tid >= n) return;
//convert global data pointer to the
int tempSum=0;
if(idx+3 * blockDim.x<=n)
{int a1=g_idata[idx];int a2=g_idata[idx+blockDim.x];int a3=g_idata[idx+2*blockDim.x];int a4=g_idata[idx+3*blockDim.x];tempSum=a1+a2+a3+a4;
}

结果如下