【CUDA 】第4章 全局内存——4.4 核函数可达到的带宽（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%

数据对比：