矩阵转置算子优化-利用padding 解决 bank conflict

相关概念sm:硬件层面的计算核心,SM 是调度的基本单位。GPU 的硬件调度器会把你的 Block “分发”给空闲的 SM 去执行grid:是 CUDA 编程模型中最高层级的线程组织结构它代表了一次 Kernel 启动所创建的所有线程的集合。你可以把它理解为整个并行计算任务的“总指挥部”或“任务总表”。block:软件层面的任务分组,它是线程协作的基本单位。同一个 Block 里的线程可以互相通信通过共享内存 SMEM也可以通过 __syncthreads() 进行同步。**warp**GPU 实际调度线程的基本小组,假如一个block包含256个thread 那么这个block里就包含 256/328 个warplanelane 就是 warp 内部线程的编号。一个warp包含32个lanebank:bank 是 shared memory 的物理存储分区和并行访问单元, 硬件根据shared memory地址把访问请求分发到对应的 bankbank id的计算方式是(字节地址 ÷ 4) % 32举个例子 我们有共享内存int sdata[16][17] 线程访问 sdata[ty][tx]那么它的线性字节地址是字节地址 (ty * 17 tx) * 4代入 Bank 公式Bank ID ((ty * 17 tx) * 4 ÷ 4) % 32 (ty * 17 tx) % 32在这里线程的二维地址直接对应sdata共享内存的二维地址 线程是(tx,ty),这个线程访问的也是sdata[tx][ty]。lane通道 就是 Warp 里的“线程编号”范围是 0 到 31。假如在一个16x16的block中我们知道一个线程的二维坐标 tx, ty那么lane的编号 (ty*16tx) % 32层级关系Block 管线程组织Warp 管 32 个线程的执行组Lane 标识 Warp 中的线程Shared Memory 是 Block 内线程共享的数据空间Bank 是 Shared Memory 在硬件上的并行存储分区。层级关系图示GPU │ ├── Block0│ │ │ ├── Thread │ ├── Thread │ ├──... │ │ │ ├── Warp0│ │ ├── lane0│ │ ├── lane1│ │ ├──... │ │ └── lane31│ │ │ ├── Warp1│ └──... │ │ └── Shared Memory │ ├── Bank0│ ├── Bank1│ ├──... │ └── Bank31│ ├── Block1├── Block2└──...bank conflict是一个 warp 中多个 lane 在同一条 shared-memory 指令中访问同一个 bank 的不同地址。由于这些请求不能像访问不同 bank 那样并行服务硬件需要将访问拆分处理从而产生额外的访问周期使相关线程产生等待降低 shared-memory 的有效吞吐量Wave调度波次): 是 GPU 硬件调度器在分配和执行线程块时形成的一个“批次”或“轮次”它代表了一组被同时调度到流多处理器上执行的线程块集合。Wave 的核心机制调度的基本单位GPU 不会逐个调度线程块而是以“波”为单位批量调度。一个 Wave 包含多个 Block这些 Block 会被同时分配到不同的 SM 上执行。Wave 的大小受限于硬件资源一个 Wave 能容纳多少 Block取决于 SM 的数量、每个 SM 能驻留的 Block 数、寄存器数量、共享内存大小等。例如如果 GPU 有 80 个 SM每个 SM 最多能驻留 16 个 Block那么理论上一个 Wave 最多可以调度 80 × 16 1280 个 Block。Wave 是性能分析的关键指标如果一个 Kernel 需要多个 Wave 才能完成意味着部分 SM 在后几个 Wave 中会空闲等待导致资源利用率下降。减少 Wave 数量是提升 GPU 利用率的重要手段。为什么 Wave 会影响性能减少调度开销每个 Wave 的启动和同步都有开销。Wave 越少总开销越低。提高 SM 利用率如果所有 Block 能在一个 Wave 内完成所有 SM 都能满负荷运行如果需要多个 Wave后几个 Wave 中部分 SM 会空闲。配合循环展开和计算强度提升如你提到的“增加每个线程处理的元素个数”这不仅能提升计算强度还能减少总 Block 数量从而可能将整个 Kernel 的执行压缩到更少的 Wave 中完成。利用padding解决bank conflic代码示例templateintBLOCK_SZ__global__voidmat_transpose_kernel_v2(constfloat*idata,float*odata,intM,intN){constintbxblockIdx.x,byblockIdx.y;constinttxthreadIdx.x,tythreadIdx.y;__shared__floatsdata[BLOCK_SZ][BLOCK_SZ1];// paddingintxbx*BLOCK_SZtx;intyby*BLOCK_SZty;if(yMxN){sdata[ty][tx]idata[y*Nx];}__syncthreads();xby*BLOCK_SZtx;ybx*BLOCK_SZty;if(yNxM){odata[y*Mx]sdata[tx][ty];}}voidmat_transpose_v2(constfloat*idata,float*odata,intM,intN){constexprintBLOCK_SZ16;dim3block(BLOCK_SZ,BLOCK_SZ);dim3grid(Ceil(N,BLOCK_SZ),Ceil(M,BLOCK_SZ));mat_transpose_kernel_v2BLOCK_SZgrid,block(idata,odata,M,N);}在示例代码中 BLOCK_SZ16 表示 一个block里开启了16x16256个线程blockIdx.x blockIdx.y 代表当前线程所在的block在grid中的坐标 x是列坐标y是行坐标threadIdx.x threadIdx.y 代表当前线程在block中的坐标x是列坐标y是行坐标这段代码中可优化的空间考虑线程0,0和线程1,15 这两个线程分别对应的bank index 为0170 %32 011715 %32 0存在 bank conflic导致效率下降解决方案 将block_size 改为32修改之后的代码templateintBLOCK_SZ__global__voidmat_transpose_kernel_v2(constfloat*idata,float*odata,intM,intN){constintbxblockIdx.x,byblockIdx.y;constinttxthreadIdx.x,tythreadIdx.y;__shared__floatsdata[BLOCK_SZ][BLOCK_SZ1];// paddingintxbx*BLOCK_SZtx;intyby*BLOCK_SZty;if(yMxN){sdata[ty][tx]idata[y*Nx];}__syncthreads();xby*BLOCK_SZtx;ybx*BLOCK_SZty;if(yNxM){odata[y*Mx]sdata[tx][ty];}}voidmat_transpose_v2(constfloat*idata,float*odata,intM,intN){constexprintBLOCK_SZ32;// 仅修改这一行即可dim3block(BLOCK_SZ,BLOCK_SZ);dim3grid(Ceil(N,BLOCK_SZ),Ceil(M,BLOCK_SZ));mat_transpose_kernel_v2BLOCK_SZgrid,block(idata,odata,M,N);}