# 从源码理解 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),
            replace<0>(
                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` 的依据是复制参考布局的行组织：

```text
一行 = 一个 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` 格式：

```text
_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) 定义：

```cpp
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>>;
```

问题形状是：

```text
(1, N, D, ((G,Hkv),B))
```

decode 只有一个 query token，原始 Q 的视图为：

```text
Q: (1, D, ((G,Hkv),B))
```

[初始化代码](/home/x/workspace/cutlass/examples/77_blackwell_fmha/77_blackwell_fmha_gen.cu:472) 将 Q 的 stride 设为：

```text
(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)：

```cpp
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 维：

```text
原问题： (1, N, D, ((G,Hkv),B))
转换后： (G, N, D, ((1,Hkv),B))
```

同时将原来的 Q head stride 搬到 M stride：

```text
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)：

```cpp
using StrideQOrig = StrideQ_;
using StrideQ = decltype(replace<0>(StrideQ_{}, 0));
```

这里使用的是普通整数 `0`，不是 `_0{}`：

```cpp
0      // 类型为 int，值存储在对象中
_0{}   // 类型为 cute::Int<0>，值编码在类型中
```

这句将第一个 stride 的类型改为 `int`。kernel 随后才能把调用方的 head stride 赋值给它。

当前示例的赋值结果是：

```text
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)：

```cpp
// 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)：

```cpp
class ThreadShape = Shape<_1, _2, _1>;

using TileShapeQK =
    decltype(shape_div(TileShape{}, ThreadShape{}));
```

本例：

```text
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)：

```cpp
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 后：

```text
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)：

```cpp
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>>>>;
```

本例：

```text
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) 依次执行：

```text
按 permutation 划分
    → 按 MMA atom 划分
    → compose AtomLayoutA_TV
    → 按 MMA thread tiling 划分
```

其输出结构为：

```text
((ThrV,(ThrM,ThrK)), (FrgV,(RestM,RestK)))
```

本例实际布局：

```text
((_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) 继续固定线程坐标：

```cpp
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 是：

```text
(FrgV, RestM, RestK)
```

最终：

```text
tSgQ.shape  = ((_128,_32),_1,_4)
tSgQ.stride = ((128,_1),_0,_32)
```

令坐标为 `((m,k),0,r)`：

```text
d = k + 32*r
源地址偏移 = 128*m + k + 32*r = D*m + d
```

这个布局的 shape 全部静态，但 stride 中仍含有一个 `int`，因此整个 layout 不是完全静态的。

## 6. 两层 `replace<0>` 分别替换哪里

首先：

```text
stride<0>(tSgQ) = (128,_1)
```

这是 `FrgV=(M_atom,K_atom)` 内的两个 stride。

内层替换：

```cpp
replace<0>(stride<0>(tSgQ), get<2>(TileShape{}))
```

结果：

```text
(128,_1) → (_128,_1)
```

外层替换：

```cpp
replace<0>(stride(tSgQ), /* 修改后的 FrgV stride */)
```

结果：

```text
((128,_1),_0,_32)
          ↓
((_128,_1),_0,_32)
```

所以 `data_layout` 为：

```text
((_128,_32),_1,_4):((_128,_1),_0,_32)
```

本例中它与原源布局的数值映射一致，但它是静态布局。

[`make_cotiled_copy()`](/home/x/workspace/cutlass/include/cute/atom/copy_atom.hpp:528) 明确要求：

```cpp
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 产生地址偏移

```cpp
auto atom_q_tv =
    Layout<Shape<Shape<_2,_32>,_16>,
           Stride<Stride<_16,_32>,_1>>{};
```

这表示 64 个线程，每个线程 16 个元素。将线程编号展平为 `t`：

```text
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 坐标编号

源码：

```cpp
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)`，参考地址偏移为：

```text
a = 128*m + k + 32*r
  = 128*m + d
```

同一坐标的默认线性编号为：

```text
q = m + 128*k + 4096*r
  = m + 128*d
```

逆映射从 `a` 得到：

```text
m = floor(a/128)
d = a % 128
q = m + 128*d
```

实际结果：

```text
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` 可直接得到：

```text
m = floor(t/8)
d = 16*(t % 8) + v
```

为对照布局打印，将线程编号进一步拆成：

```text
t = t0 + 2*t1 + 8*t2

t0 ∈ [0,1]
t1 ∈ [0,3]
t2 ∈ [0,7]
```

于是：

```text
m = t2
k = 16*t0 + v
r = t1

q = t2 + 2048*t0 + 128*v + 4096*t1
```

实际布局：

```text
layout_tv_data =
((_2,(_4,_8)),_16):((_2048,(_4096,_1)),_128)
```

这里的 `2048`、`4096`、`128` 是 data 坐标编号空间中的系数，不是新的 global memory stride。

### 7.4 逐 mode 投影并生成 `tiler`

源码：

```cpp
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(
          make_layout(
              flat_data_shape,
              replace<i>(flat_data_zeros, Int<1>{})),
          layout_tv_data));
    });
```

本例：

```text
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` |

完整结果：

```text
tiler =
(
  (_2,_8,_16):(_2048,_1,_128),
  _1:_0,
  _4:_1
)
```

它选出的 Q 坐标为：

```text
m = 0..7
d = 0..127
```

即一次分工覆盖 `8×128=1024` 个元素。整个 Q tile 仍有 `128×128` 个元素。

### 7.5 构造 `tile2data`

```cpp
auto tile2data = composition(make_layout(flat_data_shape), tiler);
```

实际结果：

```text
tile2data =
((_2,_8,_16),_1,_4):((_2048,_1,_128),_0,_4096)
```

设 copy tile 坐标为 `((t0,m,v),0,r)`：

```text
q = 2048*t0 + m + 128*v + 4096*r
```

这里组合的是 `make_layout(flat_data_shape)`，不是物理 `data_layout`：`tile2data` 生成 data 坐标编号，而不是物理地址。

### 7.6 构造最终 `layout_tv`

```cpp
auto layout_tv = composition(left_inverse(tile2data), layout_tv_data);
```

copy tile 自己的默认线性编号为：

```text
z = t0 + 2*m + 16*v + 256*r
```

代入 `m=t2`、`r=t1`：

```text
z(t,v) = t0 + 2*t2 + 16*v + 256*t1
```

因此：

```text
layout_tv =
((_2,(_4,_8)),_16):((_1,(_256,_2)),_16)
```

最终 `make_tiled_copy_impl(copy_atom, layout_tv, tiler)` 用这些静态映射构造 `TiledCopy`。

完整转换链：

```text
(tid,vid)
    │ atom_q_tv
    ↓
参考地址偏移 a
    │ left_inverse(data_layout)
    ↓
data 坐标编号 q
    │ left_inverse(tile2data)
    ↓
copy tile 坐标编号 z
```

## 8. 如何作用到真实源、共享内存和掩码

加载器接下来调用：

```cpp
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 内的坐标转换成线程和值。

核心代码：

```cpp
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 的线程数和值数将它分为：

```text
((atom_tid,atom_val),(rest_tid,rest_val))
```

其中 `rest_tid`、`rest_val` 是当前 copy tile 内相对 Atom 的重复部分，不是 Q tensor 的 `RestM`、`RestK`，也不是 tile 外的剩余坐标。

`ref2trg` 的源码变量名需要结合实际表达式理解。源端传入：

```cpp
right_inverse(AtomLayoutRef{}).compose(AtomLayoutSrc{})
```

它接收 source Atom 的线程和值坐标，先找到该值的 reference 编号，再送入 `atom_layout_TV` 的第一个 mode。目的端将 `AtomLayoutSrc` 换成 `AtomLayoutDst`。因此转换后的输入结构是：

```text
((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：

```cpp
return make_layout(zip(layout.shape()), zip(layout.stride()));
```

效果为：

```text
((trg_tid,trg_val),(rest_tid,rest_val))
                    ↓ zip
((trg_tid,rest_tid),(trg_val,rest_val))
        线程                 值
```

对应坐标仍指向同一个 tile 元素：

```text
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>>{}` 是递归控制结构：

```text
(
  终端,           // 将整个线程子布局交给普通 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。可以将其理解为：

```cpp
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))));
```

普通 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 对应：

```text
AtomNumThr = 1
AtomNumVal = 16
```

按第 7 节的 TiledLayout_TV 展开：

```text
TiledLayout_TV =
((_2,(_4,_8)),_16):((_1,(_256,_2)),_16)

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`，此时：

```text
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 个值，因此：

```text
AtomNumThr = 1
AtomNumVal = 8
rest_val extent = 16/8 = 2
```

实际结果：

```text
trg_layout_TV =
((_1,_8),((_8,_8),_2)):((_0,_8),((_128,_1),_64))

zip(trg_layout_TV) =
((_1,(_8,_8)),(_8,_2)):((_0,(_128,_1)),(_8,_64))

thrval2mn =
((_8,_8),(_8,_2)):((_128,_1),(_8,_64))
```

虽然值子布局 `(8,2):(8,64)` 满足 `8*8=64`，可以在数值上合并成 `16:8`，本句的 profile 刻意保留了这两个分组：

```text
8：一条指令内的元素
2：重复两次 Atom 覆盖的元素
```

如果改成更浅的 `Shape<_1,_1>{}`，FP16 会得到：

```text
((_8,_8),_16):((_128,_1),_8)
```

这仍能表达元素关系，但丢掉了此处后续代码需要的值分组接口。

不传 profile 则可能跨线程和值边界继续合并。FP8 本例会得到：

```text
(_2,_4,_128):(_1,_256,_2)
```

原来的线程 mode 与值 mode 已不再独立，所以不能直接替换源码中的布局。

#### 8.1.6 与真实 tensor 组合

`tensor` 是 `zipped_divide` 之后的真实 tensor 布局：

```text
((tile内坐标),(tile外剩余坐标)) → 真实 tensor 偏移
```

因此：

```cpp
auto tv_tensor = tensor.compose(thrval2mn, _);
```

将第一个 tile mode 的输入改成线程和值，同时保留 tile 外的剩余部分：

```text
((线程,(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 的实际布局为：

```text
shape  = ((_16,_1),_16,_1,_1)
stride = ((_1,_0),1024,_0,_0)
```

这是 16 个连续元素，沿 M 重复 16 次。令重复编号 `p∈[0,15]`：

```text
m = floor(t/8) + 8*p
d = 16*(t % 8) + v
```

源地址相对于当前 head group 起点的偏移是：

```text
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 |

所以：

```text
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)：

```cpp
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`，所以判定相当于：

```text
(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 的实际中间布局为：

```text
data_layout =
((_128,_16),_1,_8):((_128,_1),_0,_16)

inv_data_layout =
((_128,_128),_1):((_128,_1),_0)

layout_tv_data =
((_2,(_4,_8)),_16):((_2048,(_4096,_1)),_128)

tiler =
((_8,_16):(_1,_128),_1:_0,_8:_1)

tile2data =
((_8,_16),_1,_8):((_1,_128),_0,_2048)

layout_tv =
((_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。

从代码结构看，两者的职责不同：

```text
args.dQ：描述真实 Q 存储，作为运行时参数传入
TileShape：描述 kernel tile，由模板确定
```

若第一项使用 `_0{}`，它的类型只表示常量 0，无法再赋值为调用方的 head stride。

若直接使用 `Int<D>`，则是在源布局中用模板 D 覆盖调用方的 stride。当前示例恰好采用紧凑存储，两者一致；但保留运行时源 stride 与构造静态 copy 参考布局在职责上仍然不同。

从布局映射看，若真实源行跨度记为 P，copy 仍可先按静态 D 生成逻辑坐标，再通过真实源布局得到：

```text
source_offset = m * P + d
```

这说明静态参考地址空间不必直接充当真实物理地址空间。这里讨论的是布局组合机制，并非对整个 kernel 的所有输入格式作支持承诺。

## 11. 为什么是 `get<2>`，不是 `get<0>` 或 `get<1>`

`TileShape` 的三维分别是 M、N、D：

```text
M：Q tile 行数
N：KV 序列方向的 tile 大小
D：Q/K 特征方向的 tile 大小
```

紧凑 Q 参考布局中，一行包含 D 个元素，所以 M 方向的 stride 必须表达 D：

```text
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) | 静态分工如何作用到真实源和目的布局 |