CUDA C编程笔记
- 第五章 共享内存和常量内存
- 5.3 减少全局内存访问
- 5.3.2 使用展开的并行规约
- 思路
- reduceSmemUnroll4（共享内存）具体代码：
- 运行结果
- 意外发现书上全局加载事务和全局存储事务和ncu中这两个值相同
- 5.3.3 动态共享内存的并行规约
- reduceSmemUnroll4Dyn（共享内存）具体代码：
- 5.3.4 有效带宽
待解决的问题：意外发现书上全局加载事务和全局存储事务和ncu中这两个值相同，是否有直接相关？？？
ncu中带宽的查看方式
第五章 共享内存和常量内存
5.3 减少全局内存访问
使用共享内存的主要原因之一是要缓存片上的数据，来减少核函数中全局内存访问的次数。
第三章介绍了用全局内存的并行规约核函数，并解释了下面2个问题：
①如何重新安排数据访问模式来避免线程束分化
②如何展开循环来保证有足够的操作使指令和内存带宽饱和
本节重新使用并行规约核函数，但是这里用共享内存作为缓存来减少全局内存的访问。【并行规约+共享内存】
5.3.2 使用展开的并行规约
前面的核函数用一个线程块处理一个数据块。继续优化用第三章的思想，一次运行多个IO操作，展开线程块来提高核函数性能。
这里展开了4个线程块，即每个线程处理4个数据块的数据。
这样做的优势是：
①提高全局内存的吞吐量，因为每个线程进行了更多的并行IO。
②全局内存存储事务减少了1/4
③整体内核性能提升
思路
先重新计算全局输入数据的偏移值。
//全局索引，一次处理4个输入数据块
unsigned int idx = blockIdx.x * blockDim.x * 4 + threadIdx.x;//这里乘4
再一次性处理4个元素，每个线程读取4个数据，把这个4个数据的和放到局部变量tmpSum中，用tmpSum来初始化共享内存，而非从全局内存初始化共享内存。
//边界条件检查
if(idx < n)//在范围内的相邻块大小的元素都加起来，最多可以一次处理4个块
{
int a1, a2, a3, a4;
a1 = a2 = a3 = a4 = 0;
a1 = g_idata[idx];
if(idx + blockDim.x < n) a2 = g_idata[idx + blockDim.x];
if(idx + 2 * blockDim.x < n) a3 = g_idata[idx + 2 * blockDim.x];
if(idx + 3 * blockDim.x < n) a4 = g_idata[idx + 3 * blockDim.x];
tmpSum = a1 + a2 + a3 + a4;
}
reduceSmemUnroll4（共享内存）具体代码：
//reduceSmemUnroll4
__global__ void reduceSmemUnroll(int *g_idata, int *g_odata, unsigned int n){
//静态共享数组
__shared__ int smem[DIM];
//设置线程ID
unsigned int tid = threadIdx.x;
//全局索引，一次处理4个输入数据块
unsigned int idx = blockIdx.x * blockDim.x * 4 + threadIdx.x;//这里乘4
//展开4个块
int tmpSum = 0;//【】
//边界条件检查
if(idx < n)//在范围内的相邻块大小的元素都加起来，最多可以一次处理4个块
{
int a1, a2, a3, a4;
a1 = a2 = a3 = a4 = 0;
a1 = g_idata[idx];
if(idx + blockDim.x < n) a2 = g_idata[idx + blockDim.x];
if(idx + 2 * blockDim.x < n) a3 = g_idata[idx + 2 * blockDim.x];
if(idx + 3 * blockDim.x < n) a4 = g_idata[idx + 3 * blockDim.x];
tmpSum = a1 + a2 + a3 + a4;
}
smem[tid] = tmpSum;
__syncthreads();
//在共享内存中就地规约
if(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();
//展开warp
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] = smem[0];
}
对应的主函数调用核函数也要修改，网格除4。
这里只能给grid.x/4，不能给block/4。
如果block/4，假设原来block大小为256，调用时block/4=64，blockDim.x=64，共享内存仍分配256个空间，只有前64个有值，后面的都是未定义的有问题的值。并且归约也会崩溃，索引也有问题。
//3、reduceSmemUnroll4
cudaMemcpy(d_idata, h_idata, bytes, cudaMemcpyHostToDevice);
reduceSmemUnroll<<<grid.x / 4, block>>>(d_idata, d_odata, size);//这里要除4，因为一个线程块处理四个数据块，需要的线程块减为原来的1/4
cudaMemcpy(h_odata, d_odata, grid.x / 4 * sizeof(int), cudaMemcpyDeviceToHost);
gpu_sum = 0;
for(int i = 0; i < grid.x / 4; i++) gpu_sum += h_odata[i];
printf("reduceSmemUnroll4: %d <<<grid %d block %d>>>\n", gpu_sum, grid.x / 4,
block.x);
运行结果
[6/8] Executing 'cuda_gpu_kern_sum' stats report
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- --------- --------- -------- -------- ----------- --------------------------------------------
50.4 238,789 1 238,789.0 238,789.0 238,789 238,789 0.0 reduceGmem(int *, int *, unsigned int)
32.5 154,051 1 154,051.0 154,051.0 154,051 154,051 0.0 reduceSmem(int *, int *, unsigned int)
17.2 81,377 1 81,377.0 81,377.0 81,377 81,377 0.0 reduceSmemUnroll(int *, int *, unsigned int)
意外发现书上全局加载事务和全局存储事务和ncu中这两个值相同
存储事务：与reduceSmem相比，reduceSmemUnroll4存储事务数量减少为1/4，加载事务数量不变。
全局内存吞吐量：
加载吞吐量增加，因为大量同时加载请求。
存储吞吐量下降，较少的存储请求让总线饱和。
5.3.3 动态共享内存的并行规约
reduceSmemUnroll4中用动态共享内存代替静态共享内存
extern __shared__ int smem[];
reduceSmemUnroll4Dyn（共享内存）具体代码：
//动态共享内存+reduceSmemUnroll4
__global__ void reduceSmemUnrollDyn(int *g_idata, int *g_odata, unsigned int n){
extern __shared__ int smem[];//设置动态共享内存，其他的和reduceSmemUnroll4一样
//设置线程id
unsigned int tid = threadIdx.x;
unsigned int idx = blockIdx.x * blockDim.x * 4 + threadIdx.x;
//展开4个块
int tmpSum = 0;
if(idx < n){
int a1, a2, a3, a4;
a1 = a2 = a3 = a4 = 0;
a1 = g_idata[idx];
if(idx + blockDim.x < n) a2 = g_idata[idx + blockDim.x];
if(idx + 2 * blockDim.x < n) a3 = g_idata[idx + blockDim.x * 2];
if(idx + 3 * blockDim.x < n) a4 = g_idata[idx + blockDim.x * 3];
tmpSum = a1 + a2 + a3 + a4;
}
smem[tid] = tmpSum;
__syncthreads();
//在全局内存就地规约
if(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();
//展开warp
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[threadIdx.x] = smem[0];
}
对应main函数调用核函数的时候也要修改。
//4、reduceSmemUnroll4Dyn
cudaMemcpy(d_idata, h_idata, bytes, cudaMemcpyHostToDevice);
reduceSmemUnrollDyn<<<grid.x / 4, block, DIM * sizeof(int)>>>(d_idata, d_odata, size);//1、指定待动态分配的共享内存数量 2、除4，因为一个线程块处理四个数据块，需要的线程块减为原来的1/4
cudaMemcpy(h_odata, d_odata, grid.x / 4 * sizeof(int), cudaMemcpyDeviceToHost);
gpu_sum = 0;
for(int i = 0; i < grid.x / 4; i++) gpu_sum += h_odata[i];
printf("reduceSmemUnroll4Dyn: %d <<<grid %d block %d>>>\n", gpu_sum, grid.x / 4,
block.x);
运行效果：动态共享内存和静态共享内存运行时间差不多。
[6/8] Executing 'cuda_gpu_kern_sum' stats report
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- --------- --------- -------- -------- ----------- -----------------------------------------------
43.7 263,970 1 263,970.0 263,970.0 263,970 263,970 0.0 reduceGmem(int *, int *, unsigned int)
27.5 166,400 1 166,400.0 166,400.0 166,400 166,400 0.0 reduceSmem(int *, int *, unsigned int)
14.4 87,169 1 87,169.0 87,169.0 87,169 87,169 0.0 reduceSmemUnrollDyn(int *, int *, unsigned int)
14.4 86,720 1 86,720.0 86,720.0 86,720 86,720 0.0 reduceSmemUnroll(int *, int *, unsigned int)
5.3.4 有效带宽
规约核函数受到内存带宽的限制，因此用有效带宽来评估他们的性能指标。
有效带宽：核函数的完整执行时间内IO的数量；对于内存约束的应用程序，有效带宽是估算实际带宽利用率的好指标。
