【CUDA】第4章全局内存——4.4核函数可达到的带宽（4对角转置）【补图】

CUDA C编程笔记 第四章 全局内存4.4 核函数可达到的带宽4.4.2.4 对角转置【为每个线程分配更独立的任务】 【图】【补图+说明】

待解决的问题：

第四章 全局内存 4.4 核函数可达到的带宽 4.4.2.4 对角转置【为每个线程分配更独立的任务】

前置条件场景：启用线程块的网格时，线程块会被分配给SM。每个块有唯一的标识符bid，可以按行优先的顺序标注：

int bid = blockIdx.y * gridDim.x + blockIdx.x;//块的标识bid

当启用核函数时，线程块的ID决定分配给SM的顺序，如果所有SM都被占用，剩余的线程块等待有SM空余再分配。但由于线程块完成的速度和顺序不确定，因此可能最初相连的bid也会变得不连续。

下图是笛卡尔坐标系（直角）和对角块坐标系下的块标识顺序。

【图】

对角块坐标系用于确定一维线程块的ID，但访问数据时仍用笛卡尔坐标系。

对角坐标————笛卡尔坐标（直角）的转换

(直角坐标)block_x = (blockIdx.x + blockIdx.y) % gridDim.x;【blockIdx.x对角坐标】 (直角坐标)block_y = blockIdx.x;【blockIdx.y对角坐标】

核函数起始部分：对角坐标到直角坐标的映射计算+直角坐标计算线程索引ix、iy 借助合并读取+交叉写入——>实现转置

//3.对角转置————基于行 //①对角坐标系转直角坐标系②直角坐标系算线程索引ix、iy③转置 __global__ void transposeDiagonalRow(float *out, float *in, const int nx, const int ny){ unsigned int blk_y = blockIdx.x;//blk_y直角坐标系，blockIdx.x对角坐标系 unsigned int blk_x = (blockIdx.x+blockIdx.y) % gridDim.x;//blk_x直角坐标系，blockIdx.y对角坐标系 unsigned int ix = blockDim.x * blk_x + threadIdx.x;//用直角坐标算线程索引 unsigned int iy = blockDim.y * blk_y + threadIdx.y; if(ix < nx && iy < ny){ out[ix*ny + iy] = in[iy*nx + ix]; } } //3.对角转置————基于列 //在基于行的基础上对换in和out的下标 __global__ void transposeDiagonalRow(float *out, float *in, const int nx, const int ny){ unsigned int blk_y = blockIdx.x;//blk_y直角坐标系，blockIdx.x对角坐标系 unsigned int blk_x = (blockIdx.x+blockIdx.y) % gridDim.x;//blk_x直角坐标系，blockIdx.y对角坐标系 unsigned int ix = blockDim.x * blk_x + threadIdx.x;//用直角坐标算线程索引 unsigned int iy = blockDim.y * blk_y + threadIdx.y; if(ix < nx && iy < ny){ out[iy*nx + ix] = in[ix*ny + iy]; } } case 6:3.对角转置----基于行 kernel = &transposeDiagonalRow; kernelName = "DiagonalRow "; break; case 7:3.对角转置----基于列 kernel = &transposeDiagonalCol; kernelName = "DiagonalCol "; break;

输出结果如下：

~/cudaC/unit4$ ./4-6.1transposeNsys 6 对角转置-基于行 ./4-6.1transposeNsys starting transpose at device 0: NVIDIA GeForce RTX 3090 with matrix nx 2048 ny 2048 with kernel 6 warmup elapsed 0.000563 sec DiagonalRow elapsed 0.000075 sec <<< grid (128,128) block (16,16)>>> effective bandwidth 448.208557 GB ~/cudaC/unit4$ ./4-6.1transposeNsys 7 对角转置-基于列 ./4-6.1transposeNsys starting transpose at device 0: NVIDIA GeForce RTX 3090 with matrix nx 2048 ny 2048 with kernel 7 warmup elapsed 0.000561 sec DiagonalCol elapsed 0.000064 sec <<< grid (128,128) block (16,16)>>> effective bandwidth 525.139893 GB

查询可得，理论峰值带宽为936 GB/s 基于行的对角是理论峰值的48%↑ 基于列的对角是理论峰值的56%↓

结果：对角使得基于行性能提高↑，但使得基于列下降，基于列还是直角坐标性能好。

基于行性能提升的原因：DRAM的并行访问 DRAM分区完成发送给全局内存的请求，设备内存中连续的256字节区域分配到连续的分区。使用直角坐标把线程块映射到——>数据块时，全局内存访问无法均匀分配到整个DRAM分区，发生“分区冲突”：内存请求在部分分区内排队等待，但另一部分分区一直空闲未被调用。 对角坐标映射造成了线程块——>数据块的非线性映射，交叉访问不太可能落到一个独立的分区，这导致速度提升。

最佳性能，一般是（所有活跃warp并发访问的）全局内存被均匀地划分。

【补图+说明】