# 从源码理解 Q 的 `make_cotiled_copy`：为什么 `replace` 使用 `head_dim`
本文对应 CUTLASS 提交 `0b55a2f6`，沿 `77_blackwell_fmha_gen.cu` 的 decode/generation 路径展开。主要使用当前示例的 FP8 输入、`D=128`；最后对照 FP16。
这里讨论的是 `Sm100FmhaLoadCpAsyncWarpspecialized::load()` 中的调用，不能直接套用 prefill 的 `run_fwd_128()`、`TileShape=(256,128,128)` 或其 `ThreadShape`。
## 1. 核心结论
调用点：
```cpp
auto tiled_copy_q = make_cotiled_copy(
Copy_Atom<SM80_CP_ASYNC_CACHEALWAYS<uint128_t>, Element>{},
atom_q_tv,
make_layout(
shape(tSgQ),
replace<0>(
stride(tSgQ),
stride<0>(tSgQ),
get<2>(TileShape{})))));
```
来源：[sm100_fmha_load_cpasync_warpspecialized.hpp:176](/home/x/workspace/cutlass/examples/77_blackwell_fmha/collective/sm100_fmha_load_cpasync_warpspecialized.hpp:176)。
两层 `replace<0>` 实际替换的是 **fragment 内 M 方向的 stride**，不是 MMA 的 `ThrV` stride。
在当前紧凑存储的示例中，这个 stride 的运行时数值已经等于 `head_dim`。替换将它从普通 `int` 变成类型中编码的 `cute::Int<D>`：
```text
运行时 int，值为 128
↓
编译期 cute::Int<128>
这样构造出的 `data_layout` 才满足 `make_cotiled_copy` 对静态布局的要求。
选择 `head_dim` 的依据是复制参考布局的行组织：
一行 = 一个 Q head 的 D 个连续特征
参考地址偏移 = m * D + d
原来的 `tSgQ` 没有被修改。静态参考布局用于生成 copy 的逻辑坐标分工，真实源 tensor 仍使用其原有 stride 计算访存地址。
## 2. 本文的参数和记号
选择当前示例中的 `TileShape=(128,128,128)`。
| 记号 | 含义 | 本文取值 |
|---|---|---:|
| `D` | Q/K 的 feature dimension，即 `head_dim` | 128 |
| `G` | 同一个 KV head 对应的 Q heads 数量，即 `Hq/Hkv` | 运行时参数 |
| `Hkv` | KV heads 数量 | 运行时参数 |
| `B` | batch 数量 | 运行时参数 |
| `M_tile` | Q tile 的 M 大小 | 128 |
| `K_atom` | MMA atom 的 K 大小，FP8 下为 32 | 32 |
| `t` | 展平的 copy thread id | 0–63 |
| `v` | 一次 copy 分工内的 value id | 0–15 |
CuTe 的布局打印采用 `shape:stride` 格式：
_128：编译期常量 cute::Int<128>
128：运行时整数，其当前值为 128
后文区分三个整数空间：
| 记号 | 含义 |
|---|---|
| `a` | 参考布局的地址偏移，单位是元素 |
| `q` | data tensor 坐标的默认线性编号 |
| `z` | copy tile 坐标的默认线性编号 |
地址偏移和坐标编号不能混为一谈。
## 3. decode 为什么把 Q heads 放到 M 维
### 3.1 原始输入
[77_blackwell_fmha_gen.cu:329](/home/x/workspace/cutlass/examples/77_blackwell_fmha/77_blackwell_fmha_gen.cu:329) 定义：
using Element = cutlass::float_e5m2_t;
using ProblemShape =
Shape<_1, int, int, Shape<Shape<int, int>, int>>;
using StrideQ =
Stride<_0, _1, Stride<Stride<int, int>, int>>;
问题形状是：
(1, N, D, ((G,Hkv),B))
decode 只有一个 query token，原始 Q 的视图为：
Q: (1, D, ((G,Hkv),B))
[初始化代码](/home/x/workspace/cutlass/examples/77_blackwell_fmha/77_blackwell_fmha_gen.cu:472) 将 Q 的 stride 设为：
(0, 1, ((D, D*G), D*G*Hkv))
其中 `D`、`G`、`Hkv` 来自运行时 options。相邻同组 Q heads 之间的跨度为 `D`。
### 3.2 kernel 转换参数
[sm100_fmha_gen_kernel_warpspecialized.hpp:205](/home/x/workspace/cutlass/examples/77_blackwell_fmha/kernel/sm100_fmha_gen_kernel_warpspecialized.hpp:205)：
StrideQ dQ = replace<0>(args.dQ, 0);
get<0>(problem_shape) = get<3,0,0>(args.problem_shape);
get<3,0,0>(problem_shape) = 1;
get<0>(dQ) = get<2,0,0>(dQ);
它将同组 Q heads 的维度 `G` 移到 M 维：
原问题： (1, N, D, ((G,Hkv),B))
转换后： (G, N, D, ((1,Hkv),B))
同时将原来的 Q head stride 搬到 M stride：
dQ.M_stride = 原来的 G_stride = D
因此，进入加载器后，M 表示同组的不同 Q heads，而不是多个 query tokens。
**不能只看 `replace<0>(args.dQ,0)` 就认为最终 M stride 是 0。后面的赋值将它设成了 head stride。**
## 4. 为什么源布局的 M stride 是动态的
[mainloop 的类型定义](/home/x/workspace/cutlass/examples/77_blackwell_fmha/collective/sm100_fmha_gen_mainloop_warpspecialized.hpp:75)：
using StrideQOrig = StrideQ_;
using StrideQ = decltype(replace<0>(StrideQ_{}, 0));
这里使用的是普通整数 `0`，不是 `_0{}`：
0 // 类型为 int，值存储在对象中
_0{} // 类型为 cute::Int<0>，值编码在类型中
这句将第一个 stride 的类型改为 `int`。kernel 随后才能把调用方的 head stride 赋值给它。
当前示例的赋值结果是：
M stride 的类型 = int
M stride 的数值 = 128
赋值只改变数值，不会将 `int` 的类型变成 `cute::Int<128>`。
CuTe 对静态值的判定关注的是值是否由类型完全定义。相关定义见 [integral_constant.hpp:90](/home/x/workspace/cutlass/include/cute/numeric/integral_constant.hpp:90)：
// is_static detects if an (abstract) value is defined completely by its type
template <class T>
struct is_static : bool_constant<is_empty<T>::value> {};
因此，即使某个普通 `int` 的当前值或实际使用值始终是 128，它仍然不是 CuTe 的静态常量。
## 5. 从 `gQ` 到 `tSgQ`
### 5.1 mainloop 的实际 tile
[sm100_fmha_gen_mainloop_warpspecialized.hpp:65](/home/x/workspace/cutlass/examples/77_blackwell_fmha/collective/sm100_fmha_gen_mainloop_warpspecialized.hpp:65)：
class ThreadShape = Shape<_1, _2, _1>;
using TileShapeQK =
decltype(shape_div(TileShape{}, ThreadShape{}));
本例：
TileShape = (128,128,128)
ThreadShape = ( 1, 2, 1)
TileShapeQK = (128, 64,128)
本版本 `ExampleRunner` 虽然声明了 `ThreadShape` 模板参数，但其构造 mainloop 的代码没有转发该参数，所以这里按 mainloop 的默认 `(1,2,1)` 展开，不能直接把 `RUN` 中打印的 thread shape 当成 mainloop 实际配置。见 [77_blackwell_fmha_gen.cu:326](/home/x/workspace/cutlass/examples/77_blackwell_fmha/77_blackwell_fmha_gen.cu:326)。
### 5.2 构造 Q tensor
[加载器源码](/home/x/workspace/cutlass/examples/77_blackwell_fmha/collective/sm100_fmha_load_cpasync_warpspecialized.hpp:163)：
auto cQ = make_identity_tensor(select<0,2>(TileShape{}));
auto mQ = make_tensor(
make_gmem_ptr(params.ptr_q),
append<3>(
select<0,2>(TileShapeQK{}),
get<3>(problem_shape)),
params.dQ);
auto gQ = mQ(_, _, get<2>(blk_coord));
auto sQ = make_tensor(
make_smem_ptr(storage.smem_q.data()), SmemLayoutQ{});
固定 head group 和 batch 后：
gQ.shape = (_128,_128)
gQ.stride = (128,_1)
M tile 为 128，但有效 M 行数为运行时的 `G`，超出的行由后续掩码处理。
### 5.3 MMA atom
builder 为 FP8 选择 `SM100_MMA_F8F6F4_SS`，见 [sm100_common.inl:331](/home/x/workspace/cutlass/include/cutlass/gemm/collective/builders/sm100_common.inl:331)。
其 [trait](/home/x/workspace/cutlass/include/cute/atom/mma_traits_sm100.hpp:2485)：
constexpr static int K = 32;
using Shape_MNK = Shape<Int<M>,Int<N>,Int<K>>;
using ThrID = Layout<_1>;
using ALayout =
Layout<Shape<_1,Shape<Int<M>,Int<K>>>,
Stride<_0,Stride<_1,Int<M>>>>;
MMA atom = (128,64,32)
RestM = 128 / 128 = 1
RestK = 128 / 32 = 4
`ALayout` 内的 stride 是 MMA 的逻辑 `(M,K)` 坐标映射；组合到源 tensor 后，实际访存 stride 来自 `gQ`。
### 5.4 `thrfrg_A` 和 `partition_A` 必须分开看
[`thrfrg_A()`](/home/x/workspace/cutlass/include/cute/atom/mma_atom.hpp:291) 依次执行：
按 permutation 划分
→ 按 MMA atom 划分
→ compose AtomLayoutA_TV
→ 按 MMA thread tiling 划分
其输出结构为：
((ThrV,(ThrM,ThrK)), (FrgV,(RestM,RestK)))
本例实际布局：
((_1,(_1,_1)),((_128,_32),(_1,_4))):
((_0,(_0,_0)),((128,_1),(_0,_32)))
但是 [`partition_A()`](/home/x/workspace/cutlass/include/cute/atom/mma_atom.hpp:478) 继续固定线程坐标：
auto thr_tensor = make_tensor(
static_cast<ATensor&&>(atensor).data(),
this->thrfrg_A(atensor.layout()));
auto thr_vmk = make_coord(
get<0>(thr_vmnk_),
make_coord(get<1>(thr_vmnk_), get<3>(thr_vmnk_)));
return thr_tensor(
thr_vmk,
make_coord(_, repeat<rank<1,1>(thr_tensor)>(_)));
线程 mode 被切掉，`tSgQ` 的三个顶层 mode 是：
(FrgV, RestM, RestK)
最终：
tSgQ.shape = ((_128,_32),_1,_4)
tSgQ.stride = ((128,_1),_0,_32)
令坐标为 `((m,k),0,r)`：
d = k + 32*r
源地址偏移 = 128*m + k + 32*r = D*m + d
这个布局的 shape 全部静态，但 stride 中仍含有一个 `int`，因此整个 layout 不是完全静态的。
## 6. 两层 `replace<0>` 分别替换哪里
首先：
stride<0>(tSgQ) = (128,_1)
这是 `FrgV=(M_atom,K_atom)` 内的两个 stride。
内层替换：
replace<0>(stride<0>(tSgQ), get<2>(TileShape{}))
结果：
(128,_1) → (_128,_1)
外层替换：
replace<0>(stride(tSgQ), /* 修改后的 FrgV stride */)
((128,_1),_0,_32)
((_128,_1),_0,_32)
所以 `data_layout` 为：
((_128,_32),_1,_4):((_128,_1),_0,_32)
本例中它与原源布局的数值映射一致，但它是静态布局。
[`make_cotiled_copy()`](/home/x/workspace/cutlass/include/cute/atom/copy_atom.hpp:528) 明确要求：
static_assert(is_static<AtomTVLayout>::value);
static_assert(is_static<DataLayout>::value);
[`left_inverse()`](/home/x/workspace/cutlass/include/cute/layout.hpp:1364) 也要求静态 stride，用于在编译期合并 mode、排序 stride 和构造逆映射。
## 7. `make_cotiled_copy` 的完整变换
### 7.1 copy TV layout 产生地址偏移
auto atom_q_tv =
Layout<Shape<Shape<_2,_32>,_16>,
Stride<Stride<_16,_32>,_1>>{};
这表示 64 个线程，每个线程 16 个元素。将线程编号展平为 `t`：
a(t,v) = 16*(t % 2) + 32*floor(t/2) + v
= 16*t + v
一次参考分工覆盖地址偏移 `0..1023`。
这里的 `atom_q_tv` 是提供给构造器的线程和值偏移图；硬件 copy 指令的粒度由第一个参数 `Copy_Atom<...,Element>` 决定，二者不应混淆。
### 7.2 `left_inverse(data_layout)` 返回 data 坐标编号
源码：
auto inv_data_layout =
make_layout(left_inverse(data_layout), Layout<_1,_0>{});
auto layout_tv_data = composition(inv_data_layout, atom_tv_layout);
对坐标 `((m,k),0,r)`，参考地址偏移为：
a = 128*m + k + 32*r
= 128*m + d
同一坐标的默认线性编号为：
q = m + 128*k + 4096*r
= m + 128*d
逆映射从 `a` 得到：
m = floor(a/128)
d = a % 128
q = m + 128*d
实际结果：
coalesce(data_layout) = (_128,_128):(_128,_1)
left_inverse = (_128,_128):(_128,_1)
inv_data_layout = ((_128,_128),_1):((_128,_1),_0)
由于本例的 M 和 D 都是 128，正向布局和逆布局的打印形式碰巧一样；两者的输入输出空间不同。
追加的 `Layout<_1,_0>` 用于让超出末端的输入落到 stride-0 mode。本例 copy 地址都在正常范围内。
### 7.3 组合后得到 `layout_tv_data`
由 `a=16*t+v` 可直接得到：
m = floor(t/8)
d = 16*(t % 8) + v
为对照布局打印，将线程编号进一步拆成：
t = t0 + 2*t1 + 8*t2
t0 ∈ [0,1]
t1 ∈ [0,3]
t2 ∈ [0,7]
于是：
m = t2
k = 16*t0 + v
r = t1
q = t2 + 2048*t0 + 128*v + 4096*t1
实际布局：
layout_tv_data =
((_2,(_4,_8)),_16):((_2048,(_4096,_1)),_128)
这里的 `2048`、`4096`、`128` 是 data 坐标编号空间中的系数，不是新的 global memory stride。
### 7.4 逐 mode 投影并生成 `tiler`
auto flat_data_shape = product_each(shape(data_layout));
auto flat_data_zeros = repeat<rank(flat_data_shape)>(Int<0>{});
auto tiler = transform(
make_seq<rank(flat_data_shape)>{},
[&](auto i) {
return filter(composition(
flat_data_shape,
replace<i>(flat_data_zeros, Int<1>{})),
layout_tv_data));
});
flat_data_shape = (_4096,_1,_4)
`product_each` 将第一个顶层 mode `(128,32)` 合成 4096，但保留三个顶层 mode。
投影布局的 stride 分别为 `(1,0,0)`、`(0,1,0)`、`(0,0,1)`，用于提取各 mode 的坐标。
| mode | 提取的坐标 | `filter` 后的布局 |
|---|---|---|
| FrgV | `2048*t0 + t2 + 128*v` | `(2,8,16):(2048,1,128)` |
| RestM | `0` | `1:0` |
| RestK | `t1` | `4:1` |
完整结果：
tiler =
(
(_2,_8,_16):(_2048,_1,_128),
_1:_0,
_4:_1
)
它选出的 Q 坐标为：
m = 0..7
d = 0..127
即一次分工覆盖 `8×128=1024` 个元素。整个 Q tile 仍有 `128×128` 个元素。
### 7.5 构造 `tile2data`
auto tile2data = composition(make_layout(flat_data_shape), tiler);
tile2data =
((_2,_8,_16),_1,_4):((_2048,_1,_128),_0,_4096)
设 copy tile 坐标为 `((t0,m,v),0,r)`：
q = 2048*t0 + m + 128*v + 4096*r
这里组合的是 `make_layout(flat_data_shape)`，不是物理 `data_layout`：`tile2data` 生成 data 坐标编号，而不是物理地址。
### 7.6 构造最终 `layout_tv`
auto layout_tv = composition(left_inverse(tile2data), layout_tv_data);
copy tile 自己的默认线性编号为：
z = t0 + 2*m + 16*v + 256*r
代入 `m=t2`、`r=t1`：
z(t,v) = t0 + 2*t2 + 16*v + 256*t1
因此：
layout_tv =
((_2,(_4,_8)),_16):((_1,(_256,_2)),_16)
最终 `make_tiled_copy_impl(copy_atom, layout_tv, tiler)` 用这些静态映射构造 `TiledCopy`。
完整转换链：
(tid,vid)
│ atom_q_tv
参考地址偏移 a
│ left_inverse(data_layout)
data 坐标编号 q
│ left_inverse(tile2data)
copy tile 坐标编号 z
## 8. 如何作用到真实源、共享内存和掩码
加载器接下来调用：
auto thr_copy_q = tiled_copy_q.get_slice(thread_idx);
auto tQsQ = thr_copy_q.partition_D(sQ);
auto tQgQ = thr_copy_q.partition_S(tSgQ);
auto tQcQ = thr_copy_q.partition_S(tScQ);
[`partition_S/D()`](/home/x/workspace/cutlass/include/cute/atom/copy_atom.hpp:365) 将同一份 tiler 和 TV 映射分别组合到各个 tensor 的布局上。
| tensor | 用途 |
| `tQgQ` | 根据真实源布局计算 global memory 源地址 |
| `tQsQ` | 根据共享内存布局及 swizzle 计算目的地址 |
| `tQcQ` | 给出同一份分工对应的原始 `(m,d)` 坐标，供掩码使用 |
### 8.1 `tile2thrfrg` 中的 `zip + coalesce` 详细展开
`partition_S/D` 内部会先用 `zipped_divide` 按 tiler 划分真实 tensor，再由 [`tile2thrfrg()`](/home/x/workspace/cutlass/include/cute/atom/copy_atom.hpp:257) 将 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(_,_), _);
注意准确写法是 `Shape<_1,Shape<_1,_1>>{}`，不是 `Shape<1,...>`。`Shape` 接受类型参数，`_1` 是 CuTe 的常量类型。
#### 8.1.1 进入这一句之前的层级
`TiledLayout_TV` 描述整个 copy tile 的 `(thread,value)` 分工。`zipped_divide` 按硬件 Copy Atom 的线程数和值数将它分为：
((atom_tid,atom_val),(rest_tid,rest_val))
其中 `rest_tid`、`rest_val` 是当前 copy tile 内相对 Atom 的重复部分，不是 Q tensor 的 `RestM`、`RestK`，也不是 tile 外的剩余坐标。
`ref2trg` 的源码变量名需要结合实际表达式理解。源端传入：
right_inverse(AtomLayoutRef{}).compose(AtomLayoutSrc{})
它接收 source Atom 的线程和值坐标，先找到该值的 reference 编号，再送入 `atom_layout_TV` 的第一个 mode。目的端将 `AtomLayoutSrc` 换成 `AtomLayoutDst`。因此转换后的输入结构是：
((trg_tid,trg_val),(rest_tid,rest_val))
本例 cp.async 的 Src、Dst、Ref 布局相同，转换不改变元素关系；其他 Copy Atom 不应直接套用这一简化。
这里的 `TiledLayout_TV` 始终按 `AtomLayoutRef` 的坐标解释，同一个 `TiledCopy` 在 `partition_S/D` 时不会更换它。源端通过 `Ref^-1∘Src` 转成 reference 坐标，目的端通过 `Ref^-1∘Dst` 转成 reference 坐标，再应用统一的 reference TV 分工。
`make_cotiled_copy` 的 `data_layout` 只是 tensor 的参考坐标到地址映射；选用源 tensor 或目的 tensor 的布局构造它，不会改变 `Copy_Traits::RefLayout`。若手头已有某一侧期望的 TV 分工，需要结合 Atom 的 Ref 定义构造 reference TV，再检查最终 `get_layoutS_TV()` 和 `get_layoutD_TV()`，而不是把期望的源分工直接冒充目的参考分工。
Ref=Dst 的 ldmatrix 具体例子和两端转换公式，见 [make_tiled_copy 教程的参考端说明](/home/x/workspace/cutlass/cute_make_tiled_copy_walkthrough_zh.md)。
#### 8.1.2 `zip` 改变分组，未搬动数据
[`zip(Layout)`](/home/x/workspace/cutlass/include/cute/layout.hpp:1557) 同时转置 shape 和 stride 的 tuple：
return make_layout(zip(layout.shape()), zip(layout.stride()));
效果为：
↓ zip
((trg_tid,rest_tid),(trg_val,rest_val))
线程 值
对应坐标仍指向同一个 tile 元素：
trg_layout_TV((ta,va),(tr,vr))
== zip(trg_layout_TV)((ta,tr),(va,vr))
这不是 tensor 数据转置，也不是跨线程通信。由于输入坐标结构发生变化，不能不加坐标变换就声称 zip 前后对同一个标量输入 i 的输出总相同。
#### 8.1.3 profile 保留哪些边界
`Shape<_1,Shape<_1,_1>>{}` 是递归控制结构：
终端, // 将整个线程子布局交给普通 coalesce
(终端,终端) // 分别简化 Atom 内的值、Atom 重复的值
[`coalesce(layout,profile)`](/home/x/workspace/cutlass/include/cute/layout.hpp:880) 遇到 tuple profile 时继续逐 mode 递归，遇到标量 profile 时调用普通 `coalesce(layout)`。标量终端的数值不参与计算，所以 `(7,(7,7))` 在这里与 `(1,(1,1))` 等效。
这个 profile 不是目标 shape，也不要求每个 mode 的大小变成 1。可以将其理解为：
auto z = zip(trg_layout_TV);
auto thrval2mn = make_layout(
coalesce(get<0>(z)),
coalesce(get<1,0>(z)),
coalesce(get<1,1>(z))));
普通 coalesce 会展平指定子布局、去掉静态 size-1 mode，并在相邻 mode 满足连续条件时合并。例如 `(2,4):(1,2)` 可以变为 `8:1`，但 `(2,4):(1,256)` 不能这样合并。
因此线程子布局可能仍有多个 mode；它只是简化表示，不是强制将 64 个线程写成单一的 `64`。
#### 8.1.4 当前 FP8 的准确输出
FP8 的 16 字节 cp.async 对应：
AtomNumThr = 1
AtomNumVal = 16
按第 7 节的 TiledLayout_TV 展开：
TiledLayout_TV =
atom_layout_TV = trg_layout_TV =
((_1,_16),((_2,_4,_8),_1)):
((_0,_16),((_1,_256,_2),_0))
zip(trg_layout_TV) =
((_1,(_2,_4,_8)),(_16,_1)):
((_0,(_1,_256,_2)),(_16,_0))
thrval2mn =
((_2,_4,_8),(_16,_1)):
((_1,_256,_2),(_16,_0))
线程部分去掉 size-1 的 `trg_tid`，但 `(2,4,8):(1,256,2)` 不能按连续条件合并为单个 mode。值部分保留 `(16,1)`，表示“一条 Atom 指令内 16 个值，当前 copy tile 内每线程重复 1 次 Atom”。
将线程解码为 `t=t0+2*t1+8*t2`，此时：
thrval2mn(t,(va,0)) = t0 + 256*t1 + 2*t2 + 16*va
它的输出就是第 7 节中 copy tile 的线性坐标编号 z，还不是 global memory 地址，也不是 C++ 二元 tuple `(m,n)`。
#### 8.1.5 FP16 为什么要保留 `(8,2)`
FP16 的 16 字节 Atom 每线程处理 8 个值，整个 copy tile 的 TV 分工仍为每线程 16 个值，因此：
AtomNumVal = 8
rest_val extent = 16/8 = 2
trg_layout_TV =
((_1,_8),((_8,_8),_2)):((_0,_8),((_128,_1),_64))
((_1,(_8,_8)),(_8,_2)):((_0,(_128,_1)),(_8,_64))
((_8,_8),(_8,_2)):((_128,_1),(_8,_64))
虽然值子布局 `(8,2):(8,64)` 满足 `8*8=64`，可以在数值上合并成 `16:8`，本句的 profile 刻意保留了这两个分组：
8：一条指令内的元素
2：重复两次 Atom 覆盖的元素
如果改成更浅的 `Shape<_1,_1>{}`，FP16 会得到：
((_8,_8),_16):((_128,_1),_8)
这仍能表达元素关系，但丢掉了此处后续代码需要的值分组接口。
不传 profile 则可能跨线程和值边界继续合并。FP8 本例会得到：
(_2,_4,_128):(_1,_256,_2)
原来的线程 mode 与值 mode 已不再独立，所以不能直接替换源码中的布局。
#### 8.1.6 与真实 tensor 组合
`tensor` 是 `zipped_divide` 之后的真实 tensor 布局：
((tile内坐标),(tile外剩余坐标)) → 真实 tensor 偏移
将第一个 tile mode 的输入改成线程和值，同时保留 tile 外的剩余部分：
((线程,(Atom内值,Atom重复值)), tile外剩余坐标)
→ 真实 tensor 偏移
`tv_tensor(make_coord(_,_),_)` 再将第一个组合 mode 中的线程和值展开成两个顶层 mode。`ThrCopy::partition_S/D` 随后固定当前线程，形成实际的 per-thread fragment。
这种写法统一了源布局、共享内存布局和坐标 tensor 的线程分工；三者可以有不同的物理布局，但使用同一组逻辑 tile 坐标。
另一个 `make_tiled_copy` 例子中 `right_inverse(layout_mn)=(8,16):(16,1)` 的逐项源码推导，见 [make_tiled_copy 教程](/home/x/workspace/cutlass/cute_make_tiled_copy_walkthrough_zh.md)。不要将那个例子的 right inverse 与本例 `make_cotiled_copy` 使用的 data-layout left inverse 混淆。
### 8.2 单线程如何覆盖完整 tile
FP8 下，单线程源 partition 的实际布局为：
shape = ((_16,_1),_16,_1,_1)
stride = ((_1,_0),1024,_0,_0)
这是 16 个连续元素，沿 M 重复 16 次。令重复编号 `p∈[0,15]`：
m = floor(t/8) + 8*p
源地址相对于当前 head group 起点的偏移是：
128*m + d = 16*t + v + 1024*p
| thread | 负责的 M 行 | 每行负责的 D 范围 |
| 0 | 0、8、16、…、120 | 0–15 |
| 7 | 0、8、16、…、120 | 112–127 |
| 8 | 1、9、17、…、121 | 0–15 |
| 63 | 7、15、23、…、127 | 112–127 |
所以：
64 threads × 16 values × 16 repeats
= 16384
= 128 × 128
目的布局的物理地址可以不同，逻辑 `(m,d)` 分工保持对应。不能因为参考布局是紧凑 row-major，就认为共享内存也一定是无 swizzle 的 row-major。
### 8.3 有效行数和 `cp_async_zfill`
[加载代码](/home/x/workspace/cutlass/examples/77_blackwell_fmha/collective/sm100_fmha_load_cpasync_warpspecialized.hpp:187)：
auto limitQ = append<2>(get<0>(problem_shape), _128{});
auto src = recast<uint128_t>(tQgQ(_, _, _, _));
auto dst = recast<uint128_t>(tQsQ(_, _, _, _, state.index()));
auto c = tQcQ(_, _, _, _);
int vlen = sizeof(uint128_t) / sizeof(Element);
for (int i = 0; i < size(src); i++) {
auto cc = c(vlen*i);
bool guard = elem_less(cc, limitQ);
cutlass::arch::cp_async_zfill<16, cutlass::arch::CacheOperation::Always>(
&dst(i), &src(i), guard);
}
转换后的 `get<0>(problem_shape)=G`，所以判定相当于：
(m < G) && (d < 128)
有效 Q heads 从 global memory 加载；超出 G 的填充行清零。不是固定只加载第 0 行，也不是将同一个 Q head 广播给全部 M 行。
当前示例只支持 `D=128`，这里的 feature limit 也直接使用 `_128{}`。不能仅凭 `replace` 引用 `TileShape` 就推断整个示例已支持任意 head dimension。
## 9. FP16 对照
FP16 的 MMA trait 将 atom K 定为 `256/16=16`，见 [mma_traits_sm100.hpp:236](/home/x/workspace/cutlass/include/cute/atom/mma_traits_sm100.hpp:236)。
同样取 Q tile `128×128`：
| 布局 | FP8 | FP16 |
| MMA atom K | 32 | 16 |
| `tSgQ.shape` | `((128,32),1,4)` | `((128,16),1,8)` |
| 源 stride | `((int(128),1),0,32)` | `((int(128),1),0,16)` |
| 静态参考 stride | `((Int<128>,1),0,32)` | `((Int<128>,1),0,16)` |
| `flat_data_shape` | `(4096,1,4)` | `(2048,1,8)` |
FP16 的实际中间布局为：
data_layout =
((_128,_16),_1,_8):((_128,_1),_0,_16)
inv_data_layout =
((_128,_128),_1):((_128,_1),_0)
((_8,_16):(_1,_128),_1:_0,_8:_1)
((_8,_16),_1,_8):((_1,_128),_0,_2048)
((_2,(_4,_8)),_16):((_128,(_256,_1)),_8)
FP16 和 FP8 的 fragment 层级不同，但最后的 `(m,d)` 分工一致。
`atom_q_tv` 仍为每线程 16 个元素，不过每条 16 字节的 cp.async 对 FP16 只复制 8 个元素，因此每个这样的 16 元素分组使用两条指令；FP8 则使用一条。
## 10. 三个常见追问
### 10.1 所以替换只是为了改成静态布局？
在当前紧凑示例中，是的：源 M stride 的数值已经是 D，替换只是将其从 `int` 改成 `Int<D>`，保持数值映射一致。
更完整地说，它构造的是 copy 使用的静态参考布局，而不是修改真实源 tensor。
### 10.2 `tSgQ` 为什么是动态布局？
M stride 来自调用方的运行时参数，并且 mainloop 明确将其类型设成 `int`。`partition_A` 继承这个 stride，不会将运行时数值自动提升为类型常量。
因此它是“静态 shape + 部分动态 stride”，不是 shape 也全动态。
### 10.3 为什么不一开始直接使用静态 stride？
可以重新设计参数转换，在源布局中直接使用模板 D；但这会把真实源行跨度绑定到模板配置上，而当前代码保留调用方传入的 head stride。
从代码结构看，两者的职责不同：
args.dQ：描述真实 Q 存储，作为运行时参数传入
TileShape：描述 kernel tile，由模板确定
若第一项使用 `_0{}`，它的类型只表示常量 0，无法再赋值为调用方的 head stride。
若直接使用 `Int<D>`，则是在源布局中用模板 D 覆盖调用方的 stride。当前示例恰好采用紧凑存储，两者一致；但保留运行时源 stride 与构造静态 copy 参考布局在职责上仍然不同。
从布局映射看，若真实源行跨度记为 P，copy 仍可先按静态 D 生成逻辑坐标，再通过真实源布局得到：
source_offset = m * P + d
这说明静态参考地址空间不必直接充当真实物理地址空间。这里讨论的是布局组合机制，并非对整个 kernel 的所有输入格式作支持承诺。
## 11. 为什么是 `get<2>`，不是 `get<0>` 或 `get<1>`
`TileShape` 的三维分别是 M、N、D：
M：Q tile 行数
N：KV 序列方向的 tile 大小
D：Q/K 特征方向的 tile 大小
紧凑 Q 参考布局中，一行包含 D 个元素，所以 M 方向的 stride 必须表达 D：
reference_offset = m * D + d
M 是行数，N 是 KV tile 宽度，都不是一行 Q 的特征数。本例中 M、N、D 恰好都为 128，不代表三者可以在语义上互换。例如一个 `64×128` 的紧凑行布局，其 M stride 是 128，而不是 64。
`make_cotiled_copy` 本身并不认识 `head_dim`；是调用方用它定义了这个静态、紧凑的参考地址空间。
## 12. 应避免的错误理解
| 错误理解 | 源码支持的解释 |
| 替换的是 `ThrV` stride | `partition_A` 已切掉线程 mode；替换的是 FrgV 内的 M stride |
| 原 M stride 为 0 | kernel 参数转换将它赋值为原 Q head stride，当前示例为 D |
| 128 是运行时已知值，所以布局静态 | CuTe 静态性由类型判断；普通 `int(128)` 仍是动态项 |
| `replace` 修改了 `tSgQ` 的访存布局 | 它返回新 tuple，用于构造参考 layout；真实源 tensor 未被修改 |
| copy 的 1024 个元素就是整个 Q tile | 一次分工为 1024 个元素，重复后覆盖 16384 个元素 |
| inverse 输出可以直接当物理地址 | inverse 输出是 data 坐标编号，后续还需转换为 tile 坐标 |
| decode 只加载第 0 行 | M 维承载 G 个 Q heads，掩码判定 `m<G` |
| 静态参考 layout 就是共享内存物理布局 | 目的 tensor 使用自身的 shared-memory layout 和 swizzle |
| prefill 的 256×128 配置适用于这段代码 | 本文加载器应沿 generation/decode 调用路径分析 |
## 13. 验证范围
本文具体布局来自本次会话中对实际 mainloop/builder 的 host 端实例化，FP8 和 FP16 均核对了：
1. 当前紧凑示例中，替换前后全部 Q tile 坐标的地址映射一致。
2. `data_layout(layout_tv_data(t,v))` 恢复 `atom_q_tv(t,v)` 的参考地址。
3. `tile2data(layout_tv(t,v))` 恢复 `layout_tv_data(t,v)` 的 data 坐标编号。
4. 64 个线程的最终 partition 覆盖全部 `128×128` 元素，每个坐标恰好一次。
5. 源 partition 与共享内存 partition 在同一 pipeline stage 上对应相同的逻辑 `(m,d)` 坐标。
补充第 8.1 节时，另行枚举了 FP8、FP16 各 1024 个 thread/value 坐标，确认 `trg_layout_TV`、zip 后布局、profile coalesce 后布局对相应结构坐标的映射一致，并核对了不同 profile 的实际打印结果。
这些检查验证的是类型、布局和坐标映射；没有执行 GPU kernel。
## 14. 源码阅读顺序
| 顺序 | 源码位置 | 要确认的问题 |
|---:|---|---|
| 1 | [77_blackwell_fmha_gen.cu:329](/home/x/workspace/cutlass/examples/77_blackwell_fmha/77_blackwell_fmha_gen.cu:329) | 输入类型、shape、stride 类型 |
| 2 | [77_blackwell_fmha_gen.cu:472](/home/x/workspace/cutlass/examples/77_blackwell_fmha/77_blackwell_fmha_gen.cu:472) | Q head stride 为什么等于 D |
| 3 | [sm100_fmha_gen_kernel_warpspecialized.hpp:205](/home/x/workspace/cutlass/examples/77_blackwell_fmha/kernel/sm100_fmha_gen_kernel_warpspecialized.hpp:205) | G 如何搬到 M，head stride 如何搬到 M stride |
| 4 | [sm100_fmha_gen_mainloop_warpspecialized.hpp:75](/home/x/workspace/cutlass/examples/77_blackwell_fmha/collective/sm100_fmha_gen_mainloop_warpspecialized.hpp:75) | 为什么 M stride 的类型是 int |
| 5 | [sm100_fmha_gen_mainloop_warpspecialized.hpp:97](/home/x/workspace/cutlass/examples/77_blackwell_fmha/collective/sm100_fmha_gen_mainloop_warpspecialized.hpp:97) | TileShapeQK 和 builder 的输入 |
| 6 | [mma_traits_sm100.hpp:2485](/home/x/workspace/cutlass/include/cute/atom/mma_traits_sm100.hpp:2485) | FP8 atom K、ThrID、ALayout |
| 7 | [mma_atom.hpp:291](/home/x/workspace/cutlass/include/cute/atom/mma_atom.hpp:291) | thrfrg_A 的层级结构 |
| 8 | [mma_atom.hpp:478](/home/x/workspace/cutlass/include/cute/atom/mma_atom.hpp:478) | partition_A 如何切掉线程 mode |
| 9 | [sm100_fmha_load_cpasync_warpspecialized.hpp:163](/home/x/workspace/cutlass/examples/77_blackwell_fmha/collective/sm100_fmha_load_cpasync_warpspecialized.hpp:163) | tSgQ、replace、copy partition 和掩码 |
| 10 | [copy_atom.hpp:528](/home/x/workspace/cutlass/include/cute/atom/copy_atom.hpp:528) | inverse、投影、tiler、tile2data、layout_tv |
| 11 | [layout.hpp:1364](/home/x/workspace/cutlass/include/cute/layout.hpp:1364) | left_inverse 的静态 stride 要求 |
| 12 | [copy_atom.hpp:365](/home/x/workspace/cutlass/include/cute/atom/copy_atom.hpp:365) | 静态分工如何作用到真实源和目的布局 |