三亩地 三亩地SAN MU DI · CODE DIARY
ARTICLE DETAIL

日记详情

真实记录编程学习的某一天,欢迎挑你感兴趣的翻一翻。

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

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

相关概念

sm:硬件层面的计算核心,SM 是调度的基本单位。GPU 的硬件调度器会把你的 Block “分发”给空闲的 SM 去执行
grid:是 CUDA 编程模型中最高层级的线程组织结构,它代表了一次 Kernel 启动所创建的所有线程的集合。你可以把它理解为整个并行计算任务的“总指挥部”或“任务总表”。
block:软件层面的任务分组,它是线程协作的基本单位。同一个 Block 里的线程可以互相通信(通过共享内存 SMEM),也可以通过 __syncthreads() 进行同步。
**warp:**GPU 实际调度线程的基本小组,假如一个block包含256个thread, 那么这个block里就包含 256/32=8 个warp
lane:lane 就是 warp 内部线程的编号。一个warp包含32个lane
bank:bank 是 shared memory 的物理存储分区和并行访问单元, 硬件根据shared memory地址把访问请求分发到对应的 bank,bank 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*16+tx) % 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

代码示例:

template<intBLOCK_SZ>__global__voidmat_transpose_kernel_v2(constfloat*idata,float*odata,intM,intN){constintbx=blockIdx.x,by=blockIdx.y;constinttx=threadIdx.x,ty=threadIdx.y;__shared__floatsdata[BLOCK_SZ][BLOCK_SZ+1];// paddingintx=bx*BLOCK_SZ+tx;inty=by*BLOCK_SZ+ty;if(y<M&&x<N){sdata[ty][tx]=idata[y*N+x];}__syncthreads();x=by*BLOCK_SZ+tx;y=bx*BLOCK_SZ+ty;if(y<N&&x<M){odata[y*M+x]=sdata[tx][ty];}}voidmat_transpose_v2(constfloat*idata,float*odata,intM,intN){constexprintBLOCK_SZ=16;dim3block(BLOCK_SZ,BLOCK_SZ);dim3grid(Ceil(N,BLOCK_SZ),Ceil(M,BLOCK_SZ));mat_transpose_kernel_v2<BLOCK_SZ><<<grid,block>>>(idata,odata,M,N);}

在示例代码中 BLOCK_SZ=16 表示 一个block里开启了16x16=256个线程;
blockIdx.x, blockIdx.y 代表当前线程所在的block在grid中的坐标, x是列坐标,y是行坐标;
threadIdx.x threadIdx.y 代表当前线程在block中的坐标,x是列坐标,y是行坐标;
这段代码中可优化的空间:
考虑线程(0,0)和线程(1,15) 这两个线程分别对应的bank index 为
017+0 %32 = 0
1
17+15 %32 = 0
存在 bank conflic导致效率下降;解决方案 将block_size 改为32

修改之后的代码

template<intBLOCK_SZ>__global__voidmat_transpose_kernel_v2(constfloat*idata,float*odata,intM,intN){constintbx=blockIdx.x,by=blockIdx.y;constinttx=threadIdx.x,ty=threadIdx.y;__shared__floatsdata[BLOCK_SZ][BLOCK_SZ+1];// paddingintx=bx*BLOCK_SZ+tx;inty=by*BLOCK_SZ+ty;if(y<M&&x<N){sdata[ty][tx]=idata[y*N+x];}__syncthreads();x=by*BLOCK_SZ+tx;y=bx*BLOCK_SZ+ty;if(y<N&&x<M){odata[y*M+x]=sdata[tx][ty];}}voidmat_transpose_v2(constfloat*idata,float*odata,intM,intN){constexprintBLOCK_SZ=32;// 仅修改这一行即可dim3block(BLOCK_SZ,BLOCK_SZ);dim3grid(Ceil(N,BLOCK_SZ),Ceil(M,BLOCK_SZ));mat_transpose_kernel_v2<BLOCK_SZ><<<grid,block>>>(idata,odata,M,N);}
← 返回列表