读懂 CuTe 的 make_tiled_copy：一块矩阵怎样分给线程
读 CuTe 的数据搬运代码时，经常能看到这样的配置：
auto tiled_copy = make_tiled_copy(
Copy_Atom<UniversalCopy<uint128_t>, float>{},
make_layout(make_shape(_4{}, _8{}), make_stride(_8{}, _1{})),
make_layout(make_shape(_1{}, _4{}), make_stride(_4{}, _1{})));
三个参数分别涉及 Copy Atom、线程布局和值布局。只看名字，大概能猜到“让一批线程合作搬数据”，但很难直接回答：线程 0 搬哪几个元素？线程 8 呢？这里的 128 位，与最后那个 _4 又有什么关系？
先把结果画出来。假设要搬的是一个 4×32 的 float 矩阵，记作 A，每行连续存储。上面的配置会让 32 个线程这样分工：
列 0～3 列 4～7 列 28～31
第 0 行 线程 0 线程 1 … 线程 7
第 1 行 线程 8 线程 9 … 线程 15
第 2 行 线程 16 线程 17 … 线程 23
第 3 行 线程 24 线程 25 … 线程 31
每个线程负责同一行中相邻的 4 个 float，32 个线程一共负责 128 个元素。make_tiled_copy 建立的就是这套分工，以及如何按 Copy Atom 的要求组织这些值。
下面沿着这个例子读源码，再把矩阵扩大，看看同一套配置怎样继续工作。本文讨论的接口在源码中叫 make_tiled_copy。
先看线程最后拿到了什么
本例只使用一个 warp，也就是 32 个线程。将线程编号记为 t，它在自己那一份数据中的值编号记为 v：
t = 0～31
v = 0～3
线程 t 的值 v，对应 A(t/8, 4*(t%8)+v)
这里的除法是整数除法。把几个线程的结果列出来：
| 线程 | 值 0 | 值 1 | 值 2 | 值 3 |
|---|---|---|---|---|
| 0 | A(0,0) | A(0,1) | A(0,2) | A(0,3) |
| 1 | A(0,4) | A(0,5) | A(0,6) | A(0,7) |
| 7 | A(0,28) | A(0,29) | A(0,30) | A(0,31) |
| 8 | A(1,0) | A(1,1) | A(1,2) | A(1,3) |
| 31 | A(3,28) | A(3,29) | A(3,30) | A(3,31) |
先记住这张表。后面那些布局运算，最终都要落实到这里的矩阵位置上。
另外，创建 tiled_copy 时还没有读取 A，也没有分配保存 A 的寄存器。它甚至没有接收 A 的指针。实际使用时，还要取得线程切片，再把这套分工应用到 Tensor：
auto thr_copy = tiled_copy.get_slice(threadIdx.x);
auto tS = thr_copy.partition_S(source);
auto tD = thr_copy.partition_D(destination);
source、destination 是已有的 Tensor。partition_S/D 为当前线程创建它们的访问视图；真正的数据搬运发生在后面的 copy 中。
三个参数各自规定什么
读布局时，本文用 shape:stride 表示“形状与步长”。例如，(4,8):(8,1) 将坐标 (m,n) 换算为 8*m+n。CuTe 把一个维度称为 mode，mode 也可以嵌套。
线程布局：给线程排座位
第二个参数是：
auto thr_layout = make_layout(
make_shape(_4{}, _8{}),
make_stride(_8{}, _1{}));
它的形状是 (4,8)，表示把线程排成 4 行、8 列；步长 (8,1) 表示：
thr_layout(tm,tn) = 8*tm + tn
tm = 0～3，tn = 0～7
因此第一行是线程 0～7，第二行是线程 8～15。给定线程编号 t，可以反过来找到它的座位：
tm = t/8
tn = t%8
这里的映射方向要看清楚：thr_layout 的输入是二维坐标，输出是线程编号。它还没有规定这个线程搬哪几个矩阵元素。
值布局：每个座位放一小块数据
第三个参数是：
auto val_layout = make_layout(
make_shape(_1{}, _4{}),
make_stride(_4{}, _1{}));
形状 (1,4) 表示每个线程负责一行中的 4 个值。它将值块中的坐标 (vm,vn) 映射为该线程的值编号：
val_layout(vm,vn) = 4*vm + vn
vm 只能为 0，vn = 0～3，因此 v = vn。
沿 N 方向，线程 0 的小块占 4 列，接下来是线程 1 的 4 列，依此类推。于是整个分工覆盖的矩阵块大小为：
M = 线程行数 × 每线程值块的行数 = 4 × 1 = 4
N = 线程列数 × 每线程值块的列数 = 8 × 4 = 32
val_layout 的 stride 描述的是值编号的顺序。矩阵在内存里怎样排列，还要由 Tensor 自己的 layout 决定。两者在本例中配合起来，才得到每线程 4 个连续 float 的访问。
Copy Atom：这些值怎样搬
第一个参数是：
Copy_Atom<UniversalCopy<uint128_t>, float>{}
UniversalCopy<uint128_t> 表示按一个 128 位对象进行复制，float 表示从 CuTe Tensor 的角度，每个逻辑元素是 32 位浮点数。于是一个 Atom 处理：
128 bit / 32 bit = 4 个 float
这个 Atom 只需要一个线程，源端和目的端各有 4 个 float。可以先把它理解成“当前线程一次复制一组 16 字节的数据”。
这与 ldmatrix 那种整个 warp 协作的指令不同。这里的 32 个线程各自执行自己的 Copy Atom，合起来完成整个矩阵块的搬运。
源码中的 UniversalCopy::copy 本身很简单：
static constexpr void copy(S const& src, D& dst)
{
dst = src;
}
具体生成什么加载、存储指令，由编译器结合类型和地址空间决定。uint128_t 规定这里的复制宽度，并不绑定一条叫作“128 位 Copy”的 PTX 指令。本文的实际编译结果会在后面列出。
make_tiled_copy 的主体只有几步
当前 copy_atom.hpp 中，去掉调试打印后的实现如下：
template <class... Args,
class ThrLayout,
class ValLayout = Layout<_1>>
CUTE_HOST_DEVICE
auto constexpr
make_tiled_copy(Copy_Atom<Args...> const& copy_atom,
ThrLayout const& thr_layout = {},
ValLayout const& val_layout = {})
{
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);
}
前面两个布局的方向分别是“坐标到线程编号”和“坐标到值编号”。函数先把它们合成一张分工表，再把方向反过来，得到实际使用时更方便的查询方式：
矩阵位置 → 由哪个线程的哪个值负责
│ 反向查询
▼
给定线程编号和值编号 → 矩阵位置
这里的 TV 就是 thread、value。下面把具体数字代入。
raked_product 怎样把两张布局拼起来
第一行是：
auto layout_mn = raked_product(thr_layout, val_layout);
对本例，它将每个矩阵方向组织成“值块内部坐标在前，线程坐标在后”：
M 方向：(vm,tm)，大小为 (1,4)
N 方向：(vn,tn)，大小为 (4,8)
CuTe 用整数访问嵌套 shape 时，前面的子维度先变化。因此，矩阵坐标与它们的关系是：
m = vm + 1*tm = tm
n = vn + 4*tn
这就解释了为什么同一个线程的 4 个值会占据连续的 4 列。
再看布局的输出。32 个线程的编号占据 0～31，接着把值编号放在更高的位置：
layout_mn(m,n) = t + 32*v
这是一种将 (t,v) 编成整数的方法。它让线程编号先变化：前 32 个编号对应所有线程的值 0，再后面 32 个对应所有线程的值 1。
把 t=8*tm+tn 和 v=vn 代入：
layout_mn((vm,tm),(vn,tn)) = 128*vm + 8*tm + 32*vn + tn
vm 的范围只有 0，第一项实际总是 0。源码打印出的完整布局是：
layout_mn = ((1,4),(4,8)):((128,8),(32,1))
大小为 1 的维度，其 stride 可以保留一个看起来很大的数字；因为没有第二个合法坐标，它不会让访问越过边界。
例如 A(0,4)，有 tm=0、tn=1、vn=0，所以输出 1，表示线程 1 的值 0。A(0,1) 的输出则是 32，表示线程 0 的值 1。
从实现看，这个顺序从哪里来
对本例这种 rank 都为 2 的布局，raked_product 的主要动作可以写成：
auto result = logical_product(thr_layout, val_layout);
return zip(get<1>(result), get<0>(result));
实际源码还会将两个输入补到相同的 rank。logical_product 先保留线程和值两部分，得到：
((4,8),(1,4)):((8,1),(128,32))
线程 值
对本例，值编号的步长乘上了 32，为线程编号留出空间。随后 zip 把同一矩阵方向的子维度配起来，并让值子维度在前：
((线程 M,线程 N),(值 M,值 N))
↓ zip，值在前
((值 M,线程 M),(值 N,线程 N))
输出仍然编码 t+32*v，改变的是输入坐标怎样组织。
为什么不直接用 blocked_product
比较这两个操作，更容易看出 raked 在这里的作用。
blocked_product(thr_layout,val_layout) 会把线程子维度放在前面。其 N 方向的关系变成：
n = tn + 8*vn
于是同一个线程负责的列变为：
| 组合方式 | 线程 0 的列 | 线程 1 的列 |
|---|---|---|
raked_product | 0,1,2,3 | 4,5,6,7 |
blocked_product | 0,8,16,24 | 1,9,17,25 |
两种方式都覆盖相同的 4×32 矩阵。在本例按行连续存储的前提下，第一种让每个线程拿到连续的 16 字节，适合选定的 128 位 Atom；第二种将该线程的值隔开了。
这也是为什么仅仅算对“32 个线程，每线程 4 个值”还不够：这些值怎样分布，会影响它们能否组成一次向量访问。
right_inverse：从线程和值反查矩阵位置
接下来是：
auto layout_tv = right_inverse(layout_mn).with_shape(
make_shape(size(thr_layout), size(val_layout)));
刚才得到的是：
layout_mn：矩阵坐标 → t+32*v
而线程执行代码时，通常已经知道自己是线程 t，希望找到值 v 对应哪个矩阵位置。因此需要沿相反方向查询。
本例的布局没有重叠或空洞，正好将 128 个矩阵位置与 128 个 TV 编号一一对应。这里可以把 right_inverse 当作反向查询表。一般布局可能存在广播、空洞等情况，不能据此把它理解成任意布局都存在的完整逆函数。
逆布局输出的整数，怎样还原成坐标
这里有一个容易与内存地址混淆的约定：对于 4×32 的逻辑 shape，CuTe 用整数 q 表示坐标时，按第一个 mode 先变化的方式解码：
q = m + 4*n
m = q%4
n = q/4
这个编号由 shape 的坐标顺序决定，与 A 是否按行连续存储是两件事。
代入线程分工：
m = t/8
n = 4*(t%8)+v
q = t/8 + 16*(t%8) + 4*v
所以我们希望逆布局将 t+32*v 映射为上面的 q。
实际得到：
right_inverse(layout_mn) = (8,16):(16,1)
给它一个整数 i=t+32*v，它会先将 i 分成 (i%8,i/8)，再计算：
16*(i%8) + i/8
= 16*(t%8) + t/8 + 4*v
= q
结果与刚才按矩阵坐标推出来的一致。
为什么源码恰好构造出 (8,16):(16,1)
前面从线程分工验证了结果。这里再按 right_inverse 的实现，逐项解释 shape 中的 8、16 和 stride 中的 16、1 从哪里来。
输入是：
layout_mn = ((1,4),(4,8)):((128,8),(32,1))
第一步，coalesce(layout_mn) 去掉 shape-1，并合并可连续表达的 mode。
源码开头：
auto clayout = coalesce(layout);
auto lstride = wrap(clayout.stride());
auto lshape = wrap(clayout.shape());
先展平，得到：
shape = (1,4,4,8)
stride = (128,8,32,1)
去掉 size 为 1 的维度以后是 (4,4,8):(8,32,1)。前两个 mode 满足：
4 * 8 = 32
因此可以合并为 16:8；最后一个 mode 的 stride 为 1，无法继续按相邻连续条件合并。结果是：
clayout = (16,8):(8,1)
此处 coalesce 没有交换两个 mode。对它的坐标 (x,y)：
layout_mn(x,y) = 8*x + y
结合原来的坐标含义：
x = m + 4*v，范围 0..15
y = tn，范围 0..7
原矩阵坐标的默认编号仍是：
q = m + 4*n
= m + 4*v + 16*tn
= x + 16*y
第二步，为输入各 mode 计算默认坐标编号的步长。
源码：
auto preprod_shape = cute::fold(
lshape, cute::tuple<_1>{},
[](auto c, auto vi) { return append(c, vi*back(c)); });
对 lshape=(16,8)：
preprod_shape = (1,16,128)
其中前两项分别是输入 mode 0、mode 1 在默认坐标编号中的 stride：
q = 1*x + 16*y
最后的 128 是总大小，不是额外的输出 mode。
第三步，按原布局的地址 stride 排序，找到从 1 开始的连续覆盖链。
本例的 stride 都是静态的，排序结果如下：
| 排序后顺序 | 原 mode | extent | 原布局 stride | 默认坐标编号 stride |
|---|---|---|---|---|
| 0 | 1，即 y | 8 | 1 | 16 |
| 1 | 0，即 x | 16 | 8 | 1 |
所以 sorted_seq=(1,0)。
源码中的 fold 从以下状态开始：
result_shape = (1)
result_stride = (0)
curr_stride = 1
遍历某个原 mode 时，只有 istride == curr_stride 才接上它：
if constexpr (is_constant<decltype(istride)::value,
decltype(curr_stride)>::value) {
return make_tuple(
append(result_shape, ishape),
append(result_stride, get<i>(preprod_shape)),
ishape * istride);
}
注意新布局的 stride 取自 preprod_shape，而不是直接照抄排序后的物理/编码 stride。
本例的状态变化为：
| 步骤 | 本次 mode | 接入条件 | 新 result_shape | 新 result_stride | 新 curr_stride |
|---|---|---|---|---|---|
| 初始 | — | — | (1) | (0) | 1 |
| 接入 y | extent=8，stride=1 | 1==1 | (1,8) | (0,16) | 8 |
| 接入 x | extent=16，stride=8 | 8==8 | (1,8,16) | (0,16,1) | 128 |
最后源码调用 coalesce，去掉最前面的 size-1 mode：
(1,8,16):(0,16,1)
↓
(8,16):(16,1)
这就是最终结果。
用一行坐标代换也可以直接得到同一个布局：
正向：i = 8*x + y
逆向：y = i % 8，x = i / 8
输出：q = 16*y + x
y 先变化、范围为 8，x 后变化、范围为 16，所以逆布局的 shape 是 (8,16)；它输出 16*y+x，所以 stride 是 (16,1)。
| 输入 TV 编号 i | (y,x)=(i%8,i/8) | 逆布局输出 q | 原矩阵 (m,n)=(q%4,q/4) |
|---|---|---|---|
| 0 | (0,0) | 0 | (0,0) |
| 1 | (1,0) | 16 | (0,4) |
| 7 | (7,0) | 112 | (0,28) |
| 8 | (0,1) | 1 | (1,0) |
| 32 | (0,4) | 4 | (0,1) |
| 127 | (7,15) | 127 | (3,31) |
例如 i=32 代表线程 0 的值 1；逆布局输出 q=4，解码为矩阵位置 (0,1)。这不是物理地址 4；对本例 row-major 存储，(0,1) 的物理元素偏移为 1。
一般情况下 right_inverse 依赖能够从 stride 1 开始接起来的连续链，不是“将 stride 排序以后任意交换 shape/stride”就能得到全域逆函数。本例连续覆盖 128 个位置，所以逆布局也恰好包含 128 个输入编号。
with_shape 把输入重新分成 thread 和 value
逆布局暂时接受一个合并编号。with_shape((32,4)) 则让我们能够直接使用 (t,v)：
(t,v) → t+32*v → q
with_shape 的源码是布局复合：
return composition(*this, make_layout(shape));
这里的 make_layout((32,4)) 使用默认紧凑布局 (32,4):(1,32)，正好计算 t+32*v。
最终得到：
layout_tv = ((8,4),4):((16,1),4)
第一个顶层 mode 一共有 8×4=32 个线程。CuTe 保留了其中的嵌套结构，t 在里面拆成 (t%8,t/8)；第二个顶层 mode 是 4 个值。
因此：
layout_tv(t,v) = 16*(t%8) + t/8 + 4*v
再把结果按 (4,32) 解码，就得到文章开头的：
A(t/8, 4*(t%8)+v)
例如 layout_tv(1,0)=16，它表示逻辑坐标 (0,4)。在按行连续存储的 A 中，A(0,4) 的实际元素偏移是 4。这里的 16 是坐标编号，要经过 Tensor 的布局，才能变成存储偏移。
tiler 与 TiledCopy 中保存了什么
下一行是：
auto tiler = product_each(shape(layout_mn));
layout_mn 的 shape 为 ((1,4),(4,8))。product_each 分别计算每个顶层 mode 内部的乘积：
product_each(((1,4),(4,8))) = (4,32)
这告诉后续的分区操作：当前这套线程和值分工，每次覆盖一个 4×32 的矩阵块。
最后的 make_tiled_copy_impl 只是构造对象：
return TiledCopy<Copy_Atom<Args...>, LayoutCopy_TV, Tiler>{atom};
到这里，TiledCopy 拿到了三部分信息：
| 信息 | 本例的内容 | 用途 |
|---|---|---|
| Copy Atom | 一个线程复制 4 个 float | 规定一次 Atom 操作的要求 |
TiledLayout_TV | ((8,4),4):((16,1),4) | 从参考线程和值找到 tile 内位置 |
Tiler_MN | (4,32) | 从输入 Tensor 中划出重复的矩阵块 |
线程和值的数量也分成 Atom 与整个 tile 两层：
AtomNumThr = 1 TiledNumThr = 32
AtomNumVal = 4 TiledNumVal = 4
本例相当于沿线程方向铺开了 32 个 Atom，每个线程在这个 tile 内恰好执行一份 Atom 工作。
TiledCopy 会检查两项整除关系：
TiledNumThr % AtomNumThr == 0;
TiledNumVal % AtomNumVal == 0;
它们约束 Atom 能否按数量铺开。实际访问的连续性、对齐和有效范围，还要结合输入 Tensor 检查。
partition_S/D 怎样接上真正的数据
现在给分工表接上一个实际布局：
// source_ptr 是 float 指针；本例的行跨度为 32 个 float。
auto source = make_tensor(
make_gmem_ptr(source_ptr),
make_layout(make_shape(_4{}, _32{}),
make_stride(_32{}, _1{})));
auto thr_copy = tiled_copy.get_slice(t);
auto tS = thr_copy.partition_S(source);
source 的存储偏移为：
offset(m,n) = 32*m+n
对于线程 t 的值 v，完整的查询过程是：
(t,v)
│ TiledLayout_TV
▼
逻辑编号 q = 16*(t%8)+t/8+4*v
│ 按 tile shape (4,32) 解码
▼
矩阵坐标 (t/8,4*(t%8)+v)
│ source.layout()
▼
存储偏移 32*(t/8)+4*(t%8)+v
这条路径把“线程拿什么”与“这些元素存在哪里”连起来了。
先把 Tensor 按 tile 划开
partition_S 内部会调用 TiledCopy::tidfrg_S。它先做：
zipped_divide(stensor, Tiler_MN{})
这里内部的 stensor 是输入 Tensor 的 layout。对于当前完整的 4×32 矩阵，划分结果为：
shape = ((4,32),(1,1))
stride = ((32,1),(0,0))
tile 剩余部分
输入恰好只有一个 tile，M、N 方向的剩余次数都是 1。先保留这两个 mode，是为了让同一段代码也能处理更大的矩阵。
为什么源端和目的端各有一个 partition
tidfrg_S 还会计算：
right_inverse(AtomLayoutRef{}).compose(AtomLayoutSrc{})
这条映射先从 Atom 的源端线程和值找到它表示的数据，再找到对应的参考线程和值。TiledLayout_TV 描述的是参考分工，接上这条转换后，才能得到源端的实际分工。
目的端使用相同的过程，只是把 AtomLayoutSrc 换成 AtomLayoutDst。
本例的 UniversalCopy 没有跨线程分发：
AtomLayoutSrc = (1,4):(0,1)
AtomLayoutDst = (1,4):(0,1)
AtomLayoutRef = (1,4):(0,1)
三者相同，转换相当于保持原来的线程和值。因此源端与目的端负责的逻辑坐标相同。
上一篇的 ldmatrix 则体现了这个区分的必要性：提供源地址的 lane，与最终收到对应值的 lane 可以不同。不能把本例中 Src、Dst 相同的情况推广到所有 Copy Atom。
TiledLayout_TV 要按参考端解释，不随 partition_S/D 切换
这里有两件事要分开：构造时如何定义整块 tile 的参考分工，以及 使用时如何把参考分工转换成源端、目的端分工。
对于同一个 TiledCopy，只有一个 TiledLayout_TV。不能在调用 partition_S 时把它当源端布局、调用 partition_D 时又把它当目的端布局；也不需要为两次 partition 更换它。
但设计 TiledLayout_TV 时，必须知道当前 Copy Atom 的参考端是什么，以及自己想要约束源端还是目的端。源端与目的端的 TV 分工不同的时候，不能把其中任意一侧的布局直接当作 reference TV。
两种“参考布局”不是同一个概念
| 对象 | 输入 → 输出 | 谁规定 |
|---|---|---|
make_cotiled_copy 的 data_layout | tensor 坐标 → 参考地址偏移 | 构造器调用方 |
AtomLayoutRef | reference Atom 的 (thread,value) → Atom 元素编号 | Copy_Traits 和元素类型 |
TiledLayout_TV | reference 线程和值 → tile 坐标编号 | TiledCopy 构造过程 |
AtomLayoutSrc/Dst | 源/目的 Atom 的 (thread,value) → 同一 Atom 元素编号 | Copy_Traits 和元素类型 |
data_layout 可以用源 tensor、目的 tensor 或另一个合适的静态参考布局来组织坐标，但这不会自动改变 Copy_Traits::RefLayout。二者不能因为都称“参考布局”就互相替代。
原始 Copy_Traits 的 Src/Dst/Ref 是 bit layout，Copy_Atom<...,Element> 会通过 recast_layout 将其转换成以 Element 为单位的 AtomLayoutSrc/Dst/Ref，见 copy_atom.hpp:61。下面的式子使用转换后的元素布局。
参考端由 Copy Atom trait 规定
using AtomLayoutSrc = typename Copy_Atom::ValLayoutSrc;
using AtomLayoutDst = typename Copy_Atom::ValLayoutDst;
using AtomLayoutRef = typename Copy_Atom::ValLayoutRef;
using AtomNumThr = decltype(size<0>(AtomLayoutRef{}));
using AtomNumVal = decltype(size<1>(AtomLayoutRef{}));
using TiledLayout_TV = LayoutCopy_TV;
因此 zipped_divide(TiledLayout_TV{},make_shape(AtomNumThr{},AtomNumVal{})) 按 reference Atom 的线程数和值数 切分。
例如：
| Copy Atom | trait 中的 reference 选择 |
|---|---|
| 本文 UniversalCopy | Ref、Src、Dst 相同 |
SM80_CP_ASYNC_CACHEALWAYS | RefLayout=SrcLayout，Src/Dst 的元素关系也相同 |
SM75_U32x4_LDSM_N | RefLayout=DstLayout，Src 与 Dst 的线程和值分布不同 |
所以“参考布局选源还是目的”通常是具体 Copy Atom trait 已经做出的选择，不是每次 partition 的运行时开关。
源端、目的端怎样转换成统一 reference 坐标
设：
R：reference Atom 坐标 → Atom 元素编号
S：source Atom 坐标 → Atom 元素编号
D：destination Atom 坐标 → Atom 元素编号
partition_S 使用：
right_inverse(AtomLayoutRef{}).compose(AtomLayoutSrc{})
方向是：
source 的 (thread,value)
│ S
↓
Atom 元素编号
│ right_inverse(R)
↓
reference 的 (thread,value) 编号
partition_D 则将 S 换成 D，得到 destination → reference 的转换。
因此，虽然 tile2thrfrg 的参数名叫 ref2trg，本处传入表达式的输入实际是源/目的端坐标，输出是 reference 坐标编号。不要只按变量名推断函数方向。
将整块 reference TV 分工记为 C_R，忽略符号中未变化的 Atom 重复坐标，则：
源分工：C_S = C_R ∘ right_inverse(R) ∘ S
目的分工：C_D = C_R ∘ right_inverse(R) ∘ D
这些转换按 Atom 内部执行，tile 中 Atom 的重复部分保持不变。
如果 Ref=Src，源端转换数值上保持对应关系，目的端通过 Src^-1∘Dst 推导出来；如果 Ref=Dst，目的端保持对应关系，源端通过 Dst^-1∘Src 推导出来。对于含有广播等非一一映射的 trait，应按源码的 right_inverse 和对应合法域理解，不能假设每个布局都有双侧逆。
一个 Src 和 Dst 确实不同的 ldmatrix 例子
SM75_U32x4_LDSM_N 的 trait 定义 RefLayout=DstLayout。
取 Copy_Atom<SM75_U32x4_LDSM_N,uint32_t>，忽略大小为 1 的子维度后：
S(s,v) = 4*s + v
R(t,v) = D(t,v) = t + 32*v
其中线程范围为 0..31，值范围为 0..3。这里的编号用于解释 Atom 内部元素对应关系，不是声称两端使用相同的物理存储。
选择整块 reference 分工：
C_R = (32,4):(1,32)
它对应目的端：reference/destination lane t 的第 v 个值映射到 tile 元素 t+32*v。
源码自动计算出来的两端 TV 布局为：
get_layoutS_TV = (32,(4,1)):(4,(1,0))
get_layoutD_TV = (32,(4,1)):(1,(32,0))
例如同一个 tile 元素 32：
源：lane 8，value 0 → 4*8+0 = 32
目的：lane 0，value 1 → 0+32*1 = 32
提供源端数据的 lane 和最终收到该值的 lane 不相同。正因为有 R^-1∘S 这一步，同一个 reference TV 才能生成这两种不同的 partition。
如果我想先规定源端的分工，应该怎么做
假设当前 Atom 的 Ref=Dst，而你手头设计的是 source 分工 C_S_wanted。不能直接把这个函数填进 reference TiledLayout_TV，因为 source partition 还会再执行一次 R^-1∘S。
在相关 Atom 映射可逆的情况下，可以按以下关系转换 reference 分工：
C_R = C_S_wanted ∘ S^-1 ∘ R
即先从 reference 坐标找到 Atom 元素，再找到它的 source 坐标，最后应用期望的 source 分工。整块 tile 的实现仍需按 Atom 与重复部分的层级处理，不能只对不兼容的完整 TV 布局随意作一次标量复合。
上面的 ldmatrix 例子中，期望源端为 C_S_wanted(s,v)=4*s+v，正确 reference 分工恰好是 C_R(t,v)=t+32*v。
如果误把 4*t+v 直接当 reference 分工，源 lane 8/value 0 会先转到 reference lane 0/value 1，再被这个错误分工映射为 1；期望的元素却是 32。这不是“切分时再选择一次 source”可以解决的问题。
设计完成后，可以用：
tiled_copy.get_layoutS_TV();
tiled_copy.get_layoutD_TV();
检查最终两端的逻辑分工。它们是使用 reference TV 和各自 Atom 转换实际推导出的布局，比仅检查 TiledLayout_TV 更能直接反映源端或目的端是否符合预期。
make_tiled_copy_S/D 也不改变新 Atom 的 RefLayout
源码的两个辅助构造器分别取旧 copy 的某一侧分工：
make_tiled_copy_impl(new_atom, old_copy.get_layoutS_TV(), old_tiler);
make_tiled_copy_impl(new_atom, old_copy.get_layoutD_TV(), old_tiler);
它们将旧 copy 的 source 或 destination 分工作为新 copy 的 reference TV，适合在不同搬运阶段衔接已有分工。后缀 S/D 表示从旧 copy 的哪一侧取布局，不是将新 Atom 的 RefLayout 强行改成 Src 或 Dst；新 copy 的两端转换仍由新 Atom 的 trait 决定。
所以要同时区分三件事：你希望约束哪一侧、当前 Atom 规定哪一侧为 reference、最终 get_layoutS/D_TV() 实际得到什么。
tile2thrfrg 把矩阵坐标改写成线程和值
接着，tile2thrfrg 将参考 TV 布局按 Atom 的范围划开，转换成源端或目的端的 TV 布局，再复合到刚才的 tile 布局上。关键操作是：
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>>{});
auto tv_tensor = tensor.compose(thrval2mn, _);
return tv_tensor(make_coord(_,_), _);
此处变量 ref2trg 沿用了源码中的名字；传入的表达式实际接收源端或目的端 TV 坐标，输出参考编号，正是前面说明的转换方向。
这里需要把 zip 和带 profile 的 coalesce 分开理解。它们分别改变坐标的分组和简化的范围。
zip 将“按 Atom 分组”转为“按线程和值分组”
trg_layout_TV 此时接受：
((trg_tid,trg_val),(rest_tid,rest_val))
当前 Atom 内部 Atom 的重复部分
zip 将相同用途的坐标配对：
((trg_tid,trg_val),(rest_tid,rest_val))
↓ zip
((trg_tid,rest_tid),(trg_val,rest_val))
线程 值
zip(Layout) 对 shape 和 stride 同时做相同的 tuple 转置：
return make_layout(zip(layout.shape()), zip(layout.stride()));
底层 tuple zip 相当于：
((a,b),(c,d)) → ((a,c),(b,d))
这里重排的是布局的输入坐标层级，不会搬动 tensor 数据，也不创建跨线程通信。对应坐标的映射保持一致：
trg_layout_TV((ta,va),(tr,vr))
== zip(trg_layout_TV)((ta,tr),(va,vr))
但不应直接断言 zip 前后的 layout(i) 对同一个标量 i 一定相等，因为层级重排可能改变标量输入的解码顺序。
Shape<_1,Shape<_1,_1>>{} 是递归 profile，不是目标 shape
源码的准确写法是：
Shape<_1,Shape<_1,_1>>{}
最外层的 _1 是类型，不能写成整数模板参数 1；Shape 的参数是类型。
这个 profile 的树结构是：
(
终端, // 线程部分：在这个完整子布局内 coalesce
(终端, 终端) // 值部分：分别 coalesce Atom 内部值和重复值
)
它控制“递归到哪里，然后在该处执行普通 coalesce”，并不要求结果 shape 等于 (1,(1,1))。
coalesce(layout,profile) 的分支直接表明这一点：
if constexpr (is_tuple<IntTuple>::value) {
return transform_layout(layout, trg_profile,
[](auto const& l, auto const& t) { return coalesce(l,t); });
} else {
return coalesce(layout);
}
profile 为 tuple 时继续逐 mode 递归；profile 为标量终端时，调用不带 profile 的 coalesce(layout)。这个终端的具体数值没有参与计算。因此本例使用 (7,(7,7)) 也得到相同结果：起作用的是树结构。
对本句而言，可按下面的等价结构理解：
auto z = zip(trg_layout_TV);
auto thrval2mn = make_layout(
coalesce(get<0>(z)),
make_layout(
coalesce(get<1,0>(z)),
coalesce(get<1,1>(z))));
它保留了顶层“线程 / 值”的边界，也保留了值内部“Atom 内 / Atom 重复”的边界。
coalesce 能合并什么，不能合并什么
普通 coalesce 先展平指定子布局，再删除静态 size-1 mode，并合并符合连续条件的相邻 mode。例如：
(2,4):(1,2) → 8:1，因为 2*1 == 2
(8,4):(16,1) 保留两个 mode，因为 8*16 != 1
因此“合并线程部分”不等于“把线程部分强行变成单个 32 维”。如果 stride 不满足合并条件，线程部分仍然会有多个 mode。
profile 中保留的子布局即使 coalesce 成标量，也会由外层 make_layout 再组织起来。所以 value 的第二部分即使只有一个元素，也会以独立的 1:0 分组保留下来。
按本例的具体数值展开
本例 AtomNumThr=1、AtomNumVal=4，source、destination 与 reference 的 Atom 布局相同，转换不会改变元素关系。中间布局为：
TiledLayout_TV =
((8,4),4):((16,1),4)
trg_layout_TV =
((1,4),((8,4),1)):((0,4),((16,1),0))
zip(trg_layout_TV) =
((1,(8,4)),(4,1)):((0,(16,1)),(4,0))
分别处理三个 profile 终端：
| profile 终端 | 输入子布局 | coalesce 后 |
|---|---|---|
| 线程整体 | (1,(8,4)):(0,(16,1)) | (8,4):(16,1) |
| Atom 内的值 | 4:4 | 4:4 |
| Atom 重复的值 | 1:0 | 1:0 |
得到的中间结果是：
thrval2mn = ((8,4),(4,1)):((16,1),(4,0))
线程 值
值 mode 的 (4,1) 表示“一次 Atom 内有 4 个值，这个 tile 内每线程只有 1 次 Atom”。它与前面的 layout_tv 表达相同的参考元素关系，同时保留了 Copy 执行需要的分组。
设线程编号 t=0..31，将其在线程 mode 中解码为 (t%8,t/8)，则：
thrval2mn(t,(va,vr)) = 16*(t%8) + t/8 + 4*va
其中 va=0..3、vr=0，输出仍是本 tile 的逻辑矩阵坐标编号 q=m+4*n。变量名里的 mn 不是说这个函数一定直接返回 C++ 二元 tuple (m,n)；本例返回的是由 tile shape 解码的整数编号。
为什么不能随手省掉这个 profile
同一输入采用三种方式处理，本例实际得到：
| 调用 | 结果 | 保留的语义边界 |
|---|---|---|
本句的 (1,(1,1)) | ((8,4),(4,1)):((16,1),(4,0)) | 线程 / 值，以及 Atom 内值 / 重复值 |
更浅的 (1,1) | ((8,4),4):((16,1),4) | 仅线程 / 值 |
| 不传 profile | (8,16):(16,1) | 不保证保留线程和值边界 |
无 profile 的结果碰巧与前面的 right_inverse(layout_mn) 一样，但它已经把原来独立的线程子维度和值子维度合并了，不能再按此处的 (thread,(atom_value,rest_value)) 接口直接使用。
更浅的 (1,1) 则会把整个值部分当成一个子布局处理。虽然本例 rest_value=1，后续代码仍统一保留二元值分组；每线程重复多次 Atom 时，这一分组更直接地区分一条指令内的值与重复指令所覆盖的值。
例如 FP16 cp.async 的每线程 16 个元素，单条 16 字节指令只处理 8 个元素，需要保留 (8,2)。其具体布局和 FP8 对照见 Q copy 文档的 tile2thrfrg 说明。
为什么下一句要 tensor.compose(thrval2mn, _)
此时 tensor 已经是按 copy tiler 分开的布局：
tensor：((TileM,TileN,...),(RestM,RestN,...)) → 真实 tensor 偏移
thrval2mn 把线程和值映射到第一个 tile mode 内的逻辑坐标，所以：
auto tv_tensor = tensor.compose(thrval2mn, _);
只替换第一个输入 mode；_ 保留第二个“tile 外剩余部分”。得到：
((线程,(Atom内值,Atom重复值)), 剩余tile坐标) → 真实 tensor 偏移
最后：
return tv_tensor(make_coord(_,_), _);
将第一个组合 mode 内的线程和值展开成两个顶层 mode，tile 外的剩余坐标仍保留在后面的 mode。调用方 ThrCopy::partition_S/D 随后固定当前线程编号，得到该线程的 fragment 视图。
经过与实际存储布局的复合，tidfrg_S 的结果为：
shape = ((8,4),(4,1),(1,1))
stride = ((4,32),(1,0),(0,0))
线程 值 剩余部分
注意线程 mode 的 stride 已经从逻辑坐标编号的 (16,1) 变成实际存储偏移的 (4,32)。例如线程 1 的首地址比线程 0 多 4 个 float，线程 8 则多 32 个 float。
最后，ThrCopy::partition_S 用原来的数据构造视图，并选中当前线程：
auto thr_tensor = make_tensor(
static_cast<STensor&&>(stensor).data(),
TiledCopy::tidfrg_S(stensor.layout()));
return thr_tensor(thr_idx_, _, repeat<rank_v<STensor>>(_));
对于 rank 为 2 的输入，最后的切片可以按下面的方式理解：
thr_tensor(t, _, make_coord(_, _));
返回视图的布局为：
shape = ((4,1),1,1)
stride = ((1,0),0,0)
切片同时将数据指针推进到当前线程的起始位置。例如，线程 1 的 tS.data() 从 A(0,4) 开始，线程 8 从 A(1,0) 开始。相同的返回 layout，并不意味着这些线程访问同一处数据。
返回 tensor 的三个顶层 mode 可以读成：
((Atom 内部的值, tile 内重复的 Atom), M 方向剩余次数, N 方向剩余次数)
创建这些分区视图仍然没有执行加载。
真正的 copy 在哪里发生
将源和目的都按上述方式划分后，就可以通过一个寄存器 fragment 完成搬运：
// kernel 的约定：一个 block，32 个线程；两个设备指针均按 16 字节对齐。
__global__ void copy_4x32(float const* source_ptr, float* destination_ptr)
{
using namespace cute;
auto tiled_copy = make_tiled_copy(
Copy_Atom<UniversalCopy<uint128_t>, float>{},
make_layout(make_shape(_4{}, _8{}), make_stride(_8{}, _1{})),
make_layout(make_shape(_1{}, _4{}), make_stride(_4{}, _1{})));
auto layout = make_layout(make_shape(_4{}, _32{}),
make_stride(_32{}, _1{}));
auto source = make_tensor(make_gmem_ptr(source_ptr), layout);
auto destination = make_tensor(make_gmem_ptr(destination_ptr), layout);
auto thr_copy = tiled_copy.get_slice(threadIdx.x);
auto tS = thr_copy.partition_S(source);
auto tD = thr_copy.partition_D(destination);
auto fragment = make_fragment_like(tD);
copy(tiled_copy, tS, fragment);
copy(tiled_copy, fragment, tD);
}
每个线程创建的 fragment 保存自己的 4 个 float。第一次 copy 从源 Tensor 加载，第二次写入目的 Tensor。
copy(tiled_copy,...) 会使用 TiledCopy 继承的 Copy Atom，遍历分区视图里的剩余部分。Copy_Atom::call 再检查当前这组值能否交给底层复制操作，copy_unpack 将数据按 uint128_t 重新解释，调用 UniversalCopy::copy。
所以 make_tiled_copy 并不自动启动线程，也不在一次函数调用里调度另外 31 个线程。kernel 中的线程各自取得切片，各自执行相应的加载和存储。
本例按 SM80 目标编译后，核心 PTX 是：
ld.global.v2.u64 {r0,r1}, [source_address];
st.global.v2.u64 [destination_address], {r0,r1};
这里为了便于阅读简化了寄存器名。v2.u64 表示两个 64 位量，合计 128 位，即 4 个 float 的原始位。编译器保留了向量宽度；这两条 PTX 分别是加载和存储。PTX 是虚拟指令集，具体机器指令还要经过后续编译。
矩阵变大以后，多出来的两个 mode 才开始工作
保持三个参数不变，只把输入 Tensor 扩大成 8×128，仍然按行连续存储：
auto layout = make_layout(make_shape(_8{}, _128{}),
make_stride(_128{}, _1{}));
TiledCopy 的基本分工仍然覆盖 4×32。所以现在沿 M 方向有 2 个 tile，沿 N 方向有 4 个 tile：
8×128 的输入
= M 方向重复 2 次 × N 方向重复 4 次 × 每块 4×32
对于同样的 32 个线程，partition_S 返回：
shape = ((4,1),2,4)
stride = ((1,0),512,32)
将后两个坐标记为 rm、rn，线程 t 的访问就是：
tS(v,rm,rn) → A(t/8+4*rm, 4*(t%8)+v+32*rn)
v = 0～3
rm = 0～1
rn = 0～3
M 方向跨一个 tile，要走 4 行，每行跨度 128，所以 stride 是 4×128=512；N 方向跨一个 tile，要走 32 列，所以 stride 是 32。
例如线程 0 会负责：
第 0 行：列 0～3、32～35、64～67、96～99
第 4 行：列 0～3、32～35、64～67、96～99
每个线程一共持有 4×2×4=32 个 float。对一次源到寄存器的 copy，需要处理 8 组，每组 4 个值；随后写回目的端也处理相同的 8 组。
这也说明 Tiler_MN=(4,32) 表示基本分工覆盖的大小。一个 block 可以继续用这套分工遍历更大的输入 Tensor。若希望不同 block 各处理一个 tile，则要在外层根据 block 编号先切出相应的 Tensor；这个构造函数本身没有规定 grid 如何安排。
每线程 8 个值，会不会自动变成 256 位复制
再做一个小改动：线程布局保持不变，Copy Atom 也不变，只将值布局改成 (1,8)：
auto tiled_copy_8 = make_tiled_copy(
Copy_Atom<UniversalCopy<uint128_t>, float>{},
make_layout(make_shape(_4{}, _8{}), make_stride(_8{}, _1{})),
make_layout(make_shape(_1{}, _8{}), make_stride(_8{}, _1{})));
现在每线程在基本 tile 中负责连续 8 个 float，整个 tile 扩大为 4×64：
TiledNumVal = 8
AtomNumVal = 4
Tiler_MN = (4,64)
对于恰好一个 4×64 的按行连续存储 Tensor，分区结果为：
shape = ((4,2),1,1)
stride = ((1,4),0,0)
第一维的 (4,2) 已经把答案写出来了：每个 Atom 处理 4 个值，一个线程在 tile 内有 2 组 Atom 工作。
线程 0 的两组分别对应：
第 0 组：A(0,0)、A(0,1)、A(0,2)、A(0,3)
第 1 组：A(0,4)、A(0,5)、A(0,6)、A(0,7)
每线程负责的值数量改变了，指定的 Atom 宽度仍然是 128 位。Copy_Atom::call 会沿嵌套的第一维继续拆分，直到每组满足一次 Atom 的要求。
回头看上一篇的 retile_D，它也是在准备这种“Atom 内部的值与 Atom 分组”的寄存器视图。两篇文章的连接点正在这里：TiledCopy 给出 Copy 所需的组织方式，retile_D 再让已有的 MMA fragment 按这种方式访问。
同样的分工，换一种存储布局会怎样
前面的连续访问有一个前提：A 的 layout 为 (4,32):(32,1)，N 方向连续。
如果改为：
make_layout(make_shape(_4{}, _32{}))
CuTe 默认使用第一个 mode 连续的紧凑布局，即 (4,32):(1,4)。线程的逻辑分工仍然相同，但线程 0 的四个值在内存中的偏移变成：
逻辑元素：A(0,0) A(0,1) A(0,2) A(0,3)
元素偏移： 0 4 8 12
对应分区布局为：
((4,1),1,1):((4,0),0,0)
这四个值之间有间隔，不能直接重解释为一个连续的 uint128_t。构造分工和创建视图仍然可以完成，但使用这个固定宽度 Atom 执行复制时，会遇到布局无法满足向量化要求的编译期检查。
因此配置固定宽度 Copy 时，需要一起看三个具体条件：线程负责哪些值、这些值在源端和目的端是否按 Atom 要求组织，以及实际起始地址是否满足对齐。
本例每组 4 个 float 共 16 字节，源端和目的端的组起始地址都需要按 16 字节对齐。4×32、8×128 两种按行连续存储的例子，其行跨度和每线程起始列都保持这种对齐。
本文的尺寸也都能被基本 tile 整除。make_tiled_copy 不会自动替最后一块生成边界保护；处理尾部数据时，需要在调用处安排有效范围和相应的谓词复制，且固定宽度 Atom 的整组访问仍须有效。
与 make_tiled_copy_A/B 的关系
上一篇使用的是：
auto tiled_copy_a = make_tiled_copy_A(copy_atom, mma);
它与本文的构造函数最终都调用 make_tiled_copy_impl，区别在于参考分工从哪里来。
本文从 thr_layout 和 val_layout 出发，先自己定义线程与值的二维排列，再计算 TV 布局。make_tiled_copy_A 则直接采用 MMA 已经规定好的 A 操作数分工：
return make_tiled_copy_impl(
copy_atom,
mma.get_layoutA_TV(),
make_shape(tile_size<0>(mma), tile_size<2>(mma)));
make_tiled_copy_B 同理，采用 B 的 TV 布局，矩阵方向为 CuTe 约定的 (N,K)。
当数据要进入 MMA fragment 时，指令对哪个 lane 持有哪些值已经有明确要求，直接采用 MMA 的分工比较自然。本文的接口则适合从希望的线程排列和值块排列出发，定义一套搬运分工。
无论参考分工来自哪里，后续都需要用 Tensor 的布局找到实际数据，再按 Copy Atom 的源端、目的端约定执行搬运。
打印布局，核对一次
配套程序放在 cute_make_tiled_copy_trace.cu。在 CUTLASS 根目录运行：
nvcc -std=c++17 -arch=sm_80 -I include \
cute_make_tiled_copy_trace.cu -o /tmp/make_tiled_copy_trace
/tmp/make_tiled_copy_trace
程序会打印构造和分区过程中的布局，遍历全部 32 个线程核对坐标、覆盖范围和源目的对应关系，并在主机内存上实际执行经过 fragment 的复制。这个主机检查不需要启动 GPU kernel。
针对本文补充的两处推导，程序还会打印 right_inverse 输入的 coalesce 结果、前缀乘积和 stride 排序后的 mode 编号，逐一检查 128 个 TV 编号的逆映射；对 zip + coalesce，它会打印 zip 结果和三种 profile 的对照，并枚举对应结构坐标验证元素映射不变。
参考端部分还使用 SM75_U32x4_LDSM_N 的实际 trait 核对全部 128 个源/目的 TV 映射，以及误将 source TV 直接作为 destination reference TV 的反例。这里只检查布局，不执行 ldmatrix 指令。
构造阶段的关键输出为：
thr_layout: (_4,_8):(_8,_1)
val_layout: (_1,_4):(_4,_1)
logical_product: ((_4,_8),(_1,_4)):((_8,_1),(_128,_32))
layout_mn: ((_1,_4),(_4,_8)):((_128,_8),(_32,_1))
right_inverse: (_8,_16):(_16,_1)
layout_tv: ((_8,_4),_4):((_16,_1),_4)
tiler: (_4,_32)
分区与扩展例子的关键输出为：
partition_S: ((_4,_1),_1,_1):((_1,_0),_0,_0)
partition_D: ((_4,_1),_1,_1):((_1,_0),_0,_0)
large partition_S: ((_4,_1),_2,_4):((_1,_0),_512,_32)
eight-value partition_S: ((_4,_2),_1,_1):((_1,_4),_0,_0)
column-major partition_S: ((_4,_1),_1,_1):((_4,_0),_0,_0)
Verified base tile: 128 coordinates and 128 copied floats.
Verified repeated tile: 1024 coordinates and 1024 copied floats.
Verified two Atoms per thread: 256 coordinates and 256 copied floats.
CuTe 打印静态整数时带 _ 前缀，正文省略了这个前缀。
同一文件还提供了 cute_copy_4x32、cute_copy_8x128 两个 kernel，便于检查设备代码。生成 PTX：
nvcc -std=c++17 -arch=sm_80 -I include -ptx \
cute_make_tiled_copy_trace.cu -o /tmp/make_tiled_copy_trace.ptx
本次编译中，4×32 kernel 包含一条 ld.global.v2.u64 和一条 st.global.v2.u64；8×128 kernel 展开为各 8 条，对应每线程的 8 组数据。两个 kernel 也通过了 GPU 实际运行的结果核对。文件中的 main 只执行前述主机检查。
对照源码
- copy_atom.hpp：
make_tiled_copy的构造过程；同文件还包含TiledCopy、tile2thrfrg与ThrCopy::partition_S/D。 - layout.hpp：
raked_product；同文件可继续查看logical_product、right_inverse、with_shape和zipped_divide。 - layout.hpp:1305：
right_inverse的前缀乘积、stride 排序和连续链构造。 - layout.hpp:880：带 profile 的
coalesce如何递归到终端。 - layout.hpp:1557：Layout 的
zip同时转置 shape 和 stride。 - copy_traits.hpp：UniversalCopy 的线程与位布局，以及
copy_unpack的向量化检查。 - copy.hpp：底层操作：
UniversalCopy的类型与复制操作。 - copy.hpp：复制算法：分区视图怎样遍历并调用 Copy Atom。
- tiled_copy.cu：官方教程中的线程分区与向量复制用法。
- 上一篇：读懂 CuTe 的 retile_D：已有 MMA fragment 如何转换为 Copy 需要的寄存器访问视图。
本文依据当前工作区的 CUTLASS/CuTe C++ 实现，源码提交为 0b55a2f691d69981583568fd9eb69687b1f0de8a。