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 确实提供了一些有限机制，可将标量 类型提升为向量类型。 这同样适用于复合类型。

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

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

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

但是，各调用作用于不同的着色器阶段输入集合，其中包括 提供标识值以区分一个调用及其同级调用的内置输入。 每个调用都有自己的独立内存空间，其形式为 私有和函数地址空间中的变量。

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

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

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

对于给定特性，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

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

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 个 有效数字。

十进制 浮点字面量为：

• 一个有效数，指定为数字序列，其中可在任意位置包含可选的小数点（.）。 有效数表示以 10 为底的分数。

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

  • 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 前缀

• 然后是一个有效数，指定为十六进制数字序列，其中可在任意位置包含可选的十六进制小数点 （.）。 有效数表示以 16 为底的分数。

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

  • 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'

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 向前推进，越过空白和注释。

    • 如果 CurrentPosition 处出现 '<' (U+003C)，则：

      • 注： 此码 位可能是模板列表的起始位置。 保存足够的状态，以便将它与之后输入中出现的终止 '>' (U+003E) 匹配。

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

      • 将 CurrentPosition 向前推进到下一个码位。

      • 如果 CurrentPosition 处出现 '<' (U+003C)， 则：

        • 注： 根据 假设 1，模板参数不会以 '<' (U+003C) 开头，因此前一个码位不可能 是模板列表的起始位置。 因此当前码位和前一个码位必须构成'<<' 运算符。

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

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

      • 如果 CurrentPosition 处出现 '=' (U+003D)，则：

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

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

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

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

  • 如果 CurrentPosition 处出现 '>' (U+003E)，则：

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

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

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

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

      • 从 Pending 栈中弹出 T。

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

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

      • 将 CurrentPosition 向前推进越过此码位。

      • 如果 CurrentPosition 处出现 '=' (U+003D)，则：

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

        • 将 CurrentPosition 向前推进越过此码位。

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

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

    • 注： 进入嵌套 表达式。

    • 将 NestingDepth 加 1。

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

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

    • 注： 退出嵌套 表达式。

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

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

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

  • 如果 CurrentPosition 处出现 '!' (U+0021)，则：

    • 将 CurrentPosition 向前推进越过此码位。

    • 如果 CurrentPosition 处出现 '=' (U+003D)，则：

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

      • 将 CurrentPosition 向前推进越过此码位。

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

  • 如果 CurrentPosition 处出现 '=' (U+003D)，则：

    • 将 CurrentPosition 向前推进越过此码位。

    • 如果 CurrentPosition 处出现 '=' (U+003D)，则：

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

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

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

    • 将 NestingDepth 设置为 0。

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

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

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

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

    • 将 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 引用和指针类型。

6.1. 类型检查

WGSL 值通过对表达式求值来计算。 表达式是一段 源文本， 它被解析为名称以“expression”结尾的 WGSL 语法规则之一。 表达式 E 可以包含子表达式，这些表达式完全包含在 外层表达式 E 中。 顶层 表达式是本身不是子表达式的表达式。 请参阅§ 8.18 表达式语法摘要。

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

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

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

对特定表达式求值可能产生的值始终属于特定的 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.7 算术表达式中的规则，加法类型规则 适用于该表达式，因为：

• 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 不是引用类型，并且

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

• 对于任何其他非引用类型 T2，ConversionRank(S, T2) > ConversionRank(S, T)。

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

注： 转换为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++ 中的简单旧式数据 类型，但还 包括原子类型和抽象数值类型。

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.7 算术表达式。

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

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

6.2.8. 原子类型

原子类型封装 一种具体的整数标量类型，使得：

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

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

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

原子类型只能由workgroup 地址空间中的变量，或访问模式为read_write 的存储 缓冲区变量实例化。 对该类型执行操作的内存作用域由其 实例化所在的地址空间 决定。 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 缓冲区绑定 确定运行时大小数组的元素数量。

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

• 标量类型

• 向量类型

• 矩阵类型

• 原子类型

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

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

注： 数组元素类型不能是buffer 类型。

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

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

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

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

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

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

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

// 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>;

// 无效示例，因为可覆盖的元素数量仅
// 对工作组变量有效。
// var<private> bad_address_space: array<i32,blockSize>;

6.2.10. 结构体类型

结构体是命名的 成员 值的命名分组。

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

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

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

• 标量类型

• 向量类型

• 矩阵类型

• 原子类型

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

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

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

注： 结构体成员类型不能是buffer 类型。

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

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

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

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

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

// 具有三个成员的结构体。
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 表达式。

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

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

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

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

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

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

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

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

6.3. 缓冲区类型

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

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

缓冲区类型不是可构造的，并且只能在通过 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. 有效和无效的内存引用

引用值要么是 有效的，要么是无效的。

引用的形成方式在§ 6.5.8 形成 引用和指针值中有详细描述。 通常，有效 引用通过以下方式形成：

• 命名一个变量，或

• 对有效指针应用间接 （一元 *）操作，或

• 一个命名分量表达式，其中基是 有效的内存视图，或

• 一个索引 表达式，其中基是有效的内存视图，并使用范围内索引。

通常，无效内存引用通过以下方式形成：

• 对无效指针应用间接运算符，或

• 一个命名分量表达式，其中基是 无效内存引用，或

• 一个索引 表达式，其中基是内存视图，并且满足以下任一情况：

  • 基是无效内存引用，或

  • 索引越界。

无效指针通过 以下方式形成：

• 对无效内存引用应用取地址 运算符（一元 &），或

• 调用 bufferView 或bufferArrayView 内置函数，而所得指针的内存视图原本会包含超出 缓冲区指针实参内存视图边界之外的内存位置。

有效指针是 与有效引用相对应的指针。

6.5.5. 起源变量

指针值的起源 变量被定义为相应引用值的起源变量。

注： 起源变量是一个动态概念。 函数形式参数的起源变量取决于该函数的 调用点。 不同调用点可能提供指向不同起源变量的指针。

有效引用始终 对应某个变量的部分或全部内存位置的非空内存视图。

struct Particle {
   position: vec2f,
   velocity: vec2f,
   color_index: i32,
}

@group(0) @binding(0)
var<storage,read_write> the_particle: Particle;

fn particle_velocity_component(p: Particle, i: i32) -> f32 {
  return the_particle.velocity[i]; // 当 i 为 0 或 1 时是有效引用。
}

6.5.6. 越界访问

对无效 内存引用执行访问的操作是越界访问。

越界访问是一种程序缺陷，因为如果按照所写的方式执行，它通常会：

• 读取或写入变量之外的内存 位置，或

• 将这些位置的内容解释为错误的存储类型，或

• 导致意外的数据竞争。

因此，实现不会按照所写的方式执行该访问。 执行越界访问会产生动态错误。

注： 错误解释存储类型的一个例子出现在 上一节的示例中。 当 i 为 2 时，表达式 the_particle.velocity[i] 的 类型为 ref<storage,f32,read_write>，这意味着它是一个以f32 为 存储类型的内存视图。 但是，这些内存位置分配给了 color_index 成员，因此实际存储的值类型为i32。

6.5.7. 引用和指针的用例

引用和指针根据其使用方式进行区分：

• 变量的类型是 引用类型。

• 取地址操作 （一元 &）将引用值转换为其对应的指针值。

• 间接操作 （一元 *）将指针值转换为其对应的引用值。

• let 声明 可以具有指针类型，但不能具有引用类型。

• 形式 参数可以具有指针类型，但不能具有引用类型。

• 简单 赋值语句通过引用执行写访问来更新内存内容，其中：

  • 赋值语句的左 侧必须具有 引用类型，并且访问模式为write 或read_write。

  • 赋值语句的右 侧必须求值得到 左侧的存储类型。

• 加载规则：在 函数内部，引用会自动解引用（读取）以满足类型规则：

  • 在函数中，当存储类型为 T 的引用表达式 r 被用于 语句或表达式中，并且

  • r 的访问模式为read 或read_write，并且

  • 唯一可能匹配的类型规则要求 r 具有类型 T 的值时，

  • 该类型规则要求被视为已满足，并且

  • 在该上下文中对 r 求值的结果，是求值时 r 所引用内存位置中存储的值（类型为 T）。 也就是说，会执行读 访问来产生结果值。

以这种方式定义引用，可以简单且符合习惯地使用变量：

@compute @workgroup_size(1)
fn main() {
  // 'i' 具有引用类型 ref<function,i32,read_write>
  // 'i' 的内存位置存储 i32 值 0。
  var i: i32 = 0;

  // 'i + 1' 只能匹配 'i' 子表达式类型为 i32 的类型规则。
  // 因此表达式 'i + 1' 的类型为 i32，并且在求值时，'i' 子表达式
  // 会求值得到求值时 'i' 的内存位置中存储的 i32 值。
  // 
  let one: i32 = i + 1;

  // 更新 'i' 所引用位置中的值，使其保存值 2。
  i = one + 1;

  // 更新 'i' 所引用位置中的值，使其保存值 5。
  // 右侧的求值发生在赋值生效之前。
  i = i + 3;
}

var<private> age: i32;
fn get_age() -> i32 {
  // return 语句中表达式的类型必须为 'i32'，因为它
  // 必须与函数声明的返回类型匹配。
  // 'age' 表达式的类型为 ref<private,i32,read_write>。
  // 应用加载规则，因为引用的存储类型与
  // 表达式所需类型匹配，并且没有其他类型规则适用。
  // 在此上下文中对 'age' 求值得到的结果，是执行 return 语句时
  // 从 'age' 所引用的内存位置中加载的 i32 值。
  // 
  return age;
}

fn caller() {
  age = 21;
  // copy_age 常量将得到 i32 值 21。
  let copy_age: i32 = get_age();
}

以这种方式定义指针可以支持两个主要用例：

• 使用指针类型的 let 声明，为变量内容的一部分创建简短名称。

• 使用函数的形式参数引用调用函数可访问的变量内存。 调用函数。

  • 对此类函数的调用必须为该操作数提供指针值。 这通常需要使用取地址操作（一元 &），以获取 指向变量内容的指针。

struct Particle {
  position: vec3<f32>,
  velocity: vec3<f32>
}
struct System {
  active_index: i32,
  timestep: f32,
  particles: array<Particle,100>
}
@group(0) @binding(0) var<storage,read_write> system: System;

@compute @workgroup_size(1)
fn main() {
  // 形成一个指向 storage 内存中特定 Particle 的指针。
  let active_particle = &system.particles[system.active_index];

  let delta_position: vec3<f32> = (*active_particle).velocity * system.timestep;
  let current_position: vec3<f32>  = (*active_particle).position;
  (*active_particle).position = delta_position + current_position;
}