CuTe DSL 核心类型全解析从 Layout 代数到结构化内存的cutlass.cute编程指南【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlass导读本文围绕 cutlass.cute 及其核心类型文档 types.inc 展开系统讲解 CUTLASS Python DSL 中构成张量布局代数tensor layout algebra与结构化内存编程地基的核心类型。你将掌握IntValue、Ratio、ScaledBasis、Swizzle、Layout、ComposedLayout的语义与用法学会用cute.struct/cute.union装饰器在 Python 中精确描述 C 风格的共享内存与数据结构并理解这些类型如何经由 CuTe IRMLIR映射到 GPU 内核。文中所有 API 描述均可在 python/CuTeDSL/cutlass/cute/core.py 中找到对应实现适合希望深入 CUTLASS Python DSL 并手写高性能 GEMM/Kernel 的开发者。总览CuTe DSL 类型体系CuTe DSLcutlass.cute模块为开发者提供了一组「core types」它们共同构成了两层能力布局代数层IntValue、Ratio、ScaledBasis、Swizzle、Layout、ComposedLayout用于描述“逻辑坐标如何映射到内存线性索引”这一张量布局的核心问题结构化内存层struct/union装饰器及其辅助类型MemRange、Align用于在 Python DSL 中声明 C 风格的数据结构精确控制共享内存的布局、对齐与嵌套关系。这些类型大量复用标准 Python 协议__str__、__eq__、__getitem__、算术运算符等并作为 MLIR value caster 注册到底层编译器基础设施中从而实现“Python 书写 → CuTe IR 生成 → CUDA 内核编译”的完整链路。从源码结构看IntValue、Layout、ComposedLayout等内部类型都注册为 MLIR value caster见 python/CuTeDSL/cutlass/cute/core.py 与 python/CuTeDSL/cutlass/cute/typing.py保证与底层编译器基础设施的无缝集成。核心数值类型IntValue携带可除性约束的受限整数IntValue是 CuTe IR 中受限整数类型constrained integer type的内部代理表示其关键职责是自动追踪可除性信息divisibility。可除性约束对布局运算至关重要例如Layout的补complement、组合composition运算往往要求步长满足特定的整除关系IntValue让这些约束在类型层面显式可见。从 core.py 的实现看IntValue继承自cutlass_arith.ArithValue并在 IR 中自动发出cute.get_scalars操作。它代表的 IR 操作序列形如%0 ... : (...) - !cute.int_tuple? %1 cute.get_scalars(%0) : (!cute.int_tuple?) - i32即先将int_tuple?depth0, rank1通过cute.get_scalars展开为标量 i32再在其上进行算术运算。这样每次运算都能继续携带可除性信息。API 方法get_typed_value()— 将值还原为IntTupleType对应文档中的 IntTupleTypeget_divisibility()— 返回值的可除性约束divisibilityproperty— 直接读取可除性约束。get_typed_value()的实现在 core.py若当前值由cute.get_scalars定义且其操作数就是int_tuple则直接返回该操作数否则例如内核参数这类由ir.Block持有的标量通过_pack_int_tuple/MakeIntTupleOp重新包装为int_tuple。支持的运算全部在 core.py 中逐一实现均通过_binary_op装饰器自动传播可除性result int_val1 int_val2 # tuple_add result int_val1 - int_val2 # tuple_sub result int_val1 * int_val2 # tuple_mul result int_val1 // int_val2 # tuple_div result int_val1 % int_val2 # tuple_mod反向运算__radd__、__rsub__、__rmul__、__rfloordiv__、__rmod__同样被重载支持2 int_val这类写法运算会先把操作数包装为int_tuple再调用对应的 CuTe IR 元组算术操作。字符串表示str(int_val) # 可除性为 1 时返回 ? str(int_val) # 可除性为 4、32 位宽时返回 ?{div4}从实现core.py看可除性为 1 时输出裸?32 位宽时输出?{divN}其他位宽还会带上类型宽度如?{i16 div4}便于调试时区分符号与宽度。Ratio精确有理数Ratio用两个整数之比表示有理数用于布局运算中可除性条件不满足时产生的精确分数步长例如composition结果中出现的 1/2 步长。构造函数会对参数做类型检查ratio cute.Ratio(numerator, denominator)numerator/denominatorint类型非整数时抛出TypeError见 core.py。方法is_integral()— 分子能被分母整除时返回Truereduced()— 返回约分到最简形式的新Ratio实现中先调用父类约分逻辑再重建Ratio见 core.pyto(dtype)— 转换为Ratio原样返回、float分子/分母或int分子//分母即向下取整其他类型抛出TypeErrorcore.py。算术运算ratio1 cute.Ratio(1, 2) ratio2 cute.Ratio(3, 4) result ratio1 * ratio2 # 返回 Ratio(3, 8) ratio cute.Ratio(2, 3) result ratio * 5 # 返回 Ratio(10, 3) result 5 * ratio # 返回 Ratio(10, 3)__rmul__ 同样支持__mul__/__rmul__的实现core.py支持与Ratio或int相乘其他类型直接抛TypeError。类型转换ratio cute.Ratio(3, 2) float_val ratio.to(float) # 返回 1.5 int_val ratio.to(int) # 返回 1向下取整布局代数类型ScaledBasis带缩放因子的基元素ScaledBasis表示布局代数中的一个“缩放基元素”由一个缩放值scale和一个模式mode构成用于标识引用布局代数中的哪个基。它是 CuTe 坐标系表示的基础——布局步长stride本质上就是一组ScaledBasis的线性组合。sb cute.ScaledBasis(value, mode)valueUnion[int, Integer, Ratio, ir.Value]缩放值modeUnion[int, List[int]]标识基元素的模式既可以是单个整数也可以是整数列表对应多模式基非整数列表会抛TypeError。示例sb1 cute.ScaledBasis(2, 0) # 2 * E(0) sb2 cute.ScaledBasis(cute.Ratio(1, 2), 1) # (1/2) * E(1) sb3 cute.ScaledBasis(4, [0, 1]) # 4 * E([0, 1]) # ScaledBasis 常被用作布局步长 layout cute.make_layout((4, 8), stride(cute.ScaledBasis(2, 0), cute.ScaledBasis(1, 1))) # 该布局步长为 (20, 11)每个维度拥有各自的基 coord (2, 3) idx cute.crd2idx(coord, layout) # 将 (2, 3) 映射为 (4, 3)属性与运算value— 缩放值mode— 以整数列表形式返回模式构造时单整数会被包装成单元素列表is_static()— 缩放值是否静态已知to(dtype)— 转换为ScaledBasis或内部_ScaledBasis右乘缩放因子3 * cute.ScaledBasis(2, 0)得到ScaledBasis(6, 0)。便捷函数cute.E(mode)用于创建单位缩放基等价于ScaledBasis(1, mode)在书写布局代数表达式时更简洁。Swizzle位级重排变换Swizzle是一种对布局元素进行位级置换的变换典型用途是重排共享内存中的数据访问模式规避 bank conflict从而提升内存吞吐。它不是一个任意置换而是受控的 XOR 位操作由三个参数精确定义参数含义MBase保持不变的 least-significant 位数BBits掩码的位数参与 XOR 的位宽SShift掩码移动的距离正数右移、负数左移位模式示意来自 core.py 的文档注释0bxxxxxxxxxxxxxxxYYYxxxxxxxZZZxxxx ^--^ MBase低 bit 保持恒定 ^-^ ^-^ BBits掩码位数 ^---------^ SShiftYYY 的位移距离正右负左 给定: 0bxxxxxxxxxxxxxxxxYYxxxxxxxxxZZxxx 结果: 0bxxxxxxxxxxxxxxxxYYxxxxxxxxxAAxxx 其中 AA ZZ xor YY实践中 Swizzle 通常由cute.make_swizzle(...)工厂函数创建再通过composition与Layout组合成优化的内存访问模式可参考 include/cute/atom 与 python/CuTeDSL/cutlass/cute/core.py 中make_swizzle的实现思路。Layout布局代数的核心抽象Layout是 CuTe 表示张量布局的核心抽象它将逻辑坐标空间映射到索引空间由一对(Shape, Stride)定义。Shape描述布局的抽象维度Stride描述坐标如何映射到线性索引。CuTe 布局天然是层级化的——大布局由更小的嵌套布局拼接而成如Shape/Stride本身可以是嵌套的IntTuple。属性shape— 表示维度尺寸的IntTuplestride— 表示步长的IntTuplemax_alignment— 布局的最大对齐字节。基本用法# shape (4,8)默认列主序步长 (1,4) layout cute.make_layout((4, 8)) # 显式指定步长得到行主序 (8,1) layout cute.make_layout((4, 8), stride(8, 1)) shape layout.shape # (4, 8) stride layout.stride # (8, 1) # 坐标映射为索引: (2, 3) - 2*8 3*1 19 idx cute.crd2idx((2, 3), layout)crd2idx坐标→索引与idx2crd索引→坐标构成了布局的往返映射make_layout是创建布局的推荐入口见 core.py 附近实现。Layout 运算代数对应 CuTe C 头文件如 include/cute/layout.hpp 中的同名原语Concatenation拼接沿维度拼接多个布局Coalescence合并合并相邻模式压缩布局层级Composition组合以函数或其他布局组合布局即(f ∘ g)(x) f(g(x))Complement补计算给定子布局在父空间中的补空间Inversion求逆反转布局映射把索引空间映射回坐标空间。字符串表示以shape:stride形式打印例如layout cute.make_layout((4, 8), stride(1, 4)) print(layout) # 输出 (4,8):(1,4)ComposedLayout布局与变换的组合ComposedLayout是普通Layout的推广其“内层”可以是不再局限于线性仿射映射的任意函数映射典型如Swizzle从而表达更一般的坐标→坐标变换。它由三个组件构成inner— 内层变换Swizzle或Layoutoffset— 应用于坐标的偏移IntTupleouter— 外层布局。属性inner/offset/outer— 分别访问上述三个组件shape— 组合布局的 shapemax_alignment— 最大对齐is_normal— 是否为普通布局非一般组合。典型创建方式与访问# ComposedLayout 通常通过组合运算创建如布局与 swizzle 组合 layout cute.make_layout((8, 8)) swizzle cute.make_swizzle(...) composed cute.composition(swizzle, layout) inner composed.inner # swizzle outer composed.outer # layout offset composed.offset # 偏移 print(composed) # 以 inner o offset o outer 形式打印在 C CuTe 中ComposedLayout对应ComposedLayoutLayout, Swizzle, Offset模板Python DSL 侧则通过composition函数core.py 附近有多个重载与_ComposedLayout/_ComposedLayoutWithInnerFunc内部类实现。结构化数据类型struct 与 union在编写真实 GPU 内核尤其是持久化内核与流水线时往往需要在共享内存中精确排布多块缓冲区、屏障、计数器等数据结构。Python DSL 用cute.struct与cute.union装饰器把 C 语言的struct/union语义搬到 Python 中并允许在编译期查询大小与对齐。structC 结构体的 Python DSL 抽象cute.struct支持以下元素基础 DSL 标量类型int/float数组MemRange嵌套结构体对齐修饰Align。基本用法# 简单结构体 cute.struct class complex: real : cutlass.Float32 imag : cutlass.Float32 # 含数组与嵌套结构 cute.struct class StorageA: mbarA : cute.struct.MemRange[cutlass.Int64, stage] compA : complex intA : cutlass.Int16显式对齐控制cute.struct class StorageB: a: cute.struct.Align[ cute.struct.MemRange[cutlass.Float32, size_a], 1024 ] b: cute.struct.Align[ cute.struct.MemRange[cutlass.Float32, size_b], 1024 ] x: cute.struct.Align[cutlass.Int32, 16] compA: cute.struct.Align[complex, 16]编译期静态查询与分配访问size StorageB.__sizeof__() # 结构体字节大小 align StorageB.__alignof__() # 结构体字节对齐 storage allocator.allocate(StorageB) # 分配例如 smem 分配器 storage.a[0] ... storage.x ... ... storage.compA.real.ptr ... storage.x.ptr.load()方法__sizeof__()— 返回结构体字节大小__alignof__()— 返回结构体字节对齐size_in_bytes()— 返回结构体字节大小。struct.MemRange连续内存区间MemRange[dtype, size]描述一段具有特定元素类型与元素数量的连续内存cute.struct.MemRange[dtype, size]dtypeDSL 标量类型Type[Numeric]size元素个数int。属性size— 元素个数elem_width— 每个元素的位宽size_in_bytes— 总字节数。方法data_ptr()— 返回指向内存区间起点的指针get_tensor(layout, swizzleNone, dtypeNone)— 基于该内存区间与给定Layout可选Swizzle、dtype创建张量视图__getitem__(index)— 按索引访问元素。示例cute.struct class Buffer: data : cute.struct.MemRange[cutlass.Float32, 128] buf allocator.allocate(Buffer) ptr buf.data.data_ptr() # 数据指针 element buf.data[5] # 元素访问 layout cute.make_layout((8, 16)) tensor buf.data.get_tensor(layout) # 从内存区间构造张量MemRange与get_tensor的组合是把“原始共享内存”包装成“带布局的张量视图”的标准路径与 C CuTe 中make_tensor(SmemAtom, layout)的思路一致。struct.Align显式对齐修饰Align[dtype, alignment]为结构体成员指定显式对齐cute.struct.Align[dtype, alignment]dtype被对齐的类型标量、MemRange或结构体alignment对齐字节数必须 0。属性dtype被对齐的数据类型、align对齐值。示例cute.struct class AlignedStorage: counter: cute.struct.Align[cutlass.Int32, 16] # 标量 16 字节对齐 buffer: cute.struct.Align[ cute.struct.MemRange[cutlass.Float32, 256], 1024 # 数组 1024 字节对齐 ]unionC 联合体的 Python DSL 抽象cute.union与struct语义类似但所有成员从偏移 0 开始大小取所有成员的最大值对齐取所有成员的最大对齐所有对象起始偏移为 0对齐为所有对象对齐的最大值大小为所有对象大小的最大值。用法cute.union class value_union: as_int : cutlass.Int32 as_float : cutlass.Float32 val allocator.allocate(value_union) val.as_int 42 float_val val.as_float.ptr.load() # 将同一段位模式解释为 float方法与struct相同——__sizeof__()返回联合体字节大小__alignof__()返回字节对齐。类型层级与协议支持多数 CuTe 类型实现了标准 Python 协议以便与既有生态集成__str__()— 调试用字符串表示如Layout的(4,8):(1,4)、IntValue的?{div4}__eq__()/__ne__()— 相等性比较__getitem__()— 索引操作__add__()/__sub__()/__mul__()/__floordiv__()/__mod__()— 算术运算。MLIR 集成方面IntValue、Layout、ComposedLayout等内部类型均注册为 MLIR value caster使得这些 Python 对象可以直接作为 CuTe IR 操作的操作数参与编译期求值这是“Python 书写、MLIR 编译、CUDA 运行”流水线的关键一环。最佳实践内存对齐共享内存结构体务必显式指定对齐需求以规避 bank conflict使用struct.Align强制对齐约束通过max_alignment属性校验布局与结构化存储的对齐是否满足硬件要求。布局运算优先使用内建布局运算make_layout、composition、coalesce、complement、inverse等而非手工构造多模式布局中用ScaledBasis显式控制步长模式保证可读性与可除性信息正确传播涉及 swizzle 等复杂变换时用ComposedLayout表达保持inner o offset o outer的清晰结构。相关文档导航以下文档与本篇同属 CuTe DSL 文档体系可按需深入CuTe DSL 装饰器与调用约定入门数值、指针、向量与数组 API 参考静态与动态值的控制流静态与动态布局实战JIT 与内核参数的类型注解深度学习框架集成CuTe DSL 程序调试技巧CuTe DSL 运行时 API 与 CuTe DSL 数学运算 API【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlass创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CuTe DSL 核心类型全解析：从 Layout 代数到结构化内存的 `cutlass.cute` 编程指南
CuTe DSL 核心类型全解析从 Layout 代数到结构化内存的cutlass.cute编程指南【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlass导读本文围绕 cutlass.cute 及其核心类型文档 types.inc 展开系统讲解 CUTLASS Python DSL 中构成张量布局代数tensor layout algebra与结构化内存编程地基的核心类型。你将掌握IntValue、Ratio、ScaledBasis、Swizzle、Layout、ComposedLayout的语义与用法学会用cute.struct/cute.union装饰器在 Python 中精确描述 C 风格的共享内存与数据结构并理解这些类型如何经由 CuTe IRMLIR映射到 GPU 内核。文中所有 API 描述均可在 python/CuTeDSL/cutlass/cute/core.py 中找到对应实现适合希望深入 CUTLASS Python DSL 并手写高性能 GEMM/Kernel 的开发者。总览CuTe DSL 类型体系CuTe DSLcutlass.cute模块为开发者提供了一组「core types」它们共同构成了两层能力布局代数层IntValue、Ratio、ScaledBasis、Swizzle、Layout、ComposedLayout用于描述“逻辑坐标如何映射到内存线性索引”这一张量布局的核心问题结构化内存层struct/union装饰器及其辅助类型MemRange、Align用于在 Python DSL 中声明 C 风格的数据结构精确控制共享内存的布局、对齐与嵌套关系。这些类型大量复用标准 Python 协议__str__、__eq__、__getitem__、算术运算符等并作为 MLIR value caster 注册到底层编译器基础设施中从而实现“Python 书写 → CuTe IR 生成 → CUDA 内核编译”的完整链路。从源码结构看IntValue、Layout、ComposedLayout等内部类型都注册为 MLIR value caster见 python/CuTeDSL/cutlass/cute/core.py 与 python/CuTeDSL/cutlass/cute/typing.py保证与底层编译器基础设施的无缝集成。核心数值类型IntValue携带可除性约束的受限整数IntValue是 CuTe IR 中受限整数类型constrained integer type的内部代理表示其关键职责是自动追踪可除性信息divisibility。可除性约束对布局运算至关重要例如Layout的补complement、组合composition运算往往要求步长满足特定的整除关系IntValue让这些约束在类型层面显式可见。从 core.py 的实现看IntValue继承自cutlass_arith.ArithValue并在 IR 中自动发出cute.get_scalars操作。它代表的 IR 操作序列形如%0 ... : (...) - !cute.int_tuple? %1 cute.get_scalars(%0) : (!cute.int_tuple?) - i32即先将int_tuple?depth0, rank1通过cute.get_scalars展开为标量 i32再在其上进行算术运算。这样每次运算都能继续携带可除性信息。API 方法get_typed_value()— 将值还原为IntTupleType对应文档中的 IntTupleTypeget_divisibility()— 返回值的可除性约束divisibilityproperty— 直接读取可除性约束。get_typed_value()的实现在 core.py若当前值由cute.get_scalars定义且其操作数就是int_tuple则直接返回该操作数否则例如内核参数这类由ir.Block持有的标量通过_pack_int_tuple/MakeIntTupleOp重新包装为int_tuple。支持的运算全部在 core.py 中逐一实现均通过_binary_op装饰器自动传播可除性result int_val1 int_val2 # tuple_add result int_val1 - int_val2 # tuple_sub result int_val1 * int_val2 # tuple_mul result int_val1 // int_val2 # tuple_div result int_val1 % int_val2 # tuple_mod反向运算__radd__、__rsub__、__rmul__、__rfloordiv__、__rmod__同样被重载支持2 int_val这类写法运算会先把操作数包装为int_tuple再调用对应的 CuTe IR 元组算术操作。字符串表示str(int_val) # 可除性为 1 时返回 ? str(int_val) # 可除性为 4、32 位宽时返回 ?{div4}从实现core.py看可除性为 1 时输出裸?32 位宽时输出?{divN}其他位宽还会带上类型宽度如?{i16 div4}便于调试时区分符号与宽度。Ratio精确有理数Ratio用两个整数之比表示有理数用于布局运算中可除性条件不满足时产生的精确分数步长例如composition结果中出现的 1/2 步长。构造函数会对参数做类型检查ratio cute.Ratio(numerator, denominator)numerator/denominatorint类型非整数时抛出TypeError见 core.py。方法is_integral()— 分子能被分母整除时返回Truereduced()— 返回约分到最简形式的新Ratio实现中先调用父类约分逻辑再重建Ratio见 core.pyto(dtype)— 转换为Ratio原样返回、float分子/分母或int分子//分母即向下取整其他类型抛出TypeErrorcore.py。算术运算ratio1 cute.Ratio(1, 2) ratio2 cute.Ratio(3, 4) result ratio1 * ratio2 # 返回 Ratio(3, 8) ratio cute.Ratio(2, 3) result ratio * 5 # 返回 Ratio(10, 3) result 5 * ratio # 返回 Ratio(10, 3)__rmul__ 同样支持__mul__/__rmul__的实现core.py支持与Ratio或int相乘其他类型直接抛TypeError。类型转换ratio cute.Ratio(3, 2) float_val ratio.to(float) # 返回 1.5 int_val ratio.to(int) # 返回 1向下取整布局代数类型ScaledBasis带缩放因子的基元素ScaledBasis表示布局代数中的一个“缩放基元素”由一个缩放值scale和一个模式mode构成用于标识引用布局代数中的哪个基。它是 CuTe 坐标系表示的基础——布局步长stride本质上就是一组ScaledBasis的线性组合。sb cute.ScaledBasis(value, mode)valueUnion[int, Integer, Ratio, ir.Value]缩放值modeUnion[int, List[int]]标识基元素的模式既可以是单个整数也可以是整数列表对应多模式基非整数列表会抛TypeError。示例sb1 cute.ScaledBasis(2, 0) # 2 * E(0) sb2 cute.ScaledBasis(cute.Ratio(1, 2), 1) # (1/2) * E(1) sb3 cute.ScaledBasis(4, [0, 1]) # 4 * E([0, 1]) # ScaledBasis 常被用作布局步长 layout cute.make_layout((4, 8), stride(cute.ScaledBasis(2, 0), cute.ScaledBasis(1, 1))) # 该布局步长为 (20, 11)每个维度拥有各自的基 coord (2, 3) idx cute.crd2idx(coord, layout) # 将 (2, 3) 映射为 (4, 3)属性与运算value— 缩放值mode— 以整数列表形式返回模式构造时单整数会被包装成单元素列表is_static()— 缩放值是否静态已知to(dtype)— 转换为ScaledBasis或内部_ScaledBasis右乘缩放因子3 * cute.ScaledBasis(2, 0)得到ScaledBasis(6, 0)。便捷函数cute.E(mode)用于创建单位缩放基等价于ScaledBasis(1, mode)在书写布局代数表达式时更简洁。Swizzle位级重排变换Swizzle是一种对布局元素进行位级置换的变换典型用途是重排共享内存中的数据访问模式规避 bank conflict从而提升内存吞吐。它不是一个任意置换而是受控的 XOR 位操作由三个参数精确定义参数含义MBase保持不变的 least-significant 位数BBits掩码的位数参与 XOR 的位宽SShift掩码移动的距离正数右移、负数左移位模式示意来自 core.py 的文档注释0bxxxxxxxxxxxxxxxYYYxxxxxxxZZZxxxx ^--^ MBase低 bit 保持恒定 ^-^ ^-^ BBits掩码位数 ^---------^ SShiftYYY 的位移距离正右负左 给定: 0bxxxxxxxxxxxxxxxxYYxxxxxxxxxZZxxx 结果: 0bxxxxxxxxxxxxxxxxYYxxxxxxxxxAAxxx 其中 AA ZZ xor YY实践中 Swizzle 通常由cute.make_swizzle(...)工厂函数创建再通过composition与Layout组合成优化的内存访问模式可参考 include/cute/atom 与 python/CuTeDSL/cutlass/cute/core.py 中make_swizzle的实现思路。Layout布局代数的核心抽象Layout是 CuTe 表示张量布局的核心抽象它将逻辑坐标空间映射到索引空间由一对(Shape, Stride)定义。Shape描述布局的抽象维度Stride描述坐标如何映射到线性索引。CuTe 布局天然是层级化的——大布局由更小的嵌套布局拼接而成如Shape/Stride本身可以是嵌套的IntTuple。属性shape— 表示维度尺寸的IntTuplestride— 表示步长的IntTuplemax_alignment— 布局的最大对齐字节。基本用法# shape (4,8)默认列主序步长 (1,4) layout cute.make_layout((4, 8)) # 显式指定步长得到行主序 (8,1) layout cute.make_layout((4, 8), stride(8, 1)) shape layout.shape # (4, 8) stride layout.stride # (8, 1) # 坐标映射为索引: (2, 3) - 2*8 3*1 19 idx cute.crd2idx((2, 3), layout)crd2idx坐标→索引与idx2crd索引→坐标构成了布局的往返映射make_layout是创建布局的推荐入口见 core.py 附近实现。Layout 运算代数对应 CuTe C 头文件如 include/cute/layout.hpp 中的同名原语Concatenation拼接沿维度拼接多个布局Coalescence合并合并相邻模式压缩布局层级Composition组合以函数或其他布局组合布局即(f ∘ g)(x) f(g(x))Complement补计算给定子布局在父空间中的补空间Inversion求逆反转布局映射把索引空间映射回坐标空间。字符串表示以shape:stride形式打印例如layout cute.make_layout((4, 8), stride(1, 4)) print(layout) # 输出 (4,8):(1,4)ComposedLayout布局与变换的组合ComposedLayout是普通Layout的推广其“内层”可以是不再局限于线性仿射映射的任意函数映射典型如Swizzle从而表达更一般的坐标→坐标变换。它由三个组件构成inner— 内层变换Swizzle或Layoutoffset— 应用于坐标的偏移IntTupleouter— 外层布局。属性inner/offset/outer— 分别访问上述三个组件shape— 组合布局的 shapemax_alignment— 最大对齐is_normal— 是否为普通布局非一般组合。典型创建方式与访问# ComposedLayout 通常通过组合运算创建如布局与 swizzle 组合 layout cute.make_layout((8, 8)) swizzle cute.make_swizzle(...) composed cute.composition(swizzle, layout) inner composed.inner # swizzle outer composed.outer # layout offset composed.offset # 偏移 print(composed) # 以 inner o offset o outer 形式打印在 C CuTe 中ComposedLayout对应ComposedLayoutLayout, Swizzle, Offset模板Python DSL 侧则通过composition函数core.py 附近有多个重载与_ComposedLayout/_ComposedLayoutWithInnerFunc内部类实现。结构化数据类型struct 与 union在编写真实 GPU 内核尤其是持久化内核与流水线时往往需要在共享内存中精确排布多块缓冲区、屏障、计数器等数据结构。Python DSL 用cute.struct与cute.union装饰器把 C 语言的struct/union语义搬到 Python 中并允许在编译期查询大小与对齐。structC 结构体的 Python DSL 抽象cute.struct支持以下元素基础 DSL 标量类型int/float数组MemRange嵌套结构体对齐修饰Align。基本用法# 简单结构体 cute.struct class complex: real : cutlass.Float32 imag : cutlass.Float32 # 含数组与嵌套结构 cute.struct class StorageA: mbarA : cute.struct.MemRange[cutlass.Int64, stage] compA : complex intA : cutlass.Int16显式对齐控制cute.struct class StorageB: a: cute.struct.Align[ cute.struct.MemRange[cutlass.Float32, size_a], 1024 ] b: cute.struct.Align[ cute.struct.MemRange[cutlass.Float32, size_b], 1024 ] x: cute.struct.Align[cutlass.Int32, 16] compA: cute.struct.Align[complex, 16]编译期静态查询与分配访问size StorageB.__sizeof__() # 结构体字节大小 align StorageB.__alignof__() # 结构体字节对齐 storage allocator.allocate(StorageB) # 分配例如 smem 分配器 storage.a[0] ... storage.x ... ... storage.compA.real.ptr ... storage.x.ptr.load()方法__sizeof__()— 返回结构体字节大小__alignof__()— 返回结构体字节对齐size_in_bytes()— 返回结构体字节大小。struct.MemRange连续内存区间MemRange[dtype, size]描述一段具有特定元素类型与元素数量的连续内存cute.struct.MemRange[dtype, size]dtypeDSL 标量类型Type[Numeric]size元素个数int。属性size— 元素个数elem_width— 每个元素的位宽size_in_bytes— 总字节数。方法data_ptr()— 返回指向内存区间起点的指针get_tensor(layout, swizzleNone, dtypeNone)— 基于该内存区间与给定Layout可选Swizzle、dtype创建张量视图__getitem__(index)— 按索引访问元素。示例cute.struct class Buffer: data : cute.struct.MemRange[cutlass.Float32, 128] buf allocator.allocate(Buffer) ptr buf.data.data_ptr() # 数据指针 element buf.data[5] # 元素访问 layout cute.make_layout((8, 16)) tensor buf.data.get_tensor(layout) # 从内存区间构造张量MemRange与get_tensor的组合是把“原始共享内存”包装成“带布局的张量视图”的标准路径与 C CuTe 中make_tensor(SmemAtom, layout)的思路一致。struct.Align显式对齐修饰Align[dtype, alignment]为结构体成员指定显式对齐cute.struct.Align[dtype, alignment]dtype被对齐的类型标量、MemRange或结构体alignment对齐字节数必须 0。属性dtype被对齐的数据类型、align对齐值。示例cute.struct class AlignedStorage: counter: cute.struct.Align[cutlass.Int32, 16] # 标量 16 字节对齐 buffer: cute.struct.Align[ cute.struct.MemRange[cutlass.Float32, 256], 1024 # 数组 1024 字节对齐 ]unionC 联合体的 Python DSL 抽象cute.union与struct语义类似但所有成员从偏移 0 开始大小取所有成员的最大值对齐取所有成员的最大对齐所有对象起始偏移为 0对齐为所有对象对齐的最大值大小为所有对象大小的最大值。用法cute.union class value_union: as_int : cutlass.Int32 as_float : cutlass.Float32 val allocator.allocate(value_union) val.as_int 42 float_val val.as_float.ptr.load() # 将同一段位模式解释为 float方法与struct相同——__sizeof__()返回联合体字节大小__alignof__()返回字节对齐。类型层级与协议支持多数 CuTe 类型实现了标准 Python 协议以便与既有生态集成__str__()— 调试用字符串表示如Layout的(4,8):(1,4)、IntValue的?{div4}__eq__()/__ne__()— 相等性比较__getitem__()— 索引操作__add__()/__sub__()/__mul__()/__floordiv__()/__mod__()— 算术运算。MLIR 集成方面IntValue、Layout、ComposedLayout等内部类型均注册为 MLIR value caster使得这些 Python 对象可以直接作为 CuTe IR 操作的操作数参与编译期求值这是“Python 书写、MLIR 编译、CUDA 运行”流水线的关键一环。最佳实践内存对齐共享内存结构体务必显式指定对齐需求以规避 bank conflict使用struct.Align强制对齐约束通过max_alignment属性校验布局与结构化存储的对齐是否满足硬件要求。布局运算优先使用内建布局运算make_layout、composition、coalesce、complement、inverse等而非手工构造多模式布局中用ScaledBasis显式控制步长模式保证可读性与可除性信息正确传播涉及 swizzle 等复杂变换时用ComposedLayout表达保持inner o offset o outer的清晰结构。相关文档导航以下文档与本篇同属 CuTe DSL 文档体系可按需深入CuTe DSL 装饰器与调用约定入门数值、指针、向量与数组 API 参考静态与动态值的控制流静态与动态布局实战JIT 与内核参数的类型注解深度学习框架集成CuTe DSL 程序调试技巧CuTe DSL 运行时 API 与 CuTe DSL 数学运算 API【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlass创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考