C3 讲解：转置里的合并访存与 bank conflict

① 先看全貌：读和写不可能同时连续

y = x.T，x 是 [M, N]，y 是 [N, M]，都是行优先。y[c][r] = x[r][c]。

• 如果让相邻线程走 x 的相邻列（c 连续）：读 x[r][c] 是连续的，写 y[c][r] 时相邻线程的地址相差 M 个 float，跨步写。
• 如果反过来让相邻线程走 x 的相邻行：写连续了，读又变成跨步。

两头总有一头是跨步的。解决办法是找一个不怕跨步的中转站：shared memory。

RTX 5090 上 8192×8192 的实测（题目页的数据）：naive 805 GB/s，x.t().contiguous() 804 GB/s，tiled 1409 GB/s。

② naive kernel：二维 launch 和二维下标

drillstranspose_naive_kernel（18–24 行） · host 端 drillsdim3 block / grid（66–67 行）

dim3 block(32, 8);                                         // 每个 block 32 × 8 = 256 个线程
dim3 grid((N + 31) / 32, (M + 7) / 8);                     // x 方向覆盖列，y 方向覆盖行

const int c = blockIdx.x * blockDim.x + threadIdx.x;       // 列：快维度
const int r = blockIdx.y * blockDim.y + threadIdx.y;       // 行
if (r < M && c < N) y[(int64_t)c * M + r] = x[(int64_t)r * N + c];

• dim3 是三个分量的结构体，没写的分量默认是 1。block 和 grid 都可以是二维的，kernel 里用 .x、.y 分别取。
• 约定 threadIdx.x 对应内存中连续的那一维。warp 是按 threadIdx.x 优先切分的：一个 warp = 同一个 threadIdx.y 下的 32 个连续 threadIdx.x。所以一个 warp 的 32 个线程有相同的 r、连续的 c。
• 二维下标展开成一维地址：x[r][c] 在行优先的 [M, N] 里是 r * N + c；y[c][r] 在 [N, M] 里是 c * M + r。先转 int64_t 再乘，8192 × 8192 已经接近 int 的上限。
• 越界检查要同时检查两个维度。

同一个 warp 里，读写地址长这样（M = N = 8192）：

写入要 32 个独立的 32 字节 sector，每个只用 4 字节，有效带宽只剩 1/8。

③ tiled kernel：最绕的是坐标对调

drillstranspose_tiled_kernel（26–53 行）

每个 block 负责 x 的一个 32×32 tile，用 32 × 8 = 256 个线程，每个线程搬 4 个元素。分两个阶段，中间一个 barrier。

先看动画：一个 tile 的两个阶段

动画里最关键的是两种模式下“写 y 的事务”那一行：tiled 写 8 次，naive 写 64 次。真实的 32×32 tile 上差距是 32 倍。tiled 版本的代价是阶段 2 按列读 shared memory，这就是第 ④ 节 bank conflict 的来源。

阶段 1：把 x 的 tile 按行读进 shared memory

drills33–39 行

const int c  = blockIdx.x * TILE + threadIdx.x;   // x 的列
const int r0 = blockIdx.y * TILE + threadIdx.y;   // 这个线程负责的第一行
for (int j = 0; j < TILE; j += ROWS) {            // j = 0, 8, 16, 24
  const int r = r0 + j;
  if (r < M && c < N) tile[threadIdx.y + j][threadIdx.x] = x[(int64_t)r * N + c];
}

block 只有 8 行线程，tile 有 32 行，所以每个线程隔 8 行搬一次，共 4 次：

threadIdx.y = 1 的那一行线程依次搬 tile 的第 1、9、17、25 行。每次都是 32 个线程读 x 同一行的 32 个连续 float：合并。

为什么不直接开 32×32 = 1024 个线程、每人搬一个？可以，但每个线程只做一次 load/store，固定开销占比更高；block 太大，每个 SM 能同时驻留的 block 也更少。256 个线程、每人 4 个是常见的折中。

barrier

drills41 行 __syncthreads()：阶段 2 里一个线程要读别的线程写进 tile 的数据（它读的是 tile[threadIdx.x][...]，行列反过来了）。必须等整个 tile 写完。

阶段 2：block 坐标对调，再按行写 y

drills45–52 行

const int yc  = blockIdx.y * TILE + threadIdx.x;   // y 的列（= x 的行）
const int yr0 = blockIdx.x * TILE + threadIdx.y;   // y 的行（= x 的列）
for (int j = 0; j < TILE; j += ROWS) {
  const int yr = yr0 + j;
  if (yr < N && yc < M) y[(int64_t)yr * M + yc] = tile[threadIdx.x][threadIdx.y + j];
}

两件事同时发生：

1. tile 级别对调：x 的第 (by, bx) 个 tile 转置后是 y 的第 (bx, by) 个 tile。所以 y 的行用 blockIdx.x 算，y 的列用 blockIdx.y 算。
2. tile 内部对调：threadIdx.x 仍然走 y 的列（连续方向），保证写入合并；数据从 tile[threadIdx.x][threadIdx.y + j] 取，也就是在 shared memory 里按列读。

拿一个具体线程验算一遍：block (bx = 2, by = 5)，线程 (tx = 3, ty = 1)，只看 j = 0。

注意同一个线程在两个阶段搬的不是同一个元素：它写进 tile[1][3]，却读出 tile[3][1]。这就是为什么中间必须有 barrier。

④ 那一列 padding：bank conflict

drills27–30 行 __shared__ float tile[TILE][TILE + 1];

shared memory 被分成 32 个 bank，每个 bank 每周期能服务一个 4 字节的字。地址为 a（以 float 为单位）的元素在第 a % 32 个 bank。一个 warp 里如果有多个线程访问同一个 bank 的不同地址，就要排队串行，这叫 bank conflict。

先用 4 个 bank 看懂：二维的 tile 其实是一条一维数组

先别管 32，缩小成 4 个 bank、4×4 的 tile。关键只有一条规则：一个数在第几个 bank，只看它的地址除以 bank 数的余数。

shared memory 其实是一条长长的一维数组，地址 0、1、2、3、4……依次排下去，bank = 地址 % 4，就是 0、1、2、3、0、1、2、3……一直循环：

地址:  0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15
bank:  0  1  2  3  0  1  2  3  0  1  2  3  0  1  2  3

二维的 tile[i][j] 是按行一行接一行放进这条数组的：地址 = i × 每行长度 + j。

不加 padding（每行 4 个）：地址 = 4i + j。把每一格的 bank 号画出来：

         列0 列1 列2 列3
第0行:    0   1   2   3      ← 地址 0..3
第1行:    0   1   2   3      ← 地址 4..7
第2行:    0   1   2   3      ← 地址 8..11
第3行:    0   1   2   3      ← 地址 12..15

每一行都从 bank 0 开始：行长 4 正好等于 bank 数，转一圈又回到起点。读第 0 列时，4 个线程读地址 0、4、8、12，全是 bank 0。一个 bank 每拍只能服务一个线程，只能排队，读一次花了 4 拍。

加 padding（每行 5 个，最后一个空着不用）：地址 = 5i + j。

         列0 列1 列2 列3 (空)
第0行:    0   1   2   3   0      ← 地址 0..4
第1行:    1   2   3   0   1      ← 地址 5..9
第2行:    2   3   0   1   2      ← 地址 10..14
第3行:    3   0   1   2   3      ← 地址 15..19

每一行多占了一格，下一行的起点就往后挪了一个 bank。再读第 0 列：地址 0、5、10、15，bank 是 0、1、2、3，各不相同，4 个 bank 同时工作，1 拍读完。

下面的演示把这个过程一步步做出来：先把 tile 一行一行地放进一维内存，再读第 0 列。切换“每行 4 个 / 每行 5 个”对比。

一句话：问题不在“读一列”本身，而在于行长正好等于 bank 数时，每一行都从同一个 bank 开始，于是同一列的元素全落在同一个 bank 上。每行多加一个空位，行长从 32 变成 33，每一行的起点就错开一个 bank，同一列的 32 个元素正好分散到 32 个不同的 bank。代价是每行浪费一个 float，32 行一共 128 字节。

把上面的 4 换成 32，就是真实的硬件：

阶段 2 里一个 warp 读的是 tile[0..31][k]（同一列，tx 变化）：

• 不加 padding（每行 32 个）：tile[i][k] 的地址是 32i + k，bank = k。32 个线程全在同一个 bank，32 路冲突，读一次变成 32 次。
• 加 1 列 padding（每行 33 个）：地址是 33i + k，bank = (i + k) % 32。i = 0..31 时恰好覆盖 32 个不同的 bank，无冲突。

动画：bank conflict 是怎么排队的

试一试这四种组合：tile[8][8] + 读一列要 8 个周期（全部挤在 bank 0），tile[8][9] + 读一列只要 1 个周期，而读一行不管有没有 padding 都是 1 个周期。真实硬件上把 8 换成 32：不加 padding 是 32 路冲突。

读第 0 列时，前 8 个线程各自落在哪个 bank：

多出来的那一列永远不存数据，只是把每一行往后错开一个 bank。代价：32 × 4 = 128 字节 shared memory。

阶段 1 是按行写 tile[i][0..31]，bank = (33i + k) % 32，k 连续，本来就无冲突，padding 不影响它。

另一种做法：swizzle（异或）

padding 浪费一点空间，也会让地址不再对齐到 2 的幂，对向量化读写和 Tensor Core 不友好。生产代码更常用 swizzle：不多加列，而是把列号和行号异或，物理列 = j ^ (i % 32)。同一列 j 的 32 行被异或到 32 个不同的物理列，同样无冲突。

• CuTe 的通用实现：cutlassstruct Swizzle，核心一行是 cutlassoffset ^ shiftr(offset & yyy_msk, msk_sft)
• FlashInfer 的版本：flashinferget_permuted_offset：i * stride + (j ^ (i % 8))（以 16 字节为单位做异或）

C4 和 G2 会再遇到这个技巧。

⑤ 边界：M、N 不是 32 的倍数时

grid 是 ⌈N/32⌉ × ⌈M/32⌉，最右边和最下边的 tile 只有一部分在矩阵里。两个阶段各自做越界检查：

• 阶段 1 if (r < M && c < N)：越界的位置不读，tile 里那几格留着垃圾值。
• 阶段 2 if (yr < N && yc < M)：越界的位置不写。y 里合法的位置只会读到 tile 里合法的格子，垃圾值永远不会被写出去。

注意阶段 2 的条件是 yr < N && yc < M，因为 y 的形状是 [N, M]。M 和 N 写反是这题最常见的 bug，而方阵测试发现不了，所以测试里一定要有 M ≠ N 的形状。

⑥ 常见坑与自测

• 忘了 __syncthreads()：小矩阵上可能碰巧对，大矩阵上随机出错。
• 阶段 2 没有对调 blockIdx：写到 y 的 (by, bx) 位置，方阵上部分 tile 对、部分错；非方阵直接越界。
• 阶段 2 用了 tile[threadIdx.y + j][threadIdx.x]：这就成了原样拷贝，没有转置。
• 把 padding 加在第一维（tile[33][32]）：没用，行长度还是 32。

源码与资料索引

本仓库：drillssolutions/c3_transpose.cu · drillsmine/c3_transpose.cu · drillstests/test_c3_transpose.py · drillscontent/c3.md

对照实现

• cuda-samples 的整套演进：cudasamplestransposeNaive → cudasamplestransposeCoalesced（有 bank conflict）→ cudasamplestransposeNoBankConflicts（TILE_DIM + 1）→ cudasamplestransposeDiagonal（处理 partition camping）
• swizzle：cutlassCuTe Swizzle · flashinferFlashInfer permuted_smem

文档

• NVIDIA blog · An Efficient Matrix Transpose in CUDA C/C++
• Programming Guide · Matrix Transpose Example Using Shared Memory · Shared Memory Bank Conflicts · Coalesced Access