读懂 CuTe 的 make_tiled_copy：从线程分工到 Reference 布局
读 CuTe 的 GEMM 代码时，常会看到先创建 Copy，再为当前线程取出源、目的视图：
auto thr_copy = tiled_copy.get_slice(threadIdx.x);
auto tS = thr_copy.partition_S(src);
auto tD = thr_copy.partition_D(dst);
copy(tiled_copy, tS, tD);
这几行代码要解决的是：当前线程搬哪些元素，从哪里读，又写到哪里。 make_tiled_copy 建立线程分工，partition 将分工应用到 tensor，最后的 copy 执行搬运。
本文先用一个小矩阵解释这条调用链，再沿官方 sgemm_sm80 的配置，说明 ldmatrix 怎样把数据送入 MMA 所需的寄存器。正文先讲分工和结果，后面的源码进阶部分再展开 get_layoutA_TV()。
源码对应 CUTLASS 提交 0b55a2f691d6。示例保留官方 MMA 与加载指令参数，存储简化为二维行主序，暂不展开 swizzle 和流水线。
1. 先分清指令、线程分工与存储位置
本文涉及的数据流是：
全局显存 ──cp.async──→ 共享内存 ──ldmatrix──→ 寄存器 ──→ mma.sync
S、D 是 Source、Destination 的缩写，表示当前 Copy 的源端、目的端：
| Copy 指令 | 源端 S | 目的端 D |
|---|---|---|
cp.async | 全局显存 | 共享内存 |
ldmatrix | 共享内存 | 寄存器 |
partition_S 中的 S 因而不固定表示某一种存储。具体空间由所用指令和传入的 tensor 决定。
CuTe 又将一次搬运拆成三个层次：
| 对象 | 负责的信息 |
|---|---|
Copy_Atom | 一条指令涉及多少线程、多少值，以及两端怎样对应 |
TiledCopy | 将 Atom 扩展到一个逻辑矩阵块，安排线程与元素的分工 |
| 传入的 tensor | 元素实际存在哪里，怎样从逻辑坐标计算地址 |
这里的 tile 就是矩阵块；TV 是 thread、value 的缩写。t 表示线程编号，v 表示该线程持有的值编号。一个 warp 有 32 个线程，warp 内编号叫 lane。
后文用 shape:stride 表示布局。例如 (4,32):(32,1) 将 (m,n) 映射为元素偏移 32*m+n。CuTe 把布局的一个维度称为 mode，mode 内还可以有子维度。打印结果省略静态整数的 _ 前缀，/ 均为整数除法。
2. 从 32 个线程搬运一个小矩阵开始
选择每线程每次复制 8 字节的 cp.async。half 是 16 位浮点数，以它为单位，一次就是 4 个值：
using namespace cute;
using Atom = Copy_Atom<
SM80_CP_ASYNC_CACHEALWAYS<uint64_t>, cutlass::half_t>;
auto thr_layout = Layout<Shape<_4,_8>,Stride<_8,_1>>{};
auto val_layout = Layout<Shape<_1,_4>>{};
auto tiled_copy = make_tiled_copy(Atom{}, thr_layout, val_layout);
uint64_t 指定指令的搬运宽度，half_t 指定 CuTe 的元素单位。两张输入布局的作用分别是：
| 布局 | 含义 | 编号公式 |
|---|---|---|
thr_layout=(4,8):(8,1) | 将 32 个线程排成 4×8 网格 | t=8*tm+tn |
val_layout=(1,4):(0,1) | 每个线程持有一个 1×4 值块 | v=vn，vm=0 |
将线程网格中的每个位置展开成一个 1×4 块，便得到 4×32 的 tile：
m = vm + 1*tm = t/8
n = vn + 4*tn = v + 4*(t%8)
例如：
| 线程 t | 持有的 4 个矩阵坐标 |
|---|---|
| 0 | (0,0)～(0,3) |
| 1 | (0,4)～(0,7) |
| 7 | (0,28)～(0,31) |
| 8 | (1,0)～(1,3) |
| 31 | (3,28)～(3,31) |
设 gA、sA 分别是全局显存和共享内存中的 4×32 tensor：
auto thr_copy = tiled_copy.get_slice(threadIdx.x);
auto tS = thr_copy.partition_S(gA);
auto tD = thr_copy.partition_D(sA);
copy(tiled_copy, tS, tD);
线程 1 从 gA 的第 0 行、第 4～7 列读取，再写到 sA 的相同逻辑位置。partition 创建视图，copy 发起复制。
两端的存储 stride 可以不同。例如 (1,0) 在行跨度为 32 的源 tensor 中偏移为 32，在行跨度为 40 的目的 tensor 中偏移为 40。线程分工选择逻辑坐标，tensor 的 layout 再将坐标换成地址。
3. make_tiled_copy 怎样算出这份分工
make_tiled_copy 的主体只有几步：
auto layout_mn = raked_product(thr_layout, val_layout);
auto layout_tv = right_inverse(layout_mn).with_shape(
make_shape(size(thr_layout), size(val_layout)));
auto tiler = product_each(shape(layout_mn));
return make_tiled_copy_impl(copy_atom, layout_tv, tiler);
目标是生成 (t,v) → 矩阵位置。源码先构造反方向的映射，再求逆。
raked_product：先展开每个线程的值块
raked_product 先做逻辑乘积，再按轴配对：
auto result = logical_product(append<R>(block), append<R>(tiler));
return zip(get<1>(result), get<0>(result));
本例 block 是线程布局，tiler 是值布局。中间结果与配对顺序为：
logical_product = ((4,8),(1,4)):((8,1),(0,32))
线程 值块
zip((vm,vn),(tm,tn)) = ((vm,tm),(vn,tn))
layout_mn = ((1,4),(4,8)):((0,8),(32,1))
每个轴中 value 在前，是因为源码先将值布局传给 zip。 zip 只保留实参顺序并配对。列轴的 shape 为 (4,8)，所以列号拆成：
vn = n%4，tn = n/4
n = vn + 4*tn
列： 0 1 2 3 | 4 5 6 7 | …
值：T0v0 T0v1 T0v2 T0v3 | T1v0 T1v1 T1v2 T1v3 | …
这样一个线程持有相邻 4 列。若按 (tn,vn) 配对，则 n=tn+8*vn，线程 0 会持有第 0、8、16、24 列；blocked_product 使用的就是后一种配对顺序。
layout_mn 的 stride 则给出输出编号：
tv = 0*vm + 8*tm + 32*vn + tn = t + 32*v
tv 是合并的 TV 编号：先编号 32 个线程，再编号下一组值，因此 t 先变化。这与矩阵轴内的 (value,thread) 配对各有用途，也不表示 GPU 执行顺序。上述紧凑编码来自本例完整、无空洞的输入排列。
right_inverse：从 TV 编号找回矩阵位置
对于 shape (4,32)，CuTe 的紧凑坐标编号是 q=m+4*n。这里 q 只编码逻辑坐标，实际地址仍由 tensor 的存储 stride 计算。代入前面的线程分工：
q = t/8 + 16*(t%8) + 4*v
= 16*(tv%8) + tv/8
这意味着，将 tv 拆成 (tv%8,tv/8)，再按 stride (16,1) 计算即可：
right_inverse(layout_mn) = (8,16):(16,1)
从 right_inverse 的实现 看，也能得到相同结果：
((1,4),(4,8)):((0,8),(32,1))
→ 去掉大小为 1 的 mode：(4,4,8):(8,32,1)
→ 合并连续的前两项： (16,8):(8,1)
令 x=m+4*v、y=tn，正向是 tv=8*x+y，原坐标编号是 q=x+16*y。求逆先接入 stride 为 1 的 y，再接入 stride 为 8 的 x，于是得到 shape (8,16)、输出 stride (16,1)。一般的 right_inverse 只覆盖从 stride 1 连续接起来的范围；本例完整覆盖 128 个位置。
with_shape：恢复独立的 thread、value 输入
with_shape 将输入恢复为 (32,4)：先计算 tv=t+32*v，再查询逆布局。结果是：
TiledLayout_TV = ((8,4),4):((16,1),4)
q = 16*(t%8) + t/8 + 4*v
例如 (t,v)=(1,2) 得到 q=24，还原为 (m,n)=(24%4,24/4)=(0,6)，与线程 1 持有第 4～7 列一致。
product_each(shape(layout_mn)) 则得到 tile 大小 (4,32)。TiledCopy 保存这张 TV 与 tile 大小，之后再由 partition 接入具体 tensor。
4. Reference 怎样连接 Copy 的两端
前面的 cp.async 两端采用相同分工。对于会在线程间分发数据的指令，两端的 lane、value 编号可能不同，需要一套共同编号来连接，这就是 reference 的用途。
Copy_Atom 将 traits 中的 bit 布局换成元素单位。TiledCopy 随后使用这些名称：
| 布局 | 输入 | 输出 |
|---|---|---|
AtomLayoutSrc | 源端 Atom 的 (thread,value) | 指令内部元素编号 e |
AtomLayoutDst | 目的端 Atom 的 (thread,value) | 同一套元素编号 e |
AtomLayoutRef | reference Atom 的 (thread,value) | 同一套元素编号 e |
TiledLayout_TV | 完整 reference 的 (thread,value) | tile 坐标编号 q |
前三张布局来自 Copy traits 和元素类型，最后一张由 TiledCopy 构造过程提供。e 用来连接指令两端，q 用来定位矩阵元素。
源端 partition 的查询链为：
源端 (thread,value)
→ AtomLayoutSrc：得到 e
→ AtomLayoutRef 的逆：得到 reference TV 编号
→ TiledLayout_TV：得到 q
→ tensor：得到实际地址
对应源码是 right_inverse(AtomLayoutRef{}).compose(AtomLayoutSrc{})；目的端将 Src 换成 Dst。reference 由 Atom 的 trait 规定，调用 partition_S/D 不会重新选择它。
cp.async Atom 的线程粒度
cp.async 的封装 让每个线程指定自己的 global、shared 地址，复制一个字节块。其 trait 因而使用 ThrID=Layout<_1>，并选择 RefLayout=SrcLayout。
本例以 half 为单位，三张 Atom 布局恰好相同：
AtomLayoutSrc = AtomLayoutDst = AtomLayoutRef = (1,4):(0,1)
AtomNumThr = 1，AtomNumVal = 4
一个 Atom 描述一个线程的 4 个 half；TiledCopy 将它扩展到 32 个线程。两端到 reference 的转换均为恒等，所以它们取得相同逻辑坐标。
partition 内部：拆 Atom、转换编号、恢复线程和值
tile2thrfrg 的核心是：
auto atom_layout_TV = zipped_divide(
TiledLayout_TV{}, make_shape(AtomNumThr{}, AtomNumVal{}));
auto trg_layout_TV = atom_layout_TV.compose(ref2trg, _);
auto thrval2mn = coalesce(
zip(trg_layout_TV), Shape<_1,Shape<_1,_1>>{});
先按一个 Atom 的线程数、值数拆分 TV：
((atom_thread,atom_value),(rest_thread,rest_value))
本例：((1,4),((8,4),1)):((0,4),((16,1),0))
rest 表示 Atom 之外的重复。本例有 32 个线程方向实例，每线程只有一组 4 值。rest_thread=(8,4) 继承求逆后的 thread mode，即 (t%8,t/8)，并非原输入 thr_layout 的 (4,8)。
随后 compose(ref2trg,_) 只转换 Atom 内的编号，保留外层重复。参数虽名为 ref2trg，实际传入的是“端点 TV → reference TV”的映射。
最后的 zip 把 thread、value 各自放到一起：
((atom_thread,atom_value),(rest_thread,rest_value))
↓ zip
((atom_thread,rest_thread),(atom_value,rest_value))
本例：((1,(8,4)),(4,1)):((0,(16,1)),(4,0))
coalesce 的第二个参数是结构 profile，要求保留以下边界：
Shape<_1,Shape<_1,_1>>{}
↓
(thread,(atom_value,rest_value))
coalesce 沿 profile 递归，在每个叶子 _1 对应的子布局内删除单位维、合并连续维。_1 允许在内部合并，并不要求结果大小为 1。本例的区别是：
带 profile： ((8,4),(4,1)):((16,1),(4,0))
不带 profile：(8,16):(16,1)
前者保留“线程、Atom 内值、值重复”的接口，后者只保留扁平顺序。随后 tensor.compose(thrval2mn,_) 将这份分工接入已经切好的 tensor，再取当前线程的切片。
5. ldmatrix：先确定目的寄存器需要什么
ldmatrix 将 shared memory 数据分发到 warp 各线程的寄存器。源 lane 提供行地址，目的 lane 接收值，二者的分工开始出现差别。
先看加载 Atom，再选择它要对接的计算需求：
using LoadAtom = Copy_Atom<SM75_U32x4_LDSM_N, cutlass::half_t>;
它封装的 PTX 指令 是：
ldmatrix.sync.aligned.x4.m8n8.shared.b16 {d0,d1,d2,d3}, [addr];
一个 warp 协作加载四个 8×8 的 16 bit 块，每个 lane 接收四个 32 位寄存器，即 8 个 half。该 trait 选择 RefLayout=DstLayout，转成 half 单位后为：
Src = (32,8):(8,1)
Dst = Ref = (32,(2,4)):(2,(1,64))
因此，本次 reference TV 应当描述：目的 lane 的第 v 个 half，需要接收矩阵中的哪个元素。 实际使用中，这份需求通常来自随后执行的 MMA。
用官方 MMA 配置提供目的需求
MMA 是矩阵乘加。M、N 表示输出矩阵的行、列方向，K 是乘积的累加方向。CuTe 将 A 按 (M,K) 索引，将 B 按 (N,K) 索引；因此下文 B 的 n 表示它参与计算哪一列输出。
官方 sgemm_sm80.cu 使用：
auto mma = make_tiled_mma(
SM80_16x8x16_F16F16F16F16_TN{},
Layout<Shape<_2,_2>>{},
Tile<_32,_32,_16>{});
单条 MMA 由一个 warp 协作计算 16×8×16；TiledMMA 将 4 个 warp 排成 M、N 方向的 2×2。这 128 个线程组成一个 CTA，即 CUDA 线程块，再通过线程内重复覆盖 32×32×16 的基本 tile：
M：16 × 2 个 warp × 1 次重复 = 32
N： 8 × 2 个 warp × 2 次重复 = 32
K：16 × 1 个 warp × 1 次重复 = 16
所以 N 的四个 n8 块由“2 个 N 方向 warp × 每 warp 2 次重复”覆盖。四个 warp 的分工为：
| warp | (wm,wn) | A 的 M 行 | B 的 N 行 |
|---|---|---|---|
| 0 | (0,0) | 0～15 | 0～7、16～23 |
| 1 | (1,0) | 16～31 | 0～7、16～23 |
| 2 | (0,1) | 0～15 | 8～15、24～31 |
| 3 | (1,1) | 16～31 | 8～15、24～31 |
线程编号是 t=l+32*wm+64*wn，l 是 lane。A 不依赖 wn，B 不依赖 wm：计算不同 C 列的 warp 复用 A，计算不同 C 行的 warp 复用 B。
从 MMA TV 得到 Copy reference TV
mma.get_layoutA_TV() 已经描述每个线程的 A 值需求，可直接作为本例的 reference TV：
auto reference_tv = mma.get_layoutA_TV();
auto tiled_copy = make_tiled_copy_impl(
LoadAtom{}, reference_tv, Shape<_32,_16>{});
这里 A tile 是 32×16。将每 lane 的 8 个值写成 v=v0+2*v1+4*v2，其中各 vi 为 0 或 1，这张 TV 的结果是：
m = l/4 + 8*v1 + 16*wm
k = 2*(l%4) + v0 + 8*v2
q = m + 32*k
本例 Ref=Dst，所以 Copy 的目的 lane l、值 v，接收 MMA 同一 lane、值 v 所需的元素。这里对齐的是输入编号；MMA Atom 输出矩阵坐标编号，Copy Ref 输出指令内部元素编号，二者的数值可以不同。
6. 从源行地址，到 MMA 的同一份寄存器
有了目的需求，就可以沿 Copy Atom 的固定规则反推加载地址，再把输出接到 MMA 已有的寄存器上。
partition_S：反推每个源 lane 的行地址
将源 lane 记为 s、行内值记为 u=0…7；目的寄存器编号为 r=0…3，寄存器内的 half 编号为 p=0,1。目的值编号是 v=p+2*r。两端通过共同的 Atom 元素编号 e 相连：
Src：e = 8*s + u
Ref：e = 2*l + p + 64*r
解出 reference 编号，再代入 MMA TV：
l = 4*(s%8) + u/2，p = u%2，r = s/8
m_src = s%16 + 16*wm
k_src = u + 8*(s/16)
于是 warp 0 的源分工是：
| 源 lane | 提供的 A 位置 |
|---|---|
| 0～7 | 第 0～7 行，K=0～7 |
| 8～15 | 第 8～15 行，K=0～7 |
| 16～23 | 第 0～7 行，K=8～15 |
| 24～31 | 第 8～15 行，K=8～15 |
每个 lane 提供连续 8 个 half 的行地址。硬件随后按 ldmatrix 的规则把这些值分发给目的 lane。例如：
源 (s=0,u=2)
→ Src 得到 e=2
→ Ref 的逆得到 tv=1，即目的 (l=1,v=0)
→ MMA 构造的 TiledLayout_TV 得到 q=64
→ A(0,2)
因此，源 lane 0 提供的一行中，A(0,2) 会进入目的 lane 1 的第一个寄存器低 half。源 lane 与目的 lane 的编号承担不同角色。
retile_D：加载视图与计算视图共享寄存器
设 sA 是待加载的 A tensor，先按 MMA 的规则准备寄存器，再创建 Copy 所需的目的视图：
auto thr_mma = mma.get_slice(threadIdx.x);
auto tCrA = thr_mma.partition_fragment_A(sA);
auto thr_copy = tiled_copy.get_slice(threadIdx.x);
auto tXsA = thr_copy.partition_S(sA);
auto tXrA = thr_copy.retile_D(tCrA);
partition_fragment_A 准备当前线程的寄存器存储；retile_D 为同一份存储提供适合 Copy 的分组方式。之后的调用关系为：
copy 经 tXrA 写入 → 同一份寄存器 → gemm 经 tCrA 读取
在基本 tile 上，A 每 lane 有 8 个 half，恰好对应一次 ldmatrix.x4 的输出。warp 0、lane 0 的具体对应是：
| ldmatrix 输出 | A 坐标，低 half、高 half | mma.sync 的 A 输入 |
|---|---|---|
d0 | (0,0)、(0,1) | a0 |
d1 | (8,0)、(8,1) | a1 |
d2 | (0,8)、(0,9) | a2 |
d3 | (8,8)、(8,9) | a3 |
copy_unpack 将目的 half 视图按位重解释为四个 uint32_t，交给 ldmatrix.x4；mma_unpack 对 A 视图作相同的重解释，再交给 mma.sync。
当前 MMA 封装 声明 ARegisters=uint32_t[4]、BRegisters=uint32_t[2]。上表保证了每个 A 操作数内的两个 half 也符合指令要求。实际物理寄存器由编译器分配，表中的名称表示操作数对应关系。
同一个 Copy Atom，也可以一次准备两份 B 片段
一条 MMA 每 lane 只需要 4 个 B half；ldmatrix.x4 每 lane 输出 8 个 half。前面的 N 方向恰有两次重复，因此将 Copy 值编号分为：
v = b + 4*repeat_n
b = v%4 单条 MMA 的 B 值编号
repeat_n = v/4 两份 N 方向片段
继续看 warp 0、lane 0：
| ldmatrix 输出 | B 的坐标 (n,k) | mma.sync 的 B 输入 |
|---|---|---|
d0 | (0,0)、(0,1) | N 重复 0 的 b0 |
d1 | (0,8)、(0,9) | N 重复 0 的 b1 |
d2 | (16,0)、(16,1) | N 重复 1 的 b0 |
d3 | (16,8)、(16,9) | N 重复 1 的 b1 |
构造方式仍然相同，只是换成 B 的需求：
auto copy_b = make_tiled_copy_impl(
LoadAtom{}, mma.get_layoutB_TV(), Shape<_32,_16>{});
auto tCrB = thr_mma.partition_fragment_B(sB);
auto tXrB = copy_b.get_slice(threadIdx.x).retile_D(tCrB);
随后两条 MMA 分别读取自己的两个 B 寄存器，并复用同一份 A。MMA Atom 与 Copy Atom 各自描述指令；组合时通过 TV 和寄存器视图建立对应。线程数、值数可以整除只是尺寸条件，正确性还取决于逐元素对应。
多个 warp 需要同一块数据时，各自加载一份
warp 0 和 warp 2 使用相同 A 块，Copy 会保留这份重复需求。partition_S 只转换 Atom 内的编号，外层 wn 的零 stride 仍然保留：
warp 0：线程 0 提供 A(0,0…7) 的行地址
→ 线程 1 接收 A(0,2)
warp 2：线程 64 提供同一个行地址
→ 线程 65 接收 A(0,2)
两个 warp 分别执行自己的 ldmatrix，将同一 shared memory 元素加载到各自的线程寄存器。随后 warp 0 用它计算一组 C 列，warp 2 计算另一组 C 列。相同线程内偏移也属于不同线程的存储。B 在 M 方向的复用同理。
7. 最后再看 make_tiled_copy_A
到这里，加载过程已经建立：由 MMA 给出目的需求，按 Copy 规则求源分工，再通过 retile_D 接入寄存器。官方用便捷接口封装了同一构造：
auto s2r_copy_a = make_tiled_copy_A(LoadAtom{}, mma);
make_tiled_copy_A 的函数体可展开为：
auto layout_tv = mma.get_layoutA_TV();
auto tiler = make_shape(tile_size<0>(mma), tile_size<2>(mma));
return make_tiled_copy_impl(copy_atom, layout_tv, tiler);
它与正文的显式构造得到相同类型。两个入口的区别在于 TV 从哪里来：
| 构造器 | reference TV 的来源 | tile 大小 |
|---|---|---|
make_tiled_copy(atom,thr,val) | 线程、值布局经 raked product 求逆 | 各轴合成后的大小 |
make_tiled_copy_A(atom,mma) | mma.get_layoutA_TV() | MMA tile 的 M、K 大小 |
_A 选择 MMA 的 A 操作数；相邻 _B、_C 分别选择 B 的 (N,K)、C 的 (M,N)。源端和目的端仍由具体 Copy 指令及 partition_S/D 区分。
这也给出了设计 TiledCopy 的顺序：先确认 Atom 的 reference 规则，再用这套 lane、value 编号描述期望分工，最后将两端的视图接到实际 tensor。下面继续展开前面用到的 MMA TV，以及较大 tensor 中的寄存器分组。
源码进阶一：get_layoutA_TV 怎样生成 MMA 的需求
这一节展开第 5 节直接使用的 mma.get_layoutA_TV()。可以先记住结果，再按下列五步追踪来源。源码的最终表达式是：
return thrfrg_A(ref_A)
.compose(atile, _)
.compose(thridx_2_thrid, _);
构造时先建立左侧 A 分工；查询 (t,v) 时从右往左进入。两个 compose(...,_) 只适配 thread mode，保留 value mode。
1. 建立 A 坐标表，接入 tile 的排列
auto ref_A = make_layout(
make_shape(tile_size_mnk<0>(), tile_size_mnk<2>()));
auto t_tile = make_tile(permutation_mnk<0>(), permutation_mnk<2>());
auto t_tensor = logical_divide(ref_A, t_tile);
本例 ref_A=(32,16):(1,32)，输出坐标编号 q=m+32*k。t_tile=(32,16) 将每个轴分成基本范围和外层重复：
t_tensor = ((32,1),(16,1)):((1,0),(32,0))
结构： ((perm_m,outer_m),(perm_k,outer_k))
这里先用 logical_divide，是为了保留 M、K 两个顶层轴，下一步才能继续分别切分。zipped_divide 还会把 tile 部分与剩余部分收拢，此时的层级会不同。
ref_A 与 t_tile 虽都来自 permutation，作用也不同：tile_size_mnk 只取其大小，t_tile 保留实际排列。官方这个纯尺寸配置中，第一次 divide 不改变元素顺序；通用实现仍需保留它。例如替换为：
using PermM = Layout<Shape<_2,_16>, Stride<_16,_1>>;
Tile<PermM,_32,_16>{};
ref_A 大小仍相同，但 M 的顺序变为 0,16,1,17,…，即 P(u)=16*(u%2)+u/2。divide 后的编号是 q=P(u)+32*k；省略它就会丢掉这份排列。配套程序也核对了该配置。
2. 切出一个 MMA Atom，再接入 Atom TV
thrfrg_A 按单条 MMA 的 A 大小 16×16 切分：
auto a_tile = make_tile(make_layout(_16{}), make_layout(_16{}));
auto a_tensor = zipped_divide(t_tensor, a_tile);
auto tv_tensor = a_tensor.compose(AtomLayoutA_TV{}, _);
zipped_divide 是 tile_unzip(logical_divide(...),tiler)。在这一步，收拢恰好能让第一个 mode 完整表示一个 Atom：
((atom_m,block_m),(atom_k,block_k))
↓ 收拢
a_tensor = ((atom_m,atom_k),(block_m,block_k))
实际布局： ((16,16),(2,1)):((1,32),(16,0))
然后将第一个 mode 替换为 (lane,atom_v)。所用 MMA Atom 布局 为：
AtomLayoutA_TV = ((4,8),(2,2,2)):((32,1),(16,8,128))
atom_m = l/4 + 8*v1
atom_k = 2*(l%4) + v0 + 8*v2
它在单个 Atom 中输出 q_atom=atom_m+16*atom_k。接入 M=32 的坐标表后，与 K 有关的 stride 扩大一倍：
tv_tensor = (((4,8),(2,2,2)),(2,1)):
(((64,1),(32,8,256)),(16,0))
结构： ((lane,atom_v),(block_m,block_k))
此时已经有 Atom 内的线程分工，两个 M 块还没有分给 warp。
3. thr_tile：把块坐标分给 warp
auto thr_tile = make_tile(_,
make_tile(make_layout(_2{}), make_layout(_1{})));
auto thr_tensor = zipped_divide(tv_tensor, thr_tile);
2、1 来自 M、K 方向的 warp 份数。它们匹配 block_m、block_k；每份块坐标内部已经有 32 个 lane：
block_m = wm + 2*rest_m
block_k = wk + 1*rest_k
logical_divide 后：
((lane,atom_v),((wm,rest_m),(wk,rest_k)))
再由 tile_unzip 收拢：
((lane,(wm,wk)),(atom_v,(rest_m,rest_k)))
thread value
外层 _ 在 logical_divide 阶段保留 (lane,atom_v)；在后续收拢时，二者仍分别进入 thread、value。make_layout(_2{}) 是 2:1，make_layout(_1{}) 是 1:0，本例 wk=0。
先接入 Atom TV，再做这次收拢，是为了让第一个二元组已经是 (lane,atom_v)。若提前作用于 a_tensor，分开的会是 atom_m、atom_k，就不能继续原样复合 Atom TV。
同一规则也用于较大的 A tensor：
| A 大小 | Atom 块数 (block_m,block_k) | 分给 warp | 剩余值重复 |
|---|---|---|---|
基本 tile 32×16 | (2,1) | (2,1) | (1,1) |
官方 CTA 的 128×64 | (8,4) | (2,1) | (4,4) |
4. atile：让 N 方向的 warp 查询同一份 A
前一步的 thread 只有 (lane,wm,wk)。完整 CTA 还包含 wn，但 A 的坐标不依赖它。get_layoutA_TV 用下面的 tiler 适配：
auto atile = make_tile(_,
make_tile(Layout<Shape<_2,_2>,Stride<_1,_0>>{}, _));
其中 (2,2):(1,0) 将 (wm,wn) 映射为 wm：
(lane,((wm,wn),wk)) → (lane,(wm,wk)) → A 坐标编号
因此只改变 wn 会查询相同的 A 元素。这里从 64 种 A 线程坐标适配到 128 个完整线程坐标，表达的就是两个 N 方向 warp 各需要一份 A。CTA 始终有 128 个线程。
5. thridx_2_thrid：从实际线程编号进入这张表
get_layoutA_TV 最后接入实际线程编号。将源码拆开可写为：
auto L = thr_layout_vmnk_;
auto C = complement(L);
auto F = make_layout(L, C);
auto project = make_layout(make_shape(size(L), _1{}),
make_stride(_1{}, _0{}));
auto thridx_2_thrid = composition(project, right_inverse(F));
L 将 VMNK 坐标映射为实际线程编号；right_inverse(F) 将编号还原，project 再丢掉补全用的辅助坐标。官方布局本来就连续：
L = (32,2,2,1):(1,32,64,0)
complement(L) = 1:128
thridx_2_thrid = 128:1
所以本例 thridx_2_thrid(t)=t。它返回紧凑编号，再由左侧 thread shape 解码出 l=t%32、wm=(t/32)%2、wn=t/64、wk=0。
先求补集，是为了兼容线程编号带空洞的情况。以 L=(2,2):(1,4) 为例，有效编号是 0、1、4、5，直接 right_inverse(L)=2:1 只能覆盖前两个连续编号。补全后：
C = complement(L) = (2,1):(2,8)
F = make_layout(L,C) = ((2,2),(2,1)):((1,4),(2,8))
right_inverse(F) = (2,2,2):(1,4,2)
project = (4,1):(1,0)
C 的偏移 0、2 补齐 0…7，原有效编号经逆映射和投影后分别恢复为 0、1、2、3。若 N=size(L)，补全后的紧凑编号为 z=i+N*c，投影通过 z%N+0*(z/N) 保留 i。补集服务于编号求逆，N 方向广播则已由 atile 完成。
整条链最后可以用一个查询核对：线程 109 的值 6 解码为 l=13、wm=1、wn=1，v0=0、v1=1、v2=1，得到 (m,k)=(27,10)、q=347。线程 45 只有 wn 不同，因此取得同一个元素。
省略大小为 1 的 value mode，最终布局为：
MMA_A_TV = ((4,8,2,2),(2,2,2)):((64,1,16,0),(32,8,256))
thread： (l%4,l/4,wm,wn) value：(v0,v1,v2)
源码进阶二：核对更大的寄存器视图与 reference
完整 CTA 中的 A、B 寄存器分组
将 A、B 都扩展到官方 CTA 的二维大小 128×64，存储均简化为 (128,64):(64,1)。配套程序实际打印的寄存器布局如下：
A：tCrA = ((2,2,2),4,4):((1,2,4),32,8)
tXrA = ((8,1),4,4):((1,0),32,8)
B：tCrB = ((2,2),8,4):((1,2),16,4)
tXrB = (((4,2),1),4,4):(((1,16),0),32,4)
前三个 mode 依次表示“指令内值、M 或 N 方向重复、K 方向重复”。每线程 A、B 都有 128 个 half，但 Copy 与 MMA 使用的分组不同：
A：tXrA(v,m,k) ↔ tCrA(v,m,k)
B：tXrB(v,j,k) ↔ tCrB(v%4, 2*j+v/4, k)
B 的 MMA 视图是每片 4 个值、N 重复 8 次；Copy 视图是每片 8 个值、N 重复 4 次。j=k=0 时，Copy 前 4 个值引用原寄存器元素偏移 0…3，后 4 个引用偏移 16…19。retile_D 用视图将两份片段交给一次 Copy，无须搬动已有数据。
B 的二维位置也能从 Atom TV 直接核对：
AtomLayoutB_TV = ((4,8),(2,2)):((16,1),(8,64))
b = b0 + 2*b1
n = l/4 + 8*wn + 16*repeat_n
k = 2*(l%4) + b0 + 8*b1
其中 block_n=wn+2*repeat_n，正是 N 方向 2 个 warp 与每 warp 2 次重复的分配方式。在基本 tile 上，A、B、C 的每线程分区大小分别是 ((2,2,2),1,1)、((2,2),2,1)、((2,2),1,2)，三者都有 8 个值；C 的两个重复沿 N 展开，全部线程恰好覆盖 32×32 的输出。
已知源分工时，先转换成 Atom 的 reference 编号
设 Atom 的三张布局为 S、D、R，reference 到 tile 的布局为 C_R。对本文的双射 Atom，有：
C_S = C_R ∘ R^-1 ∘ S
C_D = C_R ∘ R^-1 ∘ D
已知源分工：C_R = C_S_wanted ∘ S^-1 ∘ R
F∘G 表示先执行 G，再执行 F；代码中使用 right_inverse，转换只作用于 Atom 内，外层重复保留。因此 R=S 时可直接填入源分工，R=D 时可直接填入目的分工。
当前 ldmatrix 的 R=D。如果直接将它的源 TV 当作 reference TV：
auto wrong = make_tiled_copy_impl(
LoadAtom{}, tiled_copy.get_layoutS_TV(),
typename decltype(tiled_copy)::Tiler_MN{});
源 (s=0,u=2) 会先转为 reference (l=1,v=0)，错误布局却将它解释为源 lane 1 的值 0，于是把 (0,2) 错放成 (1,0)。源 lane 0 的访问变为：
期望：(0,0) (0,1) (0,2) (0,3) (0,4) (0,5) (0,6) (0,7)
实际：(0,0) (0,1) (1,0) (1,1) (2,0) (2,1) (3,0) (3,1)
这个反例中线程数、值数、tile 大小全部相同，错误来自 TV 输入编号的含义与 reference 不一致。
用配套程序验证
完整程序是 cute_make_tiled_copy_reference_trace.cu。在 CUTLASS 根目录编译运行，CUDA 路径按安装位置调整：
g++ -x c++ -std=c++17 \
-Iinclude -I/usr/local/cuda/include -I/usr/local/cuda/include/cccl \
cute_make_tiled_copy_reference_trace.cu \
-o /tmp/cute_make_tiled_copy_reference_trace
/tmp/cute_make_tiled_copy_reference_trace
程序在主机端打印布局并逐元素核对：
| 配置 | 验证范围 |
|---|---|
| 小矩阵 cp.async | 128 个坐标、逆映射及带 padding 的目的地址 |
| 官方 cp.async 参数 | 1024 个源、目的坐标 |
| 官方 MMA 与 ldmatrix 的 A | TV 逐层构造、基本 tile 的两份 A 副本，以及完整 CTA 的 16384 个寄存器别名和源地址 |
| 同一配置的 B、C | B 的 N 方向分配与 M 方向复用、C 的完整覆盖，以及 B 的 16384 个寄存器别名、Copy/MMA 坐标对应和源地址 |
| 通用布局 | 带 stride 的 permutation，以及带空洞线程编号的补全与投影 |
它也检查 make_tiled_copy_A 与显式构造的类型一致，并验证错误 reference 反例。成功时输出：
PASS: make_tiled_copy layout checks only; no GPU instruction executed.
这些检查验证布局与分区，不执行 GPU 指令。检查自己的配置时，可以先对 identity tensor 分区，直接观察逻辑坐标：
auto coords = make_identity_tensor(make_shape(_32{},_16{}));
auto thr = tiled_copy.get_slice(threadIdx.x);
print_tensor(thr.partition_S(coords));
print_tensor(thr.partition_D(coords));