高性能计算-CUDA性能优化-transpose

1.介绍

对 2048 * 512 矩阵转置，使用NCU进行性能分析，并进行性能优化。测试环境 CUDA 12.8,显卡 5070。

2. Native: 二维 Block

二维block,一个线程处理一个元素

点击查看代码

//native：二维block,一个线程处理一个元素
//矩阵 M * N
template<uint32_t M,uint32_t N>
__global__ void kernel_transpose_native(float *arr,float *out)
{
    uint32_t gidx = blockIdx.x * blockDim.x + threadIdx.x;
    uint32_t gidy = blockIdx.y * blockDim.y + threadIdx.y;
    if(gidx<N && gidy<M)
        out[gidx*M+gidy] = arr[gidy*N+gidx];
}

blockSize: 64 * 8

计算吞吐量 12.8% 和存储 pipeline 吞吐量 49.5%；

global 从 L1 数据读入 4sector/request,一个warp 请求 128B,即 4sector,正常；

问题：global 写入数据 32sector/request,一个warp 有 32个 sector,逻辑写入数据量/物理写入数据量=13107232B/104857632B,写入效率只有 12.5%。

优化方案：考虑增加写回数据的访存合并度，减少 sector 数量。

目前未能解答：L2 与 DRAM 写回数据为什么没有数据流动？？

3. 二维 Block + 共享内存

介绍：考虑到一次内存事务请求 128B，使用二维block(32,32),正好一个blcok读取的数据可以增加缓存命中，数据放在共享内存

点击查看代码

//blocksize: BM * BN
template<uint32_t M,uint32_t N,uint32_t BM,uint32_t BN>
__global__ void kernel_transpose_SM(float *arr,float *out)
{
    uint32_t tidx = threadIdx.x;
    uint32_t tidy = threadIdx.y;
    uint32_t gidx = blockIdx.x * blockDim.x + threadIdx.x;
    uint32_t gidy = blockIdx.y * blockDim.y + threadIdx.y;
    __shared__ float blockShared[BM][BN];
    if(gidy<M && gidx<N)
    {
        blockShared[tidy][tidx] = arr[gidy*N+gidx];
        __syncthreads();
        out[gidx*M+gidy] = blockShared[tidy][tidx];
    }
}

备注: 设置 native kernel 为baseline

计算单元吞吐量 14.3%，提升 11.62%;存储 pipeline 吞吐量 30.2%，下降39%；SM active Cycle（SM至少有一个 active warp 的时钟周期数量）提升 60.7%，其余指标都大幅降低。

分析：native blockSize 为 64 * 8=512，tile blockSize 为 32 * 32=1024 过大会导致 SM active Cycle 指标提升。

问题1：存储吞吐指标下降。

问题2：耗时增加 63%。

共享内存数据读进来只使用一次，没有使用共享内存的必要；并且会带来 bank conflict 问题,设置 block 8 * 8的 bank conflict情况严重于32 * 32的 baseline, 如下图:
blockSize: （8, 8），使用共享内存作为 baseLine与不使用共享内存比较：

不使用共享内存耗时减少24%，计算单元吞吐量 24.4%，下降30%；存储 pipeline 吞吐量 32.9%，提升 26%；

共享内存的 bank conflict 负面影响消除；

4. 二维 Block 不使用共享内存的对比测试

blockSize: （8, 8）对比（32, 32）的计算单元吞吐量 35.4%，提升148.8%；存储 pipeline 吞吐量 34.58%，提升14.47%;性能大幅增加，应使用较小的 blocksize，增加并行度；如下图：
blockSize: （8, 8）的数据写回合并度大大增加，sectors 减少了75%，如下图：
不同 blockSize的耗时对比:

(8, 8) 效率最高

5. 二维Block + 线程Tile_float4

warp shape 对读写单个物理 sector 利用率的影响：

warp (8, 4) 排布：读可以每个sector合并访存；写每个sector 有一半数据是多余的，每个sector 无法合并访存;

warp (4, 8) 排布：写可以每个sector合并访存；读每个sector 有一半数据是多余的，每个sector 无法合并访存。

为了实现读写单个物理 sector 的合并访存，单个warp应在 X Y 两个方向上都有 8（32B）的倍数个待处理数据，
实现了三个版本：

tile + float4, 源数据直接写入目标地址: kernel_transpose_tile_FL4;
tile + float4, 使用寄存器中转，每个线程获取部分数据就写入: kernel_transpose_tile_FL4_2;
tile + float4，, 使用寄存器中转，每个线程获取全部数据后再写入，kernel_transpose_tile_FL4_3:

详见代码

点击查看代码

//强转优先级大于 []
#define CAST_FLOAT4(pointer) reinterpret_cast<float4*>(pointer)

//2维block + tile(float4 单个线程处理 4*4 的数据)
//源数据直接写入目标地址
//blocksize: BM * BN
template<uint32_t M,uint32_t N,uint32_t BM,uint32_t BN>
__global__ void kernel_transpose_tile_FL4(float *arr,float *out)
{
    uint32_t gidx = blockIdx.x * blockDim.x + threadIdx.x;
    uint32_t gidy = blockIdx.y * blockDim.y + threadIdx.y;
    if(gidy<<2 <M && gidx<<2 <N)
    {
        //取源数据地址
        float4 *srcTemp[4];
        //循环独立，可指令级并行
        #pragma unroll
        for(int i=0;i<4;i++)
            srcTemp[i] = CAST_FLOAT4(arr+((gidy<<2) + i)*N + (gidx<<2));
        
        //重组数据
        CAST_FLOAT4(out+((gidx<<2))*M + (gidy<<2))[0]  = make_float4(srcTemp[0]->x,srcTemp[1]->x,srcTemp[2]->x,srcTemp[3]->x);
        CAST_FLOAT4(out+((gidx<<2)+1)*M + (gidy<<2))[0]  = make_float4(srcTemp[0]->y,srcTemp[1]->y,srcTemp[2]->y,srcTemp[3]->y);
        CAST_FLOAT4(out+((gidx<<2)+2)*M + (gidy<<2))[0]  = make_float4(srcTemp[0]->z,srcTemp[1]->z,srcTemp[2]->z,srcTemp[3]->z);
        CAST_FLOAT4(out+((gidx<<2)+3)*M + (gidy<<2))[0]  = make_float4(srcTemp[0]->w,srcTemp[1]->w,srcTemp[2]->w,srcTemp[3]->w);
    }
}

//二维block + tile(float4 单个线程处理 4*4 的数据)
//使用寄存器中转数据的写法
template<uint32_t M,uint32_t N,uint32_t BM,uint32_t BN>
__global__ void kernel_transpose_tile_FL4_2(float *arr,float *out)
{
    uint32_t gidx = blockIdx.x * blockDim.x + threadIdx.x;
    uint32_t gidy = blockIdx.y * blockDim.y + threadIdx.y;
    if(gidy<<2 <M && gidx<<2 <N)
    {
        //取源数据
        float srcTemp[4][4];
        //循环独立，可指令级并行
        #pragma unroll
        for(int i=0;i<4;i++)
            CAST_FLOAT4(&srcTemp[i])[0] = CAST_FLOAT4(arr+((gidy<<2) + i)*N + (gidx<<2))[0];
        
        //重组数据
        float4 resultTemp[4];
        #pragma unroll
        for(int i=0;i<4;i++)
        {
            resultTemp[i] = make_float4(srcTemp[0][i],srcTemp[1][i],srcTemp[2][i],srcTemp[3][i]);
            CAST_FLOAT4(out+((gidx<<2)+i)*M + (gidy<<2))[0]  = resultTemp[i];
        }
    }
}

//二维block + tile(float4 单个线程处理 4*4 的数据)
//使用寄存器中转数据的写法,拆分循环
template<uint32_t M,uint32_t N,uint32_t BM,uint32_t BN>
__global__ void kernel_transpose_tile_FL4_3(float *arr,float *out)
{
    uint32_t gidx = blockIdx.x * blockDim.x + threadIdx.x;
    uint32_t gidy = blockIdx.y * blockDim.y + threadIdx.y;
    if(gidy<<2 <M && gidx<<2 <N)
    {
        //取源数据
        float srcTemp[4][4];
        //循环独立，可指令级并行
        #pragma unroll
        for(int i=0;i<4;i++)
            CAST_FLOAT4(&srcTemp[i])[0] = CAST_FLOAT4(arr+((gidy<<2) + i)*N + (gidx<<2))[0];
        
        //重组数据
        float4 resultTemp[4];
        #pragma unroll
        for(int i=0;i<4;i++)
            resultTemp[i] = make_float4(srcTemp[0][i],srcTemp[1][i],srcTemp[2][i],srcTemp[3][i]);

        #pragma unroll
        for(int i=0;i<4;i++)
            CAST_FLOAT4(out+((gidx<<2)+i)*M + (gidy<<2))[0]  = resultTemp[i];
    }
}

每个kernel 执行4轮，blockSize(8,8)，NC 总体数据如下：

第二种实现，每个线程获取部分数据就写入的计算单元和内存吞吐量较高。

block 处理总数据量不变: 32 * 32，kernel_transpose_tile_FL4_2 计算单元吞吐量为25.5%，提升183%；内存吞吐量为52.7，提升39%。

每个kernel 执行4轮，blockSize(32,32)，NC 总体数据如下：

仍是第二种实现的计算单元和内存吞吐量稍高。

对比kernel_transpose_tile_FL4_2，一个 block 处理数据从 1024 增加到 16384,计算单元吞吐量下降 34%,存储吞吐量下降 22%，带来了性能下降。

测试 block(8,4),结果强制分配为 block(8,8) [后经查阅，此处是默认值64，也可以修改blockSize的大小],blockSize 最小为64，结果如下图：
仍是 block(8, 8)的耗时低 8% 。

6. 二维 Block + 线程tile_float2

以上线程tile处理 4 * 4 的数据，单个线程处理数据量较多，线程数较少，导致访存延迟无法隐藏，考虑减少每个线程处理数据量；
实现：一个线程处理（2， 2）的数据；

点击查看代码

//二维block + tile(float2 单个线程处理 2*2 的数据)
//使用寄存器中转数据的写法
template<uint32_t M,uint32_t N,uint32_t BM,uint32_t BN>
__global__ void kernel_transpose_tile_FL2(float *arr,float *out)
{
    uint32_t gidx = blockIdx.x * blockDim.x + threadIdx.x;
    uint32_t gidy = blockIdx.y * blockDim.y + threadIdx.y;
    if(gidy<<1 <M && gidx<<1 <N)
    {
        //取源数据
        float srcTemp[2][2];
        //循环独立，可指令级并行
        #pragma unroll
        for(int i=0;i<2;i++)
            CAST_FLOAT2(&srcTemp[i])[0] = CAST_FLOAT2(arr+((gidy<<1) + i)*N + (gidx<<1))[0];
        
        //重组数据
        float2 resultTemp[2];
        #pragma unroll
        for(int i=0;i<2;i++)
        {
            resultTemp[i] = make_float2(srcTemp[0][i],srcTemp[1][i]);
            CAST_FLOAT2(out+((gidx<<1)+i)*M + (gidy<<1))[0]  = resultTemp[i];
        }
    }
}

与float4 版本比较相同 blockSize(8, 8)，block 处理不同数据量的对比，如下图：

float2 比 float4 计算单元吞吐量提升 33%, 存储 pipeline 吞吐量基本持平;DRAM 吞吐降低 35.8%;float2 耗时增加 27% ;

与float4 版本比较 block 处理相同数据量1024,float4 blockSize(8,8),float2 blockSize(16,16)，如下图：

单个block 处理相同数据量 float2 比 float4 耗时增加 44%;float2 计算单元和存储 pipeline 吞吐量提升；DRAM 吞吐量下降 43%。

分析 float4 提升的主要原因是 DRAM 的吞吐量增加；

float2 指令停滞指标有改善，最大影响因子 Stall Long Scoreboard 从 16 cycle 降低为 9.5 cycle。

6. 总结

使用较小的 blockSize 可以增加并行度；
使用不同的 blockSize 对读写访存的 sector 利用率不一样，应该做计算取合适的大小。
其他的优化方法：

单个线程处理 1row 2col;

posted @ 2025-08-18 18:08 安洛8 阅读(52) 评论(0) 收藏举报

刷新页面返回顶部

anluo8

高性能计算-CUDA性能优化-transpose

1.介绍

2. Native: 二维 Block

3. 二维 Block + 共享内存

4. 二维 Block 不使用共享内存的对比测试

5. 二维Block + 线程Tile_float4

6. 二维 Block + 线程tile_float2

6. 总结

公告