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

待解决的问题：

第四章 全局内存

4.4 核函数可达到的带宽

4.4.2.4 对角转置【让DRAM访问更均匀，提高性能】

前置条件场景：启用线程块的网格时，线程块会被分配给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并发访问的）全局内存被均匀地划分。

如下图所示，假设通过两个分区访问全局内存，每个分区256字节，使用32*32的线程块启动kernel。如果每个数据块128字节，需要两个分区为0、1、2、3号线程块加载数据，但实际上只用了第一个分区块，造成了分区冲突。

但如果使用对角坐标的话，使用了两个分区来给0、1、2、3号线程块，这样加载和存储请求在两个分区间均匀分配，性能更好。