是什么

ldmatrix.sync 是 按 warp 协作、从 shared memory 中一次性读取 8×8 的矩阵 tile 到寄存器的指令，用来给 Tensor Core 的矩阵乘（mma.sync / mma.sp.sync 等）准备操作数 fragment。
它是 同步（sync）且 warp 级别 的：同一 warp 的 32 个线程必须一起执行该指令。

典型语法

ldmatrix.sync.aligned.m8n8.x{1|2|4}[.trans].shared.b{16|8} {dst_list}, [addr];

• aligned：要求基地址满足对齐（通常 ≥16B，对 b16 的 8×8 刚好 16B/行）。

• m8n8：一次处理的 tile 尺寸是 8×8。

• x1|x2|x4：一次加载几个 8×8 tile 的 fragment（返回 1/2/4 份寄存器结果）。

  • 常见：x4 给 A 操作数，x2 给 B 操作数（与后续 mma.sync.m16n8k16 等形状匹配）。

• trans：在加载时对 8×8 做转置（常用于 B 操作数的列主布局需求）。

• shared：数据源必须在 shared memory。

• b16|b8：元素宽度（半精度/BF16/INT16 等用 b16；INT8 等用 b8）。

• {dst_list}：目的寄存器列表。

  • b16.x1：每个线程产出 1 个 32-bit 寄存器（里面打包 2 个 16-bit 元素）；

  • b16.x2：每线程 2 个寄存器；

  • b16.x4：每线程 4 个寄存器。
    b8 时每寄存器打包更多元素（4 个 8-bit）。

• [addr]：每个线程提供一个地址寄存器（指向 shared 内的一段），硬件按既定模式把 8×8 tile 的每行/列分配给不同线程寄存器。

小结：x 的倍数越大，一次为后续 mma 准备的 fragment 越多，也越省指令。

执行与线程/寄存器映射（直观理解）

• warp 内 32 线程被划分成 4 组（每组 8 线程），每组 8 条 lane 提供 8 条起始地址，合起来 装载一个 8×8。

• x1：warp 等价于并行处理 4 个 8×8（每 8 线程一组，各拿一个 tile）；

• x2/x4：在同一次指令中为同一目的操作数加载 2/4 份 fragment 到更多寄存器，便于后续更大的 mma 形状直接使用。

• .trans 会在装载到寄存器时做 8×8 的转置，使寄存器布局直接符合 mma.sync 所需（例如 row.col 变体里 A 用非转置、B 常用转置）。

具体 lane 到元素的映射由硬件定义，你只需按官方推荐的“行指针计算方式”把每 lane 的 [addr] 算对即可。

与 mma.sync 的常用组合（示例）

// 假设 As、Bs 是 shared memory 中准备好的 tile
uint32_t a0,a1,a2,a3;  // A fragment (b16, x4)
uint32_t b0,b1;        // B fragment (b16, x2)
uint32_t a_addr = __cvta_generic_to_shared(As_lane_ptr);

uint32_t b_addr = __cvta_generic_to_shared(Bs_lane_ptr);
// A：不转置，x4

asm volatile(

"ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%0,%1,%2,%3}, [%4];\n"

: "=r"(a0),"=r"(a1),"=r"(a2),"=r"(a3)

: "r"(a_addr));
// B：转置，x2（常见搭配）

asm volatile(

"ldmatrix.sync.aligned.m8n8.x2.trans.shared.b16 {%0,%1}, [%2];\n"

: "=r"(b0),"=r"(b1)

: "r"(b_addr));
// 然后喂给 Tensor Core 做 16x8x16 的 MMA（示意）

uint32_t d0,d1,d2,d3;   // 累加结果片段

asm volatile(

"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f32 "

"{%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%0,%1,%2,%3};\n"

: "+r"(d0),"+r"(d1),"+r"(d2),"+r"(d3)

:  "r"(a0),"r"(a1),"r"(a2),"r"(a3), "r"(b0),"r"(b1));

常见要点与坑

1. 地址必须是 shared 空间地址：用 __cvta_generic_to_shared() 或 PTX cvta.to.shared.u32 转换。

2. 对齐：确保 tile 起始地址 16B 对齐（aligned）。shared 数组建议 __align__(128)。

3. 所有 32 线程都要参与（warp 同步语义），不要在半 warp 里执行。

4. .trans 用在恰当的操作数（通常 B 操作数需要）。

5. Bank 冲突：规划好 shared 内的行跨距（stride），很多实现用 (+padding) 避免 32-bank 冲突。

6. 配套形状：x4/x2 选择应与目标 mma.sync 的 m*n*k 形状匹配，否则需要额外重排/指令。

7. 架构要求：ldmatrix 一般要求 SM 7.5+（Turing 及以后）；b8 相关在 Ampere+ 更常见。

什么时候用

• 你已经把 A、B 的 tile 从 global 拷到 shared（常配 cp.async）；

• 需要用 Tensor Core 做半精度/BF16/INT8 的 GEMM/卷积；

• 希望 一次指令 把 8×8 的片段按 硬件喜欢的布局 放进寄存器，减少显式重排。

快速对照表