当前位置：首页 > news >正文

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

news 2025/11/7 17:47:54

1.介绍

对 2048 * 512 矩阵转置，使用NCU进行性能分析，并进行性能优化

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% 和内存吞吐量 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%;存储吞吐量%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%；存储吞吐量 32.9%，提升 26%；

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

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

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

(8, 8) 效率最高

5. 二维Block + 线程Tile

实现了三个版本：

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 unrollfor(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 unrollfor(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 unrollfor(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 unrollfor(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 unrollfor(int i=0;i<4;i++)resultTemp[i] = make_float4(srcTemp[0][i],srcTemp[1][i],srcTemp[2][i],srcTemp[3][i]);#pragma unrollfor(int i=0;i<4;i++)CAST_FLOAT4(out+((gidx<<2)+i)*M + (gidy<<2))[0]  = resultTemp[i];}
}