CUDA C编程笔记
- 第四章 全局内存
- 4.4 核函数可达到的带宽
- 4.4.2.3 展开转置【为每个线程分配更独立的任务】
待解决的问题:
第四章 全局内存
4.4 核函数可达到的带宽
4.4.2.3 展开转置【为每个线程分配更独立的任务】
展开:提高转置内存带宽的利用率,来隐藏延迟
展开的目的:为每个线程分配更独立的任务,最大化当前内存请求
- 基于行+展开因子为4
展开的意思:让一个线程处理4个元素,而不是1个元素;由于展开方向沿着行(x的方向),每个线程处理连续的4个元素
为什么输出矩阵下标里面要乘ny
关键点:
①unsigned int ix = blockIdx.x * blockDim.x4 + threadIdx.x;//这里乘4
每个线程要处理4个元素的数据,每个block也对应处理4个block的范围
②if(ix+3blockDim.x < nx && iy < ny)
每个线程将处理4个元素,即当前ix,以及ix+blockDim.x,ix+2blockDim.x,ix+3blockDim.x的位置
③假设blockDim.x是B,那么每个block在x方向上的处理范围是blockIdx.x * B4到(blockIdx.x+1)B4
每个线程在block内的threadIdx.x是0到B-1
所以每个线程对应的全局ix是blockIdx.xB*4 + threadIdx.x
处理ix、ix+B、ix+2B、ix+3B这四个位置的数据
//2.展开转置————基于行(展开因子为4)
__global__ void transposeUnroll4Row(float *out, float *in, const int nx, const int ny){
unsigned int ix = blockIdx.x * blockDim.x*4 + threadIdx.x;//这里乘4
unsigned int iy = blockIdx.y * blockDim.y + threadIdx.y;
//新数组索引,分别用于输入行访问和输出列访问
unsigned int ti = iy*nx + ix;//in输入矩阵行访问
unsigned int to = ix*ny + iy;//out输出矩阵列访问
if(ix+3*blockDim.x < nx && iy < ny){
out[to] = in[ti];
out[to + ny*blockDim.x] = in[ti+blockDim.x];//转置后相邻元素距离为ny,因此输出矩阵out要乘ny
out[to + ny*2*blockDim.x] = in[ti+2*blockDim.x];
out[to + ny*3*blockDim.x] = in[ti+3*blockDim.x];
}
}
//2.展开转置————基于列(展开因子为4)
//只修改交换的下标即可
__global__ void transposeUnroll4Col(float *out, float *in, const int nx, const int ny){
unsigned int ix = blockIdx.x * blockDim.x*4 + threadIdx.x;//还是x这里乘4
unsigned int iy = blockIdx.y * blockDim.y + threadIdx.y;
//新数组索引,分别用于输入行访问和输出列访问
unsigned int ti = iy*nx + ix;//in输入矩阵行访问
unsigned int to = ix*ny + iy;//out输出矩阵列访问
if(ix+3*blockDim.x < nx && iy < ny){
out[ti] = in[to];
out[ti+blockDim.x] = in[to + ny*blockDim.x];//转置后相邻元素距离为ny,因此输出矩阵out要乘ny
out[ti+2*blockDim.x] = in[to + ny*2*blockDim.x];
out[ti+3*blockDim.x] = in[to + ny*3*blockDim.x];
}
}
输出结果如下:基于行的展开4和基于列的展开4
~/cudaC/unit4$ ./4-6.1transposeNsys 4
./4-6.1transposeNsys starting transpose at device 0: NVIDIA GeForce RTX 3090
with matrix nx 2048 ny 2048 with kernel 4
warmup elapsed 0.000560 sec
Unroll4Row elapsed 0.000077 sec <<< grid (32,128) block (16,16)>>> effective bandwidth 435.719788 GB
~/cudaC/unit4$ ./4-6.1transposeNsys 5
./4-6.1transposeNsys starting transpose at device 0: NVIDIA GeForce RTX 3090
with matrix nx 2048 ny 2048 with kernel 5
warmup elapsed 0.000564 sec
Unroll4Row elapsed 0.000059 sec <<< grid (32,128) block (16,16)>>> effective bandwidth 569.787415 GB
查询可得,理论峰值带宽为936 GB/s
基于行的4级展开是理论峰值的47%
基于列的4级展开是理论峰值的61%
数据对比: