Layout 与布局代数 / 第 1 章
读懂 CuTe 的 Layout：从 tuple 到坐标映射
访问一张按行存放的 2×3 矩阵，通常会写：
data[m * 3 + n];
这里其实做了两件事：先将坐标 (m,n) 换算成整数下标，再用下标读取数据。
CuTe 把前一件事交给 Layout：
using namespace cute;
auto layout = make_layout(make_shape(_2{}, _3{}),
make_stride(_3{}, _1{}));
layout(1,2); // 返回 5
Layout 保存坐标的形状，以及从坐标计算下标的规则。它没有接收 data，这一步也没有读取矩阵。
如果一直是普通二维矩阵，两个大小、两个步长就够用了。但在 CuTe 中，一个维度还可以拆成“组内位置、组编号”，拆出来的维度也可能继续分组。那些看起来复杂的括号，正是用 tuple 保存的层次结构。
先把括号读懂，再按 Shape 分解坐标，用 Stride 算出索引。这篇先把这些基础讲完整，Product 和 Divide 留给后续两篇。
tuple：先理解括号里装了什么
tuple 是一组有顺序的对象。它与数组的一个区别是，每个位置可以放不同类型的对象：
auto t = make_tuple(2, _3{});
第一个位置是普通 int，第二个位置是 CuTe 的编译期整数。CuTe 的 tuple 用法与 std::tuple 相近，并提供可以在 CPU 和 GPU 上使用的实现。
访问位置用 get，编号从 0 开始：
get<0>(t); // 2
get<1>(t); // 编译期整数 3
这里的 <0>、<1> 是模板参数；tuple 各位置的类型可以不同，所以不能把它当作普通数组来写 t[i]。
2 与 _2{}，数值一样，类型不同
CuTe 经常出现这种写法：
int n = 3; // 普通运行时整数
auto s = _3{}; // CuTe 编译期整数
_3 是 Int<3> 的简写，数值 3 保存在类型中。_3{} 则是创建这个类型的对象，和 Int<3>{} 等价。
这让 CuTe 在编译期知道某个维度的大小或步长，进而简化布局运算。普通 int 也可以用于 Layout，只是类型本身不携带固定数值。
例如下面两种布局，在 n=3 时都可以计算 3*m+n_coord：
auto fixed = make_layout(make_shape(_2{}, _3{}),
make_stride(_3{}, _1{}));
int n = 3;
auto mixed = make_layout(make_shape(_2{}, n),
make_stride(n, _1{}));
前者全部是静态整数；后者混用了静态和动态整数。
打印时，CuTe 用下划线区别静态数值：
fixed = (_2,_3):(_3,_1)
mixed = (_2,3):(3,_1)
下文为便于阅读，通常省略这些打印用的下划线，将前者写成 (2,3):(3,1)。
tuple 可以继续嵌套
再看这个对象：
auto t = make_tuple(_2{}, make_tuple(_3{}, _4{}));
// (2,(3,4))
它的最外层只有两个位置：第一个放 2，第二个放子 tuple (3,4)。
用文本树展开，可以看清顶层位置与叶子的区别：
t = (2,(3,4))
├─ get<0>(t) → 2
└─ get<1>(t) → (3,4)
├─ get<1,0>(t) → 3
└─ get<1,1>(t) → 4
顶层只有两个位置，但整棵树有三个整数叶子：2、3、4。
因此这两句的含义不同：
get<1>(t); // (3,4)
get<1,0>(t); // 3，等价于 get<0>(get<1>(t))
括号不能随意删。(2,(3,4)) 与 (2,3,4) 的叶子数字相同，顶层结构却不同；后面给 Layout 传坐标时，这种区别就会表现出来。
IntTuple：用同一套结构表达 Shape、Stride 和坐标
不是所有 tuple 都只装整数。CuTe 把“一个整数，或者由这类对象递归组成的 tuple”称为 IntTuple。
所以这些都属于 IntTuple：
2
(2,3)
(2,(3,4))
第一个是单个整数；后两个是有不同层次结构的整数 tuple。
对 IntTuple，有三个常用操作：
| 操作 | 回答的问题 | 对 (2,(3,4)) 的结果 |
|---|---|---|
rank | 最外层有几个位置？ | 2 |
depth | 最深有几层 tuple？ | 2 |
size | 所有叶子数字相乘是多少？ | 2×3×4=24 |
这里的 size 是 IntTuple 的数值乘积，不是 tuple 的元素数量，也不是 C++ 对象占用的字节数。
再对照几个例子：
| 写法 | rank | depth | size |
|---|---|---|---|
6 | 1 | 0 | 6 |
(6) | 1 | 1 | 6 |
(2,3,4) | 3 | 1 | 24 |
(2,(3,4)) | 2 | 2 | 24 |
CuTe 将单个整数的 rank 约定为 1。因此单个整数 6 与单元素 tuple (6) 的 rank 相同，但 depth 不同。代码中后者要用 make_tuple(_6{}) 创建。
Shape、Stride、Coord 是不同用途的名字
它们都沿用 IntTuple 的结构：
auto shape = make_shape(_2{}, _3{});
auto stride = make_stride(_3{}, _1{});
auto coord = make_coord(1, 2);
- Shape 说明每个坐标分量有多大。
- Stride 说明该分量增加 1 时，输出增加多少。
- Coord 是这次访问实际使用的坐标。
这些名字用于说明角色，并不是三种互不兼容的数据容器。在源码中，Shape、Stride、Coord 都是 cute::tuple 的别名；如果叶子的类型一样，构造出来的 C++ 类型也可以相同。
Shape 和 Stride 需要有相同的树形结构。例如：
Shape = (2,(2,3))
Stride = (3,(6,1))
每个 Shape 叶子都能找到对应的 Stride 叶子。这叫结构一致，CuTe 用 congruent(shape,stride) 检查。
把上面的 Stride 写成 (3,6,1) 就不一致了：它的顶层有三个位置，而 Shape 的顶层只有两个。数值个数相同，不代表可以丢掉层次。
Layout：把 Shape 和 Stride 放到一起
先回到普通的二维例子：
auto row = make_layout(make_shape(_2{}, _3{}),
make_stride(_3{}, _1{}));
// (2,3):(3,1)
本文约定第一个坐标 m 是行，第二个坐标 n 是列，都从 0 开始。合法坐标满足 0≤m<2、0≤n<3，计算方式为：
row(m,n) = 3*m + n
沿行方向移动一步，跳过 3 个元素；沿列方向移动一步，前进 1 个元素。这就是行主序。
如果 Shape 不变，把 Stride 改成 (1,2)：
auto column = make_layout(make_shape(_2{}, _3{}),
make_stride(_1{}, _2{}));
// column(m,n) = m + 2*n
就得到列主序：
把两种布局的输出直接填进矩阵。在相同的 (1,1) 位置，行主序返回 4，列主序返回 3：
| m/n | 0 | 1 | 2 |
|---|---|---|---|
| 0 | 0 | 1 | 2 |
| 1 | 3 | 4 | 5 |
| m/n | 0 | 1 | 2 |
|---|---|---|---|
| 0 | 0 | 2 | 4 |
| 1 | 1 | 3 | 5 |
例如同一个坐标 (1,1)：
row(1,1); // 4
column(1,1); // 3
图中格内的数字是 Layout 的输出索引。若用它访问普通数组，它就是从数组起点开始的元素偏移；不是矩阵中保存的数据值，也不是字节偏移。
Layout 也可以输出其他编号，例如线程编号。这里先以数组偏移为例，便于把映射算清楚。
不写 Stride，会生成什么
auto layout = make_layout(make_shape(_2{},_3{}));
// (2,3):(1,2)
默认使用 LayoutLeft，左边的 mode 连续，按从左向右的前缀乘积生成步长。对这个普通二维 Shape，它得到列主序。
也可以显式指定：
auto left = make_layout(make_shape(_2{},_3{}), LayoutLeft{});
auto right = make_layout(make_shape(_2{},_3{}), LayoutRight{});
// left = (2,3):(1,2)
// right = (2,3):(3,1)
因此，Shape 并不决定行主序或列主序，不能只看 (2,3) 就判断元素如何存放。
嵌套 Layout：一个 mode 里面还可以有子坐标
现在将第二个 mode 拆成两部分：
auto h = make_layout(
make_shape(_2{}, make_shape(_2{},_3{})),
make_stride(_3{}, make_stride(_6{},_1{})));
// (2,(2,3)):(3,(6,1))
将完整坐标写成 (m,(a,b))。Shape 和 Stride 对应起来后，公式就是：
0≤m<2，0≤a<2，0≤b<3
h(m,(a,b)) = 3*m + 6*a + b
例如：
h(make_coord(1,make_coord(0,2))); // 3*1+6*0+2 = 5
这里有三个叶子坐标，不能把它直接当作普通的二维行列。可以先固定 a，分别画出两张以 m 为行、b 为列的矩阵：
| m/b | 0 | 1 | 2 |
|---|---|---|---|
| 0 | 0 | 1 | 2 |
| 1 | 3 | 4 | 5 |
| m/b | 0 | 1 | 2 |
|---|---|---|---|
| 0 | 6 | 7 | 8 |
| 1 | 9 | 10 | 11 |
高亮位置都满足 m=1、b=2。改变 a 从 0 到 1，会让输出从 5 增加到 11；增加的 6 正是 a 对应的 Stride。
第二个 mode 虽然有两个子坐标，在外层仍算一个 mode：
rank(h) = 2
depth(h) = 2
size(h) = 2×2×3 = 12
shape<1>(h) = (2,3)
size<1>(h) = 6
shape<1> 返回结构，size<1> 返回这部分所有大小的乘积。它们回答的不是同一个问题。
还可以继续向下取：
shape<1,1>(h); // 3
layout<1>(h); // (2,3):(6,1)，第二个 mode 自己也是一个子布局
这样的层次结构可以保留维度的分组含义。比如一组 6 个位置，可以按两个子坐标组织，后续代码便可以整体处理这一组，也可以继续访问里面的部分。
嵌套 Shape 的默认步长也按叶子生成
如果只给出相同的 Shape：
auto shape = make_shape(_2{},make_shape(_2{},_3{}));
auto left = make_layout(shape);
auto right = make_layout(shape,LayoutRight{});
结果为：
left = (2,(2,3)):(1,(2,4))
right = (2,(2,3)):(6,(3,1))
它们分别按所有叶子从左向右、从右向左生成紧凑步长，再放回原来的树形结构。本文手动指定的 (3,(6,1)) 是第三种规则，不能因为 Shape 相同，就把它与其中任一种混为一谈。
输入一个整数时，CuTe 怎样拆坐标
Layout 可以接收完整的 tuple 坐标，也可以接收一个整数编号。
先看普通的 row=(2,3):(3,1)：
row(3);
这不是直接返回偏移 3。整数 3 是输入坐标，还要先按 Shape (2,3) 分解：
i = m + 2*n
m = i % 2
n = i / 2 // 整数除法
i=3 → (m,n)=(1,1) → row(1,1)=4
分解时，第一个 mode 优先递增。这个顺序由 Shape 的坐标编号规则决定，与实际 Stride 是否行主序无关。
将整数输入编号 i 放回同一张逻辑矩阵，会得到下面的排列。注意，这次格内写的是输入编号，不是 Layout 的输出偏移：
| m/n | 0 | 1 | 2 |
|---|---|---|---|
| 0 | 0 | 2 | 4 |
| 1 | 1 | 3 | 5 |
高亮格的坐标是 (1,1)。它在这张输入编号图中写着 3，在前面的行主序输出图中写着 4，所以 row(3)=4。
| 整数输入 i | 按 Shape 拆出的 (m,n) | row(i) |
|---|---|---|
| 0 | (0,0) | 0 |
| 1 | (1,0) | 3 |
| 2 | (0,1) | 1 |
| 3 | (1,1) | 4 |
| 4 | (0,2) | 2 |
| 5 | (1,2) | 5 |
CuTe 也允许只将某个嵌套 mode 写成整数。对前面的 h，第二个 mode 的 Shape 是 (2,3)，因此：
j = a + 2*b
a = j % 2
b = j / 2
若整个 Layout 都用一个整数 i 表示，再在外面加一层分解：
i = m + 2*j
所以以下三种写法会得到同一个偏移 5：
h(make_coord(1,make_coord(0,2))); // 完整坐标
h(1,4); // 第二个 mode 用整数 4 表示
h(9); // 整个坐标用整数 9 表示
把 j=0…5 依次排在横轴上，也能直接查到同一个结果：
| m/j | 0 | 1 | 2 | 3 | 4 | 5 |
|---|---|---|---|---|---|---|
| 0 | 0 | 6 | 1 | 7 | 2 | 8 |
| 1 | 3 | 9 | 4 | 10 | 5 | 11 |
横轴的 j 是第二个 mode 的整数坐标，不是原来完整坐标中的 b。例如 j=4 会拆成 (a,b)=(0,2)，所以高亮格是 h(1,4)=5；整个输入再合成 i=1+2×4=9。
把整数输入 9 的计算过程逐步展开如下，表中的除法均为整数除法：
| 步骤 | 计算 | 结果 |
|---|---|---|
| 按外层 Shape 拆分 | m = 9 % 2 = 1，j = 9 / 2 = 4 | (m,j) = (1,4) |
| 按子 Shape 拆分 | a = 4 % 2 = 0，b = 4 / 2 = 2 | 完整坐标 (1,(0,2)) |
| 应用 Stride | 3*m + 6*a + b = 3*1 + 6*0 + 2 | 输出偏移 5 |
可以把整个过程分成两步：Shape 决定整数输入怎样拆成坐标，Stride 决定这些坐标怎样加成输出。
不要将整数坐标直接当成内存偏移。后面阅读布局复合、线程和值的映射时，这个区别会反复出现。
size 与 cosize：逻辑位置和覆盖范围
size 数的是逻辑坐标有多少种组合，它只依赖 Shape。
但真实的索引空间可能有空隙。比如一张 2×3 的矩阵，每行起点相隔 5 个元素：
auto padded = make_layout(make_shape(_2{},_3{}),
make_stride(_5{},_1{}));
// (2,3):(5,1)
它的输出表是：
| m/n | 0 | 1 | 2 |
|---|---|---|---|
| 0 | 0 | 1 | 2 |
| 1 | 5 | 6 | 7 |
共有 6 个逻辑位置，但若从偏移 0 开始覆盖到最后一个输出 7，需要 8 个位置：
size(padded) = 2×3 = 6
cosize(padded) = (2-1)×5 + (3-1)×1 + 1 = 8
cosize 包含偏移 3、4 这样的空隙。它不是“实际访问到的不同下标有几个”；这里访问到的不同下标只有 6 个。
将物理偏移从 0 到 7 排成一行，就能看清这些空隙。下面横轴表示物理偏移，不再是逻辑列坐标；✓ 表示被访问，— 表示未被访问：
| 偏移 | 0 | 1 | 2 | 3 | 4 | 5 | 6 | 7 |
|---|---|---|---|---|---|---|---|---|
| 访问 | ✓ | ✓ | ✓ | — | — | ✓ | ✓ | ✓ |
反过来，stride 为 0 时，可以有多个逻辑坐标访问同一个位置：
auto broadcast = make_layout(make_shape(_2{},_3{}),
make_stride(_0{},_1{}));
它的两行都会输出 0、1、2，因此 size=6、cosize=3。
| m/n | 0 | 1 | 2 |
|---|---|---|---|
| 0 | 0 | 1 | 2 |
| 1 | 0 | 1 | 2 |
黄色区域中的 (0,1) 和 (1,1) 是两个逻辑位置，却都对应偏移 1。矩阵有 6 个格子，不代表访问了 6 个不同的存储位置。
这里都采用非负整数 stride。对这些例子，cosize 就是最大输出偏移加 1；带负 stride 或输出更复杂索引的布局，需要结合具体定义另看。
flatten 和 coalesce 简化的是哪一部分
知道 tuple 的层次后，可以顺便区分这两个常见操作。
flatten(h) 去掉 Shape、Stride 的嵌套层次：
(2,(2,3)):(3,(6,1))
↓ flatten
(2,2,3):(3,6,1)
它保留整数输入 i 对应的输出，但坐标接口从 (m,(a,b)) 变成了 (m,a,b)。这一步没有保证内存连续。
coalesce 则尝试合并能连续接起来的相邻 mode，例如：
(2,3):(1,2) → 6:1
因为第二个 stride 2 等于第一个大小 2 乘以第一个 stride 1，这两个 mode 对整数输入的映射可以合成一条连续规则。
同样，整数输入的映射保留了，但原来的二维坐标接口不会自动保留。是否能删除括号，要看后续代码还需不需要它们表达的分组。
从这里继续读 Product 和 Divide
到这里，可以独立读懂一张 Layout：先看 tuple 的分组，再按 Shape 分解坐标，最后用 Stride 算出索引。
后续两篇按这个顺序展开即可：
| 篇章 | 从什么问题开始 | 逐步展开的内容 |
|---|---|---|
| Product | 已有一份布局，怎样重复成更大的结构？ | 从 logical_product 的小例子开始，再比较 zipped、tiled、flat，以及 blocked、raked 的坐标分组与用途 |
| Divide | 已有一份大布局，怎样拆成块内位置与块编号？ | 先算一个一维分块，再扩展二维；比较各类 divide，最后接到 tensor 切片与线程分区 |
两篇都沿用这里的记号和坐标映射基础，不再同时引入 tuple、嵌套 mode 和整数坐标分解。Layout Product 已有完整示例；想继续看 Tensor 切片、线程和值的布局，以及 MMA 的 ALayout，可以阅读从矩阵坐标到线程分工。
配套程序与参考源码
配套程序核对了 tuple 层次、各类坐标写法、默认步长，以及 size/cosize、flatten 和 coalesce 的行为，共 106 项检查。程序在 CPU 上运行，不需要 GPU：
# 在 CUTLASS 仓库根目录执行。
nvcc -std=c++17 -I include cute_layout_basics_trace.cu \
-o /tmp/cute_layout_basics_trace
/tmp/cute_layout_basics_trace
完整打印见运行结果。文中矩阵表的偏移已与同一程序输出逐项核对；输入编号图按文中的 Shape 分解公式排列，物理偏移图则标出了 padded 布局实际访问的位置。
相关源码与文档：
- container/tuple.hpp：tuple 的构造与访问。
- int_tuple.hpp：rank、depth、size 及层次结构操作。
- layout.hpp：Shape、Stride、Coord 的别名，以及 Layout 的实现。
- stride.hpp：
crd2idx按 Shape 递归拆分坐标，再应用 Stride。 - CuTe Layout 官方教程。
本文依据本地 CUTLASS 提交 0b55a2f691d69981583568fd9eb69687b1f0de8a 的 C++ 实现。文中的精确 Shape、Stride 以该版本实际编译打印的结果为准。