着色语言

着色语言用豫言的标准语法写 GPU 着色程序：着色源经编译器的语法分析、名称解析与类型检查（配合着色内建声明）之后译成 WGSL（网页图算着色语言），也可译成 MSL（金属着色语言，见着九）或 PTX（统算，见着十）。WebGPU 创建着色模块时再检查一遍 WGSL（地址空间、一致性分析等）。

各条以“着”编号。翻译器在 库/译场/分职/着色翻译，着色内建声明在包「点染」（库/点染/），命令行工具是包「着色器翻译」（应用/百工坊/着色器翻译，产物 着色器），测试在包「点染试」（库/点染/测试）。

着一：文件后缀

着二：模块查找与构建集成

着色文件（前两种后缀）里的 寻、观、寻观 只查 名。着色器。豫 与 名。着色器。接口。豫。都找不到时报“包内模块不存在：……（着色文件只能导入着色源「名。着色器。豫」与着色内建声明「名。着色器。接口。豫」）”，所以着色源导入标准库、操作系统接口或普通模块都报错。着色内建声明需要的「整数」「小数」「爻」「有」「元类型」直接取编译器的内建类型名 「《《内建类型：…》》」。

高级文件依次查 名。豫、名。接口。豫、名。应用接口。豫、名。着色器。豫（底层文件只查底层模块，见文件后缀缀二）。找到着色源时，得到构建时由它生成的模块（虚拟路径 名。着色生成。豫），该模块导出：

• 「着色源」：WGSL 文字；
• 「金属源」：MSL 文字（见着九）；MSL 翻译失败时是一行说明原因的注释，WGSL 照常可用；
• 「并线汇编源」：PTX 文字（见着十）；PTX 翻译失败时同样是一行说明原因的注释；
• 每个入口项一个同名字符串常量，值是入口函数的名字（WGSL 与 MSL 相同）；另有“入口项名并线汇编名”常量，值是 PTX 里改写成 ASCII 的入口名。

例如 寻「应用」之「加倍」之书。 之后，「加倍」之「着色源」 是 WGSL 文字，「加倍」之「加倍入口」 是 『加倍』，「加倍」之「加倍入口并线汇编名」 是 『yy_e5_8a_a0_e5_80_8d』。高级代码导入着色内建声明时报错。

生成在编译器里完成，应用产物只含生成的字符串：

1. 虚拟模块的依赖分析返回空列表（生成的模块没有导入），只备一份空的文件依赖缓存。
2. 词法阶段调用翻译器。翻译器在内存里分析着色源及其导入的着色模块（依赖先分析），不写阶段缓存，所以并行编译的多个工作进程不会争写同一个缓存文件。生成的豫言源码与已有的词法缓存相同时不改写，后续阶段与导入它的文件照常命中缓存；着色源或着色内建改动后，生成的源码随之改变。
3. 阶段环境读取虚拟模块的源码（如语言服务）时，从它的词法缓存取；缓存有效性以词法缓存的修改时间为源码时间。
4. 虚拟模块按其着色源判断所属的包。

着三：着色内建

库/点染/点染本。着色器。接口。豫 声明着色源可用的名字，包「点染」没有依赖。着色源写 寻观「点染」之「点染本」之书。，着色源所在的包须声明依赖「点染」。

标量与定长长度：标量的宽度与有无符号由类型区分，宽度写成整数索引（见依值类型）。「定长小数」「定长整数」「定长无符号」各由一个整数宽度构造出标量类型：「浮三二」即「定长小数」于「三二」（f32），「浮十六」即「定长小数」于「一六」（f16），「整三二」即「定长整数」于「三二」（i32），「无三二」即「定长无符号」于「三二」（u32）；WGSL 没有的组合（如「定长整数」于「一六」）翻译时报错。这些是不同的类型，混用宽度或有无符号时类型检查报错。布尔用「爻」（bool）。

字面量：整数、小数字面量的类型是「整数」「小数」，经构造器「新浮三二」「新浮十六」「新整三二」「新无三二」取得标量类型：「新无三二」于「二」译成 2u，「新浮三二」于「0.5」译成 0.5f，「新整三二」于「三」译成 3i，「新浮十六」于「1.0」译成 f16(1.0)；实参不是字面量时译成类型转换，如 u32(x)。不经构造器的字面量（如「读纹素」的层级）译成 WGSL 抽象字面量。

向量与矩阵：向量是不命名元组的别名，如「二维浮」（vec2<f32>）、「三维无」（vec3<u32>）、「四维整」（vec4<i32>）、「二维半」（vec2<f16>）、「三维爻」（vec3<bool>）；向量 中「零」 是 .x。矩阵「二阶浮」「三阶浮」「四阶浮」是列向量组成的元组，矩阵 中「一」 是第二列。

内建输入输出类型：类型检查上是向量或标量的别名；写在入口参数、入口返回值或入口所用结构的字段上时，翻译器按类型名加 @builtin。

「片元颜色」即「四维浮」，按位置编号输出。

类型构造：「数组」于 T 是运行期数组 array<T>，「定长数组」于 T 于 N 是 array<T, N>（N 写成整数字面量，如「定长数组」于「浮三二」于「八一九二」），「原子」于 T 是 atomic<T>，「纹理二维」于 T 是 texture_2d<T>，「采样器」是 sampler，「子组矩阵片甲」「子组矩阵片乙」「子组矩阵片累」是子组矩阵片（见着十二，WGSL 没有），「张量片累」是张量片（见着十三，只有金属）。资源「存储缓冲」「均匀缓冲」「工作组变量」只用于顶层资源声明；存储缓冲的访问方式是「只读」或「读写」。

资源绑定：「绑定项」即「组类」【「组」乃「整数」也，「号」乃「整数」也，「访问」乃「访问方式」也】；「资源位」即「组类」【「组」乃「整数」也，「号」乃「整数」也】。「绑定」由绑定项得存储缓冲，「均匀绑定」「纹理绑定」「采样器绑定」由资源位得均匀缓冲、二维纹理、采样器，「工作组共享」于「元」得工作组变量。组号与绑定号须是整数字面量。

入口：「计算入口」「顶点入口」「片元入口」是入口项的类型；「计算选项」即「组类」【「工作组大小」乃「整数」也】。「计算入口项」乃凡「甲」而化（化「甲」而「有」）又化「计算选项」为「计算入口」者也。，顶点、片元入口项取一个函数，输入输出类型随所指函数而定。

函数：多态函数用 凡 声明。下标（读数组、写数组、读共享数组、写共享数组、取元素与原子操作的序号）是「无三二」，字面下标写成「新无三二」于「零」（译成 0u）。

着四：顶层定义的分类

翻译器读语法分析所得的定义列（名称未解析，保留源码里的名字与类型别名），名字查找依次看局部名字、本着色程序的顶层名字、着色内建。着色源的顶层定义按以下规则翻译：

1. 即 定义带字段名的组类：译成 struct。结构用作入口的参数或返回类型时，内建输入输出类型的字段加 @builtin(…)，其余字段按出现次序编 @location(0)、@location(1)……；整数类型的字段用作顶点输出或片元输入时再加 @interpolate(flat)。其余 即 定义是类型别名，只在翻译时展开。
2. 乃 签名的类型是资源（存储缓冲、均匀缓冲、工作组变量、二维纹理、采样器）：定义须是对「绑定」等的调用，译成 @group(g) @binding(b) var<storage, read_write> 名: T;（只读为 var<storage, read>，均匀缓冲为 var<uniform>，纹理与采样器为 var），工作组变量译成 var<workgroup> 名: T;。
3. 乃 签名的类型是计算、顶点或片元入口：定义须是对相应入口项的调用，所指函数加 @compute @workgroup_size(n)、@vertex 或 @fragment。入口函数的参数是内建输入输出类型时加 @builtin，是结构时不加，其余按次序编 @location（片元的整数输入加 @interpolate(flat)）；计算入口的参数只能是内建输入或由它们组成的结构。顶点、片元入口的返回类型是内建类型时写 -> @builtin(…) T，是结构时写 -> S，其余写 -> @location(0) T。
4. 乃 签名的类型是函数（化…而…）：译成 fn，参数与返回类型取自签名；类型为「有」的参数省去，返回「有」的函数没有返回类型。
5. 其余带签名的定义与不带签名的定义译成 const（值须是常量表达式）。
6. 立、裸表达式、多态函数（凡）与调用自身的顶层函数报错。

输出依源码次序，被导入的着色源在前；用到 f16 时开头加 enable f16;。

着五：函数体的翻译

函数体的值是返回值。

名字原样写成 WGSL 标识符（WGSL 允许中文标识符）；遇到 WGSL 关键字、保留字或以阿拉伯数字开头时加前缀 豫_。同一函数里重复的局部名字依次加 _2、_3……。

其他递归（局部递归函数不在尾位置调用自身、顶层函数调用自身）、闭包与匿名函数、把函数当作值、鉴、字符串与模块投影（之）都报错，报错带文件与顶层定义名。

着六：例子

计算着色：

寻观「点染」之「点染本」之书。

「数据」乃「存储缓冲」于（「数组」于「无三二」）也。
「数据」者「绑定」于（「组值」【「组」者「零」也，「号」者「零」也，「访问」者「读写」也】）也。

「加倍」乃化「全局调用号」为「有」者也。
「加倍」者得「号」而
　令「序」为「号」中「零」而
　「写数组」于「数据」于「序」于（（「读数组」于「数据」于「序」）乘（「新无三二」于「二」））
也。

「加倍入口」乃「计算入口」也。
「加倍入口」者「计算入口项」于「加倍」于（「组值」【「工作组大小」者「四」也】）也。

译成：

@group(0) @binding(0) var<storage, read_write> 数据: array<u32>;

@compute @workgroup_size(4)
fn 加倍(@builtin(global_invocation_id) 号: vec3<u32>) {
  let 序 = 号.x;
  数据[序] = 数据[序] * 2u;
}

更多例子（顶点与片元、局部尾递归、格、两种「若」、均匀缓冲、纹理、原子、定长数组）与期望的 WGSL 见 库/点染/测试/样例。

着七：工具与测试

• node 豫言操作系统/宿主/节点/宿主.cjs 着色器.wasm 译 <文件>：把着色源译成 WGSL 打印到标准输出，入口表（入口项、函数名、阶段）打印到标准错误；金属 <文件> 同样译成 MSL；统算 <文件> 同样译成 PTX（入口表同上，PTX 里的入口名由函数名按 UTF-8 字节改写成 ASCII）；检查 <文件> 只做语法分析与类型检查；生成 <文件> 打印高级代码导入它时构建所得模块的豫言源码。工具在当前目录发现包，为文件所属的包启用包上下文，所以在仓库根目录运行。
• 库/点染/测试/着色翻译。测试。豫 对样例目录里的每个着色源核对期望的 WGSL（名.wgsl）、期望的报错文字（名。错误）、期望的 MSL（名。金属）、期望的 PTX（名。统算）或生成模块的源码（名。生成）；库/点染/测试/着色导入其一。豫 是构建集成的黄金样例；库/点染/测试/金属运行。测试。豫 导入样例着色源，经「金属」库把生成的 MSL 编译、派发并核对读回的数值（没有金属时打印“跳过”）；库/点染/测试/统算运行。测试。豫 同样经「千手」库把生成的 PTX 载入、发射并核对（没有统算设备时打印“跳过”）。四者由全树测试自动运行。

着八：限制

• 隐式类型参数只由变量或定义引用形式的实参、以及期望类型推定：实参是投影、调用等复合表达式、又没有期望类型时报“函数应用的隐式参数无法完全确定”。此时先用 令 绑定子表达式，或用 也 标注类型；只出现在返回类型里的类型参数（如「转换」）用 以 给出。
• 元组字面量最右一项是括号里的元组时会被摊平（甲 与（乙 与 丙） 得三元组），矩阵的各列先定义为常量或绑定为变量。
• WGSL 的常量求值不许出现 NaN 与无穷：bitcast 一个常量位型得到 NaN 或无穷时，建模块即报错。要用 NaN 或无穷时借一个运行期的值按位合成（如 库/算阵/公用。着色器。豫 的「成非数」「成正无穷」）。
• 定长长度与标量宽度须在类型里写成整数字面量，翻译器从字面量读数；写成变量或算式时翻译报错。
• 定长数组只能在工作组变量里按元素写入；局部变量与缓冲里的定长数组按元素写入是待办事项。
• 点积、长度、距离的返回类型是浮三二；整数向量的点积、按模块取名（之）、多个着色模块间的重名检查都是待办事项。

着九：金属目标（MSL）

着色源也可以译成 MSL（金属着色语言），供金属在运行时编译（如经诺驿的 node:ffi 调用 newLibraryWithSource）。金属目标只支持计算入口；顶点、片元入口译成一行注释略去（待办事项）。类型检查与着四、着五的规则两种目标共用，下面只列 MSL 的不同写法。

• 开头写 #include <metal_stdlib> 与 using namespace metal;，其后结构、常量、函数依源码次序；资源不单独写出。
• 类型：f32、f16、i32、u32、bool 写成 float、half、int、uint、bool；向量写成 float4 等，矩阵写成 float4x4（列数×行数）；原子写成 atomic_uint、atomic_int；定长数组写成 array<T, N>；二维纹理写成 texture2d<float>，采样器写成 sampler。结构的字段写成 T 名;，不带属性。
• 资源：MSL 没有模块级资源，资源都作为函数参数。存储缓冲里的运行期数组写成 device T*（只读为 const device T*），其余存储缓冲写成 device T&，均匀缓冲写成 constant T&，纹理与采样器照写类型，工作组变量写成 threadgroup T&。普通函数把它直接或经调用用到的资源依次附在原有参数之后，调用处依次传入。
• 计算入口写成 kernel void，前面加 [[max_total_threads_per_threadgroup(N)]]（N 是工作组大小；宿主派发时每组线程数应取 N，可由计算管线的 maxTotalThreadsPerThreadgroup 读回）。参数依次是内建输入与用到的资源；由内建输入组成的结构按字段展开成名为“参数名_字段名”的参数，函数体开头再拼回结构。内建输入的属性：全局调用号 [[thread_position_in_grid]]，局部调用号 [[thread_position_in_threadgroup]]，局部调用序 [[thread_index_in_threadgroup]]，工作组号 [[threadgroup_position_in_grid]]，工作组数 [[threadgroups_per_grid]]。工作组变量在函数体开头声明为 threadgroup T 名;，不作参数。
• 绑定序号：存储缓冲与均匀缓冲加 [[buffer(n)]]，纹理加 [[texture(n)]]，采样器加 [[sampler(n)]]，n＝组号×16＋绑定号（三类各有自己的序号空间）。宿主按这个序号设置缓冲，如组 0 号 2 设在序号 2，组 1 号 0 设在序号 16。金属一个函数最多 31 个缓冲，所以组号只能用 0 与 1（未检查，待办事项）。
• 语句：let 写成 const auto 名 = 值;，var 名: T = 值; 写成 T 名 = 值;，loop 写成 while (true)，_ = 值; 写成 (void)(值);，select(乙, 甲, 条件) 写成 (条件 ? 甲 : 乙)，结构构造写成 S{…}，顶层常量写成 constant T 名 = 值;。
• 字面量：小数字面量加 f；「新浮十六」于「1.0」写成 1.0h，「新整三二」于「三」写成 int(3)，「新无三二」于「二」照样写 2u。
• 内建函数：多数与 WGSL 同名；逆平方根为 rsqrt，四舍五入为 rint（偶数舍入，同 WGSL 的 round），浮点取余为 fmod，位转换为 as_type<T>(x)；半精度解包、打包为 float2(as_type<half2>(x))、as_type<uint>(half2(v))；工作组屏障为 threadgroup_barrier(mem_flags::mem_threadgroup)，存储屏障为 threadgroup_barrier(mem_flags::mem_device)，舍弃为 discard_fragment()；原子操作为 atomic_fetch_add_explicit(&a[i], v, memory_order_relaxed)、atomic_fetch_max_explicit、atomic_fetch_min_explicit、atomic_load_explicit、atomic_store_explicit；采样纹理、读纹素、纹理尺寸写成纹理对象的 sample、read、get_width() 与 get_height()。抽象字面量与具体类型的值一起运算时先转成该类型（如 max(x, uint(1))）。
• 名字与 WGSL 目标相同（中文标识符照写）；两种目标共用一张保留字表，所以入口函数在两种目标里同名。

限制（待办事项）：顶点与片元入口；「数组长度」（MSL 没有运行期数组长度，须由宿主另传）；结构与缓冲的内存布局与 WGSL 不尽相同（如 float3 占 16 字节）；函数名与金属内建函数同名时未改写。

着十：统算目标（PTX）

着色源也可以译成 PTX（英伟达统算的虚拟指令集），由 库/千手 交给驱动即时编译（「载千手篇」），不经 C++ 与 NVRTC。统算目标只支持计算入口；顶点、片元入口写成一行注释略去（待办事项）。类型检查与着四、着五的规则各目标共用，下面只列 PTX 的写法。

• 开头写 .version 8.0、.target sm_80、.address_size 64：计算能力 8.0 及以上（安培起）的显卡都能即时编译，在 RTX 5090 上驱动编成 sm_120a。输出只含 ASCII 字符（ptxas 遇到非 ASCII 字符即报错，注释里也不行）。
• 名字：中文名字确定性地改写成 ASCII。名字的 UTF-8 字节里，ASCII 字母与数字原样保留，其余每个字节（含下划线）写成下划线加两位十六进制，前面加前缀。入口的 PTX 名由它在 WGSL 里的函数名改写、前缀 yy，如『加倍』得 yy_e5_8a_a0_e5_80_8d；工作组变量前缀 ys；寄存器只用编号（%p 谓词、%h 十六位、%r 三十二位、%rd 六十四位）。生成模块为每个入口项另导出“入口项名并线汇编名”常量（着二）。
• 入口：每个计算入口写一个 .visible .entry，用 .reqntid N, 1, 1 把每块线程数固定为工作组大小 N，宿主发射时块维须恰为 (N, 1, 1)，网格维即派发的工作组数。内建输入由特殊寄存器算出：全局调用号 = %ctaid × %ntid + %tid（逐维），局部调用号 = %tid，局部调用序 = tid.x + ntid.x × (tid.y + ntid.y × tid.z)，工作组号 = %ctaid，工作组数 = %nctaid；由内建输入组成的结构按字段分别算出。
• 参数：各是一个 .param .u64 设备指针，名为 p_g组号_b绑定号。入口（直接或经调用）用到的缓冲都在组 0 时，绑定 0 至所用最大绑定号的每个号各占一个参数（入口没用到的占位绑定也占一位），与宿主按绑定号依次传缓冲的约定一致；否则为用到的存储缓冲与均匀缓冲，按（组号，绑定号）升序。均匀缓冲也按设备指针传，宿主先把内容写进一小块显存。没有标量参数。入口开头用 ld.param.u64 取出参数，再用 cvta.to.global.u64 换成全局地址。宿主经 库/千手 发射时，参数列依次是各缓冲的「统算显存参」。
• 工作组变量在入口体里声明为 .shared .align A .b8 名[字节数]，用 ld.shared、st.shared 读写；工作组屏障与存储屏障都译成 bar.sync 0。
• 内存布局按 WGSL 存储缓冲的规则：f16 对齐 2，其余标量对齐 4；二维向量对齐为元素的两倍，三维、四维为四倍；矩阵按列向量；结构取字段对齐的最大值；数组步长为元素大小向上取整到对齐。均匀缓冲也按这套布局（WGSL 均匀缓冲另要求 16 字节对齐，待办事项）。布尔只在共享内存里按四字节存。
• 值：每个值按分量放在寄存器里（向量按分量，矩阵按列，结构按字段依次）；格是一组可重写的寄存器，「读格」时复制一份。立即数直接写进指令（f32 带小数点）；半精度立即数先按 f32 搬进寄存器再转换（PTX 不收 f16 立即数）。
• 控制流：「若」译成条件跳转与标签，值位置两支把值搬进同一组结果寄存器；局部尾递归展开成带标签的循环，自尾调用把新实参先算进临时寄存器再搬进参数寄存器，然后跳回循环开头；用户函数与顶层常量在用处内联，不单独写出 .func。
• 运算：算术、位运算、移位、比较（浮点的不等于用无序比较 neu）与逻辑运算逐分量译成 PTX 指令；f32 除法用 div.rn，浮点取余为 x − trunc(x/y)·y。数学函数用 PTX 的指令：sqrt.rn、rsqrt.approx、ex2.approx、lg2.approx、sin.approx、cos.approx、tanh.approx，指数与对数换底，幂为 ex2(y·lg2 x)，取整用 cvt.rmi、cvt.rpi、cvt.rni、cvt.rzi，饱和用 cvt.sat；半精度的除法与数学函数经 f32 算。原子加、最大、最小用 atom.global，原子读、原子存用 ld.relaxed.gpu、st.relaxed.gpu。位转换在三十二位类型之间不改变位，二维半精度与 u32 之间用 mov.b32 打包或拆开。

限制（待办事项）：顶点与片元入口；纹理与采样器；「数组长度」（须由宿主另传）；反三角函数（反正弦、反余弦、反正切、反正切二）与半精度的平滑阶跃；越界访问不检查。

着十一：子组运算

子组是一起执行的一组线程：网页图算叫 subgroup，金属叫 simdgroup，统算叫线程束（warp）。子组运算在子组内的线程之间交换数据，不经工作组内存，也不用屏障，适合替代工作组内对半归约的最后几步。

规则：

• 值只用 32 位标量（浮三二、无三二、整三二）。统算目标对别的类型报错；WGSL 与 MSL 虽也接受向量，本库不依赖。
• 须在一致控制流里调用：整个子组都执行到这一调用，不能放在只有部分线程进入的「若」分支或循环里。WGSL 由一致性分析检查，违反时建模块报 must only be called from uniform control flow；统算的 shfl.sync 以当时活跃的线程（activemask）为成员，归约时不活跃的线程代以单位元（和取 0，最大取 −∞ 或最小整数，最小取 +∞ 或最大整数），所以线程束不满（如工作组 48 个线程）时与金属一样只在活跃线程之间计算；广播与洗牌的源线程不活跃时结果未定义。翻译器把子组运算当作有副作用的调用，「若」的两支里有子组运算时不译成 select。
• 广播要求各线程给的序相同（结果是子组里第「序」个线程的值）；WGSL 的 subgroupBroadcast 要求序是常量表达式，所以广播与洗牌都译成 subgroupShuffle。
• 最大、最小遇到 NaN 的结果随目标而定（PTX 的 max.f32 返回非 NaN 的一方）；要 NaN 传播时，用异或洗牌取值后自己归并。
• 子组大小：金属与统算都是 32；网页图算随设备而定（8 至 128），可由适配器信息的 subgroupMinSize、subgroupMaxSize 查得。依赖子组的内核一般假定子组由连续的局部调用序组成、子组大小不小于 32（如按 32 个线程一段做异或洗牌归约），设备不满足时应改用不带子组运算的内核。
• 网页图算：用到子组运算或子组内建输入的着色源，译成 WGSL 时开头加 enable subgroups;（同时用到 f16 时两条都加）。设备须有 subgroups 特性，并在 requestDevice 时要求；Chrome 134 起支持，本机 Dawn 也支持。没有该特性的设备上建模块即报错，所以用子组的内核与不用子组的内核应放在不同的着色源里，由后端按设备特性选用。

限制（待办事项）：子组的投票（ballot、any、all）、前缀和与上下移洗牌；十六位与向量的子组运算；翻译器不检查子组运算是否处在一致控制流里。

着十二：子组矩阵片

预填充实验内建「载入矩阵片累」从单精度共享数组载入16×16累加片，起、步以单精度元素计；「载矩阵片乙自转置存储」从半精度位型设备缓冲载入乙的转置布局，乙[n][k] = 缓冲[起+k×步+n]，起、步以半精度元素计。两者完整区域在界内、起及步为8的倍数、全子组一致调用。金属与统算可用：统算按 mma m16n8k16 的片布局逐道读（累片用 ld.shared.v2.f32，转置乙用 ld.global.b16 再两两拼成寄存器）。

深度学习推理少行接口「载入矩阵片累前八行」只从单精度共享数组载入前8×16元素，返回16×16累加片，下八行置零。起点、行步长以单精度元素计，均须为8的倍数，仅所读八行须在界内；全子组一致调用；金属与统算可用。辅助源码只在使用该接口时生成。

深度学习缓存直读实验接口「载矩阵片甲自存储」「载矩阵片乙自存储」从无三二存储缓冲读取半精度16×16片，每字两元素，低半字在前；起、步以半精度元素计，均为8的倍数，整片在界内，乙按转置布局载入。少行实验接口「载矩阵片甲前八行」只从半精度共享数组载入前8×16元素，后八行置零；「存矩阵片累前八行」只向单精度共享数组写前8×16元素，其余元素保持原值。少行接口只要求八行在界内。所有调用须全子组一致执行；金属与统算可用（统算按 mma m16n8k16 的片布局逐道读写，自存储用 ld.global.b32），网页图算的支持留待后续版本。

「存矩阵片累到存储」接受读写单精度存储缓冲、无三二起、无三二步、子组矩阵片累，返回单位值。写出布局为缓冲[起+行×步+列]，行列各0至15；起、步须为8的倍数，完整区域由调用方检查边界。整个子组须在一致控制流里调用。Metal以设备指针调用simdgroup_store，PTX以st.global.v2.f32写出，WGSL翻译报错。

子组矩阵片是由整个子组一起持有的 16 × 16 矩阵块，用来调用硬件的矩阵乘加单元：金属译成 simdgroup_matrix，统算译成 ldmatrix 与 mma.sync。各线程持有哪几个元素由目标决定，着色源看不到，只能整块载入、乘加、存回。

规则：

• 载入的数组须是半精度（「浮十六」）的工作组数组，存回的数组须是单精度（「浮三二」）的工作组数组；「起」与「步」按元素算，须是 8 的倍数（PTX 的 ldmatrix 要求每行地址按 16 字节对齐，翻译器把工作组数组都按 16 字节对齐声明）。
• 须在一致控制流里由整个子组一起调用；翻译器把这些调用都当作有副作用的调用，不检查一致性。
• 金属：用到子组矩阵片的源码开头加一段结构与内联函数（「金矩阵片序幕」）。乙片载入时转置，所以乘加直接得 甲 · 乙ᵀ；半精度乘、单精度累加。
• 统算：甲片是 mma 的 A（行主序），乙片的前两个寄存器与后两个分别是两个 n8 块的 B（.col，即按 [n][k] 存），累片的前 4 个与后 4 个寄存器分别是两个 n8 块的 C 与 D；需要 sm_80（PTX 目标为 .target sm_80）。
• 网页图算：WGSL 没有子组矩阵片，翻译报错。高级代码导入这样的着色源时，「着色源」是一行说明原因的注释，入口表取自 MSL；用子组矩阵片的内核应放在单独的着色源里，由后端按目标选用。
• 整八片：译成 PTX，金属与网页图算的写法留待后续版本（翻译报错）。高级代码导入只能译成 PTX 的着色源时，「着色源」「金属源」都是一行说明原因的注释，入口表取自 PTX。整累片的元素可按位型当单精度读：初值取「满整矩阵片累」于 0x4B400000（1.5 × 2²³ 的位型）时，乘加后的元素当单精度读恰是 12582912 + 点积（|点积| < 2²²），由此不经整数转单精度的指令。
• 精度：甲、乙是半精度（有效位 11 位，最大 65504），累加是单精度；把单精度数据放进片之前，调用者要自己转成半精度，超出范围会成为无穷。

限制（待办事项）：只有半精度乘单精度累加（16 × 16 × 16）与整八乘三十二位整数累加（16 × 16 × 32，只有 PTX）两种；片不能直接从存储缓冲载入或存回存储缓冲。

着十三：张量片（Metal 4 张量运算）

张量片是由整个工作组（恰为 128 线程、4 个子组）一起持有的 64 × 64 单精度累加块，用 Metal 4 的张量运算（Metal Performance Primitives 的 matmul2d）计算，在 M5、A19 及以后的 GPU 上由神经加速器执行。各线程持有哪几个元素由目标决定，着色源看不到。只有金属支持；网页图算与统算翻译报错。设备是否支持由宿主判断（金属张量设备按 MTLGPUFamilyApple10 查询），着色源本身不检查。

这些调用都须在一致控制流里由整个工作组一起执行。MSL 译法：用到张量片的源码开头加 #include <metal_tensor> 与 #include <MetalPerformancePrimitives/MetalPerformancePrimitives.h>，matmul2d_descriptor(64, 64, 32, false, true, true, multiply_accumulate) 配 execution_simdgroups<4>；累加块是它的目标协作张量（yy_张量片累），甲按 tensor<device float, dextents<int32_t, 2>, tensor_inline> 带步长包住（行数作外层长度），乙按线程组半精度张量包住，存回时按（列数，行数）截断。累加块可以放进「新格」，在循环里「读格」「写格」。