搞过 CUDA 的人迟早会撞上矩阵转置这道题。很多教程会给你一个能跑出正确结果的 kernel但一测带宽就露馅Nsight 里满屏共享内存 bank conflict。这个题目把几个高频考点都串起来了共享内存、bank conflict、行主序/列主序切换、循环展开还点明了一个容易被忽略的约束——转置前后全局内存里的数据都是行主序中间那层共享内存才按列主序来用。我会把这个方案从头到尾讲透。先从 bank 冲突到底怎么算说起再给完整代码然后讲循环展开最后写一写实际排查时的坑。适合下面几类人正在学 CUDA 优化的初学者准备 GPU 编程面试的人以及做图像处理、矩阵库、神经网络算子时被转置性能卡住的工程师。1. 先把题目拆干净1.1 为什么转置会慢到让人怀疑人生矩阵转置本身计算量几乎是零就是个数据搬移问题。真正难的是访存模式。假设矩阵按行主序存在全局内存里也就是a[row][col]的地鼠是row * width col同一行内相邻元素地址连续。朴素转置核函数通常是每个线程读一个元素再写到一个矩阵转置后的位置。比如读的时候连续线程访问同一行天然合并但写的时候连续线程访问的是输出矩阵里的同一列不同行地址间隔一个width这会让全局内存的写变成一堆分散的小段访问。反过来也一样你先满足写合并读就会变成不合并。全局内存合并访问这件事单个核函数里两头都占靠朴素写法做不到所以才需要共享内存做中间缓冲。1.2 行主序、列主序和共享内存的关系题目里反复强调“转置前后都是行主序中间共享内存中是列主序”这个不要理解偏了。C/C 里的二维数组天然是行主序。共享内存本身没有二维只有一维地址所谓行主序、列主序只是你给地址起的别名。矩阵转置的完整过程可以看成先从全局内存按行主序读进一个数据块放到共享内存再从共享内存按转置后的位置取出来按行主序写回全局内存。题目里的“中间列主序”实际指的是共享内存这一层的索引方式要跟前后两边的行主序索引错开用列优先的下标来读或写。很多资料会把这一步叫“在共享内存里把矩阵转过来”。它其实不是把字节重新排成某种神秘格式而是把二维下标里的行列关系换一下。最终落到硬件上关心的只是一个线性地址除以 4 再模 32 落在哪个 bank 上。1.3 看这份笔记前需要先有的基础这篇不是从零教 CUDA 语法的入门文。我假设你已经知道怎么起一个 kernelgrid, block什么是threadIdx.x、threadIdx.y、blockIdx.x、blockIdx.y基本共享内存声明和__syncthreads()的作用全局内存合并访问的大致概念如果你这些都不熟建议先跑一个 vector add 和 matrix multiply 的简单例子再回来。下面所有分析都围绕一个核心目标让每个 warp 发出的共享内存访问尽量一次完成。2. bank conflict 的硬件细节2.1 共享内存为什么是 32 个 bank共享内存是片上存储访问延迟比全局内存低得多但设计上有一个限制它被分成 32 个 bank每个 bank 的带宽是 4 字节。也就是说一个 warp 的 32 个线程同时访问共享内存时硬件会看这 32 个地址分别落在哪些 bank。理想情况是每个线程命中一个不同的 bank那么一次内存指令就能完成。如果有两个线程落在同一个 bank但访问的不是同一个地址硬件就要把这次访问拆成多次每次都只能服务其中一个线程这个现象就叫 bank conflict。你可以把 bank 想象成 32 个并排的快递窗口。每个窗口同一时刻只能处理一个包裹只要 32 个线程每人去一个不同窗口一次就全部办完。要是 20 个人挤同一个窗口那就得排队轮 20 次。2.2 一个地址对应哪个 bank怎么算对float数组来说计算十分简单bank (数组下标的线性值) % 32因为 32 个 bank 每个 4 字节一个 float 正好占一个 bank 单位。如果是double或int4要先把字节数除以 4所以分析方式会不一样。看一个典型例子。共享内存声明成__shared__ float tile[32][32];写阶段用tile[ty][tx]一个 warp 内threadIdx.x从 0 到 31threadIdx.y固定不变。地址是ty * 32 tx因为ty是常量地址差就是tx的差所以 32 个线程分别落在 bank 0 到 31没有冲突。问题出在转置读取。如果读的时候用tile[tx][ty]地址变成tx * 32 ty对 warp 来说ty是常量tx * 32这部分模 32 永远等于 0。结果 32 个线程全部落在同一个 bankty上。这就是一次 32 路 bank conflict。原本一条共享内存读指令能完成的事要拆成 32 条性能直接掉一个数量级。2.3 为什么 padding 一列就能破局把共享内存声明改成__shared__ float tile[32][33];再读tile[tx][ty]时地址变成tx * 33 ty对同一个 warpty是常量tx从 0 到 31。因为33 % 32 1所以bank (tx ty) % 32tx每次加 1bank 也跟着每次加 1。32 个线程正好覆盖 32 个不同 bank冲突消失。这里有个容易踩的坑不是随便加几个数都行。如果声明成tile[32][34]那34 % 32 2bank 分布是2 * tx ty最后只有偶数 bank 会被用到变成 2 路冲突。padding 的列数要让步长跟 32 互质。最常用的就是 321、641 这类因为步长 33 或 65 对 32 取模都是 1。共享内存声明读tile[tx][ty]的 bank 分布冲突情况float tile[32][32]bank 恒等于ty32 路冲突float tile[32][33]bank 等于(txty)%32无冲突float tile[32][34]bank 等于(2*txty)%322 路冲突3. 完整实现padded 共享内存转置3.1 核函数的主流程整个转置核函数分成三段从全局内存按行主序合并读入一块数据。经过__syncthreads()后从共享内存按转置后的下标取数。写回全局内存时仍然保持行主序让连续线程访问连续地址。阶段 1 和阶段 3 负责解决全局内存合并访问问题阶段 2 靠 padding 解决共享内存 bank conflict 问题。线程块的形状是(32, 8)也就是blockDim.x 32这样每个 warp 恰好覆盖threadIdx.x的 0 到 31分析 bank 分布非常干净。下面是完整代码。为了顺便处理边长不是 32 整数倍的情况我加上了边界判断。实际做方阵优化时如果边长能整除 32可以把if去掉。#define TILE_DIM 32 #define BLOCK_ROWS 8 __global__ void transpose_smem_padded(float* in, float* out, int in_rows, int in_cols) { __shared__ float tile[TILE_DIM][TILE_DIM 1]; int tx threadIdx.x; int ty threadIdx.y; int bx blockIdx.x * TILE_DIM; int by blockIdx.y * TILE_DIM; // 阶段1全局内存按行主序读入共享内存 #pragma unroll for (int i 0; i TILE_DIM; i BLOCK_ROWS) { int row by ty i; int col bx tx; if (row in_rows col in_cols) { tile[ty i][tx] in[row * in_cols col]; } } __syncthreads(); // 阶段2共享内存按转置下标读出写回全局内存 #pragma unroll for (int i 0; i TILE_DIM; i BLOCK_ROWS) { int out_row bx ty i; int out_col by tx; if (out_row in_cols out_col in_rows) { out[out_row * in_rows out_col] tile[tx][ty i]; } } }主函数里的启动参数是dim3 block(TILE_DIM, BLOCK_ROWS); dim3 grid((in_cols TILE_DIM - 1) / TILE_DIM, (in_rows TILE_DIM - 1) / TILE_DIM);注意输出矩阵的列数是in_rows所以写回时用的 stride 是in_rows不是in_cols。这是新手最容易搞错的地方。3.2 为什么这样访问没有 bank conflict只看阶段 1 的写入tile[ty i][tx] in[row * in_cols col];对同一个 warptx是 0 到 31ty i固定。共享内存线性地址是(ty i) * 33 txbank 是((ty i) * 33 tx) % 32 (ty i tx) % 32这已经是无冲突的。但真正的转置读在阶段 2out[out_row * in_rows out_col] tile[tx][ty i];这里读tile[tx][ty i]线性地址是tx * 33 (ty i)同一个 warp 内ty i固定tx从 0 到 31bank 同样是(tx ty i) % 32。所以读和写两个方向都避开了 bank conflict。同时全局内存两次访问也都是合并的。阶段 1 里连续线程的col连续阶段 2 里连续线程的out_col连续。这就是这块代码最漂亮的地方两边全局访问都合并中间每一层共享内存访问都不冲突。3.3 循环展开怎么做上面代码里的两个循环都加了#pragma unroll。因为TILE_DIM / BLOCK_ROWS 4所以循环会展开成 4 段。编译器看到i 0, 8, 16, 24这 4 个常量后可以把地址计算完全摊开把共享内存访问和全局内存访问尽量排开提升指令级并行。如果你不想依赖编译器也可以手动展开。对一个 32x32 的 tile每个线程处理 4 行代码直接写成tile[ty][tx] in[(by ty) * width (bx tx)]; tile[ty 8][tx] in[(by ty 8) * width (bx tx)]; tile[ty 16][tx] in[(by ty 16) * width (bx tx)]; tile[ty 24][tx] in[(by ty 24) * width (bx tx)]; __syncthreads(); out[(bx ty) * width (by tx)] tile[tx][ty]; out[(bx ty 8) * width (by tx)] tile[tx][ty 8]; out[(bx ty 16) * width (by tx)] tile[tx][ty 16]; out[(bx ty 24) * width (by tx)] tile[tx][ty 24];这里假设是方阵width同时表示输入行数和输出列数。手动展开的好处是逻辑一眼能看穿。你在分析 bank conflict 时可以直接把 4 个读地址拎出来算。坏处是代码不通用换了TILE_DIM或BLOCK_ROWS就得重写。实际项目里我一般先写带#pragma unroll的版本性能不足再去看汇编最后才考虑手动展开。3.4 循环展开真正提升了什么很多人以为循环展开是为了减少for的跳转开销。对 GPU 来说这个收益有但不是重点。重点在于让编译器在展开后能看到多个独立访存从而调度出更好的内存级并行。比如阶段 1 展开后一个线程要发 4 条全局加载指令这 4 条之间没有数据依赖可以同时发出。如果没展开编译器也能做不少调度但展开后指令窗口更大老架构上效果更明显。要小心的是无脑加深展开。展开因子越大寄存器和指令缓存压力越大。BLOCK_ROWS从 8 改成 4循环次数从 4 变 8展开因子更大但寄存器占用上升占用率可能下降。性能往往不是单调上升而是有一个甜点值。20 系之前的卡上我常用BLOCK_ROWS8也就是每个线程处理 4 行新卡上有时BLOCK_ROWS4更好但必须实测。4. 实操过程和性能观察4.1 测试环境和基础验证我一般用 1024x1024 和 4096x4096 的 float 矩阵测编译命令nvcc -O3 -archsm_80 transpose.cu -o transpose先用小矩阵验证正确性。把 kernel 输出拷回 host和 CPU 转置结果逐元素比较。矩阵边长如果比较大随机生成数据后比对一两个关键位置是不够的必须全量比对。带宽计算用带宽 读入字节数 写出字节数转置 4096x4096 的 float 矩阵读和写都是 64 MB总共 128 MB。用cudaEvent计时算出有效带宽。这比看 kernel 执行时间更直观。4.2 第一个版本不加 padding我第一次写这个转置时共享内存声明的是tile[32][32]。结果正确但性能难看。Nsight Compute 里看共享内存访问读阶段平均冲突次数非常高。为什么前面已经算过读tile[tx][ty]时 32 个线程全部落在一个 bank 上。这等于每个共享内存读被拆成 32 次。虽然共享内存延迟已经比全局内存低很多但架不住这么拆。实测这个版本比加了 padding 的版本慢 20% 到 50%具体看 GPU 架构。4.3 第二个版本加一列 padding把声明改成tile[32][33]后bank conflict 消失性能明显上升。这个改动只多了 4 字节乘 32 行也就是 128 字节的共享内存对占用率几乎没有影响。一个 block 的共享内存用量只有33 * 32 * 4 4224 字节远小于 48 KB 的默认上限。这个版本已经能跑到比较理想的状态全局内存两侧合并共享内存中间无冲突。4.4 第三个版本循环展开同样的代码加上#pragma unroll后性能还能再上去一点。原理不是减少了循环跳转而是把 4 个独立访存暴露给编译器让长时间延迟的载入和存储任务尽量交错。这里给一个比较标准的结论在内存带宽受限的转置任务里如果全局访问已经合并、共享内存已经无冲突那再加循环展开的收益通常在 5% 到 15%。如果去掉这个收益看起来很小但转置算子往往是大矩阵处理流水线里的热点5% 可能就是整个应用的一大截提升。4.5 留意卡在哪儿我实际排查时最常用的顺序是先用Nsight Compute看全局内存合并。再看共享内存 bank conflict。最后看 warp stall 原因。如果全局内存已经合并但带宽还是不高点开 warp state 看是不是卡在long scoreboard也就是等待全局数据回来。这时候光改共享内存没用要提升并发比如让每个线程多处理几个元素或者提高占用率。如果共享内存冲突已经为零但性能还是不满意再去看指令数。转置这活儿本身不重指令数过高往往是地址计算太多可以考虑用更短的索引计算方式。5. 常见问题与排查技巧5.1 加了 padding 还是报 bank conflict先检查 padding 的步长和 32 的关系。tile[32][34]这种也会冲突原因前面讲过34 和 32 不互质。另一个常见问题是TILE_DIM不是 32却仍然用% 32来套公式。如果blockDim.x 16一个 warp 会跨两行bank 分布要重新算。最稳妥的做法是保持blockDim.x 32让一个 warp 对应共享内存的一行。这样只要步长模 32 等于 1就一定是无冲突的。5.2 结果出错先把__syncthreads()放在哪找清楚共享内存转置最典型的错误就是漏掉加载和存储之间的__syncthreads()。如果丢掉这个 barrier有些线程可能还没写完共享内存另一些线程已经去读了结果自然是花的。还有一种死锁写法把__syncthreads()放到if分支里面。因为__syncthreads()要求同一个 block 里所有线程都执行到同一位置如果一部分线程因为条件不满足跳过了 barrier整个 block 就会挂住。所以 barrier 一定要放在所有线程都会经过的地方不要放在边界检查的if内部。5.3 非方阵怎么处理如果输入是rows x cols转置输出是cols x rows那么输出全局内存的 stride 是rows不是cols。代码里的out_row来自原列out_col来自原行。只要记住这两个关系非方阵本身不复杂。另外要特别注意网格配置。grid.x根据in_cols算grid.y根据in_rows算。因为输出块的行方向对应原矩阵的列方向这样 grid 的 x 方向恰好覆盖所有输出行。5.4 循环展开后反而变慢如果你手动展开了很多段性能反而下降先查 register usage。展开会一次性生成很多局部变量可能导致编译器为每个变量分配寄存器进而降低占用率。另一个原因是共享内存指令发射已经满了。转置任务里每个数据至少要经过一次 STS 和一次 LDS共享内存带宽本身可能成为瓶颈。这时候再展开也无济于事不如减少共享内存访问次数比如尝试向量化加载和存储或者让每个线程处理一个 float4。5.5 用工具确认问题而不是靠猜Nsight Compute 是排查共享内存问题最直接的工具。建议跑ncu --set full ./transpose重点看 MemoryWorkloadAnalysis 里的 shared memory 相关指标。老的 Visual Profiler 里也可以看 shared memory conflict 统计但新项目直接用 Nsight Compute。如果指标显示冲突很高可以把错误代码简化到最小用不同 padding 值跑同一组数据。亲自对比几次后你对 bank 分布的直觉会非常准。这种东西看十篇文章不如自己在同一个矩阵上把[32][32]、[32][33]、[32][34]各跑一遍印象深刻。最后再分享一个我自己的经验转置优化最忌讳一上来就堆各种技巧。先把原始版本跑通用工具看瓶颈然后一次只改一个变量。加了 padding 冲突消失就再测一次没提升就再查是不是别的环节卡住。分享内存、bank conflict、行主序列主序这些概念本质上都是在回答同一个问题某一次访存硬件到底能不能一次干完。想明白了这一点后面再遇到各种共享内存优化题套路其实都一样。