豫言 / 技术规范

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

着色语言

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

各条以“着”编号。翻译器在 库/编译器核心/编译步骤/着色翻译/,着色内建声明在包「着色库」(库/着色库/),命令行工具是包「着色器翻译」(应用/开发工具/着色器翻译/,产物 yy着色器),测试在包「着色库测试」(库/着色库/测试/)。

着一:文件后缀

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

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

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

其余文件(高级代码)依次查 名。豫、名。接口。豫、名。应用接口。豫、名。着色器。豫。找到着色源时,得到构建时由它生成的模块(虚拟路径 名。着色生成。豫),该模块导出:

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

例如 寻「应用」之「加倍」之书。 之后,「加倍」之「着色源」 是 WGSL 文字,「加倍」之「加倍入口」 是 『加倍』,「加倍」之「加倍入口PTX名」 是 『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。

类型 别名 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)

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

类型构造:「数组」于 T 是运行期数组 array<T>,「定长数组」于 T 于 L 是 array<T, N>(N 是定长长度 L 的数),「原子」于 T 是 atomic<T>,「纹理二维」于 T 是 texture_2d<T>,「采样器」是 sampler。资源「存储缓冲」「均匀缓冲」「工作组变量」只用于顶层资源声明;存储缓冲的访问方式是「只读」或「读写」。

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

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

函数:多态函数用 承 声明。下标(读数组、写数组、读共享数组、写共享数组、取元素与原子操作的序号)是「无三二」,字面下标写成「新无三二」于「零」(译成 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]、a[i] = v;
新浮三二、新浮十六、新整三二、新无三二 带类型的字面量 0.5f、f16(1.0)、3i、2u,或类型转换
新格、读格、写格 var、读变量、赋值
采样纹理、读纹素、纹理尺寸 textureSample、textureLoad、textureDimensions
工作组屏障、存储屏障、舍弃 workgroupBarrier();、storageBarrier();、discard;
原子加、原子最大、原子最小、原子读、原子存 atomicAdd(&a[i], v) 等

着四:顶层定义的分类

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

  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 yy着色器.wasm 译 <文件>:把着色源译成 WGSL 打印到标准输出,入口表(入口项、函数名、阶段)打印到标准错误;金属 <文件> 同样译成 MSL;统算 <文件> 同样译成 PTX(入口表里是 PTX 入口名);检查 <文件> 只做语法分析与类型检查;生成 <文件> 打印高级代码导入它时构建所得模块的豫言源码。工具在当前目录发现包,为文件所属的包启用包上下文,所以在仓库根目录运行。
  • 库/着色库/测试/着色翻译。测试。豫 对样例目录里的每个着色源核对期望的 WGSL(名.wgsl)、期望的报错文字(名.错误)、期望的 MSL(名.金属)、期望的 PTX(名.统算)或生成模块的源码(名.生成);库/着色库/测试/着色导入-1。豫 是构建集成的黄金样例;库/着色库/测试/金属运行。测试。豫 导入样例着色源,经「金属」库把生成的 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 六十四位)。生成模块为每个入口项另导出“入口项名PTX名”常量(着二)。
  • 入口:每个计算入口写一个 .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 对齐 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 打包或拆开。

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