跳到正文豫言技术文档

豫言 / 技术规范

本篇描述 4cdeca1 的实现。版本与适用边界见基线说明。

着色语言

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

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

着一:文件后缀

后缀 内容
名。着色器。豫 着色源:只用豫言标准语法(乃、者、即、虑…而、复令、若…则…否则、;、于、以、中、也、「组值」【…】、「组类」【…】)
名。着色器。接口。豫 着色内建声明:写法同 。接口。豫,只有 立、即 与不带函数体的 乃 签名,语法分析给签名补占位定义。其中的函数由着色目标提供:翻译器把调用换成 WGSL、MSL 的运算符与内建函数或 PTX 指令,运行时不存在这些调用
名。着色生成。豫 保留后缀:高级代码导入着色源时,构建所得模块的虚拟路径;磁盘上没有这个文件

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

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

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

  • 「着色源」: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。

类型 别名 WGSL
「全局调用号」 三维无 @builtin(global_invocation_id)
「局部调用号」 三维无 @builtin(local_invocation_id)
「局部调用序」 无三二 @builtin(local_invocation_index)
「工作组号」 三维无 @builtin(workgroup_id)
「工作组数」 三维无 @builtin(num_workgroups)
「顶点序号」 无三二 @builtin(vertex_index)
「实例序号」 无三二 @builtin(instance_index)
「裁剪位置」 四维浮 @builtin(position)(顶点输出)
「片元坐标」 四维浮 @builtin(position)(片元输入)
「正面朝向」 爻 @builtin(front_facing)
「片元深度」 浮三二 @builtin(frag_depth)
「采样序号」 无三二 @builtin(sample_index)
「采样掩码」 无三二 @builtin(sample_mask)
「子组内序号」 无三二 @builtin(subgroup_invocation_id)(见着十一)
「子组大小」 无三二 @builtin(subgroup_size)(见着十一)

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

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

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

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

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

豫言 WGSL
加、减、乘、除以、取余 +、-、*、/、%
取负 一元 -
数乘(向量或矩阵乘标量)、矩阵乘 *
小于、大于、小于等于、大于等于、等于、不等于 <、>、<=、>=、==、!=(返回爻)
且、或者、非 &&、\|\|、!
位与、位或、位异或、位反、左移、右移(移位量是「无三二」,整三二的右移为算术移位) &、\|、^、~、<<、>>
解包半精度(一个 u32 的低、高 16 位按 binary16 解成两个 f32)、打包半精度 unpack2x16float(x)、pack2x16float(v)
平方根、逆平方根、指数、二指数、对数、二对数、幂 sqrt、inverseSqrt、exp、exp2、log、log2、pow
正弦、余弦、正切、反正弦、反余弦、反正切、反正切二、双曲正弦、双曲余弦、双曲正切 sin、cos、tan、asin、acos、atan、atan2、sinh、cosh、tanh
绝对值、符号、向下取整、向上取整、四舍五入、截尾取整、小数部分、饱和 abs、sign、floor、ceil、round、trunc、fract、saturate
最大、最小、阶跃、夹紧、混合、平滑阶跃、融合乘加 max、min、step、clamp、mix、smoothstep、fma
点积、长度、距离、叉积、归一化、反射 dot、length、distance、cross、normalize、reflect
转浮三二、转浮十六、转无三二、转整三二 f32(x) 等(向量实参得相应向量)
转换、位转换(须用 以 给出源类型与目标类型) T(x)、bitcast<T>(x)
读数组、写数组、数组长度 a[i]、a[i] = v;、arrayLength(&a)
读四浮、读四无 于 a 于 i vec4<f32>(a[i], a[(i) + 1u], a[(i) + 2u], a[(i) + 3u])(四无为 vec4<u32>);MSL 为 *(const device float4*)(&(a)[i]),PTX 为 ld.global.v4.f32(或 .u32),一次读 16 字节,要求 i 是 4 的倍数
读半字 于 a 于 i(a 为 32 位无符号存储数组) 把 a 当作 16 位单元的数组读第 i 个、零扩展成 32 位:WGSL 为 ((a[(i) >> 1u] >> (((i) & 1u) * 16u)) & 65535u),MSL 为 uint(((const device ushort*)(&(a)[0]))[i]),PTX 为 ld.global.u16 加 cvt.u32.u16;用于只按 2 字节对齐的数据(如 Q6_K)
读四有号字节 于 a 于 b(a 为 32 位无符号存储数组,b 为字节偏移) 读第 b 至 b + 3 个字节、各按 8 位有符号数转成四维浮,偏移不必对齐:WGSL 逐字节取出(f32(i32((a[(b + j) >> 2u] >> (((b + j) & 3u) * 8u)) << 24u) >> 24u)),MSL 为 float4(char4(*((const device packed_char4*)(((const device char*)(&(a)[0])) + (b))))),PTX 为四次 ld.global.s8 加 cvt.rn.f32.s32;用于只按 2 字节对齐的块量化数据(如 Q8_0)
写共享四半 于 a 于 i 于 v(a 为工作组里的半精度定长数组,v 为四维浮) 把 v 的四个分量转成半精度写进 a[i] 至 a[i + 3]:WGSL 逐个写,MSL 为 *((threadgroup half4*)(&(a)[i])) = half4(v);,PTX 为四次 cvt.rn.f16.f32 加 st.shared.v4.b16;要求 i 是 4 的倍数
读缓冲、写缓冲、读均匀、读共享、写共享 读写资源的值
读共享数组、写共享数组(工作组变量里的定长数组)、取元素(定长数组值的一项,如结构字段) a[i]、a[i] = v;
新浮三二、新浮十六、新整三二、新无三二 带类型的字面量 0.5f、40.0f、f16(1.0)、3i、2u,或类型转换
新格、读格、写格 var、读变量、赋值
采样纹理、读纹素、纹理尺寸 textureSample、textureLoad、textureDimensions
按层采样纹理 于 纹理 于 采样器 于 坐标 于 层 textureSampleLevel:按给定的层(浮三二,可带小数)采样,各阶段都能用
比较采样 于 深度纹理 于 比较采样器 于 坐标 于 参照 textureSampleCompare:按采样器的比较函数拿“参照 比较 所存深度”判定,成立得 1、不成立得 0,线性过滤时取相邻四个纹素判定的加权平均;只用于片元着色
横差、纵差 dpdx、dpdy(屏幕空间导数,只用于片元着色,须在一致控制流里)
工作组屏障、存储屏障、舍弃 workgroupBarrier();、storageBarrier();、discard;
原子加、原子最大、原子最小、原子读、原子存 atomicAdd(&a[i], v) 等
子组求和、子组最大、子组最小、子组广播、子组洗牌、子组异或洗牌 subgroupAdd、subgroupMax、subgroupMin、subgroupShuffle、subgroupShuffleXor(见着十一)
载入矩阵片甲、载入矩阵片乙、载入矩阵片乙转置、零矩阵片累、矩阵片乘加、存矩阵片累、累片元、累片换元、累片元行、累片元列、累片转甲片 WGSL 留待后续版本,翻译报错(MSL 与 PTX 的译法见着十二)
零张量片累、张量片乘加、存张量片累 WGSL 与 PTX 不支持,翻译报错(MSL 的译法见着十三)

着四:顶层定义的分类

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

  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
令「甲」为 值 而 … let 甲 = 值;;值只由字面量组成时用 const,保留抽象字面量类型;单位值不绑定名字
令「甲」为「新格」于 初值 而 … var 甲: T = 初值;,T 取 以 给出的类型,否则取初值类型的默认具体类型
「读格」于 甲、「写格」于 甲 于 值 甲、甲 = 值;
甲;乙 依次执行,前者的值有副作用时写成 _ = 值;
若 条件 则 甲 否则 乙 两支都是无副作用的简单表达式且类型是标量或向量时译成 select(乙, 甲, 条件);否则在值位置用临时 var 加 if/else,在尾位置两支各自 return
复令「环」属 类型 即 得…而 体 而 …「环」于 初值… 局部尾递归:在调用处展开成 loop,参数变成 var;自尾调用译成赋值加 continue(多个参数经临时量平行赋值),其余尾值写 return,或写进结果变量再 break
函数 于 实参… 用户函数译成 f(…),着色内建译成运算符或内建函数(见着三)
值 中「零」、值 中「字段」 向量取 .x/.y/.z/.w,矩阵取列 [i],结构取字段 .字段
「组值」【「字段」者 值 也,…】 结构构造 S(…):期望类型是结构时用它,否则按字段名找已声明的结构
甲 与 乙 … 向量或矩阵构造:期望类型给出维数与元素类型时照用,否则由元素推定(标量得向量,同维向量得矩阵)
类型 也 值 值以标注类型为期望类型翻译;值是抽象字面量时用该类型的构造器转换
不经构造器的数字字面量、「阳」、「阴」 WGSL 抽象字面量、true、false

名字原样写成 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,二维深度纹理写成 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)]]……,与片元阶段输入的同号字段按名字对上;片元返回颜色时写其类型(即颜色目标 0),返回结构时字段依次带 [[color(N)]],片元深度带 [[depth(any)]],采样掩码带 [[sample_mask]];没有返回值写 void。顶点返回别的、片元单独返回内建量报错。
  • 绑定序号:存储缓冲与均匀缓冲加 [[buffer(n)]],纹理与深度纹理加 [[texture(n)]],采样器与比较采样器加 [[sampler(n)]],n=组号×16+绑定号(三类各有自己的序号空间)。金属的采样器序号只到 15,所以采样器与比较采样器只能放在组 0(未检查,待办事项)。宿主按这个序号设置缓冲,如组 0 号 2 设在序号 2,组 1 号 0 设在序号 16。金属一个函数最多 31 个缓冲,所以组号只能用 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 占 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)。子组运算在子组内的线程之间交换数据,不经工作组内存,也不用屏障,适合替代工作组内对半归约的最后几步。

豫言 WGSL MSL PTX
「子组内序号」(内建输入) @builtin(subgroup_invocation_id) [[thread_index_in_simdgroup]] %laneid
「子组大小」(内建输入) @builtin(subgroup_size) [[threads_per_simdgroup]] WARP_SZ(32)
子组求和、子组最大、子组最小 于 值 subgroupAdd、subgroupMax、subgroupMin simd_sum、simd_max、simd_min 五步 shfl.sync.bfly.b32(掩码 16、8、4、2、1;成员取 activemask,源线程不活跃时代以单位元),每步后 add、max、min
子组广播、子组洗牌 于 值 于 序 subgroupShuffle simd_shuffle shfl.sync.idx.b32
子组异或洗牌 于 值 于 掩码 subgroupShuffleXor simd_shuffle_xor shfl.sync.bfly.b32

规则:

  • 值只用 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。各线程持有哪几个元素由目标决定,着色源看不到,只能整块载入、乘加、存回。

豫言 含义 MSL PTX
「子组矩阵片甲」「子组矩阵片乙」 半精度 16 × 16 块 结构,含 4 个 simdgroup_half8x8 4 个 32 位寄存器(每个两个半精度)
「子组矩阵片累」 单精度 16 × 16 累加块 结构,含 4 个 simdgroup_float8x8 8 个 32 位寄存器
载入矩阵片甲 于 数组 于 起 于 步 片[行][列] = 数组[起 + 行 × 步 + 列] 4 次 simdgroup_load ldmatrix.sync.aligned.m8n8.x4.shared.b16
载入矩阵片乙 于 数组 于 起 于 步 同上(乙按 [n][k] 存) 4 次转置的 simdgroup_load(存的是乙ᵀ) 同上,各线程给的行地址不同
载入矩阵片乙转置 于 数组 于 起 于 步 乙按 [k][n] 存:片[n][k] = 数组[起 + k × 步 + n](如注意力里的值矩阵) 4 次不转置的 simdgroup_load ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16
累片元 于 累 于 序 本道持有的第 序 个累片元素(0 至 7) 第 ⌊序/2⌋ 块的 thread_elements()[序 mod 2] 与 0 至 7 比较的 setp.eq.u32 与 selp.b32 链(序为常数时即时编译折掉)
累片换元 于 累 于 序 于 值 把第 序 个元素换成 值 的新累片 改 thread_elements() 后返回 8 次 setp 与 selp
累片元行、累片元列 于 序 本道第 序 个元素在 16 × 16 里的行、列 道号取 simd_prefix_exclusive_sum(1u),按 simdgroup 8 × 8 的布局求 道号取 %laneid:行 ⌊道/4⌋ + 8⌊(序 mod 4)/2⌋,列 8⌊序/4⌋ + 2(道 mod 4) + (序 mod 2)
累片转甲片 于 累 同位置转成半精度的甲片(注意力里把权重直接当甲) 逐块 half 转换 4 次 cvt.rn.f16x2.f32(千手上不必在道间搬移)
零矩阵片累 于 元 全零的累加块 make_filled_simdgroup_matrix 8 次 mov.b32 …, 0
「子组整八矩阵片甲」「子组整八矩阵片乙」 有号八位整数 16 × 32 块(乙为 32 × 16,按 [n][k] 存) 留待后续版本 4 个 32 位寄存器(每个 4 个有号字节)
「子组整矩阵片累」 三十二位整数 16 × 16 累加块,元素布局同「子组矩阵片累」 留待后续版本 8 个 32 位寄存器
载入整八矩阵片甲、载入整八矩阵片乙 于 数组 于 起 于 步 片[行][4 × 字列 + j] = 数组[起 + 行 × 步 + 字列] 的第 j 个字节(数组是无三二工作组数组,起、步按字算,须是 4 的倍数) 留待后续版本 ldmatrix.sync.aligned.m8n8.x4.shared.b16(字节布局与半精度片相同)
零整矩阵片累 于 元、满整矩阵片累 于 值 各元素为 0 或 值 的整累片 留待后续版本 8 次 mov.b32
整八矩阵片乘加 于 累 于 甲 于 乙 累 + 甲 · 乙ᵀ(三十二位整数) 留待后续版本 2 次 mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32
整累片元 于 累 于 序 本道第 序 个整数元素(0 至 7) 留待后续版本 同「累片元」
矩阵片乘加 于 累 于 甲 于 乙 累[m][n] + Σₖ 甲[m][k] × 乙[n][k](加上 甲 · 乙ᵀ) 8 次 simdgroup_multiply_accumulate 2 次 mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32
存矩阵片累 于 数组 于 起 于 步 于 累 数组[起 + 行 × 步 + 列] = 累[行][列] 4 次 simdgroup_store 4 次 st.shared.v2.f32

规则:

  • 载入的数组须是半精度(「浮十六」)的工作组数组,存回的数组须是单精度(「浮三二」)的工作组数组;「起」与「步」按元素算,须是 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 查询),着色源本身不检查。

内建 含义
「张量片累」 类型:64 × 64 单精度累加块
「零张量片累」于「元」 全零的累加块
「张量片乘加」于 累 于 甲 于 甲起 于 甲步 于 甲行数 于 乙 于 乙起 于 乙步 得 累[m][n] + Σₖ 甲[甲起 + m × 甲步 + k] × 乙[乙起 + n × 乙步 + k](m、n 取 0 至 63,k 取 0 至 31);甲是单精度存储缓冲,m 不小于甲行数的行当作零;乙是半精度工作组数组(64 行在界内)
「存张量片累」于 出 于 起 于 步 于 行数 于 列数 于 累 出[起 + m × 步 + n] = 累[m][n],只写 m < 行数、n < 列数;出是读写的单精度存储缓冲

这些调用都须在一致控制流里由整个工作组一起执行。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> 带步长包住(行数作外层长度),乙按线程组半精度张量包住,存回时按(列数,行数)截断。累加块可以放进「新格」,在循环里「读格」「写格」。