着色语言
着色语言以豫言之常法书图形处理器之着色程序:着色之源经编译器之析文、解名、检类(佐以着色内建之声),乃译为 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 之低、高十六位依 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 常量之求不许非数与无穷:
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。结构之项书为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。金属一函至多三十一缓冲,故组号惟可用 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 占十六字节);函名与金属内建之函同者未改。
着十:统算之的(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 六十四位)。所生之模为每入口项另出“入口项名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 齐二,余标量齐四;二维向量之齐倍其元,三维、四维四之;矩阵依列向量;结构取项齐之最大;数组之步为元之大进至其齐。均匀缓冲亦依此布(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包之或析之。
所限(待办事项):顶点与片元之入口;纹理与采样器;「数组长度」(须宿主另传);反三角之函(反正弦、反余弦、反正切、反正切二)与半精之平滑阶跃;越界之访不察。