CUTLASS CuTe IR 实战（四）：Accessors 与 Size/Indexing 运算——从布局中提取结构、计算尺寸并在坐标与线性索引间转换

CUTLASS CuTe IR 实战（四）：Accessors 与 Size/Indexing 运算——从布局中提取结构、计算尺寸并在坐标与线性索引间转换

【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlass

本指南是 CUTLASS 编译栈（cutlass_compiler）CuTe IR 概念系列的第四篇，聚焦于两组纯函数（Pure）运算：Accessors（访问器）——从!cute.layout、!cute.composed_layout、!cute.tile以及各类元组值中读取单个 mode、叶子或标量；以及Size / Indexing（尺寸与索引）——计算size、cosize、product等标量量，并在多维坐标与线性索引之间相互转换。读完本文，你将掌握 CuTe IR 中 18 个访问/尺寸/索引运算的语义、MLIR 写法、边界条件（错误检查），并了解它们如何被方言测试与静态折叠优化所验证，可直接用于阅读或编写 CuTe 布局代数 IR。

原始文档：04_accessors_and_size_index.rst。配套阅读：01_cute_types.rst（Shape/Stride/Layout/Tile/ComposedLayout/Swizzle 类型基础）与 cute_dialect.rst（完整的操作数类型、结果类型与汇编格式规范）。

两组运算的定位

本章涉及的运算按功能分为两组：

• Accessors（访问器）——从已有的 CuTe 值（布局、tile、复合布局、元组）中读出单个 mode、叶子或标量。它们是 01_cute_types.rst 中构造类运算（cute.make_layout、cute.make_composed_layout等）的读侧对应物。这类运算不改变语义，只是把已存在的结构"取出"为 IR 后续可以使用的 SSA 值。
• Size and indexing（尺寸与索引）——计算关于 CuTe 值的标量量（size、cosize、product），并在坐标与线性索引之间转换。

两类运算都是纯函数（无副作用），且接受任何兼容的 CuTe 值作为输入。从源码实现看，CuteOps.td 中所有相关 op 定义都带有[Pure, ...]特征（例如Cute_GetShapeOp、Cute_Crd2IdxOp均声明了Pure与InferTypeOpInterface），见 CuteOps.td。其中Pure保证无副作用，使这些运算可以被CuteFoldStatic等变换安全地常量折叠；InferTypeOpInterface则允许编译器根据输入自动推断结果类型。确切的操作数类型、结果类型和汇编格式，请以方言参考 cute_dialect.rst 为准。

---

Accessors（访问器）

访问器运算把组件从已有的 CuTe 值中抽取出来，是构造运算的读侧。下表汇总 9 个访问器（与文档保持一致）：

Op	功能
cute.get_shape	从布局、复合布局或 tile 中提取 shape。
cute.get_stride	从布局中提取 stride。
cute.composed_get_inner	提取复合布局的内部（A）组件。
cute.composed_get_offset	提取复合布局的 offset。
cute.composed_get_outer	提取复合布局的外部（B）布局。
cute.get	按层次化的 mode 索引提取子值。
cute.select	从选定的顶层 mode 子集中构造新值。
cute.get_leaves	把元组类值展平为每个叶子一个 SSA 结果。
cute.get_scalars	从 CuTe 值中提取标量整数（i32）。
cute.get_layouts_from_tile	从 tile 中提取每个布局槽位，跳过下划线。

布局 / stride / shape 提取

cute.get_shape与cute.get_stride返回布局的对应组成部分，结果为匹配类型的 CuTe 值：

%l = cute.static : !cute.layout<"(4,8):(1,4)"> %s = cute.get_shape(%l) : !cute.layout<"(4,8):(1,4)"> -> !cute.shape<"(4,8)"> %d = cute.get_stride(%l) : !cute.layout<"(4,8):(1,4)"> -> !cute.stride<"(1,4)">

cute.get_shape还接受复合布局或 tile。op 定义（CuteOps.td）明确了三种输入的行为：

• !cute.layout——返回布局自身的 shape（如从(4,8):(1,4)得到(4,8)）；
• !cute.composed_layout——返回外层 B 布局的 shape（例如S<3,5,4> o 0 o (8,4):(1,8)→(8,4)）；
• !cute.tile——返回一个 shape，其第 i 个元素是 tile 中第 i 个布局的 shape（例如[(4,8):(1,4);(2,3):(1,2)]→((4,8),(2,3))）。此时结果是各槽位 shape 的拼接。

值得注意的约束：含下划线_槽位的 tile 会被拒绝——下划线槽位没有 shape，无法贡献元素。若 tile 含下划线，应先通过cute.get_layouts_from_tile投影出布局槽位，再逐个取 shape。对应错误场景可参考 get_shape_errors.mlir 与 get_shape.mlir。

cute.get_stride仅接受!cute.layout输入，且按原样保留 scaled-basis stride（如(2,3):(1@0,1@1)→(1@0,1@1)，见 CuteOps.td）。传入复合布局或声明的结果 stride 与输入不一致都会被拒绝。

复合布局组件提取

一个!cute.composed_layout在内部表示为A ∘ offset ∘ B（坐标 c 映射到A(offset + B(c))）。三个 op 分别取出各组件：

%cl = cute.static : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)"> %a = cute.composed_get_inner(%cl) : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)"> -> !cute.layout<"(4,5):(1,4)"> %off = cute.composed_get_offset(%cl) : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)"> -> !cute.int_tuple<"2"> %b = cute.composed_get_outer(%cl) : !cute.composed_layout<"(4,5):(1,4) o 2 o (2,3):(1,2)"> -> !cute.layout<"(2,3):(1,2)">

当A是 swizzle 时，cute.composed_get_inner返回!cute.swizzle而非!cute.layout（op 定义见 CuteOps.td）：

%a = cute.composed_get_inner(%cl) : !cute.composed_layout<"S<3,5,4> o 0 o (8,4):(1,8)"> -> !cute.swizzle<"S<3,5,4>">

其余两个 op 的约束类似：composed_get_offset返回 offset!cute.int_tuple，composed_get_outer返回外层B布局；输入不是composed_layout或声明类型与实际组件不匹配时均报错。对应测试见 composed_get_inner.mlir、composed_get_offset.mlir、composed_get_outer.mlir。

模式提取：cute.get与cute.select

cute.get接收一个层次化索引列表（作为<[...]>属性），返回该路径上的子值；空列表返回输入本身：

%l = cute.static : !cute.layout<"((4,2),(3,8)):((1,4),(8,24))"> // 第一个顶层 mode %m0 = cute.get<[0]> (%l) : !cute.layout<"((4,2),(3,8)):((1,4),(8,24))"> -> !cute.layout<"(4,2):(1,4)"> // 深入到 mode (0, 1)：内部的 2:4 子布局 %m01 = cute.get<[0, 1]> (%l) : !cute.layout<"((4,2),(3,8)):((1,4),(8,24))"> -> !cute.layout<"2:4">

cute.select从顶层 mode 的选定子集构造新值，按列出顺序排列——允许重复和重排：

%l = cute.static : !cute.layout<"(4,2,8):(1,4,8)"> // 按顺序选取 mode 2 与 mode 0 %sw = cute.select<[2, 0]> (%l) : !cute.layout<"(4,2,8):(1,4,8)"> -> !cute.layout<"(8,4):(8,1)"> // 仅选 mode 0（结果仍是 rank-1 值） %m0 = cute.select<[0]> (%l) : !cute.layout<"(4,2,8):(1,4,8)"> -> !cute.layout<"(4):(1)">

注：原始文档中select<[0]>示例的类型注解与实际输入%l存在不一致（输入实为(4,2,8):(1,4,8)），本文按 op 语义修正为一致形式。模式选择语义的完整测试见 get.mlir 与 select.mlir，错误路径见 get_errors.mlir 与 select_errors.mlir。

叶子 / 标量提取

cute.get_leaves与cute.get_scalars是两个"展平"访问器：

• cute.get_leaves为输入层次结构的每个叶子返回一个 CuTe 类型的结果。当下游代码需要把叶子作为独立 SSA 值使用、而不想再从 CuTe 类型中恢复它们时，它非常有用；
• cute.get_scalars把整数叶子作为普通i32SSA 值返回。可选限定符<{only_dynamic}>把输出限制为动态叶子——当只需要运行时值（即带?的叶子）时非常方便。

%sh = cute.static : !cute.shape<"(4,(2,3))"> // 三个叶子：!cute.shape<"4">、!cute.shape<"2">、!cute.shape<"3"> %a, %b, %c = cute.get_leaves(%sh) : !cute.shape<"(4,(2,3))"> %l = cute.make_layout(%s, %d) : (!cute.shape<"(?,8)">, !cute.stride<"(1,?)">) -> !cute.layout<"(?,8):(1,?)"> // 所有标量叶子（含静态）作为 i32 %s0, %s1, %s2, %s3 = cute.get_scalars (%l) : !cute.layout<"(?,8):(1,?)"> // 仅动态叶子 %d0, %d1 = cute.get_scalars<{only_dynamic}> (%l) : !cute.layout<"(?,8):(1,?)">

从 op 定义看，get_scalars也适用于 swizzle（如cute.get_scalars(%sw) : !cute.swizzle<"S<3,5,4>">可取出num_bits、num_base、num_shift三个标量），见 CuteOps.td。测试见 get_leaves.mlir 与 get_scalars.mlir。

Tile 访问器

!cute.tile是一组布局与_通配符的序列。cute.get_layouts_from_tile把每个布局槽位作为独立 SSA 结果返回，跳过下划线：

%t = cute.static : !cute.tile<"[(4,8):(1,4);_;(2,3):(1,2)]"> %a, %b = cute.get_layouts_from_tile(%t) : !cute.tile<"[(4,8):(1,4);_;(2,3):(1,2)]">

两个结果分别是!cute.layout<"(4,8):(1,4)">和!cute.layout<"(2,3):(1,2)">。op 定义见 CuteOps.td：当输入本身是布局时返回其自身；当 tile 全部是下划线（如[_;_]）时结果为空。相关测试见 get_layouts_from_tile.mlir。

---

Size and Indexing（尺寸与索引）

这组运算计算关于 CuTe 值的标量量，或在坐标与线性索引之间转换：

Op	功能
cute.size	元组 / 布局类值的乘法尺寸（可选限制为部分 mode）。
cute.cosize	布局类值的余域尺寸（最大索引 + 1，可选限制为部分 mode）。
cute.tuple_product	把int_tuple或 shape 的所有叶子相乘。
cute.tuple_product_each	对int_tuple或 shape 按顶层 mode 分别求积。
cute.crd2idx	按 shape 把坐标转换为线性索引。
cute.idx2crd	按 shape 把线性索引转换为坐标。
cute.increment_coord	在其边界 shape 内把坐标递增一。
cute.append_to_rank	追加默认元素直到 rank 达到目标N。
cute.prepend_to_rank	前置默认元素直到 rank 达到目标N。

尺寸与乘积

cute.size与cute.cosize是布局的两个标量汇总运算。size(L)是布局走过的不同坐标数；cosize(L)是它产生的最大线性索引的后继（max-index + 1）。两者都接受可选的 mode 索引列表（<[...]>）把汇总限制到部分 mode：

%l = cute.static : !cute.layout<"(4,(16,32)):(1,(4,64))"> // 完整 size %sz = cute.size (%l) : (!cute.layout<"(4,(16,32)):(1,(4,64))">) -> !cute.int_tuple<"2048"> // 仅 mode 1 %sz1 = cute.size<[1]> (%l) : (!cute.layout<"(4,(16,32)):(1,(4,64))">) -> !cute.int_tuple<"512"> // Cosize（最大线性索引 + 1） %cs = cute.cosize (%l) : (!cute.layout<"(4,3,2):(1,4,12)">) -> !cute.int_tuple<"24">

结合 size.mlir 测试可确认cute.size的输入面非常宽：

• !cute.shape——所有 extent 的乘积（如(4,3)→12），带 mode 时取该子树（如(1,(2,4),9)的[1]→8）；
• !cute.int_tuple——如(4,(16,32))→2048；
• !cute.layout——取其定义域 shape的 size（size忽略 stride）；
• !cute.composed_layout——取外层布局定义域的 size（如S<3,4,6> o 0 o (3,4):(4,1)→12）；
• 含动态?的值——结果为动态（如(4,(16,32),(?,64))→?）；
• scaled-basis stride 布局——如(2,3):(1@0,1@1)→6，仍只取 extent 乘积。

cute.tuple_product把 int-tuple 或 shape 的所有叶子相乘归约为一个标量；cute.tuple_product_each按顶层 mode分别归约：

%t = cute.static : !cute.int_tuple<"(2,(3,4))"> %p = cute.tuple_product(%t) : (!cute.int_tuple<"(2,(3,4))">) -> !cute.int_tuple<"24"> %pe = cute.tuple_product_each(%t) : (!cute.int_tuple<"(2,(3,4))">) -> !cute.int_tuple<"(2,12)">

tuple_product也接受 shape、coord（结果!cute.int_tuple）等输入；tuple_product_each对扁平元组退化为恒等（如(4,3,2)→(4,3,2)），见 CuteOps.td。测试见 tuple_product.mlir 与 tuple_product_each.mlir。

坐标 ↔ 索引转换

坐标与列主序（column-major）线性索引之间的转换被暴露为独立运算：

• cute.crd2idx(coord, shape)使用从 shape 推导出的隐式列主序 stride打包坐标，产生线性索引；
• cute.idx2crd(idx, shape)是它的对偶——给定 shape 与线性索引，返回**字典序（colexicographic）**顺序下对应的坐标。

%shape0 = cute.static : !cute.shape<"(4,5)"> %c = cute.static : !cute.coord<"(2,1)"> %i = cute.crd2idx(%c, %shape0) : (!cute.coord<"(2,1)">, !cute.shape<"(4,5)">) -> !cute.int_tuple<"6"> %idx = cute.static : !cute.int_tuple<"9"> %shape = cute.static : !cute.shape<"(4,8)"> %crd = cute.idx2crd(%idx, %shape) : (!cute.int_tuple<"9">, !cute.shape<"(4,8)">) -> !cute.coord<"(1,2)">

crd2idx的精确公式在 op 定义（CuteOps.td）中给出：对 shape(s_0, ..., s_{n-1})与 coord(c_0, ..., c_{n-1})，结果为

c_0 + c_1*s_0 + c_2*s_0*s_1 + ... + c_{n-1}*s_0*...*s_{n-2}

嵌套的 coord/shape 逐 mode 递归。它的前置条件是$coord为!cute.coord且与$shapecongruent（层次结构一致）；结果是标量!cute.int_tuple，且当且仅当coord 的每个叶子与 shape 中每个贡献 extent 都是静态时结果才是静态的（任一?使结果变为?——因为动态 extent 会影响后续 mode 的列主序权重）。这是结构化打包：不接收 layout / stride / order 操作数；当 stride 来自具体布局时应改用cute.layout_eval。rank 不匹配（如 coord(1,2,3)对 shape(4,8)）或声明结果与推断值不一致（如声明!cute.int_tuple<"7">而推断为"9"）都会被拒绝，见 crd2idx_errors.mlir。

idx2crd的展开规则（CuteOps.td）：对一维 shapes，结果是idx mod s；对更高维 shape，索引按 shape 各 extent 的累乘逐位（digit-by-digit）拆解。$index可以是标量!cute.int_tuple（rank ≤ 1），也可以是 rank 与 shape 一致的多元素 int_tuple（后者相当于预先按 mode 分解了线性索引）。结果是 profile 与 shape 匹配的!cute.coord；仅当 index 静态且 shape 每个 extent 都静态时结果才完全静态。错误测试见 idx2crd_errors.mlir。

cute.increment_coord(coord, shape)在其边界 shape 内把坐标递增一（colexicographic 顺序），在 mode 边界处回绕：

%c = cute.static : !cute.coord<"(3,2)"> %s = cute.static : !cute.shape<"(4,8)"> %n = cute.increment_coord(%c, %s) : (!cute.coord<"(3,2)">, !cute.shape<"(4,8)">) -> !cute.coord<"(0,3)">

由于 shape 为(4,8)，(3,2)的下一个字典序坐标即(0,3)。测试见 increment_coord.mlir。

秩填充（Rank Padding）

cute.append_to_rank与cute.prepend_to_rank把输入值扩展直到其 rank 达到目标N（以<N>属性指定），用调用方提供的同类型默认元素填充。它们主要用于在组合布局前对齐各个值的 rank：

%in = cute.static : !cute.shape<"(4,8)"> %e = cute.static : !cute.shape<"1"> %r = cute.append_to_rank<4> (%in, %e) : !cute.shape<"(4,8)">, !cute.shape<"1"> // -> !cute.shape<"(4,8,1,1)"> %p = cute.prepend_to_rank<4> (%in, %e) : !cute.shape<"(4,8)">, !cute.shape<"1"> // -> !cute.shape<"(1,1,4,8)">

可接受任何 CuTe 元组类或布局类型（CuteOps.td 中注释明确"Any cute tuple-like or layout type accepted by append/prepend_to_rank"）。测试见 append_to_rank.mlir 与 prepend_to_rank.mlir。

---

测试与验证：访问器与尺寸/索引运算如何被验证

CuTe IR 方言对这两组运算提供了完整的正/反例测试覆盖，路径结构非常清晰：

• 方言级语法/类型检查：cutlass_compiler/cute_ir/test/Dialect/Cute/Accessors/下的get_shape.mlir、get_stride.mlir、composed_get_inner.mlir、composed_get_offset.mlir、composed_get_outer.mlir、get.mlir、select.mlir、get_leaves.mlir、get_scalars.mlir、get_layouts_from_tile.mlir，以及SizeIndex/下的size.mlir、cosize.mlir、tuple_product.mlir、tuple_product_each.mlir、crd2idx.mlir、idx2crd.mlir、increment_coord.mlir、append_to_rank.mlir、prepend_to_rank.mlir；
• 错误路径：每个 op 都有对应的*_errors.mlir（如get_shape_errors.mlir、crd2idx_errors.mlir、select_errors.mlir），覆盖类型不匹配、rank 不一致、tile 含下划线、声明结果与推断值冲突等失败场景；
• 展开（lowering）：cutlass_compiler/cute_ir/test/Conversion/CuteExpandOps/Accessors/与.../SizeIndex/验证访问器与尺寸运算如何展开为更基础的方言运算；cutlass_compiler/cute_ir/test/Conversion/CuteToBase/PostExpand/SizeIndex/进一步验证最终落到base方言的过程；
• 静态折叠：cutlass_compiler/cute_ir/test/Transforms/CuteFoldStatic/fold_accessors.mlir与fold_size_index.mlir验证静态输入（所有叶子为常量）下这些Pure运算被CuteFoldStatic变换折叠为常量。

测试的运行方式与其余方言测试一致（见 size.mlir 顶部的 RUN 行）：通过cute-opt工具配合-split-input-file与FileCheck断言，例如：

cute-opt %s -split-input-file | FileCheck %s

此外，类型层的单元测试（unittests 下的LayoutTypeTests.cpp、ShapeTypeTests.cpp、StrideTypeTests.cpp、IntTupleTypeTests.cpp、ComposedLayoutTypeTests.cpp、TileTypeTests.cpp等）从 C++ 侧保证了类型表示的稳定性，间接为访问器运算提供了类型级依据。

---

小结与速查

访问器（读侧，对应构造器的逆操作）

• get_shape/get_stride——取布局的 shape / stride；get_shape对复合布局返回外层 B 的 shape，对 tile 返回各槽位 shape 的拼接（tile 含_时拒绝）。
• composed_get_inner/composed_get_offset/composed_get_outer——拆解A ∘ offset ∘ B；当 A 为 swizzle 时composed_get_inner返回!cute.swizzle。
• get（<[...]>层次化索引）与select（<[...]>顶层模式子集，允许重排与重复）——按 mode 提取或重组。
• get_leaves（每叶子一个 CuTe 值结果）与get_scalars（i32标量，可用<{only_dynamic}>过滤）——展平提取。
• get_layouts_from_tile——取 tile 中每个布局槽位，跳过_。

尺寸与索引

• size= 定义域坐标数（extent 乘积，忽略 stride）；cosize= 最大索引 + 1；两者都支持<[mode]>限制。
• tuple_product（全叶子求积）与tuple_product_each（按顶层 mode 求积）。
• crd2idx（列主序隐式 stride 打包，公式c_0 + c_1*s_0 + ...）与idx2crd（按累乘逐位拆解）互为对偶；increment_coord在边界 shape 内按字典序递增并回绕。
• append_to_rank/prepend_to_rank（<N>目标 rank + 默认元素填充）用于组合布局前对齐 rank。

使用建议：这些运算全部为Pure且带类型推断，静态输入可被CuteFoldStatic完全折叠，因此适合在写布局代数变换时自由组合而不必担心副作用；遇到类型或结构错误时，优先对照各 op 的*_errors.mlir测试理解前置条件（congruent、静态性传播、tile 下划线等）。更完整的类型系统（Shape/Stride/Layout/Coord/Tile/ComposedLayout/Swizzle）见 01_cute_types.rst，下一章将介绍在提取出的结构上执行算术运算，见 05_arithmetic.rst。

【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlass