从共享内存到 MMA：读懂 A、B 的 ldmatrix 与 retile_D
读 CuTe 的 GEMM 代码时，经常会看到这样的组合：
auto tCrA = thr_mma.partition_fragment_A(sA);
auto tCrB = thr_mma.partition_fragment_B(sB);
auto tXsA = thr_copy_a.partition_S(sA);
auto tXsB = thr_copy_b.partition_S(sB);
auto tXrA = thr_copy_a.retile_D(tCrA);
auto tXrB = thr_copy_b.retile_D(tCrB);
copy(copy_a, tXsA, tXrA);
copy(copy_b, tXsB, tXrB);
其中有三个需要接起来的问题：MMA 的寄存器槽位应当装什么，ldmatrix 从哪里读、往哪里写，以及 retile_D 怎样把指令输出接到已有的寄存器存储上。
先给两个元素指定数值，作为全文的观察目标：
共享内存：A(0,9)=37 共享内存：B(0,9)=53
↓ ↓
lane 0 的 A half 槽位 5 lane 0 的 B half 槽位 3
MMA 布局确定槽位需要的元素；Copy Atom 描述指令的固定搬运路线；TiledCopy 据此计算源地址，retile_D 则将输出绑定到 MMA fragment。 下面从一个 warp 的具体例子开始，把这两条加载链路算清，再改变加载指令，最后读一遍 retile_D 的源码。
先固定矩阵、存储和指令
本文使用同一条 MMA，先为 A 选择 ldmatrix.x4，为 B 选择 ldmatrix.x2：
auto mma = make_tiled_mma(SM80_16x8x16_F16F16F16F16_TN{});
auto copy_a = make_tiled_copy_A(
Copy_Atom<SM75_U32x4_LDSM_N, half_t>{}, mma);
auto copy_b = make_tiled_copy_B(
Copy_Atom<SM75_U32x2_LDSM_N, half_t>{}, mma);
一个 warp 有 32 个线程，warp 内的线程编号叫 lane，取值为 0…31。所有编号从 0 开始。本文的 v 表示线程内一个 half 值的逻辑编号；两个 half 打包成一个 32 位指令操作数。
MMA 的计算尺寸是 M=16、N=8、K=16。按通常的数学写法，A 是 16×16，B 是 16×8。CuTe 访问 A 时使用 (M,K)，访问 B 时使用 (N,K)，因此本文约定：
A(m,k) 对应数学上的 A[m,k]
B(n,k) 对应数学上的 B[k,n]
D(m,n) = C(m,n) + sum(k=0…15, A(m,k)*B(n,k))
这只是 B 的索引顺序不同，没有额外转置数据。后面的 B(0,9)，指的是数学上的 B[9,0]。
假设数据已经准备在共享内存中，并且对整个 warp 可见。我们选择 K 连续的存储，基址满足 16 字节对齐：
auto sA = make_tensor(make_smem_ptr(smem_a),
Layout<Shape<_16,_16>, Stride<_16,_1>>{});
auto sB = make_tensor(make_smem_ptr(smem_b),
Layout<Shape<_8,_16>, Stride<_16,_1>>{});
本文用 shape:stride 表示布局。Shape 给出各维大小，stride 表示沿该维走一步增加多少个元素偏移：
| Tensor | 坐标 | 布局 | half 存储偏移 |
|---|---|---|---|
sA | (m,k) | (16,16):(16,1) | 16*m+k |
sB | (n,k) | (8,16):(16,1) | 16*n+k |
所以两者的 (0,9) 都位于各自基址之后的第 9 个 half，即字节偏移 18。本文分别用 0x1000、0x2000 作为 A、B 的示例基址。
这种存储让 K 方向的 8 个 half 连续，适合当前非转置 ldmatrix 的行加载。行主序只规定坐标如何变成地址，尚未规定哪一个 lane 提供这些地址。本文先不展开 swizzle 和流水线。
MMA 先规定每个槽位需要什么
A 的 lane 0，为什么槽位 5 对应 A(0,9)
mma_traits_sm80.hpp 中，这条 MMA 的 A 布局是：
using ALayout =
Layout<Shape <Shape < _4,_8>, Shape < _2,_2, _2>>,
Stride<Stride<_32,_1>, Stride<_16,_8,_128>>>;
输入为 (thread,value)，简称 TV。输出按 q_A=m+16*k 编码矩阵位置。这里的 q_A 是逻辑矩阵编号，与前面共享内存的偏移采用不同规则。
将线程和值的嵌套坐标写开：
t = t0 + 4*t1
v = v0 + 2*v1 + 4*v2
q_A = 32*t0 + t1 + 16*v0 + 8*v1 + 128*v2
= (t1+8*v1) + 16*(2*t0+v0+8*v2)
m = t1 + 8*v1
k = 2*t0 + v0 + 8*v2
lane 0 对应 (t0,t1)=(0,0)，值编号 5 在 (2,2,2) 中展开为 (1,0,1)：
m = 0
k = 1+8 = 9
q_A = 144
因此，lane 0 的 A 槽位 5 必须保存 A(0,9)。这条要求来自 MMA 指令的寄存器接口，CuTe 的 ALayout 将它记录下来。
B 的 lane 0，为什么槽位 3 对应 B(0,9)
同一个 traits 中，B 的布局是：
using BLayout =
Layout<Shape <Shape < _4,_8>, Shape <_2, _2>>,
Stride<Stride<_16,_1>, Stride<_8,_64>>>;
B 使用 (N,K) 坐标，输出编码为 q_B=n+8*k。将值编号展开成 v=v0+2*v1：
q_B = 16*t0 + t1 + 8*v0 + 64*v1
= t1 + 8*(2*t0+v0+8*v1)
n = t1
k = 2*t0 + v0 + 8*v1
lane 0、值编号 3 对应 t0=t1=0、v0=v1=1，因此：
n = 0
k = 1+8 = 9
q_B = 72
A、B 在 lane 0 的全部输入槽位如下。a0…a3、b0…b1 表示该 lane 的 32 位 MMA 输入操作数：
| 操作数 | half 槽位 | 低 half 对应元素 | 高 half 对应元素 |
|---|---|---|---|
a0 | 0、1 | A(0,0) | A(0,1) |
a1 | 2、3 | A(8,0) | A(8,1) |
a2 | 4、5 | A(0,8) | A(0,9) |
a3 | 6、7 | A(8,8) | A(8,9) |
b0 | 0、1 | B(0,0) | B(0,1) |
b1 | 2、3 | B(0,8) | B(0,9) |
MMA 一次需要每个 lane 提供 8 个 A half 和 4 个 B half，即 4 个 A、2 个 B 的 32 位操作数。物理寄存器的具体编号由编译器分配，这里的槽位编号描述的是指令参数中的顺序。
partition_fragment 准备的是怎样的存储
auto thr_mma = mma.get_slice(lane);
auto tCrA = thr_mma.partition_fragment_A(sA);
auto tCrB = thr_mma.partition_fragment_B(sB);
partition_fragment_A 的关键调用是：
return TiledMMA::make_fragment_A(partition_A(atensor));
partition_A 找出本线程需要的矩阵元素，得到共享内存上的视图；本例的 make_fragment_A 再通过 make_fragment_like 创建线程局部存储。B 的过程相同。创建 fragment 还没有执行共享内存加载。
打印 lane 0 的布局，可以看到存储位置发生了变化：
| 视图或 fragment | 实际布局 |
|---|---|
thr_mma.partition_A(sA) | ((2,2,2),1,1):((1,128,8),0,0) |
tCrA | ((2,2,2),1,1):((1,2,4),0,0) |
thr_mma.partition_B(sB) | ((2,2),1,1):((1,8),0,0) |
tCrB | ((2,2),1,1):((1,2),0,0) |
第一维保存一条 MMA 的线程内值，后两维是 M 或 N 方向、K 方向的 MMA 重复；本例都为 1。第一维虽然包含嵌套括号，整个 tensor 的 rank 仍然是 3。
make_fragment_like 对第一维使用 LayoutLeft 生成紧凑布局。源码中的：
compact_major<LayoutLeft>(
filter_zeros(atom_stride, atom_shape))
用原 stride 识别静态零步长对应的广播维度，保留这些别名关系；对其余维度重新生成紧凑步长。本例没有这种广播，因此 A 的 (1,128,8) 变成 (1,2,4)，B 的 (1,8) 变成 (1,2)。
这样，逻辑值编号 v 在本例 fragment 中落到 half 存储槽位 v。后面要做的事，就是把正确的共享内存数值填进这些槽位。
ldmatrix 自己规定了怎样的搬运路线
先把矩阵元素名称放在一边，单独看非转置指令：
ldmatrix.sync.aligned.x4.m8n8.shared.b16
m8n8 规定每个子矩阵有 8 行、每行 8 个 16 位元素；每行的 8 个 half 占 16 字节，由一个 lane 提供行首地址。x4 表示同时加载四个这样的子矩阵。
| 形式 | 每次加载的子矩阵数 | 提供有效行地址的 lanes | 每个 lane 的输出 |
|---|---|---|---|
.x1 | 1 | 0～7 | 1 个 32 位值，即 2 个 half |
.x2 | 2 | 0～15 | 2 个 32 位值，即 4 个 half |
.x4 | 4 | 0～31 | 4 个 32 位值，即 8 个 half |
32 个 lane 都协作执行指令并接收输出。.x1/.x2 中未提供独立行地址的上层 lanes 仍然参与；CuTe 对它们构造重复的合法源地址。
将子矩阵编号记为 b，行内坐标记为 (row,col)，硬件路线可以写成：
源地址提供 lane = 8*b + row
目的 lane = 4*row + col/2
目的输出编号 = b
输出 half 位置 = col%2 // 0 为低 half，1 为高 half
其中 row,col=0…7，除法取整数。例如，子矩阵 2 的第 0 行，由 lane 16 提供地址；行内第 1 个 half 被送到目的 lane 0 的 dst2 高 half。
指令不知道这些地址里的数据叫 A 还是 B，也不知道矩阵的 (m,k)。四个 8×8 是硬件粒度；它们是否拼成 16×16，以及对应哪几个区域，取决于软件提供的源地址。
CuTe 怎样把两套指令规则接起来
TiledCopy 采用 MMA 的 TV 分配
make_tiled_copy_A/B 的关键参数分别是：
// A
make_tiled_copy_impl(copy_atom,
mma.get_layoutA_TV(),
make_shape(tile_size<0>(mma), tile_size<2>(mma)));
// B
make_tiled_copy_impl(copy_atom,
mma.get_layoutB_TV(),
make_shape(tile_size<1>(mma), tile_size<2>(mma)));
也就是说，Copy 的 TiledLayout_TV 取自 MMA。本文把 A、B 的这张布局分别记为 F_A、F_B。在当前单 warp、单 MMA 的配置下，它们保留了前面 Atom 的元素分配，只带有大小为 1 的额外层次。
接下来区分 Copy Atom 中的三张布局：
S = AtomLayoutSrc
D = AtomLayoutDst
R = AtomLayoutRef
它们把各自的 (thread,value) 映射到同一套 Copy 内部元素编号 e。这些编号只用于连接 Copy 的源、目的和 reference。源码明确说明，这套抽象编号可以任意选择，只要映射关系一致。
本例两种 ldmatrix 都定义 RefLayout = DstLayout，因此 R=D。构造 TiledCopy 时，MMA 的 TV 分配被用来赋予 Copy 目的 TV 矩阵含义。
| 布局 | 输入 | 输出的含义 |
|---|---|---|
sA.layout() / sB.layout() | 矩阵坐标 | 实际元素存储偏移 |
MMA 的 F_A / F_B | lane、线程内值编号 | 逻辑矩阵编号 q |
Copy Atom 的 S/D/R | 各自的 thread、value | Copy 内部编号 e |
后两类布局都接收 TV，但输出不是同一套编码。不能拿 e 直接套用 q_A=m+16*k 或 q_B=n+8*k。
partition_S 用目的需求反推源地址
源码中，partition_S 经由 tidfrg_S 接入这张转换：
right_inverse(AtomLayoutRef{}).compose(AtomLayoutSrc{})
它先通过 S 找到内部编号，再通过 R 的逆映射恢复 reference TV。tile2thrfrg 随后接入 TiledLayout_TV，最后与原 tensor 的存储布局组合。
对当前 A 的 .x4、B 的 .x2，一个 Atom 恰好覆盖整个操作数 tile，因此可以简写为：
源矩阵位置 = F ∘ R⁻¹ ∘ S
目的矩阵位置 = F ∘ R⁻¹ ∘ D
∘ 表示函数组合，从右向左执行。由于 R=D，目的端组合等于 F。这就是 Copy 目的槽位与 MMA 槽位指向同一矩阵元素的原因。 较小 Atom 带来的重复组，在后面换 .x2/.x1 时再展开。
从构造方向看，MMA traits 描述 MMA 指令，Copy traits 描述 Copy 指令；CuTe 选择兼容的组合，并让本次 TiledCopy 采用 MMA 的元素分配。两种 traits 分别忠实记录各自硬件行为。
A：把偏移 9 中的 37 搬进 lane 0 的槽位 5
从目的槽位反查源位置
Copy_Atom<SM75_U32x4_LDSM_N,half_t> 将 traits 的 bit 布局转换为 half 单位后，得到：
S = (32,8):(8,1)
D = (32,(2,4)):(2,(1,64))
R = D
S(tS,p) = 8*tS+p
D(tD,v) = 2*tD + v%2 + 64*(v/2)
目的 (lane=0,v=5) 的内部编号是：
D(0,5) = 1+64*2 = 129
令 S(tS,p)=129，其中 p=0…7，解得 tS=16、p=1。这说明硬件会把源 lane 16 所提供行的第 1 个 half，送到目的 lane 0 的槽位 5。
再通过 partition_S 计算这个源位置应该读取的元素：
源 (16,1)
→ S → e=129
→ R⁻¹ → reference (0,5)
→ F_A → q_A=144
→ → A(0,9)
→ sA → half 偏移 9，字节地址 0x1012
这里出现的几个数，各有用途：
129 Copy 内部编号，用于连接搬运路线
144 A 的逻辑矩阵编号，解码为 (0,9)
9 A(0,9) 在 sA 中的 half 存储偏移
37 地址中实际保存、最终需要搬运的数值
用更小的编号也能看出两套编码的区别：A 的目的槽位 (0,1) 在 F_A 中输出 q_A=16，表示 A(0,1)；在 D 中却输出 e=1。若把这个 1 按 q_A=m+16*k 解码，会得到 A(1,0)。因此必须先通过 R⁻¹ 恢复 reference TV，再由 F_A 取得矩阵编号。
为什么是 lane 16 提供行地址
按前面的组合计算全部源位置，会得到下面四个子矩阵：
K=0…7 K=8…15
┌──────────────┬──────────────┐
M=0…7 │ 子矩阵 0 │ 子矩阵 2 │
├──────────────┼──────────────┤
M=8…15 │ 子矩阵 1 │ 子矩阵 3 │
└──────────────┴──────────────┘
对应的源行首地址为：
| 子矩阵 | 地址提供 lane | 该 lane 提供的行首 |
|---|---|---|
| 0 | r | &A(r,0) |
| 1 | 8+r | &A(8+r,0) |
| 2 | 16+r | &A(r,8) |
| 3 | 24+r | &A(8+r,8) |
其中 r=0…7。这是根据 MMA 目的分配和 Copy 路线计算出的地址安排。
所以，A 的第 0 行被拆成两段：lane 0 提供 &A(0,0)，覆盖 K 的 0…7；lane 16 提供 &A(0,8)，覆盖 K 的 8…15。矩阵行号为 0，并不意味着整行由 lane 0 提供。
lane 16 实际传入的地址是 0x1010。硬件读取该行时，第 1 个 half 位于 0x1012，数值是 37。按指令的固定分发规则：
子矩阵 2、row=0、col=1
→ 目的 lane = 4*0+1/2 = 0
→ 输出 dst2 的高 half
dst2 怎样写到已有的 MMA fragment
本例中，retile_D(tCrA) 得到：
tCrA：((2,2,2),1,1):((1,2,4),0,0)
tXrA：((8,1),1,1):((1,0),0,0)
Copy 视图第一维的 (8,1) 表示每组 8 个 half，共 1 组。它沿用 tCrA.data()，并将组内值编号接到原 fragment 的槽位上。
copy_unpack 随后做按位重解释：
Tensor rS = recast<uint128_t>(src);
Tensor rD = recast<uint32_t>(dst);
// 展开调用的含义：
CopyOp::copy(rS(0), rD(0), rD(1), rD(2), rD(3));
源端取 rS(0) 的共享内存地址作为行首；目的端的 rD(2) 引用 fragment 的 half 槽位 4、5。于是，ldmatrix 写 dst2 的高 half，就是写 tCrA(5,0,0)。
共享内存地址 0x1012 中的 37
↓ warp 协作的 ldmatrix.x4
lane 0 的 dst2 高 half
↓ 输出引用绑定到 tCrA
lane 0 的 A 槽位 5 = 37
B：把偏移 9 中的 53 搬进 lane 0 的槽位 3
B 每个 lane 需要 4 个 half，所以一次 .x2 就能准备完整输入。对应 Atom 的 half 布局是：
S = ((16,2),8):((8,0),1)
D = (32,(2,2)):(2,(1,64))
R = D
S(tS,p) = 8*(tS%16)+p
D(tD,v) = 2*tD + v%2 + 64*(v/2)
源线程维度里的零步长，使 lanes 16～31 的源视图重复 lanes 0～15。指令使用后者提供的 16 个独立行地址；目的端仍有 32 个 lane。
对目的 lane 0、槽位 3：
D(0,3) = 1+64 = 65
S(8,1) = 8*8+1 = 65
因此该槽位的数据来自 lane 8 所提供行的第 1 个 half。完整的源地址计算为：
源 (8,1)
→ S → e=65
→ R⁻¹ → reference (0,3)
→ F_B → q_B=72
→ → B(0,9)
→ sB → half 偏移 9，字节地址 0x2012
两个子矩阵沿 K 方向安排：
| 子矩阵 | B 的范围，按 (N,K) | 地址提供 lane | 行首 |
|---|---|---|---|
| 0 | N=0…7，K=0…7 | r | &B(r,0) |
| 1 | N=0…7，K=8…15 | 8+r | &B(r,8) |
lane 8 提供 &B(0,8)=0x2010。指令读取偏移一个 half 的 0x2012，把数值 53 送到 lane 0 的 dst1 高 half。
B 的目的视图为：
tCrB：((2,2),1,1):((1,2),0,0)
tXrB：((4,1),1,1):((1,0),0,0)
重新解释成两个 32 位输出后，dst1 引用 half 槽位 2、3。因此 tCrB(3,0,0)=53，正好满足前面的 B 输入分配。
这里不需要 .trans：本例的 sB 按 (N,K) 索引，K 连续；四个值经非转置 .x2 分发后，已经符合当前 MMA 的 B 槽位要求。换成另一种 B 存储时，需要重新分析地址连续性和指令形式。
MMA 不变，把 A 改成两次 x2
现在只改变 A 的 CopyAtom：
auto copy_a = make_tiled_copy_A(
Copy_Atom<SM75_U32x2_LDSM_N, half_t>{}, mma);
MMA 的 F_A 没变，每个 lane 仍需要 8 个 half。但 .x2 每次只给每个 lane 写入 4 个 half，所以 TiledCopy 需要两组指令：
第一组：填充 A 槽位 0…3，即 a0、a1
第二组：填充 A 槽位 4…7，即 a2、a3
实际的 retile_D 结果也变为：
((4,2),1,1):((1,4),0,0)
将组内 half 编号记为 a=0…3，加载组编号记为 g=0,1：
// 本例中，这两个访问指向相同存储位置。
tXrA(make_coord(a,g),0,0);
tCrA(a+4*g,0,0);
因此两次指令的输出引用绑定为：
| 加载组 | 本次 dst0 绑定 | 本次 dst1 绑定 | 本次加载的 A 区域 |
|---|---|---|---|
| 0 | a0，槽位 0、1 | a1，槽位 2、3 | 整个 M、K=0…7 |
| 1 | a2，槽位 4、5 | a3，槽位 6、7 | 整个 M、K=8…15 |
这里 dst0/dst1 是每条 .x2 的局部输出编号。第二条指令的 dst0 绑定到 a2，因此它会保留第一条指令已经填好的 a0。
A(0,9) 的源地址提供 lane 也改变了
.x2 每次使用 lanes 0～15 的行地址。本例两次的地址安排为：
| 加载组 | lanes r 提供 | lanes 8+r 提供 |
|---|---|---|
| 0 | &A(r,0) | &A(8+r,0) |
| 1 | &A(r,8) | &A(8+r,8) |
所以，A(0,9) 的数据这次来自 第二组的源 lane 0：
g=1，lane 0 提供 &A(0,8)
→ 读取行内第 1 个 half：A(0,9)=37
→ 目的 lane 0，本次 dst0 的高 half
→ 第二组局部 half 槽位 1
→ 全局 fragment 槽位 1+4*1=5
同一个矩阵元素，使用 .x4 时由 lane 16 提供行地址，使用第二次 .x2 时由 lane 0 提供行地址；最终进入的 MMA 槽位相同。
这也解释了为什么不能把 Copy 内部编号当作矩阵编号。第二次 .x2 中，该元素的 Atom 内部编号是 e=1：
源 Atom TV (0,1)
→ S → e=1
→ R⁻¹ → reference Atom TV (0,1)
→ 加入重复组 g=1 → reference tile TV (0,5)
→ F_A → q_A=144，即 A(0,9)
内部编号只标识一次 Atom 调用中的路线；重复组确定这次调用在整个 tile 中的位置。.x4 中的 e=129 和这里的 e=1，都可以在各自的组合下对应同一个 A(0,9)。
B 也可以改成两次 x1
B 的 F_B 同样可以保持不变，改选：
auto copy_b = make_tiled_copy_B(
Copy_Atom<SM75_U32x1_LDSM_N, half_t>{}, mma);
此时一次指令给每个 lane 写入 2 个 half，retile_D(tCrB) 的结果是：
((2,2),1,1):((1,2),0,0)
| 加载组 | 行地址安排，r=0…7 | lane 0 收到的元素 | 最终 B 槽位 |
|---|---|---|---|
| 0 | lane r 提供 &B(r,0) | B(0,0)、B(0,1) | 0、1 |
| 1 | lane r 提供 &B(r,8) | B(0,8)、B(0,9) | 2、3 |
B(0,9)=53 由第二组的 lane 0 提供行地址，经本次 dst0 高 half，写入最终 B 槽位 1+2*1=3。
四种配置可以放在一起看：
| 操作数与 CopyAtom | 每 lane 的总 half 数 | 每次写入 half 数 | 指令组数 | retile_D 第一维 |
|---|---|---|---|---|
A / .x4 | 8 | 8 | 1 | (8,1):(1,0) |
A / .x2 | 8 | 4 | 2 | (4,2):(1,4) |
B / .x2 | 4 | 4 | 1 | (4,1):(1,0) |
B / .x1 | 4 | 2 | 2 | (2,2):(1,2) |
第一维中第二个大小为 1 的 mode 使用零步长，因为没有第二组需要访问。
当前单个 B tile 每 lane 只有 4 个 half，不能直接给它套一个每 lane 需要 8 个 reference half 的 .x4 Atom；TiledCopy 的数量断言就不满足。若想用更大的 CopyAtom，需要重新设计覆盖的 tile 和 fragment 范围。
带着 A 的 8/4、B 的 4/2，读一遍 retile_D
前面已经知道了所需结果：A 的 .x2 要把 8 个槽位分成两组，每组 4 个；B 的 .x1 要把 4 个槽位分成两组，每组 2 个。现在看源码如何从输入布局得到它们。
外层沿用存储，内层计算布局
ThrCopy::retile_D 的关键代码是：
return make_tensor(
static_cast<DTensor&&>(dtensor).data(),
TiledCopy::retile(dtensor.layout()));
它将已有 fragment 的 data() 与新布局放在一起。因此 tXrA 引用的是 tCrA 已有的存储，tXrB 也一样。
但“共用 data()”只说明存储相同，还需要解释每个 Copy 坐标为什么能落到正确的原槽位。这个对应关系由内部 retile 计算。下面保留其主体，省略修饰：
template <class Tensor>
auto retile(Tensor&& tensor)
{
constexpr int R = remove_cvref_t<Tensor>::rank;
auto V = size<0>(tensor);
auto frg_layout_mn = upcast<TiledNumThr{} * V>(
right_inverse(TiledLayout_TV{})
.with_shape(shape(Tiler_MN{})));
auto frg_layout_v = zipped_divide(
logical_product(make_layout(V), right_inverse(frg_layout_mn)),
make_layout(AtomNumVal{}));
auto t_tensor = zipped_divide(
tensor, prepend(product_each(shape(frg_layout_mn)), V));
auto v_tensor = t_tensor.compose(frg_layout_v, _);
return v_tensor(_, append<R>(Int<0>{}, _));
}
注意，沿当前调用路径，这里的形参 tensor 实际接收 fragment 的 Layout，并不读取 fragment 的数值。
先定位 fragment 的范围
| 参数 | A / .x2 | B / .x1 |
|---|---|---|
R | 3 | 3 |
V=size<0>(tensor) | 8 | 4 |
TiledNumThr | 32 | 32 |
AtomNumVal | 4 | 2 |
Tiler_MN | (16,16)，实际为 M、K | (8,16)，实际为 N、K |
V 只统计第一维中一条 MMA 的线程内值数量，不是任意大 fragment 的总元素数。AtomNumVal 则来自 Copy 的 reference 布局；本例 reference 就是目的端，所以它也等于一次加载给每个 lane 写入的 half 数。
第一段计算：
right_inverse(TiledLayout_TV{}).with_shape(shape(Tiler_MN{}))
先将矩阵位置反查为展平的 TV 编号。线程维度在前且大小为 32，因此 tv=t+32*v。例如：
A(0,9) → lane 0、值 5 → tv=160
B(0,9) → lane 0、值 3 → tv=96
with_shape 将输入重新组织成矩阵的两维。接着 upcast<32*V> 以所有线程前 V 个值的规模为单位压缩布局，用来确定 Copy 还需要从哪些 fragment 位置收集值。
本例中，A 的 32*8=256 个值已经覆盖整个 16×16 tile，B 的 32*4=128 个值也覆盖整个 8×16 tile，所以两者都没有额外的收集位置：
product_each(shape(frg_layout_mn)) = (1,1)
right_inverse(frg_layout_mn) = 1:0
1:0 只有一个合法坐标 0，并映射到 0。这里的 (1,1) 表示只需要当前两个矩阵方向上的一个 fragment 位置。
再生成 Copy 的组内编号与组号
现在代入第二段：
auto frg_layout_v = zipped_divide(
logical_product(make_layout(V), right_inverse(frg_layout_mn)),
make_layout(AtomNumVal{}));
make_layout(V) 为已有的 V 个值生成编号；logical_product 接上 fragment 收集位置。本例收集位置只有一个，所以这个过程为：
| 阶段 | A / .x2 | B / .x1 |
|---|---|---|
make_layout(V) | 8:1 | 4:1 |
logical_product | (8,1):(1,0) | (4,1):(1,0) |
按 AtomNumVal 做 zipped_divide | (4,2):(1,4) | (2,2):(1,2) |
用 (a,g) 表示“组内 half、加载组”，结果分别生成：
A：v = a+4*g
B：v = a+2*g
此时 frg_layout_v 的输出是输入 tile 的逻辑编号。要找到实际存储槽位，还要接回输入 fragment 的布局。
将 Copy 编号组合到原来的存储布局
第三段先划分输入：
auto t_tensor = zipped_divide(
tensor, prepend(product_each(shape(frg_layout_mn)), V));
代入 (1,1) 后，tiler 分别是 (8,1,1) 和 (4,1,1)。它们取出当前第一维的全部值，同时把外层剩余部分留在第二个 mode：
A 的 t_tensor shape = ((8,1,1),(1,1,1))
B 的 t_tensor shape = ((4,1,1),(1,1,1))
本次 tile 剩余部分
再做组合：
auto v_tensor = t_tensor.compose(frg_layout_v, _);
以 A 为例，坐标如何落到存储的过程是：
Copy 分组坐标 (a,g)
↓ frg_layout_v
输入 tile 的逻辑编号 a+4*g
↓ t_tensor：沿用原 fragment 的布局
原 fragment 的 half 存储槽位 a+4*g
B 同理，使用 a+2*g。因此，这里的正确对应不只是来自共用 data()：分组编号经过原 fragment 布局求值，明确选中了它应当引用的槽位。
最后，本例 R=3，append<R>(Int<0>{},_) 生成 (0,_,_)。它固定剩余部分中第一维的坐标为 0，保留后面的两个矩阵重复方向：
return v_tensor(_, make_coord(Int<0>{}, _, _));
实际结果正是：
A / x2：((4,2),1,1):((1,4),0,0)
B / x1：((2,2),1,1):((1,2),0,0)
外层再接上原 data()，就得到加载使用的 tXrA/tXrB。
这段源码以输入 fragment 遵循 reference 的值顺序为前提；本例 R=D，适用于 retile_D。由于所有外层范围都是 1，结果看起来就是连续分组。更大的 tile 中，frg_layout_mn 还可能要求跨 fragment 收集，不能把通用 retile 一概替换成简单 reshape。想继续看这种情况，可以接着读原 retile 文章中的跨 fragment 例子。
两份 fragment 最后怎样交给 MMA
下面把基础配置放回一段完整的调用流程。假设 sA/sB 已准备好，gC 是按 (M,N) 索引的输出 tensor：
auto mma = make_tiled_mma(SM80_16x8x16_F16F16F16F16_TN{});
auto copy_a = make_tiled_copy_A(
Copy_Atom<SM75_U32x4_LDSM_N,half_t>{}, mma);
auto copy_b = make_tiled_copy_B(
Copy_Atom<SM75_U32x2_LDSM_N,half_t>{}, mma);
auto thr_mma = mma.get_slice(threadIdx.x);
auto thr_a = copy_a.get_slice(threadIdx.x);
auto thr_b = copy_b.get_slice(threadIdx.x);
auto tCrA = thr_mma.partition_fragment_A(sA);
auto tCrB = thr_mma.partition_fragment_B(sB);
auto tCrC = thr_mma.partition_fragment_C(gC);
auto tXsA = thr_a.partition_S(sA);
auto tXsB = thr_b.partition_S(sB);
auto tXrA = thr_a.retile_D(tCrA);
auto tXrB = thr_b.retile_D(tCrB);
copy(copy_a, tXsA, tXrA);
copy(copy_b, tXsB, tXrB);
clear(tCrC);
gemm(mma, tCrA, tCrB, tCrC);
copy(tCrC, thr_mma.partition_C(gC));
这段例子由一个 warp 执行。整个 warp 先执行 A 的加载，再执行 B 的加载；每条加载指令都由所有 lane 协作，并为每个 lane 写入自己的输出寄存器。A、B 分别准备在不同的 fragment 中。
mma_unpack 把两份 fragment 重新解释成打包操作数：
Tensor rA = recast<uint32_t>(A); // 4 个操作数
Tensor rB = recast<uint32_t>(B); // 2 个操作数
再调用指令封装。对本文这条 MMA，实际 PTX 形式为：
mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16
{d0,d1},
{a0,a1,a2,a3},
{b0,b1},
{c0,c1};
lane 0 的 a2 高 half 已经是 37，b1 高 half 已经是 53。MMA 按固定接口分别将它们作为 A(0,9) 和 B(0,9) 参与 warp 的矩阵乘加。
改用 A 的 .x2、B 的 .x1 时，前面的源分工和加载分组随之改变；这里 MMA 接收的操作数数量、槽位含义保持相同。
打印布局，再核对实际生成的指令
配套程序为 cute_mma_ldmatrix_ab_retile_trace.cu。它包含四种配置的主机端布局验证，以及三个用于编译检查的 GPU kernel。
在 CUTLASS 根目录编译并运行主机验证：
g++ -x c++ -std=c++17 \
-Iinclude -I/usr/local/cuda/include -I/usr/local/cuda/include/cccl \
cute_mma_ldmatrix_ab_retile_trace.cu \
-o /tmp/cute_mma_ldmatrix_ab_retile_trace
/tmp/cute_mma_ldmatrix_ab_retile_trace
将四个配置的关键目标输出整理在一起：
A/x4: group0 source(16,1) -> e129 -> q144 -> lane0 slot5 = 37
A/x2: group1 source(0,1) -> e1 -> q144 -> lane0 slot5 = 37
B/x2: group0 source(8,1) -> e65 -> q72 -> lane0 slot3 = 53
B/x1: group1 source(0,1) -> e1 -> q72 -> lane0 slot3 = 53
PASS: 4 configurations, 768 fragment slots and 384 packed bindings; host only.
程序按非转置 ldmatrix 的行、列分发规则独立计算路线，核对全部 lane 的源坐标、16 字节行对齐、连续 half 地址、目的槽位覆盖，以及 retile_D 与原 fragment 的指针别名。随后在主机模拟写入，核对实际数值和打包后的 half 位模式。
还可以只编译 GPU kernel，检查 PTX：
nvcc -std=c++17 -Iinclude -arch=sm_80 --ptx \
cute_mma_ldmatrix_ab_retile_trace.cu \
-o /tmp/cute_mma_ldmatrix_ab_retile_trace.ptx
实际生成的指令数量如下，每一项都是一个 warp 的一条 MMA 所需的加载：
| Kernel | A 的加载 | B 的加载 | MMA |
|---|---|---|---|
ab_x4_x2 | 1 条 .x4 | 1 条 .x2 | 1 条 |
ab_x2_x2 | 2 条 .x2 | 1 条 .x2 | 1 条 |
ab_x2_x1 | 2 条 .x2 | 2 条 .x1 | 1 条 |
主机程序验证布局、地址和存储绑定；PTX 验证真实模板实例化后的指令形式与数量。本次未在 GPU 上启动 kernel。
对照源码
本文依据 CUTLASS 提交 0b55a2f691d6。以下链接固定到该版本，避免后续 main 更新造成行号漂移。
| 要核对的内容 | 源码 |
|---|---|
| MMA 的 A、B TV 布局 | mma_traits_sm80.hpp |
| MMA 指令及 A、B 操作数类型 | mma_sm80.hpp |
thrfrg_A/B 如何采用 Atom 布局 | mma_atom.hpp |
get_layoutA_TV/get_layoutB_TV | mma_atom.hpp |
partition_fragment_A/B | mma_atom.hpp |
make_fragment_like 的布局构造 | layout.hpp |
| Copy 抽象编号的约定 | copy_traits.hpp |
.x1/.x2/.x4 的 Copy traits | copy_traits_sm75.hpp |
| bit 布局转换为元素布局 | copy_atom.hpp |
tidfrg_S、tile2thrfrg 与内部 retile | copy_atom.hpp |
partition_S、retile_D 与 make_tiled_copy_A/B | copy_atom.hpp |
copy_unpack 的打包输出绑定 | copy_traits.hpp |
三种非转置 ldmatrix 的封装 | copy_sm75.hpp |
| Copy 分组如何展开成多次 Atom 调用 | copy_atom.hpp、copy.hpp |
| MMA fragment 转换为指令操作数 | mma_traits.hpp |