读懂 CuTe 的 Layout：从矩阵坐标到线程分工

假设有一个 3×4 的矩阵，按行连续放在内存中。访问第 1 行、第 2 列时，通常会写：

float value = data[1 * 4 + 2];

如果矩阵改成按列连续存放，同一个坐标就要换一种算法：

float value = data[1 + 2 * 3];

前一个下标是 6，后一个是 7。矩阵坐标没有变，变的是“这个坐标怎样对应到存储位置”。

CuTe 的 Layout 就是把这一步计算单独拿出来。先用小矩阵把它算清楚，再看嵌套括号、Tensor 切片，以及 GEMM 里那些线程和值组成的布局。

先把一张矩阵的地址表写出来

先固定行主序，也就是每一行内部连续存放：

逻辑坐标                    从 data 开始的元素偏移

      n=0  n=1  n=2  n=3          n=0  n=1  n=2  n=3
m=0   ·    ·    ·    ·      m=0    0    1    2    3
m=1   ·    ·    ·    ·      m=1    4    5    6    7
m=2   ·    ·    ·    ·      m=2    8    9   10   11

这里把行坐标叫 m，列坐标叫 n，都从 0 开始。沿着行坐标走一步，要跳过 4 个元素；沿着列坐标走一步，只前进 1 个元素。因此：

offset(m,n) = 4*m + n

用 CuTe 表达：

using namespace cute;

auto row = make_layout(
    make_shape(3,4),
    make_stride(4,1));

row(1,2);  // 6

Shape 表示坐标各维有多大，Stride 表示沿这一维走一步时，索引增加多少。CuTe 通常将布局打印成 shape:stride：

(3,4):(4,1)
  ↑     ↑
shape stride

对这组普通二维坐标，计算就是坐标与 stride 的逐项乘加：

坐标     (1,2)
stride   (4,1)

结果 = 1*4 + 2*1 = 6

此时 row(1,2) 只返回整数 6。它没有接收 data，也没有读取矩阵元素。

如果要访问数据，再把指针与 Layout 放到一个 Tensor 中：

float data[12] = {};
auto A = make_tensor(data + 0, row);

A(1,2);  // 访问 data[6]

对这里基于普通指针的 Tensor，这次访问可以读成：

A(m,n) = data[row(m,n)]

Layout 返回的偏移以元素为单位。float 的偏移 6 是 6 个 float，对应 24 字节；同一个 Layout 配上 half 指针，偏移 6 就是 6 个 half。

同一个 Shape，可以有不同的存放方式

保留 (3,4)，把 stride 改成 (1,3)：

auto column = make_layout(
    make_shape(3,4),
    make_stride(1,3));

现在沿行坐标前进只需要移动一个元素，得到列主序：

      n=0  n=1  n=2  n=3
m=0    0    3    6    9
m=1    1    4    7   10
m=2    2    5    8   11

同一个坐标 (1,2)，两种布局给出：

row(1,2)    = 1*4 + 2*1 = 6
column(1,2) = 1*1 + 2*3 = 7

如果让两个 Tensor 使用同一根指针，却配上不同的 Layout，它们会以不同坐标读取原有存储。创建这种视图时，内存里的元素没有被重新排列。若要得到一份重新排布的矩阵，还需要实际执行读写。

Shape 也没有规定坐标名称。本文约定第一个坐标是行，第二个坐标是列，才有上面的行主序、列主序说法。后面它们也可以表示线程和值。

size 和 cosize，分别在数什么

真实矩阵的行与行之间，经常留有额外空间。例如，逻辑上仍是 4 列，但每行的起点相隔 6 个元素：

auto padded = make_layout(
    make_shape(3,4),
    make_stride(6,1));

映射变成：

      n=0  n=1  n=2  n=3
m=0    0    1    2    3
m=1    6    7    8    9
m=2   12   13   14   15

size(padded) 数的是合法逻辑坐标，共 3*4=12 个。它没有计算行间空隙。

cosize(padded) 则为 16。对于本文这些普通整数 stride 非负的布局，从偏移 0 到最大可达偏移 15，需要覆盖 16 个位置：

size   = 3*4 = 12
cosize = (3-1)*6 + (4-1)*1 + 1 = 16

cosize 包含偏移 4、5、10、11 这样的空隙。因此，不能按 size=12 分配数组后，就用这张 padded 布局访问全部矩阵元素。若实际分配 3*6=18 个元素也可以，只是最后一行末尾的两个位置不会被这张布局访问。

再看另一个方向。把行 stride 改成 0：

auto broadcast = make_layout(
    make_shape(3,4),
    make_stride(0,1));

三行现在映射到相同的位置：

0  1  2  3
0  1  2  3
0  1  2  3

size   = 12
cosize = 4

12 个逻辑坐标只访问 4 个存储位置。(0,2)、(1,2)、(2,2) 都对应偏移 2。这种零 stride 可以表达广播；通过其中一个坐标修改数据，其他别名坐标读到的也是修改后的值。

所以，看到 Layout 时，应当分别看逻辑位置有多少、索引覆盖多大，以及是否有不同坐标指向同一个位置。

row(5) 为什么不是偏移 5

二维坐标会算之后，还有一个很容易误读的用法：

row(5);

这里仍然是在给 Layout 输入一个坐标，只不过把二维坐标写成了整数编号。CuTe 先按 Shape 将它拆开，再乘上 Stride。

对 Shape (3,4)，整数编号先让最左边的 mode 变化：

i = m + 3*n

m = i % 3
n = i / 3       // 整数除法

取 i=5：

5 → (5%3,5/3) → (2,1) → 2*4+1 → 9

因此：

row(5) == row(2,1);  // 都返回 9

将全部编号列出来，会更清楚：

整数坐标 i	等价的 (m,n)	row(i)
0	(0,0)	0
1	(1,0)	4
2	(2,0)	8
3	(0,1)	1
4	(1,1)	5
5	(2,1)	9
6	(0,2)	2
7	(1,2)	6
8	(2,2)	10
9	(0,3)	3
10	(1,3)	7
11	(2,3)	11

右边这串结果不是按物理地址从小到大排列。原因是整数坐标的拆解顺序由 Shape 决定，与 stride 是否行主序无关。

这一步最好在脑中分成两次计算：

整数坐标 i
    │ 根据 Shape 拆坐标：取余、整除
    ▼
完整坐标 (m,n)
    │ 根据 Stride 算索引：逐项乘加
    ▼
元素偏移 offset

对于配上 row 的 Tensor，A(5) 也会走这条路径，访问 data[9]。如果你要直接读取底层偏移为 5 的值，应当读 data[5]。

反过来，从物理偏移查坐标是另一个问题。本例的紧凑行主序布局中，偏移 5 对应 (1,1)，整数坐标为 4。CuTe 的 get_hier_coord、get_1d_coord 提供相关接口，但源码注明这些接口适用于 compact layout，不能把它们直接当成任意 padded 或广播布局的通用逆映射。

多一层括号，怎样改变坐标的含义

GEMM 代码里很少只看到 (3,4)，更常见的是嵌套 Shape。例如：

auto h = make_layout(
    make_shape(2, make_shape(2,2)),
    make_stride(4, make_stride(2,1)));

把它写成布局记号：

shape  = (2,(2,2))
stride = (4,(2,1))

Shape 与 Stride 的括号结构一一对应：第一个大小 2 配步长 4，后面的 (2,2) 配步长 (2,1)。

CuTe 将一个维度叫作 mode。这里最外层有两个 mode：第一个大小是 2，第二个由两个大小为 2 的子 mode 组成。第二个顶层 mode 一共容纳 2*2=4 个位置，所以仍然可以把它画成一张 2×4 的表。

先用最完整的坐标表示：

坐标   (m,(a,b))
索引 = 4*m + 2*a + b

其中 m、a、b 都只能取 0 或 1。

如果仍想写 h(m,j)，就要把 j 拆到第二个 mode 的 (2,2) 中。这里继续采用最左子 mode 先变化的顺序：

a = j % 2
b = j / 2

j	(a,b)	第二个 mode 贡献的偏移 2*a+b
0	(0,0)	0
1	(1,0)	2
2	(0,1)	1
3	(1,1)	3

于是整张表是：

      j=0  j=1  j=2  j=3
m=0    0    2    1    3
m=1    4    6    5    7

第二列与第三列的顺序交换了。整数坐标 j=1 先拆成 (1,0)，再乘上 (2,1)，因此得到 2；它不会根据“哪条 stride 更小”重新选择拆解顺序。

下面三种坐标写法访问同一个位置：

h(make_coord(1, make_coord(1,0)));  // 1*4 + 1*2 + 0*1 = 6
h(1,1);                             // j=1 拆成 (1,0)，得到 6
h(3);                               // 3 拆成 (m=1,j=1)，得到 6

嵌套 Shape 的用途就在这里：既能保留子结构，也能将某个子结构临时当成一个整数 mode 使用。后面 MMA 的“一个线程持有 8 个值，但这 8 个值又分成 (2,2,2)”也是相同的表达方式。

rank、depth、shape 与 size

对上面的 h：

表达式	结果	含义
rank(h)	2	顶层有两个 mode
depth(h)	2	最深的 tuple 嵌套有两层
shape(h)	(2,(2,2))	完整的大小结构
size(h)	8	全部逻辑位置数量
shape<1>(h)	(2,2)	第二个 mode 的内部结构
size<1>(h)	4	第二个 mode 的逻辑位置数量
size<1,0>(h)	2	第二个 mode 中，第一个子 mode 的大小
layout<1>(h)	(2,2):(2,1)	取出第二个 mode 的 Shape 和 Stride

shape<1> 保留括号，size<1> 则将里面的大小相乘。rank 只数最外层，所以这里不会因为有三个整数叶子就变成 3。

flatten 会去掉括号，数据仍在原处

对同一个 h 调用：

auto flat = flatten(h);

得到：

h    = (2,(2,2)):(4,(2,1))
flat = (2,2,2):(4,2,1)

原来的完整坐标 (m,(a,b))，现在写成 (m,a,b)。逐项乘加的公式仍是 4*m+2*a+b。使用单个整数坐标遍历时，两者也有相同的结果：

flat(i) == h(i);  // i=0..7

这里的 flatten 是展开嵌套结构，结果的 rank 为 3。它没有把数组重新排成连续的一维数据，也没有把所有 mode 合成一个 mode。

反过来，把 flat 的第 1、2 个 mode 分组：

auto again = group<1,3>(flat);  // 区间 [1,3)，包含 mode 1、2

就恢复了 (2,(2,2)):(4,(2,1))。这些操作让同一组索引关系可以按不同层级组织。

省略 Stride，CuTe 会替你生成什么

常见代码只写 Shape：

auto L = make_layout(make_shape(3,4));

默认使用 LayoutLeft。对这个二维 Shape，它从左向右生成紧凑 stride：

shape  = (3,4)
stride = (1,3)

也就是前面的列主序。若希望生成行主序：

auto R = make_layout(make_shape(3,4), LayoutRight{});
// (3,4):(4,1)

对嵌套 Shape，这两个标签会跨过括号，按全部整数叶子的顺序生成 stride。继续用 (2,(2,2))：

LayoutLeft  → (2,(2,2)):(1,(2,4))
LayoutRight → (2,(2,2)):(4,(2,1))

后者就是刚才的 h，画成二维表时却是 0,2,1,3。这并不矛盾：LayoutRight 决定如何生成嵌套坐标的 stride，而整数 j 仍按 Shape 的固定顺序拆成 (a,b)。

因此，阅读嵌套布局时，直接看 Shape 与 Stride 的对应关系，会比只记“Left 是列主序、Right 是行主序”更可靠。

打印里的 _3，为什么带一条下划线

到目前为止，例子都用普通整数构造。CuTe 还提供把数值写入类型的静态整数：

Int<3>{}   // 编译期已知的 3
_3{}       // 同一个静态整数类型的简写

因此，下面两个布局具有相同的数值映射，但类型不同：

auto dynamic_row = make_layout(
    make_shape(3,4), make_stride(4,1));

auto static_row = make_layout(
    Shape<_3,_4>{}, Stride<_4,_1>{});

打印结果为：

dynamic_row = (3,4):(4,1)
static_row  = (_3,_4):(_4,_1)

Shape<_3,_4> 表达的仍然是 (3,4)；它将两个维度大小写在模板类型中。Stride 同理。

两种形式也可以混用：

int columns = 4;

auto mixed = make_layout(
    make_shape(_3{}, columns),
    make_stride(columns, _1{}));
// (_3,4):(4,_1)

这张布局的行数和列步长在类型中已知，列数与行步长则由运行时值提供。静态信息使 CuTe 能在编译期进行布局推导和约束检查；普通整数即使刚好写成字面量 4，类型仍然是 int，不会自动变成 _4。

本文其他布局记号省略 _，方便看数字。配套程序主要使用静态布局，因此实际打印会带这个前缀。

切出一行时，原来的行偏移去了哪里

回到 padded 矩阵：

(3,4):(6,1)

第 1 行的元素偏移应当是 6,7,8,9。现在分别对 Layout 和 Tensor 切片：

auto row_layout = padded(1,_);

int storage[16] = {};
auto A = make_tensor(storage + 0, padded);
auto row_view = A(1,_);

_ 表示保留这一维，1 表示固定行坐标。row_layout 与 row_view.layout() 都会得到：

shape  = (4)
stride = (1)

但直接调用 row_layout(0) 返回 0，并不是第 1 行的原始偏移 6。原因是普通 Layout 只保存 Shape 与 Stride，固定坐标带来的常量偏移由切片的另一部分处理。

slice_and_offset 同时返回这两部分：

auto [new_layout, offset] = slice_and_offset(
    make_coord(1,_), padded);

// new_layout = (4):(1)
// offset     = 6

Tensor 的切片源码正好把它们接起来：

auto [sliced_layout, offset] = slice_and_offset(coord, layout());
return make_tensor(data() + offset, sliced_layout);

于是 row_view 的指针已经前进到 storage+6，新布局只需要计算行内偏移：

row_view(j)
  = (storage+6)[j]
  = storage[6+j]

切出第 2 列也一样：

auto column_view = A(_,2);

// 指针   = storage + 2
// layout = (3):(6)

它访问 storage[2]、storage[8]、storage[14]。两个视图都使用原来的数组；通过视图写入，也会修改原数组。

这个例子中没有执行数据搬运。变化的是视图指针和剩余坐标的索引规则。

带着这几个数字，看 Layout 的调用源码

现在再看 Layout 的定义，主要信息已经可以对应上了。省略访问器后，它的结构是：

template <class Shape, class Stride = LayoutLeft::Apply<Shape>>
struct Layout : private cute::tuple<Shape,Stride>
{
  // ...
};

布局保存的就是 Shape 与 Stride。它没有保存矩阵数据的指针。

多个坐标参数会先组合成一个坐标 tuple：

row(1,2);
// 转给 operator()(make_coord(1,2))

接着进入：

template <class Coord>
auto operator()(Coord const& coord) const {
  if constexpr (has_underscore<Coord>::value) {
    return slice(coord, *this);
  } else {
    return crd2idx(coord, shape(), stride());
  }
}

有 _ 就切布局，没有 _ 就调用 crd2idx 计算索引。后者位于 stride.hpp，可以按三条规则读：

当前处理的情况	运算	本文例子
坐标、Shape、Stride 都是整数	coord*stride	行坐标 1、行步长 4，贡献 4
三者都是 tuple	各 mode 递归计算，再求和	(1,2) 配 (4,1)，得到 4+2
坐标是整数，Shape/Stride 是 tuple	先按子 Shape 的大小取余、整除，再递归	row(5) 先将 5 拆成 (2,1)

嵌套坐标只是让这三条规则多走几层。h(1,1) 的实际计算可以展开为：

h 的 Shape  = (2,(2,2))
h 的 Stride = (4,(2,1))
输入坐标    = (1,1)

第一 mode：1*4 = 4
第二 mode：1 拆成 (1,0)，得到 1*2 + 0*1 = 2

最终结果：4+2 = 6

Shape 在这里参与坐标解释。它没有替你分配数据，也没有让一次越界的 Tensor 读取自动变得有效；kernel 的边界处理仍要另外安排。

Layout 为什么还会用来描述线程和值

到这里看到的布局，都是“矩阵坐标→存储偏移”。但 Layout 本身不认识矩阵或内存：它接收坐标，按 Shape 与 Stride 返回一个索引。

把输入换成 (thread,value)，它就可以描述线程分工。thread 是线程编号，value 是当前线程负责的第几个值。CuTe 将这类布局叫作 TV 布局。

先用四个逻辑线程分担一个 2×4 tile，每个线程负责两个值：

auto tv = make_layout(
    make_shape(4,2),
    make_stride(1,4));

// tv(thread,value) = thread + 4*value

这里约定它返回的整数 q 按 q=m+2*n 编码 tile 的逻辑坐标。因为 Shape 是 (2,4)，因此：

m = q % 2
n = q / 2

各线程对应的元素为：

thread	value 0	value 1
0	q=0 → (0,0)	q=4 → (0,2)
1	q=1 → (1,0)	q=5 → (1,2)
2	q=2 → (0,1)	q=6 → (0,3)
3	q=3 → (1,1)	q=7 → (1,3)

tile 的实际存储，则交给另一张布局：

auto data_layout = make_layout(
    make_shape(2,4),
    make_stride(4,1));

现在可以将两层关系连起来。例如，线程 0 的第 1 个值：

(thread=0,value=1)
        │ tv
        ▼
q=4，表示矩阵坐标 (0,2)
        │ data_layout
        ▼
物理元素偏移 2

q=4 并不表示存储偏移一定是 4。它是这套 TV 规则采用的逻辑位置编号；数据 Layout 再决定该位置实际放在哪里。

CuTe 可以用 composition 接上这两层：

auto tv_to_address = composition(data_layout, tv);

tv_to_address(0,1);  // data_layout(tv(0,1)) = data_layout(4) = 2

前面解释过整数坐标的拆解，因此这里的 data_layout(4) 能将 4 解释为 (0,2)。composition 的具体构造算法可以留到布局代数一篇再展开；此处先跟住这一条映射即可。

这张例子里的 tv 只是定义分工规则，没有创建 CUDA 线程，也没有发出 copy 指令。

从小例子走到 MMA 的 ALayout

再看一张真实的 TV 布局，就能与 retile 那篇接起来了。那篇使用：

auto mma = make_tiled_mma(SM80_16x8x16_F32F16F16F32_TN{});

这条 SM80 MMA 由一个 warp 的 32 个线程协同执行，完成 M=16、N=8、K=16 的矩阵乘加。A tile 是 16×16，共 256 个 half；每个线程提供其中的 8 个 half。

哪些 A 元素由哪个线程提供，由指令接口规定。CuTe 在 mma_traits_sm80.hpp 中，用下面的布局记录这套分工；本文的 FP32 累加版本继承了这张 ALayout：

using ALayout = Layout<
    Shape <Shape < _4,_8>, Shape < _2,_2,  _2>>,
    Stride<Stride<_32,_1>, Stride<_16,_8,_128>>>;

Shape 的两个顶层 mode 分别是：

((4,8), (2,2,2))
  线程     每线程的值

第一个 mode 容纳 4*8=32 个线程，第二个容纳 2*2*2=8 个值。按照前面的整数拆解规则：

lane → (lane%4, lane/4)

v → (a,b,c)
a = v%2
b = (v/2)%2
c = v/4

再与 stride 逐项乘加：

q = 32*(lane%4) + lane/4 + 16*a + 8*b + 128*c

这里的 q 按 A 的 (M,K) 逻辑坐标编码：q=m+16*k。因此可以读出：

m = lane/4 + 8*b
k = 2*(lane%4) + a + 8*c

把 lane 0、lane 1 的 8 个值列出来：

v	(a,b,c)	lane 0	lane 1
0	(0,0,0)	A(0,0)	A(0,2)
1	(1,0,0)	A(0,1)	A(0,3)
2	(0,1,0)	A(8,0)	A(8,2)
3	(1,1,0)	A(8,1)	A(8,3)
4	(0,0,1)	A(0,8)	A(0,10)
5	(1,0,1)	A(0,9)	A(0,11)
6	(0,1,1)	A(8,8)	A(8,10)
7	(1,1,1)	A(8,9)	A(8,11)

回头看那几层括号，就有具体含义了：a 选择相邻的两个 K 值，b 选择上下相隔 8 行的位置，c 选择左右相隔 8 列的位置。

这里仍然要区分指令分工与实际存储。若 A 在 shared memory 中采用 (16,16):(16,1)，lane 0 的 v=1 对应：

ALayout(0,1) = 16
      │ 按 q=m+16*k 解码
      ▼
A(0,1)
      │ 按实际数据布局算偏移
      ▼
shared memory 中的元素偏移 1

换一张数据 Layout，会改变这个元素的存储地址；指令要求 lane 0 提供 A(0,1) 这件事仍然不变。加载代码需要将对应元素送到 lane 0 的指定寄存器中。

至于一个线程内部，这 8 个值怎样放进它的寄存器 fragment，又是一张局部存储布局。在 retile 文章的单个 MMA 例子里，它是：

tCrA.layout() = ((2,2,2),1,1):((1,2,4),0,0)

它把 (a,b,c) 映射到当前线程的 half 槽位 a+2*b+4*c。这与上面的 ALayout 回答的问题不同：前者定位线程自己的存储，后者定位整个 A tile 中的矩阵元素。

同一套 Layout 结构就这样用在了三个地方：

布局	输入	输出的含义
A 数据布局	矩阵坐标 (m,k)	相对于数据指针的元素偏移
MMA ALayout	(lane,v)	A tile 内的逻辑位置编号
tCrA.layout()	当前线程的 fragment 坐标	当前线程内部的 half 存储槽位

看 GEMM 时，先确认一张布局的输入和输出分别是什么，再代入数字。那些相似的括号就不容易混在一起了。

用程序把这些映射跑一遍

完整验证程序放在 cute_layout_trace.cu，正文只保留了相关代码片段。它检查全部矩阵坐标、嵌套坐标、切片别名，以及线程分工的对应关系。

在当前仓库根目录编译：

cd /home/x/workspace/cutlass
nvcc -std=c++17 -I include cute_layout_trace.cu -o /tmp/cute_layout_trace
/tmp/cute_layout_trace

程序只在 host 上执行，不启动 CUDA kernel。MMA 部分读取的是指令的布局元信息，没有执行 GPU MMA 指令。

几个关键输出是：

row(i): 0 4 8 1 5 9 2 6 10 3 7 11
hierarchical: (_2,(_2,_2)):(_4,(_2,_1))
   0   2   1   3
   4   6   5   7
flatten(h): (_2,_2,_2):(_4,_2,_1)
slice_and_offset offset: 6
t=0 v=1 -> q=4 -> (0,2) -> offset=2
lane=0 v=1 -> A(0,1) q=16 physical_offset=1
Verified 951 checks; 8 tiny-tile values and 256 MMA A values covered.

课后实践：自己改出一张布局

前面的程序用来核对正文。下面换几组参数，留出六处代码自己填写：前五题从地址计算走到数据操作，第六题选做，用真实 MMA 布局检查线程分工。

练习文件是 cute_layout_exercises.cu。需要填写的函数集中在文件开头，分别标了 TODO(1) 到 TODO(6)；后面的验收代码不需要修改。它们也都在 host 上运行，第五题里的“线程”只是循环变量，不会启动 GPU 线程。

先编译练习版，再单独运行第一题：

cd /home/x/workspace/cutlass
nvcc -std=c++17 -I include cute_layout_exercises.cu -o /tmp/cute_layout_exercises
/tmp/cute_layout_exercises 1

文件里已经放了能编译的初始实现，但六题都有一处关键逻辑需要修改。刚开始出现 FAIL 是预期行为，提示会同时打印实际值和期望值。改完一题后重新编译，把最后一个参数换成下一题的编号；不带参数会运行全部六题。

建议每题都先在纸上算一个坐标，再改代码。只看 PASS 容易漏掉最值得想清楚的一步：输入的数字究竟是什么，输出的数字又表示什么。

第一题：给矩阵的每行留出空位

有一张 3×4 矩阵，要求第 0 行从偏移 0 开始，第 1 行从偏移 8 开始，第 2 行从偏移 16 开始。每行内部的四个元素仍然连续存放。

修改 exercise1_layout()，保持 Shape 为 (3,4)，填写合适的 Stride。然后回答三个问题：

1. 三行的地址表分别是什么？
2. size 和 cosize 分别是多少？
3. 最少需要多少个 int 才能访问所有有效元素？最后一行末尾是否必须再留四个空位？

验收代码会逐个检查十二个坐标，而不只是检查 padded(0,0)。后者无论 stride 填多少都是 0，无法说明布局是否正确。

第二题：把行主序矩阵复制成列主序

这次真的搬一次数据。源 Tensor 和目标 Tensor 都是 3×4，布局分别为：

src.layout() = (3,4):(4,1)
dst.layout() = (3,4):(1,3)

源矩阵已经按 src(m,n)=10*m+n 填好了：

 0   1   2   3
10  11  12  13
20  21  22  23

修改 exercise2_copy()，使复制后每个 dst(m,n) 都等于 src(m,n)。初始代码直接按底层数组下标复制，能够运行，但结果不满足这个要求。

先用双层 (m,n) 循环完成，再试着只用一个 i，通过 Tensor 的 operator() 完成。这里两个 Tensor 的 Shape 完全相同，可以使用相同的整数逻辑坐标枚举它们。

验收时检查的是完整矩阵的值。程序还会打印目标底层数组，观察它与上面的矩阵逐行展开有什么不同。顺便算一下：src.layout()(5) 与 dst.layout()(5) 为什么可以不同，但 src(5) 与 dst(5) 应该相等？

第三题：把一个 mode 拆成两层

构造这样一张布局：

Shape = (2,(2,3))
h(m,(a,b)) = 6*m + 3*a + b

其中 m、a 的范围都是 [0,2)，b 的范围是 [0,3)。修改 exercise3_layout() 中的嵌套 Stride。

这题先别急着填代码，先做三次同一个位置的计算：

h(1, (0,2))
h(1, 4)
h(9)

它们应该得到同一个偏移。第二种写法需要把 4 拆到第二个 mode 的 (2,3) 中；第三种写法还要先拆出最外层的 m。

验收程序会遍历全部十二个位置，检查自然嵌套坐标、二维坐标和整数坐标是否一致，还会检查 flatten(h) 保留了整数坐标的映射。自己再打印一下 flatten(h)，确认 group<1,3> 为什么能够恢复原来的层次。

第四题：不搬数据，做一个转置视图

现在手上是一张带空位的 Tensor A：

A.layout() = (3,4):(8,1)
A(m,n) = 10*m+n

修改 exercise4_transpose()，返回一个 4×3 的 Tensor B，要求：

B(n,m) == A(m,n)
B.data() == A.data()

只允许构造新视图，函数里不写复制循环。Shape 已经交换好了，剩下的是填写新 Stride。

接下来，验收程序会执行：

auto row = B(2,_);
row(1) = 777;

先预测这次写入会改动 A 的哪个元素、底层数组的哪个位置。再检查 row 的起始指针和 stride：它是 B 的一行，却不一定在物理存储上连续。

验收条件除了转置后的全部值，还包括：原来的数据指针没有变；写入能通过 A 读到；其他元素和行间空位保持原值。这题完成后，应该能分清第二题的“复制成另一种存储布局”和这一题的“换一套坐标看同一份存储”。

第五题：让八个线程把 tile 填满

把一个 4×4 tile 分给八个逻辑线程，每个线程持有两个值。实际数据布局带有行间空位：

data_layout = (4,4):(6,1)

线程分工则规定为：

q = t + 8*v
0 <= t < 8
0 <= v < 2

修改 exercise5_tv()，构造 Shape 为 (8,2) 的 TV Layout。q 使用 q=m+4*n 的逻辑编码，不是上面那张带空位布局的物理偏移。

然后沿着下面的路径，手算 (t,v)=(1,1)：

(t,v) -> q -> (m,n) -> data_layout(m,n)

验收程序用 composition(data_layout,tv) 得到物理偏移，再向那里写入 100*t+v。运行后把打印的矩阵读一遍：每个数都能反推出是哪个线程的第几个值写入的。

这题必须同时满足三个条件：十六个逻辑位置恰好各写一次；组合布局与手动解码的物理地址一致；行间空位没有被写入。只检查“没有重复”还不够，也要检查有没有遗漏。

第六题，选做：给 lane 4 找齐八个 A 元素

沿用正文中的 SM80_16x8x16_F32F16F16F32_TN。修改 exercise6_mma_coord(lane,v)，返回每个 (lane,v) 对应的矩阵坐标 (m,k)，不要在这个函数里直接查询 MMA 的 TV Layout。

从 ALayout 的分层 stride 出发，先把 v 拆成三个二进制分量，再算出 m 和 k。可以先只做 lane 4：列出它的八个元素，并为每个元素分别计算：

逻辑位置编号 q = m + 16*k
行主序存储偏移 = 16*m + k

验收程序会拿你的公式与 mma.get_layoutA_TV() 对照，检查全部 32×8 个位置，并确认 A tile 的 256 个元素各覆盖一次。这里只读取真实指令的布局元信息，不执行 MMA 指令。

做完再对照参考结果

下面给出关键结果。完整参考实现也保存在同一个练习文件的 CUTE_LAYOUT_EXERCISES_REFERENCE 分支里。编译时开启这个宏，可以独立运行参考版，不必覆盖自己填写的代码：

cd /home/x/workspace/cutlass
nvcc -std=c++17 -I include -DCUTE_LAYOUT_EXERCISES_REFERENCE \
  cute_layout_exercises.cu -o /tmp/cute_layout_exercises_reference
/tmp/cute_layout_exercises_reference

第一题：地址表与最小存储范围

make_layout(Shape<_3,_4>{}, Stride<_8,_1>{})

 0   1   2   3
 8   9  10  11
16  17  18  19

size=12，cosize=20。最大有效偏移为 19，分配二十个 int 就能容纳全部有效元素。可以额外分配到二十四个，让最后一行也占满八个槽位，但这不是当前布局访问所必需的范围。

第二题：复制的是同一个逻辑坐标

可以使用双层循环，执行 dst(m,n)=src(m,n)；本题也可以直接写：

for (int i = 0; i < int(size(src)); ++i) {
  dst(i) = src(i);
}

这里 i=5 对应 (m,n)=(2,1)。源和目标分别通过自己的 Layout 找到偏移 9 和 5，读写的矩阵值都是 21。直接用 data()[i] 会跳过这一步映射。

目标底层数组应为：

0 10 20 1 11 21 2 12 22 3 13 23

第三题：先拆坐标，再乘 stride

make_layout(Shape<_2,Shape<_2,_3>>{},
            Stride<_6,Stride<_3,_1>>{})

按 (m,j) 打印的地址表是：

0 3 1 4 2 5
6 9 7 10 8 11

第二个 mode 中，a=j%2，b=j/2。因此 j=4 拆成 (0,2)；整个 Shape 中，整数 9 先拆成 m=9%2=1、j=9/2=4。三种写法最终都是 6*1+3*0+2=8。

这张布局的 rank=2、depth=2、size=12。flatten(h) 是 (2,2,3):(6,3,1)，随后 group<1,3> 把后两个 mode 重新放进一组。

第四题：转置视图的 stride 也要交换

B 的布局应为 (4,3):(1,8)。练习函数可以从输入 Tensor 中取出 stride：

auto transposed_stride = make_stride(stride<1>(a), stride<0>(a));

B(2,_) 的指针为 A.data()+2，剩余布局为 (3):(8)，初始值是 2,12,22。

写 row(1)=777 就是写 B(2,1)，也就是 A(1,2)。它的物理偏移是 8*1+2=10。发生变化的是同一份存储中的一个元素。

第五题：线程分工与物理地址分别计算

make_layout(Shape<_8,_2>{}, Stride<_1,_8>{})

对于 (t,v)=(1,1)，先得到 q=9，再得到 (m,n)=(1,2)，最后得到物理偏移 6*1+2=8。用 q=9 直接当底层数组下标，会写到另一个位置。

矩阵中的 100*t+v 应为：

  0 400   1 401
100 500 101 501
200 600 201 601
300 700 301 701

例如第一行第三列的 1，表示线程 0 的第 1 个值。逻辑位置的访问计数最终全是 1，物理行间的空位仍然是初始值 -99。

第六题：lane 4 的完整结果

先拆 v：

int a = v % 2;
int b = (v / 2) % 2;
int c = v / 4;
int m = lane / 4 + 8*b;
int k = 2*(lane % 4) + a + 8*c;

代入 lane 4，得到：

v	A 坐标 (m,k)	逻辑编号 q	行主序物理偏移
0	(1,0)	1	16
1	(1,1)	17	17
2	(9,0)	9	144
3	(9,1)	25	145
4	(1,8)	129	24
5	(1,9)	145	25
6	(9,8)	137	152
7	(9,9)	153	153

q 和物理偏移偶尔相等，不能据此把它们当成一回事。比如 v=5 时，两者分别是 145 和 25。

上述参考版已在当前仓库编译运行，六题全部通过，最终输出：

Passed 6/6 exercises.

对照源码

本文基于当前仓库提交 0b55a2f6，只使用 C++ CuTe。

• layout.hpp：Layout、operator()、make_layout、size/cosize、slice_and_offset、flatten/group。
• stride.hpp：crd2idx 怎样递归解释坐标，并计算索引。
• int_tuple.hpp：嵌套整数 tuple 的 rank、depth 与 size。
• tensor_impl.hpp：Tensor 索引，以及切片时指针偏移与新 Layout 的组合。
• mma_traits_sm80.hpp：SM80 MMA 的线程和值布局。
• CuTe Layout 官方教程：本文基础概念的完整参考。
• retile_D 文章：继续看同一份寄存器怎样分别用于 Copy 写入和 MMA 读取。