着色语言
着色语言用豫言的标准语法写 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』。高级代码导入着色内建声明时报错。
生成在编译器里完成,应用产物只含生成的字符串:
- 虚拟模块的依赖分析返回空列表(生成的模块没有导入),只备一份空的文件依赖缓存。
- 词法阶段调用翻译器。翻译器在内存里分析着色源及其导入的着色模块(依赖先分析),不写阶段缓存,所以并行编译的多个工作进程不会争写同一个缓存文件。生成的豫言源码与已有的词法缓存相同时不改写,后续阶段与导入它的文件照常命中缓存;着色源或着色内建改动后,生成的源码随之改变。
- 阶段环境读取虚拟模块的源码(如语言服务)时,从它的词法缓存取;缓存有效性以词法缓存的修改时间为源码时间。
- 虚拟模块按其着色源判断所属的包与编译模式。
着三:着色内建
库/着色库/着色内建。着色器。接口。豫 声明着色源可用的名字,包「着色库」没有依赖。着色源写 寻观「着色库」之「着色内建」之书。,着色源所在的包须声明依赖「着色库」。
标量与定长长度:标量的宽度与有无符号由名义类型区分。「定长长度」即「元类型」,每个长度是一个类型:「定长一」「定长二」「定长三」「定长四」「定长八」「定长十六」「定长三十二」「定长六十四」「定长一百二十八」「定长二百五十六」「定长五百一十二」「定长一千零二十四」「定长二千零四十八」「定长四千零九十六」。「定长小数」「定长整数」「定长无符号」各由一个定长长度构造出标量类型:「浮三二」即「定长小数」于「定长三十二」(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) 等 |
着四:顶层定义的分类
翻译器读语法分析所得的定义列(名称未解析,保留源码里的名字与类型别名),名字查找依次看局部名字、本着色程序的顶层名字、着色内建。着色源的顶层定义按以下规则翻译:
即定义带字段名的组类:译成struct。结构用作入口的参数或返回类型时,内建输入输出类型的字段加@builtin(…),其余字段按出现次序编@location(0)、@location(1)……;整数类型的字段用作顶点输出或片元输入时再加@interpolate(flat)。其余即定义是类型别名,只在翻译时展开。乃签名的类型是资源(存储缓冲、均匀缓冲、工作组变量、二维纹理、采样器):定义须是对「绑定」等的调用,译成@group(g) @binding(b) var<storage, read_write> 名: T;(只读为var<storage, read>,均匀缓冲为var<uniform>,纹理与采样器为var),工作组变量译成var<workgroup> 名: T;。乃签名的类型是计算、顶点或片元入口:定义须是对相应入口项的调用,所指函数加@compute @workgroup_size(n)、@vertex或@fragment。入口函数的参数是内建输入输出类型时加@builtin,是结构时不加,其余按次序编@location(片元的整数输入加@interpolate(flat));计算入口的参数只能是内建输入或由它们组成的结构。顶点、片元入口的返回类型是内建类型时写-> @builtin(…) T,是结构时写-> S,其余写-> @location(0) T。乃签名的类型是函数(化…而…):译成fn,参数与返回类型取自签名;类型为「有」的参数省去,返回「有」的函数没有返回类型。- 其余带签名的定义与不带签名的定义译成
const(值须是常量表达式)。 立、裸表达式、多态函数(承)与调用自身的顶层函数报错。
输出依源码次序,被导入的着色源在前;用到 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打包或拆开。
限制(待办事项):顶点与片元入口;纹理与采样器;「数组长度」(须由宿主另传);反三角函数(反正弦、反余弦、反正切、反正切二)与半精度的平滑阶跃;越界访问不检查。