着色语言

着色语言以豫言之常法书图形处理器之着色程序：着色之源经编译器之析文、解名、检类（佐以着色内建之声），乃译为 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)；值为整之浮类构造字面亦有点，「新浮三二」于「四零」译为 40.0f。实参非字面则译为类之转，如 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，「深度纹理二维」为 texture_depth_2d，「比较采样器」为 sampler_comparison（以为影图），「子组矩阵片甲」「子组矩阵片乙」「子组矩阵片累」为子组矩阵片（见着十二，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 常量之求不许非数与无穷：bitcast 一常量之位型而得非数或无穷，建模即报误。须用非数或无穷者，借运行之值依位合成之（如 库/算阵/公用。着色器。豫 之「成非数」「成正无穷」）。
• 定长之度与标量之宽须于类中书为整数之字面，译器依字面读其数；书为变量或算式则译时报误。
• 定长数组惟于工作组变量可逐元书之；局部之变与缓冲中之定长数组逐元书之，待办事项。
• 点积、长度、距离返浮三二；整向量之点积、按模取名（之）、诸着色模块间重名之察，皆待办事项。

着九：金属之的（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，二维深纹书为 depth2d<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 名;，不为参。
• 顶点、片元之入口书为 vertex、fragment 之函（不带工作组之大）。参：内建之入加属，顶点序号 [[vertex_id]]，实例序号 [[instance_id]]，片元坐标 [[position]]，正面朝向 [[front_facing]]，采样序号 [[sample_id]]，采样掩码 [[sample_mask]]；余入（顶点之属、片元插值之入）收于阶段之入之构 yy_函名_入，为参 yy_入 [[stage_in]]，顶点之项带 [[attribute(N)]]（N 即 WGSL 之 @location，宿主之顶点之描依之取属），片元之项带 [[user(locN)]]，整类复加 [[flat]]；结构之参依项展之（内建之项为参，名“参名_项名”，余项入阶段之入之构），函体之首复拼为构；常位之参于函首取之；所用之资为带序之参，工作组之变量报误。返：顶点惟返裁剪之位则书 float4（金属以 float4 之返为位），返构则书构名，构之裁剪位置之项带 [[position]]，余项依次带 [[user(loc0)]]、[[user(loc1)]]……，与片元阶段之入同号之项以名相接；片元返色则书其类（即色鹄零），返构则项依次带 [[color(N)]]，片元深度带 [[depth(any)]]，采样掩码带 [[sample_mask]]；无返书 void。顶点返他物、片元独返内建者报误。
• 绑之序：存储缓冲与均匀缓冲加 [[buffer(n)]]，纹理与深纹加 [[texture(n)]]，采样器与比较之采样器加 [[sampler(n)]]，n＝组号×16＋绑号（三类各有其序之域）。金属采器之序止于十五，故二采样器惟可置于组零（未察，待办事项）。宿主依此序置其缓冲，如组 0 号 2 置于序 2，组 1 号 0 置于序 16。金属一函至多三十一缓冲，故组号惟可用 0 与 1（未察，待办事项）。
• 句：let 书为 const auto 名 = 值;，var 名: T = 值; 书为 T 名 = 值;，loop 书为 while (true)，_ = 值; 书为 (void)(值);，select(乙, 甲, 条件) 书为 (条件 ? 甲 : 乙)，结构之造书为 S{…}，顶层之常量书为 constant T 名 = 值;。
• 字面：小数之字面加 f，值为整之浮字补其点（如 40.0f、40.0h）；「新浮十六」于「1.0」书为 1.0h，「新整三二」于「三」书为 int(3)，「新无三二」于「二」仍书 2u。
• 内建之函：多与 WGSL 同名；逆平方根为 rsqrt，四舍五入为 rint（舍入于偶，同 WGSL 之 round），横差、纵差为 dfdx、dfdy，浮之取余为 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()，按层采样纹理书为 sample(采器, 坐标, level(层))，比较采样书为深纹之 sample_compare(采器, 坐标, 参照)。抽象之字面与具体之值共算，则先化为其类（如 max(x, uint(1))）。
• 名同 WGSL 之的（中文之标识照书）；二的共一保留字之表，故入口之函于二的同名。

所限（待办事项）：「数组长度」（MSL 无运行数组之长，须宿主另传）；结构与缓冲之存布与 WGSL 不尽同（如 float3 占十六字节）；函名与金属内建之函同者未改。

着十：千手之的（PTX）

着色之源亦可译为 PTX（英伟达千手之虚指令之集），由 库/千手 付驱动即时编之（「载千手篇」），不经 C++ 与 NVRTC。千手之的惟支计算之入口；顶点、片元之入口书一行注而略之（待办事项）。类型之检与着四、着五之规，诸的共之；下惟列 PTX 之书。

• 首书 .version 8.0、.target sm_80、.address_size 64：算力八点零以上（安培以降）之卡皆可即时编之，于 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绑号。入口（直或经召）所用之缓皆在组零，则绑零至所用最大绑号每号各占一参（入口所不用之占位之绑亦占一位），合于宿主依绑号次第传缓之约；否则为所用之存储缓冲与均匀缓冲，依（组号，绑号）升序。均匀缓冲亦以器之指传之，宿主先书其容于一小块显存。无标量之参。入口之首以 ld.param.u64 取其参，复以 cvta.to.global.u64 化为全局之址。宿主经 库/千手 发射，参列依次为诸缓之「千手显存参」。
• 工作组变量于入口体中声为 .shared .align A .b8 名[字数]，以 ld.shared、st.shared 读写之；工作组屏障与存储屏障皆译为 bar.sync 0。
• 存之布依 WGSL 存储缓冲之规：f16 齐二，余标量齐四；二维向量之齐倍其元，三维、四维四之；矩阵依列向量；结构取项齐之最大；数组之步为元之大进至其齐。均匀缓冲亦依此布（WGSL 均匀缓冲另须齐于十六字，待办事项）。布尔惟于共享之存依四字存之。
• 值：每值依其分置于寄存器（向量依分，矩阵依列，结构依项之序）；格为一组可重书之寄存器，「读格」则抄之。立即之数径书于令中（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）。子组之算于子组之线间相易其数，不经工作组之存，亦不用屏障，宜以代工作组内对半而归之末数步。

其规：

• 值惟用三十二位之标量（浮三二、无三二、整三二）。千手之的遇他类则报误；WGSL 与 MSL 虽亦受向量，本库不恃之。
• 须于一致之控流召之：全子组皆行至此召，不得置于惟部分之线所入之「若」之支或环中。WGSL 以一致之析察之，违则建模报 must only be called from uniform control flow；千手之 shfl.sync 以其时活线（activemask）为员，归时不活之线代以单位之元（和取零，大取 −∞ 或最小之整，小取 +∞ 或最大之整），故线束不满（如工作组四十八线）则同金属，惟于活线间算之；广播与洗牌之源线不活则其果无定。译器以子组之算为有副作用之召，「若」之二支有子组之算则不译为 select。
• 广播须诸线所予之序同（其果为子组中第「序」线之值）；WGSL 之 subgroupBroadcast 须序为常量之式，故广播与洗牌皆译为 subgroupShuffle。
• 最大、最小遇非数，其果随的而定（PTX 之 max.f32 还非非数之一方）；欲非数相传，则以异或洗牌取值而自并之。
• 子组之大：金属与千手皆三十二；网页图算随器而定（八至百二十八），可由适配器之信之 subgroupMinSize、subgroupMaxSize 察之。恃子组之核，率设子组由连续之局部调用序而成、子组之大不小于三十二（如依三十二线一段行异或洗牌之归），器不合则宜改用不带子组之算之核。
• 网页图算：用子组之算或子组内建之入之着色源，译为 WGSL 则首加 enable subgroups;（并用 f16 则二者皆加）。器须有 subgroups 之特，且于 requestDevice 时求之；Chrome 一百三十四起支之，本机 Dawn 亦支之。无此特之器建模即报误，故用子组之核与不用子组之核宜置于异之着色源，由后端依器之特而择之。

所限（待办事项）：子组之投（ballot、any、all）、前缀之和与上下移之洗牌；十六位与向量之子组之算；译器不察子组之算是否处于一致之控流。

着十二：子组矩阵片

预填之试内建「载入矩阵片累」自单精共享列载十六乘十六累片，起、步以单精元计；「载矩阵片乙自转置存储」自半精位型设备缓载乙之转置布，乙[n][k] = 缓冲[起+k×步+n]，起、步以半精元计。二者全域在界、起与步为八之倍、全子组一致召之。金属与千手可用：千手依 mma m16n8k16 之片布逐道读之（累片用 ld.shared.v2.f32，转置之乙用 ld.global.b16 而两两拼为寄存器）。

深学推理少行之口「载入矩阵片累前八行」惟自单精共享列载前八乘十六元，还十六乘十六累片，下八行置零。起与行步以单精元计，皆八倍，惟所读八行须在界；全子组一致召之；金属与千手可用。惟用此口乃生辅源。

深度学习缓存直读之试「载矩阵片甲自存储」「载矩阵片乙自存储」自无三二存储缓读半精十六乘十六片，每字二元，低半字在先；起、步以半精元计，皆为八之倍，全片在界，乙以转置之布载之。少行之试「载矩阵片甲前八行」惟自半精共享列载前八乘十六元，后八行置零；「存矩阵片累前八行」惟向单精共享列书前八乘十六元，余元保旧值。少行之召惟须八行在界。诸召皆须全子组一致行之；金属与千手可用（千手依 mma m16n8k16 之片布逐道读写，自存储用 ld.global.b32），网页图算之支待后续之版。

「存矩阵片累到存储」受读写单精存储缓冲、无三二起、无三二步、子组矩阵片累，而返单位值。书之布局为缓冲[起+行×步+列]，行列各零至十五；起、步须为八之倍，召者察全域之界。全子组须于一致控流召之。Metal以设备指针召simdgroup_store，PTX以st.global.v2.f32书之，WGSL译则报误。

子组矩阵片者，全子组共持之十六 × 十六之阵块，以用器之阵乘加之单元：于金属译为 simdgroup_matrix，于千手译为 ldmatrix 与 mma.sync。诸线各持何元由的定之，着色之源不得见，惟可全块载之、乘加之、存回之。

其规：

• 所载之列须为半精（「浮十六」）之工作组之列，所存之列须为单精（「浮三二」）之工作组之列；「起」与「步」依元计，须为八之倍（PTX 之 ldmatrix 须每行之址齐于十六字节，译器声诸工作组之列皆齐于十六字节）。
• 须于一致之控流由全子组共召之；译器以此诸召皆为有副作用之召，不察其一致。
• 金属：用子组矩阵片之源首加一段结构与内联之函（「金矩阵片序幕」）。乙片载时转之，故乘加径得 甲 · 乙ᵀ；半精乘、单精累。
• 千手：甲片为 mma 之 A（行主序），乙片之前二与后二寄存器各为二 n8 块之 B（.col，即依 [n][k] 存），累片之前四与后四寄存器各为二 n8 块之 C 与 D；须 sm_80（PTX 之的为 .target sm_80）。
• 网页图算：WGSL 无子组矩阵片，译则报误。高级之代导此着色源，则「着色源」为一行注明其故，入口之表取于 MSL；用子组矩阵片之核宜置于别之着色源，由后端依的而择之。
• 整八片：译为 PTX，金属与网页图算之书法待后之版（译则报误）。高级之代导惟能译为 PTX 之着色源，则「着色源」「金属源」皆为一行注明其故，入口之表取于 PTX。整累片之元可依位型读为单精：初值取「满整矩阵片累」于 0x4B400000（1.5 × 2²³ 之位型），则乘加后之元以单精读之恰为 12582912 + 点积（|点积| < 2²²），由是不经整数转单精之令。
• 精度：甲、乙为半精（有效之位十一，最大 65504），累加为单精；单精之数入片之前，召者须自转为半精，逾其界则为无穷。

所限（待办事项）：惟半精乘单精累（十六 × 十六 × 十六）与整八乘三十二位整数累（十六 × 十六 × 三十二，惟 PTX）二者；片不可径自存储缓冲载之或存回存储缓冲。

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

张量片者，全工作组（恰百二十八线、四子组）共持之六十四 × 六十四单精之累块，以 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> 带步而裹之（行数为外层之长），乙依线组半精张量裹之，存回依（列数，行数）截之。累块可置于「新格」，于环中「读格」「写格」。