WebGPU 着色语言

1. 简介

WebGPU 着色语言（WGSL）是 [WebGPU] 的着色器语言。 也就是说，使用 WebGPU API 的应用程序使用 WGSL 来表达在 GPU 上运行的程序， 这些程序称为着色器。

// 一个使用点光源为带纹理几何体进行光照的片元着色器。

// 来自存储缓冲区绑定的光源。
struct PointLight {
  position : vec3f,
  color : vec3f,
}

struct LightStorage {
  pointCount : u32,
  point : array<PointLight>,
}
@group(0) @binding(0) var<storage> lights : LightStorage;

// 纹理和采样器。
@group(1) @binding(0) var baseColorSampler : sampler;
@group(1) @binding(1) var baseColorTexture : texture_2d<f32>;

// 函数参数是来自顶点着色器的值。
@fragment
fn fragmentMain(@location(0) worldPos : vec3f,
                @location(1) normal : vec3f,
                @location(2) uv : vec2f) -> @location(0) vec4f {
  // 从纹理中采样表面的基础颜色。
  let baseColor = textureSample(baseColorTexture, baseColorSampler, uv);

  let N = normalize(normal);
  var surfaceColor = vec3f(0);

  // 遍历场景中的点光源。
  for (var i = 0u; i < lights.pointCount; i++) {
    let worldToLight = lights.point[i].position - worldPos;
    let dist = length(worldToLight);
    let dir = normalize(worldToLight);

    // 确定此光源对表面颜色的贡献。
    let radiance = lights.point[i].color * (1 / pow(dist, 2));
    let nDotL = max(dot(N, dir), 0);

    // 将光源贡献累积到表面颜色。
    surfaceColor += baseColor.rgb * radiance * nDotL;
  }

  // 返回累积后的表面颜色。
  return vec4(surfaceColor, baseColor.a);
}

1.1. 概述

WebGPU 以 GPU 命令的形式向 GPU 发出一个工作单元。 WGSL 涉及两种 GPU 命令：

• 绘制命令 在输入、输出和所附加资源的上下文中执行渲染 管线。

• 分派 命令在输入和所附加资源的上下文中执行计算管线。

两种管线都使用以 WGSL 编写的着色器。

着色器是 WGSL 程序中在管线内执行着色器 阶段的部分。 一个着色器包括：

• 一个入口点函数。

• 从入口点开始的所有被调用函数的传递闭包。 此集合既包括用户定义函数，也包括内置函数。 （有关更严格的定义，请参阅“着色器阶段中的函数”。）

• 由所有这些函数静态访问的变量和常量集合。

• 用于定义或分析所有这些函数、变量和常量的类型集合。

注： WGSL 程序不要求具有入口点；但是，这样的 程序无法由 API 执行，因为创建 GPUProgrammableStage 时需要入口点。

执行着色器阶段时，实现：

• 计算在模块作用域声明的常量的值。

• 将资源绑定到 着色器的资源接口中的变量， 从而使这些资源的内容在着色器执行期间可供着色器使用。

• 为其他模块作用域变量分配内存， 并使用指定的初始值填充该内存。

• 如果入口点存在形式参数，则用着色器阶段的输入填充这些形式参数。

• 如果入口点存在返回值，则将其连接到着色器阶段的输出。

• 然后调用入口点。

WGSL 程序由以下内容组成：

• 指令，用于指定模块级行为控制。

• 函数，用于指定执行行为。

• 语句，它们是声明或可执行行为的单元。

• 字面量，它们是纯数学值的文本表示。

• 常量，每个常量都为在特定时间计算出的值提供一个名称。

• 变量，每个变量都为保存某个值的内存提供一个名称。

• 表达式，每个表达式都组合一组值以产生一个结果 值。

• 类型，每种类型描述：

  • 一组值。

  • 对所支持表达式的约束。

  • 这些表达式的语义。

• 属性，用于修改对象以指定额外信息，例如：

  • 指定入口 点的接口。

  • 指定诊断过滤器。

注： WGSL 程序目前由单个 WGSL 模块组成。

WGSL 是一种命令式语言：行为被指定为要执行的一系列语句。 语句可以：

• 声明常量或变量。

• 修改变量的内容。

• 使用结构化编程构造修改执行顺序：

  • 选择性执行：if（带有可选的 else if 和 else 子句）、switch。

  • 重复：loop、while、for。

  • 跳出嵌套执行构造：continue、break、break if。

  • 重构：函数调用和return。

• 对表达式求值，以作为上述行为的一部分计算值。

• 在着色器创建时检查假设，其对象为常量表达式。

WGSL 是静态类型语言：特定表达式计算出的每个值都具有特定类型， 该类型仅通过检查程序源代码即可确定。

WGSL 具有描述布尔值和数字 （整数和浮点数）的类型。 这些类型可以聚合成复合类型 （向量、矩阵、 数组和结构体）。 WGSL 具有提供独特操作的特殊类型（例如原子类型）。 WGSL 将可存储在内存中的类型描述为内存视图。 WGSL 以 纹理和采样器的形式提供常用的渲染类型。 这些类型具有相关的内置函数， 用于公开通常由 GPU 硬件提供的图形渲染功能。

WGSL 的隐式转换和提升非常有限。 它只为作为参数传递给函数调用的抽象 类型和缓冲区指针提供隐式转换和提升。 将值从一种具体数值或布尔类型转换为 另一种类型，需要显式转换、 值构造器或 位的重新解释；但是，WGSL 确实提供了一些有限机制，可将标量 类型提升为向量类型。 这也适用于复合类型。

一个着色器阶段的工作被划分为一个或多个调用， 每个调用都会执行入口点，但所处条件略有不同。 着色器阶段中的调用共享对某些变量的访问：

• 该阶段中的所有调用都共享着色器接口中的资源。

• 在计算 着色器中，同一 工作组中的调用共享 workgroup 地址空间中的变量。 不同工作组中的调用不共享这些变量。

但是，这些调用操作的是不同的着色器阶段输入集合，包括内置输入， 这些内置输入提供用于将某个调用与其同级调用区分开的标识值。 每个调用都有自己独立的内存空间，其形式为 private 和function 地址空间中的变量。

着色器阶段内的调用并发执行，并且通常可能并行执行。 着色器作者负责确保着色器阶段中调用的动态行为：

• 满足某些原始操作的一致性要求，包括 纹理采样和控制屏障。

• 协调对共享变量可能发生冲突的访问，以避免数据竞争。

WGSL 有时允许某个给定特性具有多种可能的行为。 这会造成可移植性风险，因为不同实现可能表现出不同的行为。 WGSL 的设计旨在尽量减少这种情况，但受到可行性的约束， 同时还要实现跨广泛设备获得高性能的目标。

行为 要求是实现处理或执行 WGSL 程序时将 执行的操作。它们描述了实现与程序员之间契约中 实现所承担的义务。 当这些义务可能并非显而易见时，本规范会明确说明这些义务。

1.2. 语法表示法

以下语法表示法描述 WGSL 句法语法的约定：

• 规则两侧的斜体文本表示一个语法规则。

• 规则右侧以单引号（'）开头和结尾的粗体等宽文本表示关键字和词法单元。

• 普通文本中的冒号（:）定义一个语法规则。

• 普通文本中的竖线（|）表示可选项。

• 普通文本中的问号（?）表示前面的关键字、词法单元、规则 或分组出现零次或一次（即可选）。

• 普通文本中的星号（*）表示前面的关键字、词法单元、规则或分组出现 零次或多次。

• 普通文本中的加号（+）表示前面的关键字、词法单元、规则或分组 出现一次或多次。

• 普通文本中一对匹配的左括号（(）和右括号 （)）表示一组元素。

1.3. 数学术语和表示法

角度：

• 按照约定，角度以弧度计量。

• 用于测量角度的参考射线是从原点 (0,0) 指向 (+∞,0) 的射线。

• 令 θ 为比较射线与参考射线之间所张的角。 当比较射线逆时针移动时，θ 增大。

• 一个完整圆周包含 2 π 弧度。

• 示例：

  • 角度 0 从原点向右指向 (1,0)

  • 角度 2π 从原点向右指向 (1,0)

  • 角度 π/4 从原点指向点 (1,1)

  • 角度 π/2 从原点指向点 (0,1)

  • 角度 π 从原点指向点 (-1,0)

  • 角度 (3/2)π 从原点指向点 (0,-1)

双曲角是 一个无量纲面积，而不是传统 意义上的角。具体而言：

• 考虑双曲线 x² - y² = 1，其中 x > 0。

• 令 R 为从原点到双曲线上某点 (x, y) 的射线。

• 令 a 为由 R、x 轴和双曲线 本身围成面积的两倍。

• 当 R 位于 x 轴上方时，将 a 视为正值，位于 下方时视为负值。

于是，面积 a 是一个双曲角，使得 x 是 a 的双曲余弦，而 y 是 a 的双曲正弦。

正无穷， 记作 +∞，是一个严格大于所有实数的唯一值。

负无穷， 记作 −∞，是一个严格小于所有实数的唯一值。

扩展实数 （也称为仿射扩展实数）是实数集合再加上 +∞ 和 −∞。 计算机使用浮点类型近似表示 扩展实数，包括两个无穷值。 参见§ 15.7 浮点求值。

区间是具有下界和上界的连续 数字集合。 根据上下文，它们可以是整数、浮点数、实数或扩展实数的集合。

• 闭区间 [a,b] 是满足 a ≤ x ≤ b 的数 x 的集合。

• 半开区间 [a,b) 是满足 a ≤ x < b 的数 x 的集合。

• 半开区间 (a,b] 是满足 a < x ≤ b 的数 x 的集合。

向下取整表达式 对扩展实数 x 定义如下：

• ⌊ +∞ ⌋ = +∞

• ⌊ −∞ ⌋ = −∞

• 对于实数 x，⌊x⌋ = k，其中 k 是满足 k ≤ x < k+1 的唯一整数

向上取整 表达式对扩展实数 x 定义如下：

• ⌈ +∞ ⌉ = +∞

• ⌈ −∞ ⌉ = −∞

• 对于实数 x，⌈x⌉ = k，其中 k 是满足 k-1 < x ≤ k 的唯一整数

truncate 函数 对扩展实数 x 定义如下：

• truncate(+∞) = +∞

• truncate(−∞) = −∞

• 对于实数 x，计算绝对值小于或 等于 x 绝对值的最近整数：

  • 如果 x ≥ 0，则 truncate(x) = ⌊x⌋；如果 x < 0，则 = ⌈x⌉。

roundUp 函数 对正整数 k 和 n 定义为：

• roundUp(k, n) = ⌈n ÷ k⌉ × k

roundDown 函数 对正整数 k 和 n 定义为：

• roundDown(k, n) = ⌊n ÷ k⌋ × k

c 列 r 行矩阵 A 的转置 是一个 r 列 c 行矩阵 A^T，通过将 A 的各行复制为 A^T 的各列形成：

• transpose(A) = A^T

• transpose(A)_i,j = A_j,i

列向量的转置通过将列向量解释为 1 行矩阵来定义。 类似地，行向量的转置通过将行向量解释为 1 列矩阵来定义。

2. WGSL 模块

WGSL 程序由单个 WGSL 模块组成。

模块是由可选指令序列，后跟模块作用域声明和断言组成的序列。 模块由以下内容组成：

• 指令，用于指定模块级行为控制。

• 函数，用于指定执行行为。

• 语句，它们是声明或可执行行为的单元。

• 字面量，它们是纯数学值的文本表示。

• 变量，每个变量都为保存某个值的内存提供一个名称。

• 常量，每个常量都为在特定时间计算出的值提供一个名称。

• 表达式，每个表达式都组合一组值以产生一个结果 值。

• 类型，每种类型描述：

  • 一组值。

  • 对所支持表达式的约束。

  • 这些表达式的语义。

• 属性，用于修改对象以指定额外信息，例如：

  • 指定入口 点的接口。

  • 指定诊断过滤器。

2.1. 着色器生命周期

WGSL 程序及其可能包含的着色器的生命周期中有四个关键事件。 前两个对应于用于准备 WGSL 程序以供执行的 WebGPU API 方法。 后两个是着色器执行的开始和结束。

这些事件是：

1. 着色器模块 创建

   • 当调用 WebGPU createShaderModule() 方法时会发生此事件。 此时会提供 WGSL 程序的源文本。

2. 管线 创建

   • 当调用 WebGPU createComputePipeline() 方法 或 WebGPU createRenderPipeline() 方法 时会发生此事件。 这些方法使用一个或多个先前创建的着色器模块以及其他 配置信息。

   • 在创建管线期间，仅考虑指定 GPUProgrammableStage 的指定入口点所构成的着色器代码。 也就是说，与该入口点无关的代码在编译前实际上会被移除。

   • 注： 每个着色器阶段被视为 单独编译，因此可能 包含模块的不同部分。

3. 着色器 执行开始

   • 当向 GPU 发出绘制或分派命令， 开始执行管线， 并调用着色器阶段入口点函数时，会发生此事件。

4. 着色器执行 结束

   • 当着色器中的所有工作完成时，会发生此事件：

     • 其所有调用都终止，并且

     • 对资源的所有访问均完成，并且

     • 如果存在输出，则将其传递给下游管线阶段。

这些事件按顺序发生，原因是：

• 数据依赖关系：着色器执行需要管线，而管线需要着色器模块。

• 因果关系：着色器必须先开始执行，之后才能结束执行。

2.2. 错误

WebGPU 实现可能因两个原因无法处理着色器：

• 如果着色器不满足 WGSL 或 WebGPU 规范的要求，则会发生程序错误。

• 即使满足了所有 WGSL 和 WebGPU 要求，也可能发生未分类 错误。 可能的原因包括：

  • 着色器过于复杂，超出了实现的能力， 但这种情况难以通过规定的限制加以描述。 简化着色器可能可以绕过该问题。

  • WebGPU 实现中的缺陷。

在着色器生命周期的三个阶段中可能发生处理错误：

• 着色器创建 错误 是在着色器模块创建时可以切实检测到的错误。 检测仅依赖 WGSL 模块源文本 以及 createShaderModule API 方法可获得的其他信息。 本规范中描述程序必须 执行某事的语句，在违反这些断言时通常会产生着色器创建错误。

• 管线创建错误 是在管线创建时可以检测到的错误。 检测依赖 WGSL 模块源文本 以及特定管线创建 API 方法可获得的其他信息。 这些错误只会由正在为 GPUProgrammableStage 编译的入口点的 着色器中存在的代码触发。

• 动态错误是 着色器执行期间发生的错误。 这些错误可能可以检测，也可能无法检测。

注： 例如，数据竞争可能无法检测。

每项要求将在最早的机会进行检查。 即：

• 如果未满足在着色器创建时可检测的要求，则产生着色器创建错误。

• 如果未满足在管线创建时可检测、但无法更早检测的要求，则产生管线创建错误。

当从上下文无法明确判断时，本规范会指出 未满足某项特定要求 会导致着色器创建、管线创建还是动态错误。

错误的后果如下：

• 具有着色器创建错误或管线创建错误 的 WGSL 模块不会被纳入 管线，因此 不会被执行。

• 可检测的错误将触发诊断。

• 如果发生动态错误：

  • 内存访问 将被限制为：

    • 着色器阶段输入，

    • 着色器阶段输出，

    • 绑定到 WGSL 模块中某个变量的资源的任何部分，以及

    • WGSL 模块中声明的其他变量。

  • 除此之外，程序可能不会按照本规范其余部分所描述的方式运行。 注：这些影响可能是非局部的。

2.3. 诊断

实现可以在着色器模块创建或管线创建期间生成诊断。 诊断是实现为帮助应用程序作者而生成的消息。

当满足某个特定条件时，会创建或触发一个诊断， 该条件称为触发规则。 源文本中满足该条件的位置，以源文本中的点或范围表示， 称为 触发位置。

诊断具有以下 属性：

• 一个严重程度。

• 一个触发规则。

• 一个触发位置。

诊断的严重程度是以下值之一，按从高到低排序：

错误

该诊断是错误。 这对应于着色器创建错误或管线创建 错误。

警告

该诊断描述值得应用程序开发者关注的异常，但 并非错误。

信息

该诊断描述值得应用程序开发者关注的显著情况，但 既不是错误也不是警告。

关闭

该诊断已禁用。它不会传递给应用程序。

触发规则的名称是以下两种形式之一：

• 一个diagnostic_name_token，或

• 两个diagnostic_name_token 词法单元，由句点 '.'（U+002E）分隔。

2.3.1. 诊断处理

被触发的诊断将按如下方式处理：

1. 对于每个诊断 D，找到包含 D 的触发 位置、具有最小受影响范围，且具有相同触发规则的诊断过滤器。

   • 如果存在这样的过滤器，则将其应用于 D，更新 D 的严重程度。

   • 否则，D 保持不变。

2. 丢弃严重程度为关闭的诊断。

3. 如果至少有一个剩余诊断 DI 的严重程度为信息，则：

   • 具有相同触发规则的其他信息 诊断可以被丢弃，只保留原始 诊断 DI。

4. 如果至少有一个剩余诊断 DW 的严重程度为警告，则：

   • 具有相同触发规则的其他信息或 警告 诊断可以被丢弃，只保留原始 诊断 DW。

5. 如果至少有一个剩余诊断具有错误严重程度，则：

   • 其他诊断可以被丢弃，包括其他具有错误 严重程度的诊断。

   • 生成程序错误。

     • 如果诊断是在着色器模块创建时触发的，则该错误是着色器创建错误。

     • 如果诊断是在管线创建时触发的，则该错误是管线创建错误。

6. 如果在着色器模块创建时进行处理，则剩余诊断 会填充 WebGPU GPUCompilationInfo 对象的 messages 成员。

7. 如果在管线创建期间进行处理，则错误诊断 会导致验证 GPUProgrammableStage 时 WebGPU 验证失败。

注： 这些规则允许实现一旦检测到错误就停止 处理 WGSL 模块。 此外，针对某种特定警告的分析可以在出现第一个警告时停止， 针对某种特定信息诊断的分析也可以在首次出现时停止。 WGSL 不规定执行不同种类分析的顺序，也不规定单次 分析内部的顺序。 因此，对于同一个 WGSL 模块，不同实现可能报告具有相同严重程度的不同诊断实例。

2.3.2. 可过滤的触发规则

大多数诊断都会无条件 报告给 WebGPU 应用程序。 某些类型的诊断可以过滤，其中一种方式是指定其触发规则的名称。 下表列出了可以过滤的标准触发规则集合。

使用由单个诊断名称词法单元组成的未识别触发规则时，用户代理应 触发警告。

实现可以支持此处未指定的触发规则， 前提是它们使用diagnostic_rule_name 的多词法单元形式拼写。 使用多词法单元形式拼写的未识别触发规则本身可以触发 诊断。

本规范的未来版本可以移除某个特定规则，或减弱其默认严重程度 （即用严重程度较低的默认值替换当前默认值），并仍被视为满足向后 兼容性。 例如，WGSL 的未来版本可以将derivative_uniformity 的默认严重程度从 error 更改为 warning 或 info。 在规范作出此类更改之后，先前有效的程序仍将保持有效。

2.3.3. 诊断过滤

一旦具有可过滤 触发规则的诊断被触发，WGSL 就提供机制来丢弃该 诊断，或修改其严重程度。

诊断过滤器 DF 有三个参数：

• AR：称为受影响范围的源文本范围

• NS：新的严重程度

• TR：一个触发规则

将诊断过滤器 DF(AR,NS,TR) 应用于诊断 D 会产生 以下效果：

• 如果 D 的触发位置位于 AR 中，且 D 的触发规则是 TR，则 将 D 的严重程度属性设置为 NS。

• 否则，D 保持不变。

范围诊断 过滤器是一种诊断过滤器，其受影响范围是指定的 源文本范围。 范围诊断过滤器被指定为受影响源范围开始处的 @diagnostic属性， 如下表所示。 @diagnostic 属性不得出现在任何其他位置。

注： 以下内容也属于复合语句： 函数体、case 子句、独立 default 子句、 while 和for 循环的循环体， 以及if_clause、else_if_clause 和else_clause 的主体。

var<private> d: f32;
fn helper() -> vec4<f32> {
  // 在 "if" 的主体中禁用 derivative_uniformity 诊断。
  //
  if (d < 0.5) @diagnostic(off,derivative_uniformity) {
    return textureSample(t,s,vec2(0,0));
  }
  return vec4(0.0);
}

全局 诊断过滤器可用于将诊断过滤器应用于整个 WGSL 模块。

diagnostic(off,derivative_uniformity);
var<private> d: f32;
fn helper() -> vec4<f32> {
  if (d < 0.5) {
    // 此处的 derivative_uniformity 诊断已由
    // 全局诊断过滤器禁用。
    return textureSample(t,s,vec2(0,0));
  } else {
    // derivative_uniformity 诊断被设置为 'warning' 严重程度。
    @diagnostic(warning,derivative_uniformity) {
      return textureSample(t,s,vec2(0,0));
    }
  }
  return vec4(0.0);
}

两个诊断过滤器 DF(AR1,NS1,TR1) 和 DF(AR2,NS2,TR2) 在以下情况下发生冲突：

• (AR1 = AR2)，并且

• (TR1 = TR2)，并且

• (NS1 ≠ NS2)。

诊断过滤器必须不冲突。

注： 当多个全局诊断过滤器 不发生冲突时，允许同时存在。

WGSL 的诊断过滤器设计为使其受影响范围完全嵌套。 如果 DF1 的受影响范围与 DF2 的受影响范围重叠，则 要么 DF1 的受影响范围完全包含在 DF2 的受影响范围内，要么反之。

对于源位置 L 和触发规则 TR，如果存在最近的外围诊断过滤器， 则它是诊断过滤器 DF(AR,NS,TR)，其中：

• L 落在受影响范围 AR 中，并且

• 如果存在另一个过滤器 DF'(AR',NS',TR)，且 L 落在 AR' 中，则 AR 包含在 AR' 中。

由于受影响范围是嵌套的，最近的外围诊断：

• 是唯一的范围诊断过滤器，

• 否则，是一组重复的（不冲突的）全局诊断 过滤器之一，

• 否则不存在。

2.4. 限制

WGSL 实现将支持满足以下限制的着色器。 WGSL 实现可以支持超出指定限制的着色器。

注： 如果 WGSL 实现不支持 超出指定限制的着色器，则应发出错误。

3. 文本结构

text/wgsl 媒体类型用于将内容标识为 WGSL 模块。 参见附录 A：text/wgsl 媒体类型。

WGSL 模块是使用 UTF-8 编码且不带字节顺序标记（BOM）的 Unicode 文本。

WGSL 模块文本由 Unicode 码点序列组成，这些码点被分组为连续的非空集合，形成：

• 注释

• 词法单元

• 空白

程序文本不得包含空码点（U+0000）。

3.1. 解析

要解析 WGSL 模块：

1. 移除注释：

   • 将第一个注释替换为空格码点（U+0020）。

   • 重复此操作，直到不再有注释。

2. 使用§ 3.9 模板列表中的算法查找模板列表。 此步骤用于消除 '<'（U+003C）和 '>'（U+003E）码点 作为模板列表分隔符与作为比较运算符等其他用途之间的歧义。

3. 解析整个文本，尝试匹配translation_unit 语法规则。 解析使用 LALR(1) 解析器（一个词法单元的前瞻）[DeRemer1969]，并进行以下 定制：

   • 词法单元化与解析交错进行，并且具有上下文感知能力。 当解析器请求下一个词法单元时：

     • 消耗并忽略起始的一系列空白码点。

     • 如果下一个码点是模板列表的开始，则消耗 它并返回_template_args_start。

     • 如果下一个码点是模板列表的结束，则消耗 它并返回_template_args_end。

     • 否则：

       • 词法单元候选是由剩余未消耗码点的非空前缀构成的任何 WGSL 词法单元。

       • 返回的词法单元是最长的词法单元 候选，且它同时是当前 解析器状态的有效前瞻词法单元。[VanWyk2007]

如果出现以下情况，则会产生着色器创建 错误：

• 整个源文本无法转换为有效词法单元的有限序列，或

• translation_unit 语法规则无法匹配 整个词法单元序列。

3.2. 空白和换行

空白是 Unicode Pattern_White_Space 属性中的一个或多个码点的任意组合。 以下是Pattern_White_Space 中的码点集合：

• 空格（U+0020）

• 水平制表符（U+0009）

• 换行符（U+000A）

• 垂直制表符（U+000B）

• 换页符（U+000C）

• 回车符（U+000D）

• 下一行（U+0085）

• 从左到右标记（U+200E）

• 从右到左标记（U+200F）

• 行分隔符（U+2028）

• 段落分隔符（U+2029）

换行是一段连续的 空白码点序列， 用于表示一行的结束。 它被定义为 UAX14 第 6.1 节不可定制的换行规则中的LB4 和LB5 所定义的表示“强制换行”的空白。 即，换行是以下任一种：

• 换行符（U+000A）

• 垂直制表符（U+000B）

• 换页符（U+000C）

• 后面没有紧跟换行符（U+000A）的回车符（U+000D）

• 后面紧跟换行符（U+000A）的回车符（U+000D）

• 下一行（U+0085）

• 行分隔符（U+2028）

• 段落分隔符（U+2029）

注： 以行号形式报告源文本位置的诊断应使用换 行来 计算行数。

3.3. 注释

注释是一段 不影响 WGSL 程序有效性或含义的文本，但注释可以分隔词法单元。 着色器作者可以使用注释记录其程序。

行尾 注释是一种注释， 由两个码点 //（U+002F 后跟 U+002F）以及之后的码点 组成，直到但不包括：

• 下一个换行，或

• 程序末尾。

块注释是一种 注释，由以下内容组成：

• 两个码点 /*（U+002F 后跟 U+002A）

• 然后是以下内容的任意序列：

  • 一个块注释， 或

  • 不包含 */（U+002A 后跟 U+002F） 或 /*（U+002F 后跟 U+002A）的文本

• 然后是两个码点 */（U+002A 后跟 U+002F）

注： 块注释可以嵌套。 由于块注释要求起始和结束文本序列匹配，并且允许任意嵌套， 因此不能使用正则表达式识别块注释。 这是正则语言抽引引理的结果。

const f = 1.5; // 这是行尾注释。
const g = 2.5; /* 这是一个块注释
                跨越多行。
                /* 块注释可以嵌套。
                 */
                但所有块注释都必须终止。
               */

3.4. 词法单元

词法单元是一段连续码点序列， 构成以下之一：

• 一个字面量。

• 一个关键字。

• 一个保留字。

• 一个语法词法单元。

• 一个标识符。

• 一个上下文相关名称。

3.5. 字面量

字面量是以下之一：

• 布尔 字面量：true 或 false。

• 数值 字面量：整数字面量或浮点字面量， 用于表示一个数字。

3.5.1. 布尔字面量

const a = true;
const b = false;

3.5.2. 数值字面量

数值 字面量的形式通过模式匹配定义。

整数字面量是：

• 以下任一形式指定的整数：

  • 0

  • 十进制数字序列，其中第一位不是 0。

  • 0x 或 0X 后跟十六进制数字序列。

• 然后是可选的 i 或 u 后缀。

注： 非零整数字面量的前导零（例如 012）是禁止的， 以避免与其他语言中“前导零表示八进制”的记法混淆。

const a = 1u;
const b = 123;
const c = 0;
const d = 0i;

const a = 0x123;
const b = 0X123u;
const c = 0x3f;

浮点 字面量是十进制浮点字面量 或十六进制浮点字面量。

浮点 字面量在逻辑上分为两部分：用于表示小数的有效数，以及可选的指数。 粗略地说，字面量的值等于有效数乘以某个底数的指定指数次幂。 如果某个有效数数字非零， 或其左侧和右侧都存在非零有效数数字，则该数字是有效的。 有效数字从左到右计数：第 N 个有效数字的左侧有 N-1 个 有效数字。

十进制 浮点字面量是：

• 一个有效数，由一系列数字指定，其中可以在某处包含一个可选的小数点（.）。 有效数表示十进制记法中的小数。

• 然后是一个可选的指数后缀，由以下内容组成：

  • e 或 E。

  • 然后是一个以十进制数指定的指数，可带可选的前导符号（+ 或 -）。

  • 然后是可选的 f 或 h 后缀。

• 小数点、指数或 f 或 h 后缀中至少有一个 必须 存在。 如果都不存在，则该词法单元是一个整数字面量。

const a = 0.e+4f;
const b = 01.;
const c = .01;
const d = 12.34;
const f = .0f;
const g = 0h;
const h = 1e-3;

注： 十进制有效数在 20 个 十进制数字后截断，在小数部分保留约 log(10)/log(2)×20 ≈ 66.4 个有效位。

十六进制浮点字面量是：

• 0x 或 0X 前缀

• 然后是一个有效数，由一系列十六进制数字指定，其中可以在某处包含一个可选的十六进制小数点 （.）。 该有效数表示十六进制记法中的小数。

• 然后是一个可选的指数后缀，由以下内容组成：

  • p 或 P

  • 然后是一个以十进制数指定的指数，可带可选的前导符号（+ 或 -）。

  • 然后是可选的 f 或 h 后缀。

• 十六进制小数点或指数中至少有一个必须存在。 如果两者都不存在，则该词法单元是一个整数字面量。

const a = 0xa.fp+2;
const b = 0x1P+4f;
const c = 0X.3;
const d = 0x3p+2h;
const e = 0X1.fp-4;
const f = 0x3.2p+2h;

注： 十六进制有效数在 16 个 十六进制数字后截断，在小数部分保留约 4 ×16 = 64 个有效位。

当数值字面量具有 后缀时，该字面量表示特定具体标量类型中的值。 否则，该字面量表示下面定义的抽象数值类型之一中的值。 无论哪种情况，字面量所表示的值都是其数学值按照 § 15.7.6 浮点转换中的规则转换为目标 类型后的值。

如果出现以下情况，则会产生着色器创建 错误：

• 带有 i 或 u 后缀的整数字面量 无法由目标类型表示。

• 带有 f 或 h 后缀的十六进制浮点字面量 溢出或无法由目标 类型精确表示。

• 带有 f 或 h 后缀的十进制浮点字面量 溢出目标类型。

• 在未启用f16 扩展时使用了带有 h 后缀的浮点字面量。

注： 十六进制浮点值 0x1.00000001p0 需要 33 个有效数位才能精确表示， 但f32 只有 23 个显式有效数位。

注： 如果要使用 f 后缀将 十六进制浮点字面量强制设为某种类型，则该字面量还必须 使用二进制指数。例如，写作 0x1p0f。相比之下，0x1f 是一个 十六进制整数字面量。

3.6. 关键字

关键字是一个词法单元，用于指代预定义的语言 概念。 WGSL 关键字列表参见§ 16.1 关键字摘要。

3.7. 标识符

标识符是一种 用作名称的词法单元。 参见§ 5 声明和作用域。

WGSL 使用两个语法非终结符来区分使用场景：

• ident 用于命名 已声明对象。

• member_ident 用于命名结构体类型的成员。

标识符的形式基于 Unicode 标准附录 #31中 Unicode 版本 14.0.0的规定， 并作如下补充。

标识符使用以下以UAX31 语法描述的配置：

<Identifier> := <Start> <Continue>* (<Medial> <Continue>+)*

<Start> := XID_Start + U+005F
<Continue> := <Start> + XID_Continue
<Medial> :=

这意味着包含非 ASCII 码点的以下标识符是 有效的：Δέλτα、réflexion、Кызыл、𐰓𐰏𐰇、 朝焼け、سلام、검정、שָׁלוֹם、गुलाबी、 փիրուզ。

但有以下例外：

• 标识符不得与关键字或保留字具有相同拼写。

• 标识符不得为 _（单个下划线， U+005F）。

• 标识符不得以 __（两个下划线， U+005F 后跟 U+005F）开头。

Unicode 版本 14.0.0 的 Unicode 字符数据库包含 一份非规范性列表， 其中列出了 XID_Start 和 XID_Continue 的所有有效码点。

注： 某些内置函数的返回类型是结构体类型，其名称不能在 WGSL 源代码中使用。 这些结构体类型被描述为仿佛是以两个下划线开头的名称预声明的一样。 结果值可以通过类型推断保存到新声明的 let 或 var 中， 或者立即按名称提取其中一个成员。有关使用示例，请参见 frexp 和 modf 的说明。

3.7.1. 标识符比较

两个 WGSL 标识符相同，当且仅当它们由完全相同的码点序列组成。

注： 本规范不允许为了比较而对值进行 Unicode 规范化。 视觉上和语义上相同但使用不同 Unicode 字符序列的值不会 匹配。 建议内容作者始终使用相同的编码序列，或者在选择值时避免可能 引发问题的字符。有关更多信息，请参见[CHARMOD-NORM]。

注： 当把某个标识符的所有实例替换为该标识符的某个同形异码字符后会改变 WGSL 模块的含义时，用户代理应发出开发者可见的 警告。 （同形异码字符是一段码点序列，对读者而言可能看起来与另一段码点 序列相同。 用于检测同形异码字符的映射示例包括上一段中提到的转换、映射和匹配算法。 如果可以通过 反复将子序列替换为其同形异码字符，使一个标识符转换成另一个序列，则这两个码点序列互为同形异码。）

3.8. 上下文相关名称

上下文相关 名称是一个词法单元，用于命名一个 概念，但仅用于特定的语法上下文中。 该词法单元的拼写可能与标识符相同，但该词法单元不会解析到已声明的对象。 本节列出了用作上下文相关名称的词法单元。 该词法单元不得是关键字或保留字。

3.8.1. 属性名称

参见§ 12 属性。

attribute 名称包括：

• 'align'

• 'binding'

• 'builtin'

• 'compute'

• 'const'

• 'diagnostic'

• 'fragment'

• 'group'

• 'id'

• 'interpolate'

• 'invariant'

• 'location'

• 'blend_src'

• 'must_use'

• 'size'

• 'vertex'

• 'workgroup_size'

• 'subgroup_size'

3.8.2. 内置值名称

内置值 名称词法单元是一个词法单元，用于 内置值的名称中。

参见§ 13.3.1.1 内置输入和输出。

内置值的名称 包括：

• 'vertex_index'

• 'instance_index'

• 'position'

• 'front_facing'

• 'frag_depth'

• 'sample_index'

• 'sample_mask'

• 'local_invocation_id'

• 'local_invocation_index'

• 'global_invocation_id'

• 'global_invocation_index'

• 'workgroup_id'

• 'workgroup_index'

• 'num_workgroups'

• 'subgroup_invocation_id'

• 'subgroup_size'

• 'primitive_index'

• 'subgroup_id'

• 'num_subgroups'

3.8.3. 诊断规则名称

诊断 名称词法单元是一个词法单元，用于诊断触发规则的名称中。

参见§ 2.3 诊断。

预定义的诊断规则名称包括：

• 'derivative_uniformity'

• 'subgroup_uniformity'

3.8.4. 诊断严重程度控制名称

有效的诊断 过滤器严重程度控制名称列于§ 2.3 诊断中，但其 形式与标识符相同：

诊断过滤器 严重程度控制名称包括：

• 'error'

• 'warning'

• 'info'

• 'off'

3.8.5. 扩展名称

有效的启用扩展 名称列于§ 4.1.1 启用扩展中，但通常具有与标识符相同的 形式：

启用扩展 名称包括：

• 'f16'

• 'clip_distances'

• 'dual_source_blending'

• 'subgroups'

• 'primitive_index'

• 'subgroup_size_control'

有效的语言 扩展名称列于§ 4.1.2 语言扩展中，但 通常具有与标识符相同的形式：

语言 扩展名称包括：

• 'readonly_and_readwrite_storage_textures'

• 'packed_4x8_integer_dot_product'

• 'unrestricted_pointer_parameters'

• 'pointer_composite_access'

• 'uniform_buffer_standard_layout'

• 'subgroup_id'

• 'subgroup_uniformity'

• 'texture_and_sampler_let'

• 'texture_formats_tier1'

• 'linear_indexing'

• 'immediate_address_space'

• 'fragment_depth'

• 'buffer_view'

• 'swizzle_assignment'

3.8.6. 插值类型名称

插值类型名称词法单元是一个词法单元，用于 interpolate_type_name 的插值类型 名称中。

参见§ 13.3.1.4 插值。

插值 类型名称包括：

• 'perspective'

• 'linear'

• 'flat'

3.8.7. 插值采样名称

插值采样名称词法单元是一个词法单元，用于插值 采样的名称中。

参见§ 13.3.1.4 插值。

插值 采样名称包括：

• 'center'

• 'centroid'

• 'sample'

• 'first'

• 'either'

3.8.8. 重排名称

重排名称用于向量访问表达式和重排视图 表达式中：

3.8.9. 深度模式名称

深度模式 名称词法单元是一个词法单元，在支持fragment_depth 特性时，用于 frag_depth 内置值中。

• 'less'

• 'greater'

3.9. 模板列表

模板 参数化是一种指定用于修改一般概念的参数的方式。 要编写模板参数化，先写一般概念，后跟一个模板列表。

忽略注释和空白时，模板列表是：

• 起始的 '<'（U+003C）码点，然后

• 一个或多个模板参数组成的、由逗号分隔的 列表，然后

• 一个可选的末尾逗号，然后

• 终止的 '>'（U+003E）码点。

模板 参数的形式由下面的模板列表发现算法隐式定义。 通常，它们是名称、表达式或类型。

注： 例如，短语 vec3<f32> 是一个模板参数化，其中 vec3 是被修改的一般概念， 而 <f32> 是包含一个参数，即 f32 类型的模板列表。 它们组合在一起，vec3<f32> 表示一个特定的向量类型。

注： 例如，短语 var<storage,read_write> 使用模板参数 storage 和 read_write 修改一般的 var 概念。

用于界定模板列表的 '<'（U+003C）和 '>'（U+003E）码点 也用于拼写：

• relational_expression 中的比较运算符。

• shift_expression 中的移位运算符。

• 用于执行移位操作后再赋值的compound_assignment_operator。

该语法歧义按有利于模板列表的方式解决：

• 模板列表在解析的早期阶段发现，早于声明、表达式、语句的解析。

• 在后续阶段进行词法单元化时， 模板列表起始的 '<'（U+003C）被映射到_template_args_start 词法单元，而 模板列表终止的 '>'（U+003E）被映射到_template_args_end 词法单元。

模板 列表发现算法如下所示。 它使用以下假设和性质：

1. 模板 参数是一个表达式，因此不会以 '<'（U+003C）或 '='（U+003D）码点开头。

2. 表达式不包含码点 ';'（U+003B）、'{'（U+007B）或 ':'（U+003A）。

3. 表达式不包含赋值。

4. '='（U+003D）码点出现的唯一情况是作为比较操作的一部分， 即 以下之一： '<='、 '>='、 '==' 或 '!='。 否则，'='（U+003D）码点会作为赋值的一部分出现。

5. 模板列表分隔符会遵循由括号 '(...)' 和数组索引 '[...]' 形成的嵌套表达式。 模板列表的开始和结束必须出现在同一嵌套层级。

算法： 模板列表发现

输入： 程序源文本。

记录类型：

令 UnclosedCandidate 为包含以下内容的记录类型：

• position，源文本中的一个位置

• depth，一个整数，即 position 处的表达式嵌套深度

令 TemplateList 为包含以下内容的记录类型：

• start_position，开始此模板列表的 '<'（U+003C）码点在源文本中的位置。

• end_position，结束此模板列表的 '>'（U+003E）码点在源文本中的位置。

输出： DiscoveredTemplateLists，一个 TemplateList 记录列表。

过程：

• 将 DiscoveredTemplateLists 初始化为空列表。

• 将 Pending 变量初始化为一个空的 UnclosedCandidate 记录栈。

• 将整数变量 CurrentPosition 初始化为 0。 它通过源文本开始处之后的码点数量，编码当前正在检查的码点的位置。

  • 执行算法时，该变量将沿文本向前推进。 到达文本末尾时，立即终止算法并返回 DiscoveredTemplateLists。

• 将整数变量 NestingDepth 初始化为 0。

• 重复以下步骤：

  • 推进 CurrentPosition，越过空白、注释和字面量。

  • 如果 ident_pattern_token 与 CurrentPosition 处的文本匹配，则：

    • 推进 CurrentPosition，越过ident_pattern_token。

    • 如果存在空白和注释，则推进 CurrentPosition 越过它们。

    • 如果 '<'（U+003C）出现在 CurrentPosition 处，则：

      • 注： 此码 点可能是模板列表的开始。 保存足够的状态，以便将其与稍后在输入中出现的终止 '>'（U+003E）相匹配。

      • 将 UnclosedCandidate(position=CurrentPosition,depth=NestingDepth) 压入 Pending 栈。

      • 将 CurrentPosition 推进到下一个码点。

      • 如果 '<'（U+003C）出现在 CurrentPosition 处，则：

        • 注： 根据 假设 1，没有模板参数以 '<'（U+003C）开头，因此前一个码点不可能 是模板列表的开始。 因此当前码点和前一个码点必须是'<<' 运算符。

        • 从 Pending 栈中弹出顶部条目。

        • 推进 CurrentPosition 越过此码点，并 开始循环的下一次迭代。

      • 如果 '='（U+003D）出现在 CurrentPosition 处，则：

        • 注： 根据 假设 1，没有模板参数以 '=' （U+003C）开头，因此前一个码点不可能是 模板列表的开始。 假定当前码点和前一个码点构成'<=' 比较运算符。 跳过 '='（U+003D）码点，以免后续 步骤将其误认为赋值。

        • 从 Pending 栈中弹出顶部条目。

        • 推进 CurrentPosition 越过此码点，并 开始循环的下一次迭代。

    • 开始循环的下一次迭代。

  • 如果 '>'（U+003E）出现在 CurrentPosition 处，则：

    • 注： 此码点可能是 模板列表的结束。

    • 如果 Pending 非空，则令 T 为其顶部条目，并且 如果 T.depth 等于 NestingDepth，则：

      • 注： 此码 点结束当前模板列表，其开始位置记录在 T 中。

      • 将 TemplateList(start_position=T.position, end_position=CurrentPosition) 添加到 DiscoveredTemplateLists。

      • 从 Pending 栈中弹出 T。

      • 推进 CurrentPosition 越过此码点，并开始 循环的下一次迭代。

    • 否则，此码点不会结束模板列表：

      • 推进 CurrentPosition 越过此码点。

      • 如果 '='（U+003D）出现在 CurrentPosition 处，则：

        • 注： 假定 当前码点和前一个码点构成'>=' 比较运算符。 跳过 '='（U+003D）码点，以免后续 步骤将其误认为赋值。

        • 推进 CurrentPosition 越过此码点。

      • 开始循环的下一次迭代。

  • 如果 '('（U+0028）或 '['（U+005B）出现在 CurrentPosition 处，则：

    • 注： 进入嵌套 表达式。

    • 将 NestingDepth 加 1。

    • 推进 CurrentPosition 越过此码点，并开始下一次 循环迭代。

  • 如果 ')'（U+0029）或 ']'（U+005D）出现在 CurrentPosition 处，则：

    • 注： 退出嵌套 表达式。

    • 从 Pending 栈中弹出条目，直到栈为空，或者直到其 顶部条目的 depth < NestingDepth。

    • 将 NestingDepth 设置为 0 或 NestingDepth − 1 中 较大的值。

    • 推进 CurrentPosition 越过此码点，并开始下一次 循环迭代。

  • 如果 '!'（U+0021）出现在 CurrentPosition 处，则：

    • 推进 CurrentPosition 越过此码点。

    • 如果 '='（U+003D）出现在 CurrentPosition 处，则：

      • 注： 假定 当前码点和前一个码点构成'!=' 比较 运算符。 跳过 '='（U+003D）码点，以免后续步骤 将其误认为赋值。

      • 推进 CurrentPosition 越过此码点。

    • 开始循环的下一次迭代。

  • 如果 '='（U+003D）出现在 CurrentPosition 处，则：

    • 推进 CurrentPosition 越过此码点。

    • 如果 '='（U+003D）出现在 CurrentPosition 处，则：

      • 注： 假定 当前码点和前一个码点构成'==' 比较运算符。 跳过 '='（U+003D）码点，以免后续步骤 将其误认为赋值。

      • 推进 CurrentPosition 越过此码点，并开始 循环的下一次迭代。

    • 注： 假定此码点 是赋值的一部分，而赋值不能作为表达式的一部分出现，因此 也不能出现在模板列表中。 清除待处理的未闭合候选项。

    • 将 NestingDepth 设置为 0。

    • 从 Pending 栈中移除所有条目。

    • 推进 CurrentPosition 越过此码点，并开始下一次 循环迭代。

  • 如果 ';'（U+003B）或 '{'（U+007B）或 ':'（U+003A） 出现在 CurrentPosition 处，则：

    • 注： 这些不能出现在 表达式中间，因此也不能出现在模板列表中。 清除待处理的未闭合候选项。

    • 将 NestingDepth 设置为 0。

    • 从 Pending 栈中移除所有条目。

    • 推进 CurrentPosition 越过此码点，并开始下一次 循环迭代。

  • 如果 '&&' 或'||' 与 CurrentPosition 处的文本匹配，则：

    • 注： 这些运算符 的优先级低于比较运算。拒绝当前表达式层级中所有待处理的未闭合 候选项。

    • 注： 依据此规则，程序片段 a<b || c>d 中不会找到 模板列表。 相反，它会被识别为两个比较表达式的短路析取。

    • 从 Pending 栈中弹出条目，直到栈为空，或者直到其 顶部条目的 depth < NestingDepth。

    • 推进 CurrentPosition 越过这两个码点，并开始下一次 循环迭代。

  • 推进 CurrentPosition 越过当前码点。

注： 该算法会显式跳过字面量， 因为某些数值字面量以字母结尾，例如 1.0f。 结尾的 f 不应被误认为是ident_pattern_token 的开始。

注： 在短语 A ( B < C, D > ( E ) ) 中，片段 < C, D > 是一个模板列表。

注： 该算法遵循表达式嵌套：特定 模板列表的开始和结束不能出现在不同的表达式嵌套层级。 例如，在 array<i32,select(2,3,a>b)> 中，模板列表有三个参数，其中 最后一个是 select(2,3,a>b)。 a>b 中的 '>' 不会终止模板列表，因为它被包含在 调用 select 函数的表达式的括号部分中。

注： 模板列表的两端必须出现在 同一个索引 表达式中。例如 a[b<d]>() 不包含有效的模板列表。

注： 在短语 A<B<<C> 中， 短语 B<<C 被解析为 B，后跟左移运算符'<<'，再后跟 C。 模板发现算法开始检查 B，然后是 '<'（U+003C），但随后 发现下一个 '<'（U+003C）码点不能开始模板参数，因此 紧接在 B 后面的 '<' 不是模板列表的开始。 最初的 '<' 和最后的 '>' 是唯一的模板列表分隔符，并且它具有 模板参数 B<<C。

注： 短语 A<B<=C> 的分析 与上一条注释类似，因此短语 B<=C 被解析为 B，后跟 小于或等于运算符'<='，再后跟 C。 模板发现算法开始检查 B，然后是 '<'（U+003C），但随后 发现下一个 '='（U+003D）码点不能开始模板参数，因此 紧接在 B 后面的 '<' 不是模板列表的开始。 最初的 '<' 和最后的 '>' 是唯一的模板列表分隔符，并且它具有 模板参数 B<=C。

注： 检查短语 A<(B>=C)> 时，存在一个模板列表，它从第一个 '<'（U+003C） 码点开始，到最后一个 '>'（U+003E）码点结束，并具有模板参数 B>=C。 检查第一个 '>'（U+003C）码点（位于 B 之后）后，需要特别识别 '='（U+003D）码点，以免将其假定为赋值的一部分。

注： 检查短语 A<(B!=C)> 时，存在一个模板列表，它从第一个 '<''（U+003C） 码点开始，到最后一个 '>'（U+003E）码点结束，并具有模板参数 B!=C。 检查 '!'（U+0021）码点（位于 'B' 之后）后，需要特别识别 '=' （U+003D）码点，以免将其假定为赋值的一部分。

注： 检查短语 A<(B==C)> 时，存在一个模板列表，它从第一个 '<'（U+003C） 码点开始，到最后一个 '>'（U+003E）码点结束，并具有模板参数 B==C。 检查第一个 '='（U+003D）码点（位于 'B' 之后）后，需要特别识别第二个 '='（U+003D）码点，以免将它们中的任何一个假定为 赋值的一部分。

模板 列表发现完成后， 解析将尝试将每个模板列表与template_list 语法规则匹配。

4. 指令

指令是一个词法单元序列，用于修改 WebGPU 实现处理 WGSL 程序的方式。

指令是可选的。 如果存在，所有指令必须出现在任何声明或const 断言之前。

4.1. 扩展

WGSL 预计会随着时间推移而演进。

扩展是对 WGSL 规范的一组 具有名称且相互协调的 修改集合，由以下内容的任意组合组成：

• 通过新语法添加新的概念和行为，包括：

  • 声明、语句、属性和内置函数。

• 移除当前规范或先前发布的扩展中的限制。

• 用于缩小允许行为集合的语法。

• 用于限制程序某一部分可用特性的语法。

• 描述扩展如何与现有规范交互，以及可选地如何与 其他扩展交互。

假设而言，扩展可以：

• 添加数值标量类型，例如不同位宽的整数。

• 添加用于约束浮点舍入模式的语法。

• 添加用于表明着色器不使用原子类型的语法。

• 添加新的语句类型。

• 添加新的内置函数。

• 添加用于约束着色器调用执行方式的语法。

• 添加新的着色器阶段。

扩展有两种：启用扩展和语言扩展。

4.1.1. 启用扩展

启用扩展是 一个扩展，其功能仅在满足以下条件时 可用：

• 实现支持它，并且

• 着色器通过enable 指令显式请求它，并且

• 创建 GPUDevice 时请求的必需特性之一，是相应的 WebGPU GPUFeatureName。

启用扩展旨在 公开并非普遍受支持的硬件功能。

enable 指令是 一种指令，用于开启一个或多个启用扩展的支持。 如果实现不支持列出的所有启用扩展，则会产生着色器创建 错误。

与其他指令一样，如果存在enable 指令，它必须出现在所有声明和const 断言之前。 扩展名称不是标识符： 它们不会解析到声明。

有效的启用扩展列于下表。

// 启用用于任意精度浮点类型的假设扩展。
enable arbitrary_precision_float;
enable arbitrary_precision_float; // 冗余的 enable 指令也是允许的。

// 启用用于控制舍入模式的假设扩展。
enable rounding_mode;

// 假定 arbitrary_precision_float 启用以下内容：
//    - 类型 f<E,M>
//    - 作为函数返回值、形式参数和 let 声明中的类型
//    - 作为从 AbstractFloat 构造值的值构造器
//    - 除法运算符的操作数：/
// 假定 @rounding_mode 属性由 rounding_mode enable 指令启用。
@rounding_mode(round_to_even)
fn halve_it(x : f<8, 7>) -> f<8, 7> {
  let two = f<8, 7>(2);
  return x / 2; // 使用舍入到偶数的舍入模式。
}

4.1.2. 语言扩展

语言扩展 是一种扩展，只要实现支持它，就会自动 可用。 程序无需显式请求它。

语言扩展 包含理论上任何 WebGPU 实现都可以支持的功能。 如果某个特性尚未得到普遍支持，那是因为某些 WebGPU 实现尚未 实现它。

注： 例如，do-while 循环可以作为语言 扩展。

WebGPU GPU 对象的 wgslLanguageFeatures 成员列出实现所支持的 语言扩展 集合。

requires 指令 是一种指令，用于记录 程序对一个或多个语言扩展的使用。 它不会改变实现公开的功能。 如果实现不支持其中一个必需扩展，则会产生着色器创建 错误。

WGSL 模块可以使用requires 指令表明可能存在不可移植性， 并表明预期的最低可移植性要求。

注： WebGPU 实现之外的工具可以 检查程序使用的所有语言扩展是否都由程序中的requires 指令覆盖。

与其他指令一样，如果存在requires 指令，它必须出现在所有声明和const 断言之前。 扩展名称不是标识符： 它们不会解析到声明。

注： 其意图是，随着时间推移，WGSL 将定义 包含当时普遍支持的语言扩展中全部功能的语言扩展。 在requires 指令中，它们可以作为列出所有 这些常见特性的简写。 它们代表逐步增加的功能集合，可以被视为某种意义上的语言版本。

4.2. 全局诊断过滤器

全局诊断 过滤器是一种诊断过滤器，其受影响范围是整个 WGSL 模块。 它是一种指令，因此出现在 任何模块作用域 声明之前。 它的拼写方式类似属性形式，但没有前导的 @（U+0040）码点，并以 分号终止。

5. 声明和作用域

声明将一个 标识符与以下 类型的对象之一关联：

• 一个类型

• 一个类型生成器

• 一个值

• 一个变量

• 一个函数

• 一个形式 参数

• 一个枚举项

换句话说，声明为对象引入一个名称。

如果声明出现在程序源代码中，但位于任何其他声明的文本之外，则该声明位于模块 作用域。

函数声明 出现在模块作用域。 如果函数声明具有形式参数，则包含这些形式参数的声明， 并且它可以在其主体内包含变量和值声明。 因此，这些包含的声明不位于模块作用域。

注： 唯一可以包含另一个 声明的声明类型是函数声明。

某些对象由 WebGPU 实现提供，并且被视为 仿佛在 WGSL 模块源代码开始之前就已声明。 我们称这些对象为预声明对象。 例如，WGSL 预声明：

• 内置 函数，

• 内置类型，例如 i32 和f32，

• 内置类型生成器，例如 array、ptr 和 texture_2d，以及

• 枚举项，例如read_write、 workgroup 和rgba8unorm。

声明的作用域是 已声明标识符可能表示其关联对象的 程序源位置集合。 我们称该标识符在这些源位置处位于 作用域内 （相对于该声明）。

声明出现的位置决定其作用域：

• 预声明对象以及在模块作用域声明的对象，在整个程序源代码中都位于作用域内。

• 用户声明函数的每个形式 参数在相应的函数主体范围内都位于作用域内。 参见§ 11.1 声明用户定义的函数。

• 否则，作用域是一段紧接声明结束之后开始的文本。 有关详细信息，请参见§ 7 变量和值声明。

同一 WGSL 源程序中的两个声明不得同时：

• 引入相同的标识符名称，并且

• 具有相同的作用域结束位置。

注： 预声明对象在 WGSL 源代码中没有声明。 因此，用户在模块作用域或函数内部指定的声明可以与预声明 对象具有相同的名称。

标识符按如下方式使用，并由语法上下文区分：

• 匹配ident 语法元素的词法单元：

  • 用于声明中，作为正在声明的对象的名称，或

  • 用作名称，表示在其他位置声明的对象。这是常见情况。

• 匹配member_ident 语法元素的词法单元：

  • 用于结构体 类型声明中， 作为成员的名称，或

  • 用作名称，表示结构体值的成员，或表示对结构体成员的引用。 参见§ 8.5.4 结构体访问表达式。

当ident 词法单元作为 表示在其他位置声明的对象的名称出现时， 对于某个声明，它必须位于作用域内。 标识符词法单元表示的对象按如下方式确定：

• 如果该词法单元至少位于一个非模块作用域声明的作用域内， 则该词法单元表示这些声明中最近声明所关联的对象。

  注： 最近的此类声明出现在 标识符词法单元之前。

• 否则，如果存在具有该名称的模块作用域声明，则该词法单元 表示该已声明对象。

  注： 模块作用域声明可以出现在 标识符词法单元之前或之后。

• 否则，如果存在具有该名称的预声明对象，则该词法单元表示该 对象。

当使用上述算法将标识符映射到声明时，我们 称该标识符解析 到该声明。 类似地，我们也称该标识符解析到已声明对象。

如果任何模块作用域 声明是递归的，则会产生着色器创建错误。 也就是说，声明之间不能存在循环：

考虑以下有向图：

• 每个节点对应一个声明 D。

• 当 D 的定义中提到一个解析到 T 的标识符时， 从声明 D 到声明 T 存在一条边。

该图不得包含环。

注： 函数主体是函数声明的一部分，因此 函数不得直接或间接递归。

注： 非模块作用域标识符声明必须在文本中先于其使用位置。

// 有效，用户定义的变量可以与内置函数具有相同名称。
var<private> modf: f32 = 0.0;

// 有效，foo_1 在整个程序中都位于作用域内。
var<private> foo: f32 = 0.0; // foo_1

// 有效，bar_1 在整个程序中都位于作用域内。
var<private> bar: u32 = 0u; // bar_1

// 有效，my_func_1 在整个程序中都位于作用域内。
// 有效，foo_2 的作用域持续到函数结束。
fn my_func(foo: f32) { // my_func_1, foo_2
  // 对 'foo' 的任何引用都会解析到函数参数。

  // 无效，modf 解析到模块作用域变量。
  let res = modf(foo);

  // 无效，foo_2 的作用域在函数结束处结束。
  var foo: f32; // foo_3

  // 有效，bar_2 的作用域持续到函数结束。
  var bar: u32; // bar_2
  // 对 'bar' 的引用解析到 bar_2
  {
    // 有效，foo_4 的作用域持续到复合语句结束。
    var foo : f32; // foo_4

    // 有效，bar_3 的作用域持续到复合语句结束。
    var bar: u32; // bar_3
    // 对 'bar' 的引用解析到 bar_3

    // 无效，bar_4 与 bar_3 具有相同的作用域结束位置。
    var bar: i32; // bar_4

    // 有效，i_1 的作用域持续到 for 循环结束
    for ( var i: i32 = 0; i < 10; i++ ) { // i_1
      // 无效，i_2 与 i_1 具有相同的作用域结束位置。
      var i: i32 = 1; // i_2.
    }
  }

  // 无效，bar_5 与 bar_2 具有相同的作用域结束位置。
  var bar: u32; // bar_5

  // 有效，later_def 是模块作用域声明，在整个程序中都位于作用域内。
  var early_use : i32 = later_def;
}

// 无效，bar_6 与 bar_1 具有相同的作用域。
var<private> bar: u32 = 1u; // bar_6

// 无效，my_func_2 与 my_func_1 具有相同的作用域结束位置。
fn my_func() { } // my_func_2

// 有效，my_foo_1 在整个程序中都位于作用域内。
fn my_foo( //my_foo_1
  // 有效，my_foo_2 的作用域持续到函数结束。
  my_foo: i32 // my_foo_2
) { }

var<private> later_def : i32 = 1;

// 此声明遮蔽预声明的 'min' 内置函数。
// 由于此声明位于模块作用域，因此它在整个
// 源代码中都位于作用域内。内置函数不再可访问。
fn min() -> u32 { return 0; }

const rgba8unorm = 12; // 此声明遮蔽预声明的 'rgba8unorm' 枚举项。

6. 类型

程序计算值。

在 WGSL 中，类型是一组 值，并且每个值恰好属于一个类型。 值的类型决定可以对该值执行的操作的语法和语义。

例如，数学数字 1 对应于 WGSL 中以下不同的值：

• 32 位有符号整数值 1i，

• 32 位无符号整数值 1u，

• 32 位浮点值 1.0f，

• 如果启用了f16 扩展，则为 16 位浮点值 1.0h，

• AbstractInt 值 1， 以及

• AbstractFloat 值 1.0

WGSL 将这些视为不同的值，因为它们的机器表示和操作不同。

某些类型是不透明的，这意味着 其值具有由实现定义的表示， 无法在 WGSL 程序中直接检查、分解或操作。

类型要么是预声明的，要么 通过 WGSL 源代码中的声明创建。

某些类型表示为模板参数化。 类型生成器是一个 预声明对象，当使用模板 列表进行参数化时，它表示一个类型。 例如，类型 atomic<u32> 将类型生成器 atomic 与 模板列表 <u32> 组合。

我们区分类型的概念与在 WGSL 中表示该类型的语法。 在许多情况下，本规范中类型的拼写与其 WGSL 语法相同。 例如：

• 32 位无符号整数值集合在本规范中写作 u32， 在 WGSL 模块中也是如此。

• 对于结构体类型或包含结构体的类型，其拼写方式不同。

某些 WGSL 类型仅用于分析源程序以及 确定程序的运行时行为。 本规范会描述这些类型，但它们不会出现在 WGSL 源文本中。

注： 引用类型不会写入 WGSL 模块中。参见§ 6.5.3 引用和指针类型。

注： 重排视图类型不会写入 WGSL 模块中。参见§ 6.5.4 重排视图类型。

6.1. 类型检查

WGSL 值通过对表达式求值来计算。 表达式是一段 源文本， 被解析为名称以 "expression" 结尾的 WGSL 语法规则之一。 表达式 E 可以包含子表达式，这些表达式被适当地包含在 外层表达式 E 中。 顶层 表达式是自身并非子表达式的表达式。 参见§ 8.19 表达式语法摘要。

表达式求值所产生的具体值取决于：

• 静态上下文： 表达式周围的源文本，以及

• 动态 上下文： 对表达式求值的调用的状态， 以及该调用正在运行的执行上下文。

对特定表达式求值可能得到的值始终属于某个特定的 WGSL 类型， 称为该表达式的静态类型。 WGSL 的规则被设计为使表达式的静态类型仅取决于 表达式的静态 上下文。

类型断言是 从某个 WGSL 源表达式到 WGSL 类型的映射。 表示法

e : T

是一个类型断言，表示 T 是 WGSL 表达式 e 的静态类型。

注： 类型断言是关于程序 文本的事实陈述。 它不是运行时检查。

语句通常使用表达式，并且可能对这些表达式的静态类型提出要求。 例如：

• if 语句的条件表达式必须为bool 类型。

• 在显式指定类型的 let 声明中，初始化器 表达式必须求值得到该类型。

对成功解析的 WGSL 模块进行类型检查，是将 每个表达式映射到其静态类型， 并验证每个语句的类型要求是否满足的过程。 如果类型检查失败，则会产生一种特殊的着色器创建错误，称为类型错误。

类型检查可以通过递归地将类型规则 应用于语法短语来执行，其中语法短语是表达式或语句。 类型规则描述 语法短语的静态上下文 如何确定该短语中所含表达式的静态类型。 类型规则有两个部分：

• 一个结论。

  • 如果短语是表达式，则结论是该 表达式的类型断言。

  • 如果短语是语句，则结论是一组类型断言， 该语句的每个顶层表达式各有一个。

  • 在这两种情况下，语法短语都以示意方式指定， 使用斜体名称表示子表达式 或其他由语法确定的参数。

• 前置条件，包括：

  • 对于表达式：

    • 如果存在子表达式，则为子表达式提供类型断言。 每个断言都可以直接满足，或者通过可行的自动转换满足（定义见§ 6.1.2 转换等级）。 WGSL将对确定程序静态类型所需的所有const 表达式求值。

    • 该表达式在语句中的使用方式。

  • 对于语句：

    • 语句的语法形式，以及

    • 语句中顶层表达式的类型断言。

  • 对其他示意参数的条件（如果有）。

  • 可选地，还包括其他静态上下文。

类型规则可以在其前置条件和结论中具有类型 参数。 当类型规则的结论或前置条件包含类型参数时， 我们称它是参数化的。 如果不包含，则称该规则是完全具体化的。 我们可以通过为参数化类型规则的每个类型参数替换一个类型， 并对规则中给定参数的所有出现位置使用同一类型， 从而得到完全具体化的类型规则。 将类型分配给规则类型参数的过程 称为替换。

例如，以下是逻辑非的类型规则 （形式为 !e 的表达式）：

这是一个参数化规则，因为它包含类型参数 T， 该参数可以表示以下四种类型中的任一种： bool、vec2<bool>、 vec3<bool> 或 vec4<bool>。 应用将 T 映射到 vec3<bool> 的替换 会产生以下完全具体化的类型规则：

通过应用某种满足规则其他条件的替换， 从参数化规则能够产生的每条完全具体化规则 称为该参数化规则的一个重载。 例如，布尔非规则有四个重载， 因为可以用四种方式之一为其类型参数 T 分配类型。

注： 换句话说，参数化类型规则为 一组完全具体化的类型规则提供模式， 每个规则都是通过对参数化规则应用不同替换得到的。

当满足以下条件时，类型规则适用于某个语法短语：

• 规则的结论与语法短语的一种有效解析相匹配，并且

• 规则的前置条件得到满足。

如果存在一个替换， 产生适用于该表达式的完全 具体化类型规则，则参数化类型规则 适用于该表达式。

考虑表达式 1u+2u。 它有两个字面量子表达式：1u 和 2u， 两者的类型都是 u32。 顶层 表达式是加法。 参照§ 8.8 算术表达式中的规则，加法的类型规则 适用于该表达式，因为：

• 1u+2u 与形式为 e1+e2 的解析相匹配，其中 e1 表示 1u，e2 表示 2u，并且

• e1 的类型为 u32，并且

• e2 的类型为 u32，并且

• 我们可以用 u32 替换类型规则中的类型参数 T， 从而得到适用于整个表达式的完全 具体化规则。

分析语法 短语时，可能出现三种情况：

• 没有类型规则适用于该表达式。这会导致类型错误。

• 恰好有一条完全 具体化类型规则适用于该表达式。 在这种情况下，会断言该规则的结论，从而确定 该表达式的静态类型。

• 有多条类型规则适用。也就是说，多个重载的前置条件得到满足。 在这种情况下，使用§ 6.1.3 重载解析中描述的决胜过程。

  • 如果重载解析成功，则确定单个重载适用于该表达式。 会断言该重载结论中的类型 断言， 因而确定语法 短语中一个或多个表达式的类型。

  • 如果重载解析失败，则会产生类型错误。

继续上面的示例，只有一条类型规则适用于表达式 1u+2u，因此类型 检查 接受该类型规则的结论，即 1u+2u 的类型为 u32。

当满足以下条件时，WGSL 源程序是类型良好的：

• 可以通过应用类型规则确定程序中每个表达式的静态类型，并且

• 每个语句的类型要求都得到满足。

否则会产生类型错误， 并且该源程序不是有效的 WGSL 模块。

WGSL 是一种静态类型语言， 因为对 WGSL 模块进行类型检查时，仅需检查程序源文本，就将要么成功，要么 发现类型错误。

6.1.1. 类型规则表

用于表达式的 WGSL 类型规则组织在类型规则 表中， 每条类型规则占一行。

表达式的 语义是对该表达式求值所产生的效果， 其主要效果是产生一个结果值。 适用于某个表达式的类型规则中的描述列将指定该表达式的 语义。 该语义通常取决于类型规则参数的值，包括 任何子表达式的假定值。 有时表达式的语义还包括产生 结果值以外的效果，例如其子表达式的非结果值效果。

fn foo(p : ptr<function, i32>) -> i32 {
  let x = *p;
  *p += 1;
  return x;
}

fn bar() {
  var a: i32;
  let x = foo(&a); // 对 foo 的调用返回一个值
                   // 并更新 a 的值
}

6.1.2. 转换等级

当类型断言 e:T 用作类型规则前置条件时，在以下情况下满足该断言：

• e 已经是 T 类型，或

• e 是 S 类型，并且类型 S 可以按照下面的定义自动 转换为类型 T。

该规则通过下面表中定义的、作用于类型对的ConversionRank 函数进行形式化。 ConversionRank 函数 表达将一种类型（Src）的值自动转换为 另一种类型（Dest）时的偏好程度和可行性。 等级越低越优先。

可行的 自动转换将值从类型 Src 转换为类型 Dest，并且在 ConversionRank(Src,Dest) 为有限值时允许进行。 此类转换会保留值，但受§ 15.7 浮点求值中所述限制的约束。

注： 自动转换只发生在两种 情况下。 第一种，是将const 表达式转换为其对应的、可 在 GPU 上使用的有类型数值。 第二种，是从内存引用加载时，得到存储在该内存中的值。

注： 转换等级为无穷大的转换不可行，即 不允许进行。

注： 不执行转换时，转换 等级为零。

1. 这些转换要求unrestricted_pointer_parameters 和buffer_view 语言特性。

如果满足以下条件，则类型 T 是类型 S 的具体化：

• T 是具体类型，并且

• T 不是引用类型，并且

• T 不是重排视图类型，并且

• ConversionRank(S, T) 为有限值，并且

• 对于满足上述条件的任何其他类型 T2，ConversionRank(S, T2) > ConversionRank(S, T)。

类型为 T 的值 e 的具体化 值，是将把 T 映射到 T 的具体化类型的可行转换 应用于 e 后得到的值。

注： 转换为 f32 始终比转换为 f16 更优先，因此 自动转换只会在模块中启用了 f16 扩展时产生 f16 值。

6.1.3. 重载解析

当有多个类型规则适用于某个语法短语时，会使用 决胜过程 来确定应采用哪一个。 此过程称为重载解析， 并假定类型检查已经成功找到了子表达式的静态类型。

考虑一个语法 短语 P，以及所有适用于 P 的类型规则。 重载解析算法将这些类型规则称为重载候选。 对于每个候选：

• 其前置条件已经直接满足，或通过自动 转换满足。

• 其结论具有：

  • 与 P 的有效解析匹配的语法形式，以及

  • 与 P 中每个顶层表达式相对应的类型 断言。

对 P 的重载解析按如下步骤进行，目标是找到唯一一个最优先的重载 候选：

1. 对于每个候选 C，枚举语法短语中各子表达式的转换等级。 候选的前置条件已经满足，因此对于 P 中第 i 个子表达式：

   • 其静态类型已经计算出来。

   • 存在从该表达式的静态类型到前置条件中相应类型断言所要求类型的可行自动转换。 令 C.R(i) 为该转换的ConversionRank。

2. 如果某个候选的某个子表达式在可行自动转换后解析为 抽象类型， 但该候选的另一个子表达式不是const 表达式，则淘汰该候选。

   注： 因此，如果短语中的任何子表达式 不是const 表达式， 那么该短语中的所有子表达式都必须具有具体类型。

3. 对候选排序：给定两个重载候选 C1 和 C2，如果满足以下条件，则 C1 比 C2 更优先：

   • 对于 P 中每个表达式位置 i，C1.R(i) ≤ C2.R(i)。

     • 也就是说，将 C1 应用于 P 所需的每个表达式转换 至少与将 C2 应用于 P 所需的相应表达式转换一样优先。

   • 至少存在一个表达式位置 i，使得 C1.R(i) < C2.R(i)。

     • 也就是说，至少存在一个应用 C1 所需的表达式转换， 严格优先于应用 C2 所需的相应转换。

4. 如果存在唯一候选 C，它比所有其他候选都更优先，则重载解析成功，并得到候选类型规则 C。 否则，重载解析失败。

6.2. 普通类型

普通类型是用于布尔值、数字、向量、 矩阵或这些值的聚合的机器 表示类型。

普通类型是标量类型、原子 类型或复合类型中的一种。

注： WGSL 中的普通类型类似于 C++ 中的 Plain-Old-Data 类型，但还 包括原子类型和抽象数值类型。

6.2.1. 抽象数值类型

这些类型无法在 WGSL 源代码中写出。它们仅用于类型检查。

某些表达式在着色器创建时求值， 并且其数值范围和精度可能大于 GPU 直接实现的范围和精度。

WGSL 为这些求值定义了两种抽象数值类型：

• AbstractInt 类型是可用 64 位二进制补码格式表示的整数集合， 其中符号位位于最高有效位位置。

• AbstractFloat 类型是可用 IEEE-754 binary64（双精度）格式表示的有限浮点数集合。

在这些类型之一中对表达式进行求值不得溢出，也不得产生无穷值或 NaN 值。

如果一个类型是抽象 数值类型，或包含抽象数值类型，则该类型是抽象的。 如果一个类型不是抽象的，则它是具体的。

没有 后缀的数值字面量表示抽象数值类型中的值：

• 没有 i 或 u 后缀的整数字面量 表示一个AbstractInt 值。

• 没有 f 或 h 后缀的浮点字面量表示一个AbstractFloat 值。

示例：表达式 log2(32) 按如下方式分析：

• log2(32) 被解析为对 log2 内置函数的函数调用，其 操作数为AbstractInt 值 32。

• log2 不存在具有整数标量形式参数的重载。

• 相反，会应用重载 解析，考虑三个可能的重载和可行 自动转换：

  • AbstractInt 到AbstractFloat。（转换等级 4）

  • AbstractInt 到f32。（转换等级 5）

  • AbstractInt 到f16。（转换等级 6）

• 最终计算以AbstractFloat 形式进行（例如 log2(32.0)）。

示例：表达式 1 + 2.5 按如下方式分析：

• 1 + 2.5 被解析为加法操作，其子表达式分别为AbstractInt 值 1 和AbstractFloat 值 2.5。

• 不存在 e+f 的重载，其中 e 是整数类型而 f 是浮点类型。

• 但是，使用可行自动转换后，存在三个潜在重载：

  • 1 转换为AbstractFloat 值 1.0（等级 4），而 2.5 保持为AbstractFloat（等级 0）。

  • 1 转换为f32 值 1.0f（等级 5），而 2.5 转换为f32 值 2.5f（等级 1）。

  • 1 转换为f16 值 1.0f（等级 6），而 2.5 转换为f16 值 2.5h（等级 2）。

• 第一个重载是更优先的候选，类型检查成功。

• 最终计算以AbstractFloat 形式 1.0 + 2.5 进行。

示例：let x = 1 + 2.5;

• 此示例与上面的示例类似，不同之处在于 x 不能解析为抽象数值类型。

• 因此，有两个可行的重载候选：使用 f32 或f16 的加法。

• 更优先的候选使用f32。

• 该声明的效果如同写成 let x : f32 = 1.0f + 2.5f;。

示例：1u + 2.5 会产生着色器创建错误：

• 1u 项是类型为u32 的表达式。

• 2.5 项是类型为AbstractFloat 的表达式。

• 不存在有效的重载候选：

  • 不存在从 GPU 实体化的整数标量类型到 浮点类型的可行自动转换。

  • 没有类型规则匹配 e+f，其中 e 属于整数 标量类型，而 f 属于浮点类型。

示例：-1 * i32(-2147483648) 不会产生着色器创建错误：

• -1 项是类型为AbstractInt 的表达式。

• i32(-2147483648) 项是类型为i32 的表达式。

• 乘法运算符不存在针对这两个类型的重载，并且i32 项不能向上转换为AbstractInt。

• 唯一可行的自动转换是将AbstractInt 转换为i32，因此：

  • 最终计算是两个 i32 之间的乘法操作。

  • i32 操作使用二进制补码 算术，并且具有定义明确的溢出行为。

  • 会发生回绕。

// 显式指定类型的无符号整数字面量。
var u32_1 = 1u; // 变量保存一个 u32

// 显式指定类型的有符号整数字面量。
var i32_1 = 1i; // 变量保存一个 i32

// 显式指定类型的浮点字面量。
var f32_1 = 1f; // 变量保存一个 f32

// 显式指定类型的无符号整数字面量不能取负。
var u32_neg = -1u; // 无效：一元负号不支持 u32

// 当要求具体类型，但语句或表达式的任何部分都
// 没有强制使用某个特定具体类型时，整数字面量会被
// 解释为 i32 值：
//   let 声明的初始化器必须是可构造类型（或指针）。
//   从 AbstractInt 到可构造类型最优先的自动转换
//   是 AbstractInt 到 i32，转换等级为 2。因此 '1' 被推断为 i32。
let some_i32 = 1; // 类似 let some_i32: i32 = 1i;

// 从声明类型推断。
var i32_from_type : i32 = 1; // 变量保存 i32。AbstractInt 到 i32，转换等级 2
var u32_from_type : u32 = 1; // 变量保存 u32。AbstractInt 到 u32，转换等级 3

// 无后缀整数字面量可以在需要时转换为浮点：
//   自动将 AbstractInt 转换为 f32，转换等级为 5。
var f32_promotion : f32 = 1; // 变量保存 f32

// 无效：不存在从浮点到整数的可行转换
var i32_demotion : i32 = 1.0; // 无效

// 从表达式推断。
var u32_from_expr = 1 + u32_1; // 变量保存 u32
var i32_from_expr = 1 + i32_1; // 变量保存 i32

// 值必须可以表示。
let u32_too_large   : u32 = 1234567890123456890; // 无效，溢出
let i32_too_large   : i32 = 1234567890123456890; // 无效，溢出
let u32_large : u32 = 2147483649; // 有效
let i32_large : i32 = 2147483649; // 无效，溢出
let f32_out_of_range1 = 0x1p500; // 无效，超出范围
let f32_hex_lost_bits = 0x1.0000000001p0; // 无效，无法用 f32 精确表示

// 最小整数：对 AbstractInt 进行一元取负，然后推断为 i32。
// 从 AbstractInt 到可构造类型最优先的转换（转换
// 等级最低）是 AbstractInt 到 i32。
let i32_min = -2147483648;  // 类型为 i32

// 无效。与上面一样选择 AbstractInt 到 i32，但该值超出
// 范围，因此产生着色器创建错误。
let i32_too_large_2 = 2147483648; // 无效。

// 子表达式可以解析为 AbstractInt 和 AbstractFloat。
// 以下示例都有效，并且变量的值都是 6u。
var u32_expr1 = (1 + (1 + (1 + (1 + 1)))) + 1u;
var u32_expr2 = 1u + (1 + (1 + (1 + (1 + 1))));
var u32_expr3 = (1 + (1 + (1 + (1u + 1)))) + 1;
var u32_expr4 = 1 + (1 + (1 + (1 + (1u + 1))));

// 根据内置函数参数进行推断。

// 最优先的候选是 clamp(i32,i32,i32)->i32
let i32_clamp = clamp(1, -5, 5);
// 最优先的候选是 clamp(u32,u32,u32)。
// 字面量使用从 AbstractInt 到 u32 的自动转换。
let u32_clamp = clamp(5, 0, u32_from_expr);
// 最优先的候选是 clamp(f32,f32,f32)->f32
// 字面量使用从 AbstractInt 到 f32 的自动转换。
let f32_clamp = clamp(0, f32_1, 1);

// 以下示例全部提升为 f32，初始值为 10f。
let f32_promotion1 = 1.0 + 2 + 3 + 4;
let f32_promotion2 = 2 + 1.0 + 3 + 4;
let f32_promotion3 = 1f + ((2 + 3) + 4);
let f32_promotion4 = ((2 + (3 + 1f)) + 4);

// 类型规则违规。

// 无效，初始化器只能解析为 f32：
// 不存在从 AbstractFloat 到 u32 的可行自动转换。
let mismatch : u32 = 1.0;

// 无效。clamp 不存在允许混合符号参数的重载。
let ambiguous_clamp = clamp(1u, 0, 1i);

// 推断在语句层级完成。

// let 声明的初始化器必须是可构造类型（或指针）。
// 从 AbstractInt 到可构造类型最优先的自动转换
// 是 AbstractInt 到 i32，转换等级为 2。因此 '1' 被推断为 i32。
let some_i32 = 1; // 类似 let some_i32: i32 = 1i;

let some_f32 : f32 = some_i32; // 类型错误：i32 不能赋值给 f32

// 另一个溢出情况
let overflow_u32 = (1 -2) + 1u; // 无效，-1 超出 u32 的范围

// 理想值超出 32 位范围，但又回到范围内
let out_and_in_again = (0x1ffffffff / 8);

// 类似，但无效
let out_of_range = (0x1ffffffff / 8u); // 要求以 32 位进行计算，
                                       // 使 0x1ffffffff 超出范围。

6.2.2. 布尔类型

bool 类型包含值 true 和 false。

6.2.3. 整数类型

u32 类型是 32 位 无符号整数的集合。

i32 类型是 32 位 有符号整数的集合。 它使用二进制补码表示，符号位位于最高有效位位置。

对表达式中的具体整数类型进行会溢出的运算时，产生的结果按 2^bitwidth 取模

注： AbstractInt 也是整数类型。

6.2.4. 浮点类型

f32 类型是 IEEE-754 binary32（单精度）格式的 32 位 浮点值集合。 有关详细信息，请参见§ 15.7 浮点求值。

f16 类型是 IEEE-754 binary16（半精度）格式的 16 位 浮点值集合。如果程序未包含 enable f16; 指令来启用 f16 扩展，却使用了 f16 类型，则会产生着色器创建 错误。有关详细信息，请参见§ 15.7 浮点求值。

下表列出了浮点类型的某些极值。 每个值都有对应的负值。

注： AbstractFloat 也是浮点类型。

6.2.5. 标量类型

标量类型包括bool、AbstractInt、 AbstractFloat、i32、u32、f32 和f16。

数值标量类型 包括AbstractInt、 AbstractFloat、i32、u32、f32 和f16。

整数标量类型 包括AbstractInt、i32 和 u32。

标量转换 将一种标量类型中的值映射到另一种标量类型中的值。 通常，结果值在目标类型限制范围内接近原始值。 标量转换通过以下任一种方式发生：

• 显式调用值构造器，或

• 通过可行自动转换，将const 表达式从抽象数值类型转换为 另一种类型。

6.2.6. 向量类型

向量是由 2、3 或 4 个标量 分量组成的分组序列。

如果向量的分量类型是数值标量，则该向量是数值 向量。

向量的主要使用场景包括：

• 同时表示方向和大小。

• 表示空间中的位置。

• 表示某个颜色空间中的颜色。 例如，各分量可以分别是红、绿、蓝的强度， 而第四个分量可以是 alpha（不透明度）值。

许多对向量（以及矩阵）的操作按逐分量方式执行，即 结果通过对 每个标量分量独立执行操作形成。

vec2<f32>  // 是由两个 f32 组成的向量。

let x : vec3<f32> = a + b; // a 和 b 都是 vec3<f32>
// x[0] = a[0] + b[0]
// x[1] = a[1] + b[1]
// x[2] = a[2] + b[2]

WGSL 还预声明以下 类型别名：

6.2.7. 矩阵类型

矩阵是由 2、3 或 4 个浮点向量组成的分组序列。

矩阵的主要使用场景是表示线性变换。 在这种解释中，矩阵的向量被视为列向量。

乘积运算符（*）用于以下任一种情况：

• 用标量大小缩放变换。

• 将变换应用于向量。

• 将变换与另一个矩阵组合。

参见§ 8.8 算术表达式。

mat2x3<f32>  // 这是一个由 32 位浮点数组成的 2 列 3 行矩阵。
             // 等价地，它是两个 vec3<f32> 类型的列向量。

WGSL 还预声明以下 类型别名：

6.2.8. 原子类型

原子类型封装一个具体整数 类型，使得：

• 原子对象向并发观察者提供某些保证，并且

• 对原子对象唯一有效的操作是原子 内置函数。

表达式不得求值得到原子类型。

原子类型只能由workgroup 地址空间中的变量，或具有read_write 访问模式的存储 缓冲区变量实例化。 如果存储类型是或包含 atomic<vec2<u32>>，则该变量必须 位于storage地址空间中。

对该类型执行操作时的内存作用域由实例化该类型所在的地址空间 决定。 workgroup 地址空间中的原子类型具有 Workgroup 的内存作用域，而 storage 地址空间中的原子类型具有 QueueFamily 的内存作用域。

原子 修改是任何 会设置原子对象内容的操作。 即使新值与 对象现有值相同，该操作仍计为一次修改。

在 WGSL 中，对于每个对象，原子修改彼此有序。 也就是说，在着色器阶段执行期间，对于每个原子对象 A，所有 代理都观察到应用于 A 的修改操作具有相同的顺序。 不同原子对象之间的顺序可能完全无关；这并不 暗示任何因果关系。 请注意，workgroup 空间中的变量在一个 工作组内共享，但不同 工作组之间不共享。

6.2.9. 数组类型

数组是可通过索引访问的元素值序列。

数组中的第一个元素位于索引 0，后续每个元素位于下一个整数索引。 参见§ 8.5.3 数组访问表达式。

表达式不得求值得到运行时大小数组类型。

固定大小数组的元素数量表达式 N 受以下约束：

• 它必须是一个override 表达式。

• 它必须求值得到一个具体整数标量。

• 如果 N 不大于零：

  • 如果 N 是const 表达式，则产生着色器创建错误。

  • 否则，产生管线创建错误。

注： 如果 N 依赖任何override 声明，则元素数量值在管线创建 时完全确定；否则在着色器模块创建时确定。

注： 为了满足类型等价条件，任何不是 const 表达式的 override 表达式都必须是标识符。 参见由可重写 常量确定大小的工作组变量

运行时大小数组中的元素数量 由与相应存储缓冲区变量关联的缓冲区绑定大小决定。 参见§ 13.3.4 缓冲区绑定 确定运行时大小数组元素数量。

数组元素类型必须是以下之一：

• 标量类型

• 向量类型

• 矩阵类型

• 原子类型

• 具有创建时固定占用空间的数组类型

• 具有创建时固定占用空间的结构体类型。

注： 数组元素类型不能是缓冲区类型。

注： 元素类型必须是普通类型。

两个数组类型相同，当且仅当以下所有条件均为真：

• 它们具有相同的元素类型。

• 它们的元素数量规范匹配，即以下条件之一为真：

  • 它们都是运行时大小数组。

  • 它们都是具有创建时固定占用空间的固定大小数组，并且 元素数量的值相等，即使一个是有符号值而另一个是无符号值。 （在这种情况下，有符号值和无符号值可以比较，因为元素数量 始终为正。）

  • 它们都是固定大小数组，并且其元素数量被指定为解析 到同一个管线可重写常量声明的标识符。

// array<f32,8> 和 array<i32,8> 是不同的类型：
// 元素类型不同
var<private> a: array<f32,8>;
var<private> b: array<i32,8>;
var<private> c: array<i32,8u>;  // array<i32,8> 和 array<i32,8u> 是相同类型

const width = 8;
const height = 8;

// array<i32,8>、array<i32,8u> 和 array<i32,width> 是相同类型。
// 它们的元素数量都求值得到 8。
var<private> d: array<i32,width>;

// array<i32,height> 和 array<i32,width> 是相同类型。
var<private> e: array<i32,width>;
var<private> f: array<i32,height>;

注： 由可重写常量确定大小的数组类型，唯一有效的用途是 作为workgroup 地址空间中的内存视图。 这包括工作组变量的存储类型。 参见§ 7 变量和值声明。

override blockSize = 16;

var<workgroup> odds: array<i32,blockSize>;
var<workgroup> evens: array<i32,blockSize>; // 相同类型

// 以下各项均与 'odds' 和 'evens' 的类型不同。

// 不同类型：不是标识符 'blockSize'
var<workgroup> evens_0: array<i32,16>;
// 不同类型：使用算术表达式表示元素数量。
var<workgroup> evens_1: array<i32,(blockSize * 2 / 2)>;
// 不同类型：使用括号，而不仅仅是标识符。
var<workgroup> evens_2: array<i32,(blockSize)>;

// 一个无效示例，因为可重写的元素数量只能出现在
// 最外层。
// var<workgroup> both: array<array<i32,blockSize>,2>;

// 一个无效示例，因为可重写的元素数量只对
// workgroup 变量有效。
// var<private> bad_address_space: array<i32,blockSize>;

6.2.10. 结构体类型

结构体是由具有名称的成员 值组成的具名分组。

结构体类型在模块作用域声明。 在程序源代码的其他位置，结构体类型由其标识符名称表示。 参见§ 5 声明和作用域。

两个结构体类型相同，当且仅当它们具有相同名称。

结构体成员类型必须是以下之一：

• 标量类型

• 向量类型

• 矩阵类型

• 原子类型

• 具有创建时固定占用空间的固定大小 数组类型

• 运行时大小数组 类型，但仅当它是结构体的最后一个成员时

• 具有创建时固定占用空间的结构体类型

注： 结构体成员类型不能是缓冲区类型。

注： 所有用户声明的结构体类型都是具体类型。

注： 每个成员类型必须是普通类型。

结构体成员和数组元素类型限制的一些结果包括：

• 指针、纹理或采样器不得出现在数组或结构体内部的任何嵌套层级中。

• 当运行时大小 数组是较大类型的一部分时，它只能作为结构体的最后一个元素出现，而该结构体本身不能成为外层数组或结构体的一部分。

// 一个具有三个成员的结构体。
struct Data {
  a: i32,
  b: vec2<f32>,
  c: array<i32,10>, // 最后的逗号是可选的
}

// 声明一个存储 Data 类型值的变量。
var<private> some_data: Data;

以下属性可以应用于结构体成员：

• align

• builtin

• location

• blend_src

• interpolate

• invariant

• size

属性 builtin、location、blend_src、 interpolate 和invariant 是 IO 属性。 结构体 S 成员上的IO 属性仅在 S 被用作入口点的形式参数类型或返回类型时有效。 参见§ 13.3.1 阶段间输入和输出接口。

属性 align 和size 是布局属性， 如果结构体类型用于定义uniform 缓冲区或 存储缓冲区， 则可能需要这些属性。 参见§ 14.4 内存布局。

struct my_struct {
  a: f32,
  b: vec4<f32>
}

// 运行时数组
alias RTArr = array<vec4<f32>>;
struct S {
  a: f32,
  b: f32,
  data: RTArr
}
@group(0) @binding(0) var<storage> buffer: S;

6.2.11. 复合类型

如果一个类型具有由其他类型组合表示的 内部结构，则它是复合类型。 内部各部分彼此不重叠，称为分量。 复合值可以分解为其分量。参见§ 8.5 复合值分解表达式。

复合类型包括：

• 向量类型

• 矩阵类型

• 数组类型

• 结构体类型

对于复合类型 T，T 的嵌套深度，记作 NestDepth(T)，定义为：

• 向量类型为 1

• 矩阵类型为 2

• 元素类型为 E 的数组类型为 1 + NestDepth(E)

• 如果 T 是成员类型为 M₁,...,M_N 的结构体类型，则为 1 + max(NestDepth(M₁),..., NestDepth(M_N))

6.2.12. 可构造类型

许多类型的值都可以被创建、加载、存储、传入函数 以及从函数返回。 我们称这些值为可构造的。

如果一个类型是以下之一，则它是可构造类型：

• 标量类型

• 向量类型

• 矩阵类型

• 固定大小 数组类型，如果它具有创建时固定占用空间，并且其元素类型 可构造。

• 结构体类型，如果其所有 成员均可构造。

注： 所有可构造类型都具有创建时固定占用空间。

注： 原子类型和运行时大小数组类型不可 构造。 包含原子类型和运行时大小数组的复合类型也不可构造。

6.2.13. 固定占用空间类型

变量的内存占用空间 是用于存储变量内容的内存位置数量。 变量的内存占用空间取决于其存储类型，并会在着色器生命周期中的某个时刻 最终确定。 大多数变量的大小很早就在着色器创建时确定。 某些变量可能稍后在管线创建时确定大小， 还有一些最晚在着色器执行开始时确定。

如果一个类型的具体化大小在着色器创建时已经完全确定，则该类型具有创建时固定占用空间。

如果一个类型的大小在管线创建 时已经完全确定，则该类型具有固定 占用空间。

注： 所有具体的创建时固定占用空间类型和固定占用空间类型都是可存储的。

注： 管线创建依赖于着色器创建，因此具有创建时固定占用空间的类型也具有固定占用空间。

具有创建时固定占用空间的类型包括：

• 标量类型

• 向量类型

• 矩阵类型

• 原子类型

• 固定大小 数组类型，当：

  • 其元素 数量是const 表达式。

• 结构体类型，如果其所有 成员都具有创建时固定占用空间。

• 固定大小 缓冲区类型，当：

  • 其大小是const 表达式。

注： 可构造类型具有创建时固定占用空间。

具有固定 占用空间的类型包括以下任一种：

• 具有创建时固定占用空间的类型

• 固定大小 数组类型（不进一步约束其元素数量）

• 固定大小 缓冲区类型（不进一步约束其大小）

注： 元素数量是一个 非const 表达式的override 表达式的固定大小数组，其唯一有效用途是作为 workgroup 地址空间中的内存视图。 这包括工作组变量的存储类型。

注： 固定占用空间类型可以直接或 间接包含原子类型，而可构造类型不能。

注： 固定占用空间类型不包括运行时大小数组、 运行时大小 缓冲区，以及任何包含运行时大小数组的结构体。

6.3. 缓冲区类型

这些类型要求buffer_view 语言特性。

buffer 值表示其内容可以解释为 另一种主机可共享类型的内存。 在其他方面，缓冲区值是不透明的。 当变量中的数据需要以多种 不同类型访问，或着色器作者希望将一个变量细分为 多个逻辑变量时，它们很有用。

缓冲区类型不是可构造的，并且只能通过 buffer_view 内置函数将其数据解释为非不透明的主机可共享类型后进行访问。

表达式不得求值得到缓冲区 类型。

固定大小缓冲区只能由storage、uniform 和workgroup 地址空间中的变量实例化。 在 storage 和 uniform 地址空间中，大小必须是const 表达式。 运行时大小缓冲区只能由 storage 地址空间中的变量实例化。

如果 N 不大于 0：

• 如果 N 是const 表达式，则产生着色器创建错误。

• 否则，产生管线创建错误。

如果启用了f16 扩展 且 N 不能被 2 整除：

• 如果 N 是const 表达式，则产生着色器创建错误。

• 否则，产生管线创建错误。

如果未启用f16 扩展 且 N 不能被 4 整除：

• 如果 N 是const 表达式，则产生着色器创建错误。

• 否则，产生管线创建错误。

两个缓冲区类型相同，当且仅当以下任一条件成立：

• 它们都是运行时大小缓冲区。

• 它们都是具有创建时固定占用空间的固定大小缓冲区，并且大小的值 相等，即使一个是有符号值而另一个是无符号值。（在这种情况下，有符号值和 无符号值可以比较，因为大小始终 为正。）

• 它们都是固定大小缓冲区，并且其大小被指定为解析 到同一个管线可重写常量声明的标识符。

6.4. 枚举类型

枚举类型是一组 有限的具名值。 枚举用于区分某个特定概念的一组可能情况，例如一组 有效的纹素格式。

枚举项是枚举中的一个 具名值。 每个枚举项都不同于所有 其他枚举项，也不同于所有其他种类的值。

WGSL 源代码中没有声明新枚举项或新枚举类型的机制。

注： 枚举项用作模板参数。

6.4.1. 预声明枚举项

下表列出了 WGSL 中的枚举类型、它们预声明的枚举项，以及这些枚举类型所需的语言扩展。 枚举类型确实存在，但不能在 WGSL 源代码中写出。

6.5. 内存视图

除了使用普通 值进行计算之外，WGSL 程序还经常会 通过内存访问操作从内存读取值或向内存写入值。 每次内存访问都通过内存视图执行。

内存视图包括：

• 特定地址空间中的一组内存 位置，

• 一个内存模型引用，

• 将这些位置的内容解释为某种 WGSL 类型，称为存储类型，以及

• 一个访问模式。

内存视图的访问模式必须受该地址空间支持。参见§ 7 变量和值声明。

6.5.1. 可存储类型

变量中包含的值必须属于 可存储类型。 可存储类型可以具有由 WGSL 定义的显式表示， 如§ 14.4.4 值的内部布局中所述， 也可以是不透明的，例如纹理和采样器。

如果一个类型既是具体类型，又是以下之一，则它是可存储类型：

• 标量类型

• 向量类型

• 矩阵类型

• 原子类型

• 数组类型

• 结构体类型

• 纹理类型

• 采样器类型

• 缓冲区类型

注： 也就是说，可存储类型是具体的普通类型、 纹理类型、采样器类型和缓冲区类型。

6.5.2. 主机可共享类型

主机可共享类型用于描述主机与 GPU 之间共享的缓冲区内容， 或在主机与 GPU 之间复制而无需格式转换的内容。 用于此目的时，该类型还可以按照§ 14.4 内存布局中的说明应用布局属性。 如§ 7.3 var 声明中所述，uniform 缓冲区和存储缓冲区 变量的存储类型必须是主机可共享的。

如果一个类型既是具体类型，又是以下之一，则它是主机可共享的：

• 数值标量类型

• 数值向量类型

• 矩阵类型

• 原子类型

• 固定大小 数组类型，如果它具有创建时固定占用空间，且其元素类型 是主机可共享的

• 运行时大小数组 类型，如果其元素类型是主机可共享的

• 结构体类型，如果其所有 成员都是主机可共享的

• 缓冲区类型

注： 阶段间输入 和输出类型的限制 在§ 13.3.1 阶段间输入和输出接口及 后续各节中说明。 这些类型也有大小，但计算方式不同。

注： 纹理和 采样器也可以在主机和 GPU 之间共享， 但其内容是不透明的。 本节中的主机可共享类型专用于storage 和uniform 缓冲区。

6.5.3. 引用和指针类型

WGSL 有两种用于表示内存视图的类型： 引用类型和指针类型。

两个指针类型相同，当且仅当它们具有相同的地址空间、存储类型和访问模式。

在分析 WGSL 模块时，引用类型和指针类型都由 地址空间、可存储类型和访问模式完整参数化。 在本规范的代码示例中，注释会展示这种完整参数化的形式。

但是，在 WGSL 源文本中：

• 引用类型不得出现。

• 指针类型可以出现。

  • 指针类型的拼写通过以下参数化：

    • 地址 空间，

    • 存储类型， 以及

    • 有时还包括访问模式，具体见§ 14.3 地址空间。

  • 如果指针类型出现在程序源代码中， 则在程序中的某个位置，使用该指针类型的 地址 空间、存储 类型和访问模式来声明变量或从 函数调用返回该类型，也必须是有效的。

    注： 此限制禁止声明某些 永远无法在运行时使用的类型别名 和函数形式参数。 如果没有此限制，声明指针类型的别名将是有效的， 但永远无法创建该类型的指针值。 类似地，声明具有指针形式参数的函数将是有效的， 但永远无法调用该函数。

fn my_function(
  /* 'ptr<function,i32,read_write>' 是一种指针值类型，它引用
     用于保存 'i32' 值的内存，并使用 'function'
     地址空间中的内存位置。这里 'i32' 是存储类型。
     隐含的访问模式是 'read_write'。
     默认值参见“地址空间”一节。 */
  ptr_int: ptr<function,i32>,

  // 'ptr<private,array<f32,50>,read_write>' 是一种指针值类型，它
  // 引用用于保存 50 个 'f32' 类型元素数组的内存，并使用
  // 'private' 地址空间中的内存位置。
  // 这里的存储类型是 'array<f32,50>'。
  // 隐含的访问模式是 'read_write'。
  // 默认值参见“地址空间”一节。
  ptr_array: ptr<private, array<f32, 50>>
) { }

引用类型和指针类型都是内存视图的集合： 特定内存视图与唯一的引用值以及唯一的指针值关联：

类型为 ptr<AS,T,AM> 的每个指针值 p 对应 类型为 ref<AS,T,AM> 的唯一引用值 r， 反之亦然， 其中 p 和 r 描述相同的内存视图。

6.5.4. 重排视图类型

重排视图可用于 访问内存中向量分量的一个子集， 读取时对其进行排列，写入时将其分散写入。

重排视图由swizzle_assignment 语言扩展启用。

算法： SwizzleViewRead

生成一个向量，该向量通过选择并可选地重新排列经由指针 p 读取的向量分量而形成。

注： 通过重排视图读取时，会从内存中读取整个 向量，即使某些 向量分量未出现在结果中。

输入：类型为 swizzle<AS,S,N,K> 的重排 视图 (p, IndexList)

输出：类型为 vecK<S> 的向量值 v，即读取 (p, IndexList) 的结果。

过程：

1. 按照约定，IndexList = « Idx₀, ..., Idx_K−1 »

2. 从 p 所引用的内存中读取向量 u。u 的类型为 vecN<S>。

3. 形成类型为 vecK<S> 的值 v，其中：

   • v 的分量 i 等于 u 的分量 Idx_i 的值。 使用索引语法： v[i] = u[Idx_i]

算法： SwizzleViewWrite

将 update 向量的分量写入指针 p 所引用内存中的向量的指定分量。

注： 通过重排视图写入时，会写入内存中的整个向量， 即使某些分量值在此过程中未发生变化。

有关重排视图写入如何与一致性分析和 内存模型交互的讨论，请参见下文。

输入：

• 类型为 swizzle<AS,S,N,K> 的重排视图 (p, IndexList)

• 类型为 vecN<S> 的向量 old_contents

  • old_contents 旨在表示过去某个时刻存储在 p 所引用内存中的向量。

• 类型为 vecK<S> 的向量 update

IndexList 中的所有索引必须彼此不同。 此要求意味着 K ≤ N。

变量：

• 类型为 vecN<S> 的 new

过程：

1. 按照约定，IndexList = « Idx₀, ..., Idx_K−1 »

2. 将变量 new 初始化为 old_contents。

3. 对于 0..K-1 中的每个 i：

   • 将 new 的分量 Idx_i 设置为 update 的分量 i 的值。 使用索引语法： new[Idx_i] = update[i]。

4. 将 new 写入 p 所引用的内存。

注： 要求所有 Idx_i 都彼此不同 可保证步骤 3 最多只更新 new 的每个分量一次。

注： 步骤 2 和步骤 3 可以替换为以下声明式规则：

形成 new，它是类型为 vecN<S> 的向量，其中 new 的分量 i 为：

• 如果对于 0...K−1 中的某个 j，i = Idx_j， 则为 update 的分量 j。

• 否则，为 old_contents 的分量 i。

重排视图具有重排视图类型：

重排视图通过以下方式形成：

• 将多字母重排 名称应用于访问模式为read_write、 且存储类型为向量的引用或指针值，或

• 将多字母重排 名称应用于重排视图。

根据这些规则，重排视图中的 IndexList 部分在 着色器 创建时完全确定。

有关详细信息，请参见§ 8.6 重排视图表达式。

fn swizzle_read_and_write() {
  var v: vec4u;

  // 等同于：v.y = 1; v.z = 2;
  // 但是，向量只会从内存读取一次并写入内存一次。
  v.yz = vec2u(1,2);

  // 等同于：let u = vec3(v.a, v.g, v.b);
  // 表达式 v.agb 是类型为 swizzle<function,u32,4,3> 的重排视图。
  // let 声明的有效值类型必须是具体可构造类型
  // 或指针类型。因此，会通过调用重排视图加载规则将重排视图自动转换为 vec3u
  //。
  let u = v.agb;

  // 重排可以链接。
  // 等同于：v.zy = vec2(99,100);
  v.yz.yx = vec2(99,100);

  // 在重排视图上应用单个字母会得到对
  // 底层向量某个标量分量的引用。
  // 等同于：v.z = 99;
  v.yz.y = 99;

  // 对重排视图进行索引会得到对向量某个
  // 分量的引用。其行为如同形成单字母向量访问，
  // 但索引会选择重排名称中的哪个字母。
  v.zy[1] = 50; // 等同于 v.y = 50;
}

// 此示例假定除 swizzle_assignment 外还支持 pointer_composite_access。
requires pointer_composite_access;

fn swizzle_read_and_write_via_pointer(p: ptr<function,vec4u>) {
  // 等同于：(*p).y = 1; (*p).z = 2;
  p.yz = vec2u(1,2);

  // 等同于：let u = vec3((*p).a, (*p).g, (*p).b);
  let u = p.agb;

  // 重排可以链接。
  // 等同于：p.zy = vec2(99,100);
  p.yz.yx = vec2(99,100);

  // 可以从指针构造重排视图。
  // 在重排视图上应用单个字母会得到对
  // 底层向量某个标量分量的引用。
  // 等同于：(*p).z = 99;
  p.yz.y = 99;

  // 对重排视图进行索引会得到对向量某个
  // 分量的引用。其行为如同形成单字母向量访问，
  // 但索引会选择重排名称中的哪个字母。
  p.zy[1] = 50; // 等同于 (*p).y = 50;
}