WebGPU Shading Language

1. 介绍

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

@fragment
fn main() -> @location(0) vec4<f32> {
    return vec4<f32>(0.4, 0.4, 0.8, 1.0);
}

1.1. 技术概览

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

• draw command在inputs、outputs和附加resources的上下文中执行render pipeline。

• dispatch command在compute pipeline和附加resources的上下文中执行compute pipeline。

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

shader是WGSL程序的一部分，用于在管线中执行shader stage。着色器包括：

• 一个入口点 entry point function。

• 所有被调用函数的传递闭包，从入口点开始。该集合包括user-defined和built-in函数。（有关更严格的定义，请参阅“functions in a shader stage”。）

• 所有这些函数静态访问（statically accessed）的一组变量和常量。

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

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

• 计算在 module-scope 中声明的常量的值。

• 将资源（resources）绑定到着色器资源接口(resource interface)中的变量，使这些资源的内容在执行期间可供着色器使用。

• 为其他模块范围(module-scope)的变量分配内存，并使用指定的初始值填充该内存。

• 使用阶段的管线输入填充入口点的形式参数（如果存在）。

• 将入口点返回值(return value)（如果存在）连接到阶段的管线输出。

• 然后它调用入口点。

一个 WGSL 程序被组织成：

• 函数，指定执行行为。

• 文字，是纯数学值的文本表示。

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

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

• 类型，每个类型描述：

  • 对受支持表达式的限制。

  • 这些表达式的语义。

WGSL 是一种命令式语言：行为被指定为要执行的语句序列。声明：

• 声明常量或变量

• 修改变量内容

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

  • 选择：if/else, switch

  • 循环: loop, while, for.

  • 跳出嵌套命令结构：break, continue

  • 重构：函数调用和返回

  • 丢弃（仅限片元着色器）：终止调用并丢弃输出

• 评估表达式以计算值作为上述行为的一部分。

WGSL 是静态类型的：由特定表达式计算的每个值都属于特定类型，仅通过检查程序源来确定。

WGSL 具有以数组和结构的形式描述布尔值、数字、向量、矩阵和由它们组成的数组和结构。其他类型描述内存。

WGSL 没有数字或布尔类型之间的隐式转换或提升。 将值从一种数字或布尔类型转换为另一种需要显式转换(conversion)、构造(construction)或重新解释位(reinterpretation of bits)。这也适用于向量类型。

WGSL 有纹理和采样器类型。连同相关的内置函数，这些支持通常用于图形渲染的功能，并且通常由GPU提供。

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

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

• 在计算着色器（compute shader）中，同一工作组（workgroup）中的调用共享工作组（workgroup）存储类（address class）中的变量。不同工作组中的调用不共享这些变量。

但是，调用作用于不同的管线输入集，包括提供识别值以将调用与其对等调用区分开的内置输入。 此外，每个调用在 private 和 function 地址空间中以变量的形式拥有自己的独立内存空间。

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

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

• 协调对共享变量的潜在冲突访问，以避免竞争条件。

WGSL 有时允许给定特征的几种可能行为。这是一种可移植性危害，因为不同的实现可能会表现出不同的行为。WGSL的设计旨在最大限度地减少此类情况，但受到可行性和在各种设备上实现高性能的目标的限制。

1.2. 符号

floor expression在实数x上定义：

• ⌊x⌋ = k, 其中k是一个唯一整数，使得k ≤ x < k+1

ceiling expression在实数x上定义：

• ⌈x⌉ = k, 其中k是一个唯一整数，使得k-1 < x ≤ k

roundUp函数在正整数k和n上定义：

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

c列r行矩阵A的转置（transpose）是将A的行复制为A^T的列形成的r列c行矩阵A^T：

• transpose(A) = A^T

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

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

2. 着色器生命周期

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

事件为：

1. Shader module creation

   • 当WebGPU createShaderModule方法被调用后发生。此时WGSL程序的源文件被提供。

2. Pipeline creation

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

3. Shader execution start

   • 当向GPU发出draw或dispatch command命令、开始执行管线并调用shader stage entry point函数时发生。

4. Shader execution end

   • 当所有着色器中的工作被完成时发生：

     • 所有调用（invocations）完结。

     • 所有对资源（resources）的访问完结。

     • 输出（如果有）已经传递给下游管线阶段。

事件按照以下方式排序：

• 数据依赖：着色器执行需要一个管线，一个管线同样需要一个着色器模块。

• 因果关系：着色器必须开始执行才能完成执行。

2.1. 处理错误

一个WebGPU实现可能由于两个原因对着色器处理失败：

• 程序错误（program error）当着色器不满足WGSL或WebGPU规范要求时发生。

• 未分类错误（uncategorized error）当所有WGSL和WebGPU需求都被满足时可能发生，可能的原因包括：

  • 着色器太复杂，超出了实现的能力，但在某种程度上不容易被规定的限制（limits）捕获。 简化着色器可能会解决此问题。

  • WebGPU实现中的缺陷。

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

• shader-creation error是在着色器模块创建时可行检测到的错误。检测必须仅依赖于WGSL程序源文本和其他可用于 createShaderModule API 方法的信息。

• pipeline-creation error是在管线创建时可检测到的错误。检测必须仅依赖于特定管线创建API方法可用的WGSL程序源文本和其他信息。

• dynamic error是在着色器执行期间发生的错误。这些错误可能会也可能不会被检测到。

注意：比如，一个竞争条件可能不会被检测到。

每个要求将被尽早检查。也就是：

• 当未能满足着色器创建时可检测到的要求，会导致shader-creation错误。

• 当未能满足管线创建时可检测到的要求，但更早无法检测到，会导致pipeline-creation错误。

当上下文不清楚时，本规范会指出未能满足特定要求是否会导致着色器创建、管线创建或动态错误。

WebGPU规范描述了每种错误的后果。

TODO：更新WebGPU规范，参考这里定义的三种错误。

3. 文本结构

WGSL程序为文本。本规范并未规定该文本的特定编码。但是，UTF-8始终是WGSL程序的有效编码。

注意：像这样推广 UTF-8 的目的是为了简化WGSL程序的交换并鼓励工具之间的互操作性。

WGSL程序文本由一系列双字节对字符 code points组成，分组为连续的非空集形成：

• comments

• tokens

• blankspaces

程序文本不得包含空字符。

3.1. 解析

1. 去掉注释comments：

   • 用一个space字符编码(U+0020)替换第一个注释。

   • 重复直到没有注释。

2. 从头到尾扫描，将剩余的字符编码分组为标记tokens和空格blankspaces：

   • 下一组由剩余未分组字符编码的最长非空前缀组成

3. 删除空格，只保留标记。

4. 解析token序列，尝试匹配 translation_unit 语法规则。

在以下情况下会导致shader-creation error：

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

• translation_unit 语法规则与整个token序列不匹配。

3.2. 空格和换行符

Blankspace 是 Unicode 中的一个或多个代码点的任意组合。 以下是 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)

line break 是一个连续的 blankspace 代码点序列，指示行尾。 它被定义为表示“强制中断”的空格，定义为 UAX14 第 6.1 节不可定制的换行规则 LB4 和 LB5。 也就是说，换行符是以下任何一种：

• 换行 (U+000A)

• 垂直制表符 (U+000B)

• 换页 (U+000C)

• 回车符 (U+000D) 后不跟换行符 (U+000A)

• 回车 (U+000D) 后跟换行符 (U+000A)

• 下一行 (U+0085)

• 行分隔符 (U+2028)

• 段落分隔符 (U+2029)

注意：根据行号报告源文本位置的诊断应使用 line breaks数行。

3.3. 注释

注释（comment）是一块不影响WGSL程序有效性和含义的文本，只是注释可以分隔tokens。着色器作者可以使用注释去给程序添加文档信息。

行尾注释（line-ending comment）是一种由两个字符编码 // (U+002F followed by U+002F) 和后面的字符编码组成的注释，直到但不包括：

• 下一个line break，或

• 程序结尾

block comment是一种comment，由以下内容组成：

• 两个字符编码 /* (U+002F 跟着 U+002A)

• 然后是任何序列：

  • A block comment，或

  • 不包含 */ (U+002A 跟着 U+002F) 或 /* (U+002F followed by U+002A) 的文本

• 然后两个字符编码 */ (U+002A 跟着 U+002F)

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

const f = 1.5;   // 这是行尾注释。
const g = 2.5;   /* 这是一个块注释
                  跨越几行。
                  /* 块注释可以嵌套
                   */
                  但是所有的块注释必须完结。
                */

3.4. 标记

一个标记（token）是一系列连续字符编码构成以下内容：

• 一个literal.

• 一个keyword.

• 一个reserved word.

• 一个syntactic token.

• 一个identifier.

3.5. 文字

文字（literal）为一下之一：

• 布尔文字boolean literal：true或 false。

• 数字文字numeric literal：integer literal 或 floating point literal，用于表示一个数字。

const a = true;
const b = false;

numeric literal的形式是通过模式匹配定义 integer literal 是：

• 指定为以下任意项的整数： *0

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

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

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

const a = 0x123;
const b = 0X123u;
const c = 1u;
const d = 123;
const e = 0;
const f = 0i;
const g = 0x3f;

floating point literal 是 十进制浮点文字l 或 十六进制浮点文字。

• 十进制浮点文字是：

  • 尾数，指定为数字序列，其中包含可选的小数点 (.)。

  • 然后是一个可选的指数后缀，包括：

    • e 或 E。

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

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

  • 小数点、指数、f 或 h 后缀中至少有一个必须存在。 如果没有，则标记是 integer literal。

  • 字面量的值是尾数乘以 10 的指数幂。 如果未指定指数，则假定指数为 0。

• 十六进制浮点字面量 是：

  • 一个 0x 或 0X 前缀

  • 然后是尾数，指定为十六进制数字序列，其中有一个可选的十六进制点 (.)。

  • 然后是一个可选的指数后缀，包括：

    • p 或 P

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

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

  • 至少一个十六进制点或指数必须存在。 如果两者都不是，则标记是 integer literal。

  • 文字的值是尾数乘以 2 的指数次方的值。 如果未指定指数，则假定指数为 0。

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;
const i = 0xa.fp+2;
const j = 0x1P+4f;
const k = 0X.3;
const l = 0x3p+2h;
const m = 0X1.fp-4;
const n = 0x3.2p+2h;

当 numeric literal 具有后缀时，该文字表示特定 scalar 类型中的值。 否则，文字表示下面定义的 abstract numeric types 之一的值。

如果出现以下情况，则会导致 shader-creation error：

• 带有 i 或 u 后缀的 integer literal 不能由目标类型表示。

• 带有 f 或 h 后缀的 hexadecimal floating point literal 溢出或无法由目标类型精确表示。

• 带有 f 或 h 后缀的 decimal floating point literal 会溢出目标类型。

• 使用带有 h 后缀的 floating point literal 而 f16 extension 未启用。

Note: The hexadecimal float value 0x1.00000001p0 requires 33 mantissa bits to be represented exactly, but f32 only has 23 explicit mantissa bits.

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

3.6. 关键字

一个关键字(keyword)是一个token，总是指代被预定义的语言概念。有关WGSL关键字列表，请参阅§ 14.1 关键字摘要。

3.7. 标识符

标识符（identifier）是一种用作名称的token，请参阅§ 3.10 声明和范围和§ 3.9 指令。

标识符的形式基于 Unicode 标准附件 #31 用于 Unicode 版本 14.0.0, 以下详细说明。

标识符使用根据 UAX31 Grammar 描述的以下配置文件：

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

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

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

除以下例外：

• 标识符的拼写不得与 keyword 或 reserved word 相同。

• 标识符不能是 _ （单下划线， U+005F ）。

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

Unicode 版本 14.0.0 的 Unicode 字符数据库 包括非规范列表，其中包含 XID_Start 和 XID_Continue。

注意：某些 built-in functions 的 return type 是其名称不能用于WGSL源的结构类型。 这些结构类型被描述为 predeclared，名称以两个下划线开头。 结果值可以使用类型推断保存到新声明的 let 或 var 中，或者立即通过名称立即提取其成员之一。 参见 frexp 和 modf 描述中的示例用法。

3.7.1. 标识符比较

当且仅当它们由相同的代码点序列组成时，两个 WGSL 标识符才是相同的。

特别是，两个标识符在 WGSL 中可能是不同的，但在常规规范化下被认为是相同的， 映射和匹配算法，例如：

• 标准化表格 C (NFC),

• 规范化形式 D (NFD),

• 规范化形式 KC (NFKC),

• 规范化形式 KD (NFKD),

• UTS46 映射，和

• charmod-norm 匹配。

注意：当 WGSL 程序的含义发生改变时，用户代理应该发出开发人员可见的警告，如果 标识符的所有实例都替换为该标识符的同形异义词之一。 （同形文字是一个代码点序列，在读者看来可能与另一个代码点序列相同。 用于检测同形文字的映射示例是中提到的转换、映射和匹配算法 上一段。如果标识符可以通过以下方式将一个代码点转换为另一个，则两个代码点序列是同形异义词 反复用同形文字替换一个子序列。）

3.8. 属性

属性（attribute）修改对象或类型。WGSL为应用属性提供了统一的语法。属性用于多种目的，例如指定带有API的接口。一般来说，从语言的角度来看，出于类型和语义检查的目的，可以忽略属性。

对每个对象或类型，属性不能被多次指定。

下面的管道阶段属性 将函数指定为特定 shader stage 的 entry point。 这些属性只能应用于 function declarations， 并且在给定的函数上最多可以存在一个。 它们不带参数。

3.9. 指令

指令（directive）是一个token序列，它修改了WebGPU实现对WGSL程序的处理方式。见§ 10.1 启用指令。

3.10. 声明和范围

声明（declaration）将标识符（identifier）与以下类型的对象之一相关联：

• type

• value

• variable

• function

• formal parameter

换句话说，声明引入了对象的名称（name）。

声明的范围（scope）是一组程序位置，其中使用声明的标识符可能表示其关联对象。我们说标识符在那些源位置的（声明的）范围内（in scope）。

使用标识符时，对于某些声明，它必须是 in scope 或作为指令的一部分。 当标识符用于该名称的一个或多个声明的范围内时，该标识符将表示最接近该用途的非 module-scope 声明的对象，或者如果是模块范围声明 范围内没有其他声明。 我们说标识符使用 resolves 到那个声明。

声明出现的位置决定了它的范围。 通常，范围是在声明结束后立即开始的一段文本。 module scope 处的声明是例外，如下所述。

当该标识符已经在范围内且与该名称的另一个实例具有相同的范围结束时，声明不得引入该名称。

某些对象由 WebGPU 实现提供，并被视为已被每个 WGSL 程序声明。 我们说这样的对象是预先声明的（predeclared）。他们的范围是整个WGSL程序。预先声明的对象的示例为：

• built-in functions, and

• built-in types.

如果声明出现在任何其他声明的文本之外，则声明在模块范围。 整个程序的模块范围声明是 in scope。 也就是说，模块范围内的声明可以被该声明之后或之前的源文本引用。

如果任何模块范围声明是递归的，则为 shader-creation error。 也就是说，声明之间不能有循环：

考虑有向图，其中：

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

• 声明 D 有一条边 到声明 T 当定义 对于 D 提到 resolves 到 T 的标识符。

此图不得有环。

注意：function body 是 function declaration 的一部分，因此 函数不得直接或间接递归。

注意：非module scope 标识符的使用必须遵循文本中该标识符的声明。 然而，对于 module scope 声明来说，情况并非如此，它们可能在文本中被乱序引用。

注意：只有 function declaration 可以包含其他声明。

// 无效，无法重用内置函数名称。
var<private> modf: f32 = 0.0;

// 有效，foo_1 在整个程序的范围内。
var<private> foo: f32 = 0.0; // foo_1

// 有效，my_func_1 在整个程序的范围内。
var<private> bar: u32 = 0u; // bar_1

// 有效，my_func_1 在程序结束前一直在作用域内。
// 有效，foo_2 一直在作用域内直到函数结束。
fn my_func(foo: f32) { // my_func_1, foo_2
  // 对 'foo' 的任何引用都解析为函数参数。

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

  // 有效，bar_2 一直在作用域内直到函数结束。
  var bar: u32; // bar_2
  // 对 'bar' 的引用解析为 bar_2
  {
    // 有效，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++ ) { // i_1
      // 无效，i_2 与 i_1 具有相同的结束范围。
      var i: i32 = 1; // i_2.
    }
  }

  // 无效，bar_5 与 bar_2 具有相同的结束范围。
  var bar: u32; // bar_5

  // 有效的模块范围声明在整个程序的范围内。
  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;

4. 类型

程序计算值。

在WGSL中，类型（type）是一组值，每个值都属于一种类型。值的类型决定了可以对该值执行的操作的语法和语义。

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

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

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

• 32 位浮点值 1.0f,

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

• AbstractInt 值 1，和

• AbstractFloat 值 1.0

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

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

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

• 32-位无符号整数值的集合在本规范中拼写为 u32，在WGSL程序中也一样。

• 结构类型或包含结构的类型的拼写不同。

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

注意：WGSL reference types 未写在WGSL程序中。见§ 4.5 内存视图类型。

4.1. 类型检查

WGSL值是通过计算表达式来计算的。表达式（expression） 是解析为名称以“_expression”结尾的 WGSL 语法规则之一。 表达式 E 可以包含 subexpressions，它们是正确包含在外部表达式 E 中的表达式。 顶级表达式 是一个本身不是子表达式的表达式。 参见 § 6.17 表达式语法总结。

表达式求值产生的特定值取决于：

• static context: 表达式周围的源文本，以及

• dynamic context: 计算表达式的调用状态，以及调用在运行时的执行上下文。

计算特定表达式可能产生的值将始终属于特定的WGSL类型，称为表达式的静态类型（static type）。WGSL的规则被设计为表达式的静态类型仅取决于表达式的静态上下文。

类型断言（type assertion）是从一些WGSL源表达式到WGSL类型的映射。符号

e : T

是一个意思为T为WGSL表达式e的静态类型的断言。

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

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

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

• 在指定显式类型的 let 声明中，初始化表达式必须计算为该类型。

类型检查 一个成功解析的WGSL程序是映射过程 每个表达式为其静态类型， 并验证是否满足每个语句的类型要求。 如果类型检查失败，则会产生一个 shader-creation error 的特殊情况，称为 type error。

可以通过递归应用 type rules 来执行类型检查 到句法短语，其中 句法短语 是 expression 或 statement。 类型规则 描述 syntactic phrase 的 static context 如何 确定该短语中包含的表达式的静态类型。 type rule 有两个部分：

• 一个结论。

  • 如果短语是表达式，则结论是表达式的 type assertion。

  • 如果短语是陈述，则结论是一组 type assertions， 每个语句的 顶级表达式 一个。

  • 在这两种情况下，syntactic phrases 都是示意性指定的， 使用 italicized 名称来表示子表达式 或其他语法确定的参数。

• 先决条件，包括：

  • 对于表达式：

    • 为子表达式键入断言，当它有子表达式时。 每个都可以直接满足，或者通过 feasible automatic conversion（如 § 4.1.2 转化排名 中定义）。

    • 如何在语句中使用表达式。

  • 对于语句：

    • 语句的句法形式，以及

    • 在语句中为 顶级表达式 键入断言。

  • 其他原理图参数的条件（如果有）。

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

类型规则的每个不同类型参数化称为 overload。 例如，unary negation（-e 形式的表达式）有十二个重载，因为它的类型规则由类型 T 参数化。 可以是以下任何一种：

• i32

• vec2<i32>

• vec3<i32>

• vec4<i32>

• f32

• vec2<f32>

• vec3<f32>

• vec4<f32>

• f16

• vec2<f16>

• vec3<f16>

• vec4<f16>

在以下情况下，类型规则适用于句法短语：

• 规则的结论匹配 syntactic phrase 的有效解析，并且

• 规则的先决条件被满足。

考虑表达式，1u+2u。 它有两个 literal 子表达式：1u 和 2u，都是 u32 类型。 顶级表达式 是一个附加项。 参考 § 6.8 算术表达式 规则，标量 u32 加法的类型规则适用于表达式，因为：

• 1u+2u 匹配 e1+e2 形式的解析，带有 e1代表 1u 和e2代表“2u”，和

• e1是 u32 类型，并且

• e2是u32类型。

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

• 没有类型规则适用于表达式。这会导致 type error。

• 只有一种类型规则适用于表达式。 在这种情况下，规则的 conclusion 被断言，确定表达式的静态类型。

• 不止一种类型规则适用。即满足多个overload的前提条件。 在这种情况下，使用 § 4.1.3 过载分辨率 中描述的平局过程。

  • 如果重载解析成功，则确定将单个类型规则应用于表达式。 该重载的 结论 中的 type assertions 被断言， 并因此确定 syntactic phrase 中一个或多个表达式的类型。

  • 如果重载解析失败，则结果为 type error。

继续上面的例子，只有一个类型规则适用于表达式 1u+2u，所以类型检查 接受该类型规则的结论，即 1u+2u 是 u32 类型。

一个WGSL源程序为well-typed当：

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

• 满足每个语句的类型要求

否则会出现类型错误（type error）并且源程序不是有效的WGSL程序。

WGSL是一种静态类型语言（statically typed language），因为 WGSL 程序的类型检查要么成功要么发现类型错误，而只需要检查程序源文本。

4.1.1. 类型规则表

对于表达式WGSL类型规则（type rules）被组织成类型规则表（type rule tables），每个类型规则占一行。

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

TODO：示例：非结果值效应是函数调用子表达式的任何副作用。

4.1.2. 转化排名

当一个类型断言 e:T用作 type rule precondition，满足以下条件：

• e已经是 T 类型，或者

• e 的值可以自动转换为 T 类型的值。

该规则由 ConversionRank 函数对类型对进行编码，在下表中定义。 ConversionRank 函数表示自动将一种类型 (Src) 的值转换为另一种类型 (Dest) 的偏好和可行性。 较低的等级更可取。

可行的自动转换 将值从 Src 类型转换为 Dest 类型，并且当 ConversionRank(Src,Dest) 是有限的时允许。 这种转换是保值的，受 § 12.5 浮点计算 中描述的限制。

注意：自动转换仅在两种情况下发生。 首先，将 creation-time constant 转换为可在 GPU 上使用的相应类型数值时。 其次，当发生来自内存引用的加载时，会产生存储在该内存中的值。

注意：无限等级的转换是不可行的，即不允许。

注意：当不进行转换时，转换等级为零。

4.1.3. 过载分辨率

当多个 type 规则适用于句法短语 时，使用平局程序 来决定哪一个应该生效。 此过程称为重载解析， 并假设类型检查已经成功找到 subexpressions 的静态类型。

考虑一个 syntactic phrase P，并且所有 type 规则，这些规则将 应用于 P。 重载解析算法将这些类型规则称为重载候选。 对于每个候选人：

• 其先决条件已直接或通过自动转换得到满足。

• 它的 conclusion 有：

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

  • 一个 type assertion 对应于 P 中的每个 top-level expression。

P 的重载分辨率如下进行，目标是找到一个最 preferable overload Candidate：

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

   • 其静态类型已计算。

   • 从表达式的静态类型到前置条件中对应类型断言所需要的类型有一个feasible automatic conversion。 令 C.R(i) 为该转换的 ConversionRank。

2. 给候选者排序：给定两个重载候选者C1和 C2、C1 首选优于C2如果：

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

     • 即每次表达式转换都需要应用C1到 P至少与应用 C2 所需的相应表达式转换一样可取到 P。

   • 至少有一个表达式位置 i其中C1.R(i) < C2.R(i)。

     • 即应用 C1 至少需要进行一次表达式转换这比应用 C2 所需的相应转换更可取。

3. 如果只有一个候选人 C这是 preferred 高于所有其他，然后重载解析成功，产生候选类型规则 C。 否则，重载解析失败。

   TODO: Examples

4.2. 创建时间常量的类型

某些表达式在 shader-creation time 时计算， 并且具有可能比 GPU 直接实现的数值范围和精度更大的数值范围和精度。

WGSL 为这些评估定义了两个 抽象数字类型：

• AbstractInt 类型是整数集i，具有-2⁶³ ≤ 我 < 2⁶³。

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

对其中一种类型的表达式的求值不得溢出或产生未定义的结果。 否则，结果是 shader-creation error。

这些类型不能在 WGSL 源代码中拼写。它们仅由 type checks 使用。

不是抽象数字类型也不包含抽象数字的类型 类型称为 concrete。

没有后缀的 numeric literal 表示 abstract numeric type 中的值：

• 没有 i 或 u 后缀的 integer literal 表示 AbstractInt 值。

• 没有 f 后缀的 floating point literal 表示 AbstractFloat 值。

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

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

• 没有带有整数标量形式参数的 log2 重载。

• 考虑到两个可能的重载和 可行的自动转换，而是应用 overload resolution：

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

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

• 结果计算以 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）。

• 第一个重载是 preferable Candidate 并且类型检查成功。

• 结果计算以 AbstractFloat 1.0 + 2.5 的形式出现。

示例：让 x = 1 + 2.5;

• 这个例子和上面的例子类似，只是 x 不能解析为 abstract numeric type。

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

• 声明的效果就像是写成 let x : f32 = 1.0f + 2.5f;。

示例：1u + 2.5 导致 shader-creation error：

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

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

• 没有有效的重载候选：

  • 从 GPU 物化的整数类型到浮点类型的自动转换是不可行的。

  • 没有类型规则匹配 e+f 与 e 为整数类型和 f 为浮点类型。

  EXAMPLE: 文字类型推断

       // 显式类型的无符号整数文字。
       var u32_1 = 1u； // 变量持有一个 u32
  
       // 显式类型的有符号整数文字。
       var i32_1 = 1i; // 变量保存一个 i32
  
       // 显式类型的浮点字面量。
       var f32_1 = 1f; // 变量包含一个 f32
  
       // 显式类型的无符号整数文字不能被否定。
       var u32_neg = -1u; // 无效：一元减法不支持 u32
  
       // 没有后缀的整型字面量往往会被推断为 i32：
       // let-declaration 的初始化程序必须是可构造的（或指针）。
       // 最优选的从 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 to i32，但是值out of
       // 范围，产生着色器创建错误。
       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);
  
       // TODO: 当 AbstractFloat 获得加法支持时，这些将变为有效，
       // 通过提升。
      // let f32_promotion1 = 1.0 + 2 + 3 + 4; // TODO: like let f32_promotion1:f32 = 10f;
      // let f32_promotion2 = 2 + 1.0 + 3 + 4; // TODO: like let f32_promotion1:f32 = 10f;
      // let f32_promotion3 = 1f + ((2 + 3) + 4); // TODO: like let f32_promotion1:f32 = 10f;
      // let f32_promotion4 = ((2 + (3 + 1f)) + 4); // TODO: like let f32_promotion1:f32 = 10f;
  
       // 类型规则违规。
  
       // 无效，初始化器只能解析为 f32：
       // 从 AbstractFloat 到 u32 的自动转换不可行。
          let mismatch : u32 = 1.0;
  
       // 无效的。没有允许混合符号参数的钳位过载。
       let ambiguous_clamp = clamp(1u, 0, 1i);
  
       // 推理在语句级别完成。
  
       // let-declaration 的初始化程序必须是可构造的（或指针）。
       // 最优选的从 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 超出范围。
  

4.3. 普通类型

普通类型（Plain types） 是布尔值、数字、向量、矩阵或这些值的机器表示类型，。

普通类型（plain type）是一个标量（scalar）类型，一个原子（atomic）类型，或一个复合（composite）类型。

注意：WGSL中的普通类型类似于C++中的Plain-Old-Data类型，但也包括原子类型。

4.3.1. 布尔类型

布尔（bool）类型包括值 true 和 false。

4.3.2. 整数类型

u32类型为一组32-位无符号整数的集合。

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

4.3.3. 浮点类型

f32类型是一组IEEE-754 binary32 （单精度）格式的32-位浮点型数值。详见§ 12.5 浮点计算。 f16 类型是 IEEE-754 binary16（半精度）格式的 16 位浮点值集。 如果使用 f16 类型，则为 shader-creation error，除非程序包含“enable f16;” 启用 f16 extension 的指令。 有关详细信息，请参阅 § 12.5 浮点计算。

4.3.4. 标量类型

标量（scalar）类型有bool, i32, u32, f32, 和 f16。

数字标量（numeric scalar）类型有i32, u32, f32, 和 f16。

整数标量（integer scalar）有i32和u32。

4.3.5. 向量类型

向量（vector）是一组由2，3，或4个标量（scalar）或者abstract numeric type组件构成的序列。

如果向量的组件类型是 numeric scalar，则向量是 numeric vector。

一个向量是一个抽象向量，如果它的组件类型是一个abstract numeric type。

向量的关健用例包括：

• 同时表示方向和大小。

• 表示空间中的位置。

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

向量上的许多操作都是按组件（component-wise）进行的，即结果向量是通过对每个组件独立操作而形成的。

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]

4.3.6. 矩阵类型

矩阵（matrix）是一组由2，3，或4个浮点型向量组成的序列，

矩阵的关键用例是体现线性变换。在这种解释中，矩阵的向量被视为列向量。

乘积操作符(*)可用于：

• 按标量大小缩放变换。

• 对一个向量应用变换。

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

见§ 6.8 算术表达式。

mat2x3<f32>  // 此为 2 列 3 行的 32-位浮点数矩阵。
             // 等价地，此为 2 列类型 vec3<f32> 构成的向量。

4.3.7. 原子类型

原子类型（atomic type）封装了一个整数标量integer scalar类型，使得：

• 原子对象为并发观察者提供了一定的保证，并且

• 对原子对象唯一有效的操作是原子内置函数（atomic builtin functions）。

表达式不得计算为原子类型。

原子类型只能由 workgroup 地址空间中的变量或具有 read_write 访问模式的 storage buffer 变量实例化。 对类型的操作的 memory scope 由实例化的 address space 确定。 workgroup 地址空间中的原子类型的内存范围为 Workgroup，而 storage 地址空间中的原子类型的内存范围为 QueueFamily。

atomic modification是原子对象上的任何operation，它设置对象的内容。 即使新值与对象的现有值相同，该操作也算作修改。

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

4.3.8. 数组类型

数组（array）是一组可索引的元素值。

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

表达式不得计算为运行时大小的数组类型。

元素计数表达式 N 固定大小的数组必须：

• override expression，以及

• 计算为一个大于零的整数标量（integer scalar）。

注意：元素计数在pipeline creation时已被完全确定。

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

• scalar 类型

• vector 类型，带有 concrete 组件

• matrix 类型，带有 concrete 组件

• atomic 类型

• 一个数组类型具有 creation-fixed footprint

• 有 creation-fixed footprint 的 structure 类型 。

注意：元素类型必须是 plain type。

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

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

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

  • 它们都是runtime-sized。

  • 它们都是固定大小的 creation-fixedfootprint，并且 等值元素计数，即使一个是有符号的，另一个是无符号的。（在这种情况下，有符号和无符号值是可比较的，因为元素计数必须大于零。）

  • 它们都是固定大小的，元素计数指定为相同的 pipeline-overridable 常量。

// array<f32,8> and array<i32,8> are different types:
// different element types
var<private> a: array<f32,8>;
var<private> b: array<i32,8>;
var<private> c: array<i32,8u>;  // array<i32,8> and array<i32,8u> are the same type

const width = 8;
const height = 8;

// array<i32,8>, array<i32,8u>, and array<i32,width> are the same type.
// Their element counts evaluate to 8.
var<private> d: array<i32,width>;

// array<i32,height> and array<i32,width> are the same type.
var<private> e: array<i32,width>;
var<private> f: array<i32,height>;

注意：由可覆盖常量确定大小的数组的唯一有效使用是作为 workgroup 空间中变量的存储类型。

override blockSize = 16;

var<workgroup> odds: array<i32,blockSize>;
var<workgroup> evens: array<i32,blockSize>;

 // 一个无效的例子，因为可覆盖的元素计数可能只发生
 // 在外层。
 // var<workgroup> both: array<array<i32,blockSize>,2>;
 // 一个无效的例子，因为可覆盖的元素数量只有
 // 对工作组变量有效。
 // var<private> bad_storage_space: array<i32,blockSize>;

4.3.9. 结构类型

structure是一组已命名的成员值。

一个结构成员类型必须为以下之一：

• 标量（scalar）类型

• 向量vector类型

• 矩阵matrix类型

• 原子（atomic）类型

• 一个 fixed-size array 类型带有 creation-fixedfootprint

• 一个 runtime-sized 数组类型，但前提是它是结构的最后一个成员

• 一个 structure 类型，具有 creation-fixed footprint

注意：任何成员类型必须为普通类型（plain type）。

限制结构成员和数组元素类型的一些后果是：

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

• 当 runtime-sized 数组是较大类型的一部分时，它只能作为结构的最后一个元素出现，它本身不能是封闭数组或结构的一部分。

// 拥有四个成员的结构
struct Data {
  a: i32,
  b: vec2<f32>,
  c: array<i32,10>,
  d: array<f32>, // 最后这个逗号可以不写
};

WGSL定义了以下可应用于结构成员的属性：

• builtin

• location

• align

• size

注意：如果结构类型用于定义统一缓冲区（uniform buffer）或存储缓冲区（storage buffer），则可能需要布局属性。见§ 4.4.7 内存布局。

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

// TODO: runtime-sized array syntax may have changed
// Runtime Array
type RTArr = array<vec4<f32>>;
struct S {
  a: f32,
  b: f32,
  data: RTArr
}
@group(0) @binding(0) var<storage> buffer: S;

4.3.10. 复合类型

如果一个类型内部结构表达为其他类型的复合，则其为复合的（composite）。内部部分不重叠，并称为组件（components）。

复合类型为：

• vector类型

• matrix类型

• array类型

• structure类型

对于一个复合类型T，T的nesting depth，写作NestDepth(T)为：

• 对向量类型，为1

• 对矩阵类型，为2

• 对类型为E的数组类型，1 + NestDepth(E)

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

4.3.11. 可构造类型

多种类型的数值可以被创建，加载，存储，传递至函数，以及作为函数返回值。 我们称之为可构造的（constructible）。

一个类型为可构造的（constructible）如果其为以下之一：

• scalar类型

• vector类型，带有 concrete 组件

• matrix类型，带有 concrete 组件

• 一个 fixed-size array 类型，如果它有 creation-fixedfootprint 并且它的元素类型是可构造的。

• structure类型, 如果它的全部成员都为可构造的。

注意：所有可构造类型为普通类型（plain）并有 creation-fixed footprint。

注意：原子类型和runtime-sized数组类型是不可构造的。包含原子和runtime-sized数组的复合类型是不可构造的。

4.3.12. 固定占用类型

变量的内存占用memory footprint是用于存储变量内容的memory locations的数量。 变量的内存占用取决于它的 store type，并在 shader 生命周期 中的某个时刻最终确定。 大多数变量在 着色器创建 时间很早就确定了大小。 一些变量可能会在 pipeline creation 时间稍后调整大小， 而其他变量可能会在 start of shader execution 时调整大小。

如果 plain type 的大小在 shader creation 时完全确定， 则 plain type 具有创建固定的占用空间creation-fixed footprint。

如果 plain type 的大小在 pipeline creation 时完全确定，则其具有固定占用空间 fixed footprint。

注意：管道的创建依赖于着色器的创建，因此具有 creation-fixedfootprint 的类型也具有 fixedfootprint。

带有 creation-fixedfootprint 的普通类型是：

• scalar 类型

• vector 类型，带有 concrete 组件

• matrix 类型，带有 concrete 组件

• atomic 类型

• fixed-size array 类型, 当:

  • 它的 element count 是文字，或者是 module-scope 的名称 let declaration

• structure 类型，如果它的所有成员都有 creation-fixed footprint。

注意：constructible 类型具有 creation-fixed footprint。

带有 fixed footprint 的普通类型是以下任何一种：

• 具有 creation-fixed footprint 的类型

• 一个 fixed-size array 类型

注意：固定大小数组的唯一有效用途是元素计数是 override expression 而不是 creation-time expression 是 store type for a workgroup 变量。

注意：固定足迹类型可以直接或间接包含 atomic 类型，而 constructible 类型不能。

注意：固定足迹类型不包括 runtime-sized 数组，以及任何包含 runtime-sized 数组的结构或数组，递归。

4.4. 内存

WGSL中，可存储类型的值可以存储在内存中，以供以后检索。本节介绍内存的结构，以及如何使用WGSL类型来描述内存的内容。

4.4.1. 内存位置

内存由一组不同的内存位置（memory locations）组成。每个内存位置的大小为8位。影响内存的操作与一组由一个或多个组成的内存位置交互。

如果两组存储器位置的交集非空，则两组存储器位置重叠（overlap）。每个变量声明都有一组内存位置，不会与任何其他变量声明的内存位置集重叠。对结构和数组的内存操作可以访问元素之间的填充，但不得访问结构或数组末尾的填充。

4.4.2. 内存访问模式

内存访问（memory access）是一种作用于内存位置的操作。

• 读操作（read access）观察内存位置的内容。

• 写操作（write access）设置内存位置的内容。

单个操作可以为读，写，或读与写。

特定的内存位置可能只支持某些类型的访问，表示为内存的访问模式（access mode）：

read

支持读访问，但不是写。

write

支持写访问，但不是读。

read_write

读写访问都支持。

4.4.3. 可存储类型

变量（variable）中包含的值必须是可存储（storable）类型。可存储类型可能具有WGSL定义的显式表示，如§ 4.4.7.4 值的内部布局中所述，或者它可能是不透明的，例如纹理和采样器。

如果一个类型为以下之一，其为可存储的（storable）

• 标量（scalar）类型

• 向量（vector）类型，带有 concrete 组件

• 矩阵（matrix）类型，带有 concrete 组件

• 原子（atomic）类型

• 数组（array）类型

• 结构（structure）类型

• 纹理（texture）类型

• 采样器（sampler）类型

注意：也就是说，可存储类型是普通类型（plain types）、纹理类型和采样器类型。

4.4.4. IO可共享类型

管线输入和输出值必须为IO可共享类型。

一个类型是IO可共享的（IO-shareable）当其为以下之一：

• 标量（scalar）类型

• 数值向量（numeric vector）类型

• 结构（structure）类型，如果其成员全部为标量（scalars）或数值向量（numeric vectors）。

以下类型的值必须是IO可共享类型：

• 从内置变量读取或写入的值。

• 接受作为来自上游管线阶段的输入的值。

• 作为管线中下游处理的输出或输出附件写入的值。

注意：只有内置管线输入可能具有布尔类型。用户输入或输出数据属性不得为bool类型或包含bool类型。见§ 9.3.1 内置输入和输出。

4.4.5. 主机可共享类型

主机可共享类型用于描述在主机和GPU之间共享的缓冲区内容，或者在主机和GPU之间复制而无需格式转换的内容。 用于此目的时，该类型必须额外使用布局属性进行修饰，如§ 4.4.7 内存布局中所述。 我们将在§ 5.3 模块作用域变量中看到统一缓冲区（uniform buffer）和存储缓冲区（storage buffer）变量的存储类型（store type）必须是主机可共享的。

类型为主机可共享的（host-shareable）当其为以下之一：

• 数值标量（numeric scalar）类型

• 数值向量（numeric vector）类型

• 矩阵（matrix）类型

• 原子（atomic）类型

• fixed-size array类型，如果其元素类型为主机可共享的，并且有creation-fixed footprint

• runtime-sized数组类型，如果其元素类型为主机可共享的

• 结构（structure）类型，如果其所有元素为主机可共享的。

WGSL定义了以下影响内存布局的属性：

• stride

• align

• size

注意：当IO-shareable类型T不是bool并且不包含bool时，其为主机可共享的。许多类型为主机可共享的，但不是IO可共享的，包括原子类型（atomic types），runtime-sized数组，以及任何包含它们的复合类型。

注意：IO可共享类型和主机可共享类型都具有具体大小，但是各自计数。IO可共享类型的大小由位置计数指标决定，见§ 9.3.1.3 输入输出位置。主机可共享类型的大小由字节计数指标决定，见§ 4.4.7 内存布局。

4.4.6. 地址空间

内存位置被划分为存储类（address spaces）。每个地址空间都具有确定可变性、可见性、它可能包含的值以及如何使用变量的单独属性。

注意：标记 handle 被保留：它在WGSL程序中从不被使用。

注意：纹理变量包含一个不透明的句柄，用于访问底层的纹素网格。句柄本身始终是只读的。在大多数情况下，底层纹素是只读的。对于只写存储纹理，底层纹素是只写的。

4.4.7. 内存布局

统一缓冲区（Uniform buffer）和存储缓冲区（storage buffer）变量用于共享在内存中组织为字节序列的批量数据。缓冲区在 CPU 和 GPU 之间共享，或在管线中的不同着色器阶段之间，或在不同管线之间共享。

由于缓冲区数据无需重新格式化或转换即可共享，缓冲区生产者和消费者必须就memory layout达成一致， 这是如何将缓冲区中的字节组织成类型化的WGSL值的描述。

缓冲区变量的存储类型（store type）必须是主机可共享的（host-shareable），具有完全详细的内存布局，如下所述。

每个缓冲区变量必须在uniform或storage地址空间中声明。

类型的内存布局仅在评估具有以下内容的表达式时才有意义：

• uniform或storage存储类中的变量，或

• 指向uniform或storage存储类的指针。

一个8位字节是host-shareable内存的最基本单位。本节中定义的术语表示 8 位字节的计数。

我们将使用以下符号：

• AlignOf(T) 是主机可共享类型 T 的 alignment。

• AlignOfMember(S, i) 是主机可共享结构 S 的第 i 个成员的对齐方式。

• SizeOf(T) 是主机可共享类型 T 的 byte-size。

• SizeOfMember(S, i) 是主机可共享结构 S 的第 i 个成员的大小。

• OffsetOfMember(S, i) 是第 i 个成员从主机可共享结构 S 开始的偏移量。

• StrideOf(A) 是主机可共享数组类型A的元素跨度，定义 作为从一个数组元素的开头到下一个元素的开头的字节数。 它等于数组元素类型的大小，向上舍入到元素类型的对齐：

  StrideOf(数组<E, N>) = roundUp(AlignOf(E), SizeOf(E))
  StrideOf(array<E>) = roundUp(AlignOf(E), SizeOf(E))

4.4.7.1. 对齐和大小

每个host-shareable数据类型T都有对齐和大小。

类型的 alignment 是对该类型的值可以放置在内存中何处的约束，以整数表示： 一个类型的对齐方式必须平均划分该类型值的起始 memory location 的字节地址。 对齐允许使用更有效的硬件指令来访问这些值，或者满足对某些地址空间的更严格的硬件要求。 （参见 address space layout constraints）。

注意：根据构造，每个对齐值始终是2的幂。

类型或结构成员的 byte-size 是为存储类型或结构成员的值而保留在主机可共享内存中的连续字节数。

主机可共享类型的对齐方式和大小在下表中递归定义：

4.4.7.2. 结构成员布局

结构 S 的第 i 个成员 具有大小和对齐方式，分别用 SizeOfMember(S, i) 和 AlignOfMember(S, i) 表示。 成员大小和对齐方式用于计算每个成员从结构开始的字节偏移量，如 § 4.4.7.4 值的内部布局 中所述。

SizeOfMember(S, i) 是 k 如果 S 的第 i 个成员 具有属性 size(k)。 否则，它是 SizeOf(T) 其中 T 是成员的类型。

AlignOfMember(S, i) 是 k 如果第 i 个成员具有属性 align(k)。 否则，它是 AlignOf(T) 其中 T 是成员的类型。

如果结构成员使用 size 属性修饰，则该值必须至少与成员类型的大小一样大：

SizeOfMember(S, i) ≥ SizeOf(T)
哪里T 是 S 的第 i 个成员的类型。

第一个结构成员总是从结构的开头有一个零字节偏移：

OffsetOfMember(S, 1) = 0

每个后续成员都放置在满足成员类型对齐的最低偏移量处，并且避免与前一个成员重叠。 对于每个成员索引 i > 1:

OffsetOfMember(S, i) = roundUp(AlignOfMember(S, i ), OffsetOfMember(S, i-1) + SizeOfMember(S, i-1))

struct A {                                     //             align(8)  size(24)
    u: f32,                                    // offset(0)   align(4)  size(4)
    v: f32,                                    // offset(4)   align(4)  size(4)
    w: vec2<f32>,                              // offset(8)   align(8)  size(8)
    x: f32                                    // offset(16)  align(4)  size(4)
    // -- implicit struct size padding --      // offset(20)            size(4)
}

struct B {                           //             align(16) size(160)
    a: vec2<f32>,                              // offset(0)   align(8)  size(8)
    // -- implicit member alignment padding -- // offset(8)             size(8)
    b: vec3<f32>,                              // offset(16)  align(16) size(12)
    c: f32,                                    // offset(28)  align(4)  size(4)
    d: f32,                                    // offset(32)  align(4)  size(4)
    // -- implicit member alignment padding -- // offset(36)            size(4)
    e: A,                                      // offset(40)  align(8)  size(24)
    f: vec3<f32>,                              // offset(64)  align(16) size(12)
    // -- implicit member alignment padding -- // offset(76)            size(4)
    g: array<A, 3>,    //element stride 24     // offset(80)  align(8)  size(72)
    h: i32                                    // offset(152) align(4)  size(4)
    // -- implicit struct size padding --      // offset(156)           size(4)
}

@group(0) @binding(0)
var<storage,read_write> storage_buffer: B;

struct A {                                     //             align(8)  size(32)
    u: f32,                                    // offset(0)   align(4)  size(4)
    v: f32,                                    // offset(4)   align(4)  size(4)
    w: vec2<f32>,                              // offset(8)   align(8)  size(8)
    @size(16) x: f32                       // offset(16)  align(4)  size(16)
};

struct B {                           //             align(16) size(208)
    a: vec2<f32>,                              // offset(0)   align(8)  size(8)
    // -- implicit member alignment padding -- // offset(8)             size(8)
    b: vec3<f32>,                              // offset(16)  align(16) size(12)
    c: f32,                                    // offset(28)  align(4)  size(4)
    d: f32,                                    // offset(32)  align(4)  size(4)
    // -- implicit member alignment padding -- // offset(36)            size(12)
    @align(16) e: A,                        // offset(48)  align(16) size(32)
    f: vec3<f32>,                              // offset(80)  align(16) size(12)
    // -- implicit member alignment padding -- // offset(92)            size(4)
    g: array<A, 3>,    //element stride 32       // offset(96)  align(8)  size(96)
    h: i32                                    // offset(192) align(4)  size(4)
    // -- implicit struct size padding --      // offset(196)           size(12)
}

@group(0) @binding(0)
var<uniform> uniform_buffer: B;

4.4.7.3. 数组布局示例

// Array where:
//   - alignment is 4 = AlignOf(f32)
//   - element stride is 4 = roundUp(AlignOf(f32),SizeOf(f32)) = roundUp(4,4)
//   - size is 32 = stride * number_of_elements = 4 * 8
var small_stride: array<f32, 8>;

// Array where:
//   - alignment is 16 = AlignOf(vec3<f32>) = 16
//   - element stride is 16 = roundUp(AlignOf(vec3<f32>), SizeOf(vec3<f32>))
//                          = roundUp(16,12)
//   - size is 128 = stride * number_of_elements = 16 * 8
var bigger_stride: array<vec3<f32>, 8>;

// Array where:
//   - alignment is 4 = AlignOf(f32)
//   - element stride is 4 = roundUp(AlignOf(f32),SizeOf(f32)) = 4
// 如果 B 是绑定的有效缓冲区绑定大小
// 绘制或调度命令，元素个数为:
//   N_runtime = floor(B / element stride) = floor(B / 4)
@group(0) @binding(0)
var<storage> weights: array<f32>;

// Array where:
//   - alignment is 16 = AlignOf(vec3<f32>) = 16
//   - element stride is 16 = roundUp(AlignOf(vec3<f32>), SizeOf(vec3<f32>))
//                          = roundUp(16,12)
// 如果 B 是绑定的有效缓冲区绑定大小
// 绘制或调度命令，元素个数为:
//   N_runtime = floor(B / element stride) = floor(B / 16)
var<uniform> directions: array<vec3<f32>>;

4.4.7.4. 值的内部布局

本节描述了如何将值的内部结构放置在缓冲区的字节位置中，给定一个假设的整个值的位置。 这些布局取决于值的类型，以及结构成员的 align 和 size 属性。

放置值的缓冲区字节偏移量必须满足类型对齐要求： 如果 T 类型的值 被放置在缓冲区偏移 k，然后 k = c × AlignOf(T)，对于一些非负整数c。

无论地址空间如何，数据都会以相同的方式出现。

当u32或i32类型的值V放置在主机共享缓冲区的字节偏移量k处时，则：

• 字节k包含V的0-7位

• 字节k+1包含V的8-15位

• 字节k+2包含V的16-23位

• 字节k+3包含V的24-31位

注意：回想一下，i32使用二进制补码表示，因此符号位位于第31位。

f32类型的值V以IEEE-754 binary32格式表示。它有1个符号位、8个指数位和23个分数位。当V被放置在主机共享缓冲区的字节偏移 k处时，则：

• 字节k包含分数位的0-7位。

• 字节k+1包含分数位的8-15位。

• 字节k+2的0-6位包含分数位的16-22位。

• 字节k+2的第7位包含指数位的第0位。

• 字节k+3的第0-6位包含指数位的1-7位。

• 字节k+3的第7位包含符号位。

一个值 V f16 类型的以 IEEE-754 binary16 格式表示。 它有 1 个符号位、5 个指数位和 10 个小数位。 当 V 被放置在字节偏移 k 主机共享缓冲区，然后：

• 字节 k 包含分数的 0 到 7 位。

• 字节 k+1 的位 0 到 1 包含小数的位 8 到 9。

• 字节 k+1 的位 2 到 6 包含指数的位 0 到 4。

• 字节k+1 的第7 位包含符号位。

注意：上述规则隐含表明主机共享缓冲区中的数值以little-endian格式存储。

当原子类型（atomic type） atomic<T> 的值V放置在主机共享缓冲区中时，它具有与基础类型T的值相同的内部布局。

当向量类型 vector type vecN<T> 的值 V 放置在主机共享缓冲区的字节偏移量 k 处时，则：

• V.x 被放置在字节偏移量 k 处

• V.y 被放置在字节偏移量 k + SizeOf(T) 处

• 如果 N ≥ 3，则V.z被放置在字节偏移量 k + 2 × SizeOf(T) 处

• 如果 N ≥ 4，则V.w被放置在字节偏移量 k + 3 × SizeOf(T) 处

当一个值 V vector type matCxR<T> 被放置在字节偏移 k 主机共享缓冲区，然后：

• 列向量 i 的 V 被放置在字节偏移 k + i × AlignOf(vecR<T>)

当数组类型 array type A 的值放置在主机共享内存缓冲区的字节偏移量 k 处时，则：

• 数组的元素 i 被放置在字节偏移量 k| + i × StrideOf(A) 处

当结构类型 structure type S 的值被放置在主机共享内存缓冲区的字节偏移量 k 处时，则：

• 结构值的i’^th 成员被放置在字节偏移k 处 + OffsetOfMember(S,i)。 请参阅 § 4.4.7.2 结构成员布局。

4.4.7.5. 地址空间布局约束

地址空间（storage）和（uniform）具有不同的缓冲区布局约束，本节将对此进行介绍。

变量直接或间接引用的所有结构和数组类型都必须遵守变量地址空间的约束。 违反地址空间约束会导致 shader-creation error。

在本节中，我们将 RequiredAlignOf(S, C) 定义为主机可共享类型 S 的值的字节偏移 alignment 要求 在地址空间 C 中使用时。

T 类型的结构成员 必须有一个字节偏移量从结构的开头开始，该结构是地址空间 C 的 RequiredAlignOf(T, C) 的倍数：

OffsetOfMember(S, M) = k × RequiredAlignOf(T, C)
其中 k 为一个正整数，且 M 为具有类型 T 的结构 S 的成员。

元素类型 T 的数组的元素步幅（element stride）必须是对地址空间C 的 RequiredAlignOf(T, C) 的倍数：

StrideOf(array<T, N>) = k × RequiredAlignOf(T, C)
StrideOf(array<T>) = k × RequiredAlignOf(T, C)
其中k 是一个正整数

注意：RequiredAlignOf(T, C) 不会对对齐（align）修饰允许的值施加任何额外限制，也不会影响 AlignOf(T) 的规则。数据按照前面部分中定义的规则进行布局，然后根据 RequiredAlignOf(T, C) 规则验证生成的布局。

统一（uniform）存储类还要求：

• 数组元素对齐到 16 字节边界。也就是说，对于某个正整数 k'，StrideOf(array<T,N>) = 16 × k’。

• 如果结构成员本身具有结构类型 S，则该成员的开头和任何后续成员的开头之间的字节数必须至少为 roundUp(16, SizeOf(S))。

注意：以下示例展示了如何在结构成员上使用 align 和 size 属性来满足统一缓冲区的布局要求。 特别是，可以使用这些技术将具有 std140 布局的 GLSL 缓冲区机械地转换为 WGSL。

struct S {
  x: f32,
}
struct Invalid {
  a: S,
  b: f32 // 无效：a 和 b 之间的偏移量为 4 字节，但必须至少为 16
}
struct Valid {
  a: S,
  [[align(16)]] b: f32 // 有效：a 和 b 之间的偏移量为 16 字节
}
@group(0) @binding(0) var<uniform> invalid: Invalid;
@group(0) @binding(1) var<uniform> valid: Valid;

struct small_stride {
  a: array<f32,8> // stride 4
}
@group(0) @binding(0) var<uniform> invalid: small_stride; // Invalid

struct wrapped_f32 {
  @size(16) elem: f32
}
struct big_stride {
  a: array<wrapped_f32,8> // stride 16
}
@group(0) @binding(1) var<uniform> valid: big_stride;     // Valid

4.5. 内存视图类型

除了使用普通（plain）值进行计算之外，WGSL程序还经常通过内存访问（memory access）操作从内存中读取值或将值写入内存。 每个内存访问都是通过内存视图（memory view）执行的。

内存视图（memory view）由以下内容组成：

• 特定存储类（address space）中的一组内存位置（memory locations），

• 将这些位置的内容解释为 WGSL 类型（type），以及

• 一个访问模式（access mode）。

内存视图的访问方式必须被存储类支持。见§ 4.4.6 地址空间。

WGSL 有两种表示内存视图的类型：引用类型（reference types）和指针类型（pointer types）。

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

但是，在 WGSL 源文本中：

• 引用类型必须不能出现。

• 指针类型可能出现。一个参数化的指针类型通过以下内容拼写：

  • 地址空间address space，

  • 存储类型store type，以及

  • 有时通过访问模式access mode，如§ 4.5.1 访问模式默认值中所述。

fn my_function(
  // 'ptr<function,i32,read_write>' 是指针值的类型，它使用’function' 地址空间中的内存位置来引用内存以保持’i32' 值。
  // 这里 'i32' 是指针类型。
  // 隐含的访问模式为 'read_write' ，访问模式的默认值见下文。
  ptr_int: ptr<function,i32>,

  // 'ptr<private,array<f32,50>,read_write>' 是指针值的类型，
  // 它指的是使用 'private' 存储类中的内存位置保存 50 个类型为 'f32' 元素的数组的存储。
  // 这里的指针类型是’array<f32,50>'。 隐含的访问模式是“read_write”。见下面的访问模式默认值。
  ptr_array: ptr<private, array<f32, 50>>
) { }

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

每个类型ptr<S,T,A>的指针值 p 对应于一个类型ref<S,T,A>的单独参考值r，反之亦然，其中 p 和 r 描述相同的内存视图。

4.5.1. 访问模式默认值

内存视图的访问模式通常由上下文决定：

• storage 存储类支持 read 和 read_write 两种访问模式。

• 每个其他地址空间仅支持一种访问模式，如地址空间表中所述。

当在 WGSL 源码中写入变量声明（variable declaration）或指针类型（pointer type）时：

• 对 storage 地址空间，访问模式是可选的，默认为 read。

• 对其他地址空间，访问模式必须不填写。

4.5.2. 原始变量

在 WGSL 中，参考值始终对应于某些变量的部分或全部内存位置的内存视图。这定义了参考值的原始变量（originating variable）。

一个指针值总是对应一个引用值，因此指针的原始变量与相应引用的原始变量相同。

注意：原始变量是一个动态概念。函数形式参数的原始变量取决于函数的调用位置（call sites）。 不同的调用位置可以提供指向不同起始变量的指针。

如果一个引用或指针访问出界，则产生一个无效内存引用（invalid memory reference）。

从无效引用中加载（Loads）返回以下内容之一：

• 来自绑定到原始变量（originating variable）的 WebGPU 缓冲区（WebGPU buffer）的任何内存位置（memory location(s)）的值

• 引用的存储类型的零值（zero value）。

• 如果加载的值是向量，则值 (0, 0, 0, x)，其中 x 是：

  • 0、1 或最大正值对整数分量

  • 0.0 或 1.0 对浮点分量

• 将值存储到绑定到原始变量的 WebGPU 缓冲区（WebGPU buffer）的任何内存位置（memory location(s)）

• 没有被执行

对无效内存引用进行操作的读-修改-写原子（Read-modify-write atomics）必须从相同的内存位置（memory locations）加载和存储，如果它们访问内存。

4.5.3. 引用和指针用例

引用和指针通过使用方法来区分：

• 一个变量（variable）的类型为一个引用类型。

• 寻址（address-of）操作（unary &）将一个引用值转换为其对应的指针值。

• 间接寻址（indirection）操作（unary *）将指针值转换为其对应的引用值。

• let声明（let declaration）可以为指针类型，但是不能为引用类型。

• 形式参数（formal parameter）可以是指针类型，但是不能为引用类型。

• 赋值语句（simple assignment）执行写访问（write access）以通过引用更新内存内容，其中：

  • 赋值语句的左侧left-hand side必须是引用类型，访问模式为 write 或 read_write。

  • 赋值语句的右侧right-hand side必须计算为左侧的存储类型。

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

  • 在函数中，当在语句或表达式中使用存储类型为 T 的引用表达式 r 时，其中

  • r 具有访问模式 read 或 read_write，且

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

  • 此类型规则的要求视作被满足，且

  • 在该上下文中，表达式 r 的计算结果为存储在计算时被 r 引用的内存位置的值（类型为T）。也就是说，执行读取访问（read access）以产生结果值。

以这种方式定义引用可以简单地惯用变量：

@compute
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 声明，为变量的部分内容形成一个短名称。

• 使用函数的形式参数来引用调用函数（calling function）可访问变量的内存。

  • 对此类函数的调用必须为该操作数提供一个指针值。这通常需要使用寻址（address-of）操作（unary &）来获取指向变量内容的指针。

注意：以下示例使用了本规范稍后解释的 WGSL 功能。

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
fn main() {
  // 在存储内存中形成一个指向特定粒子的指针。
  let active_particle: ptr<storage,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;
}

fn add_one(x: ptr<function,i32>) {
  // 更新 'x' 的位置以包含下一个更高的整数值，（或环绕到最大的负 i32 值）。
  // 在左侧，一元“*”将指针转换为随后可以分配给的引用。
  // 默认情况下，它具有 read_write 访问模式。
  // 在右侧：
  // - 一元'*'将指针转换为引用，具有 read_write 访问模式。
  // - 唯一匹配的类型规则是加法 (+) 并且要求 '*x' 具有类型 i32，这是 '*x' 的存储类型。
  //   所以负载规则应用并且 '*x' 计算为在计算时为 '*x' 存储在内存中的值，即 0 的 i32 值。
  // - 将 1 添加到 0，以生成右侧的最终值 1 - 手边。 将 1 存储到 '*x' 的内存中。
  *x = *x + 1;
}

@compute
fn main() {
  var i: i32 = 0;

  // 修改’i' 的内容，使其包含1。使用一元'&' 获取’i' 的指针值。
  // 这是一个明确的信号，表明被调用的函数可以访问 'i' 的内存，并且可以修改它。
  add_one(&i);
  let one: i32 = i;  // 'one' has value 1.
}

4.5.4. 形成引用和指针值

引用值通过以下方式之一形成：

• 解析为范围内（in-scope）变量 v 的标识符（identifier）表示 v 内存的引用值。

  • 解析成的变量是对引用的原始变量（originating variable）。

• 对一个指针使用间接寻址（indirection (unary *)）操作。

  • 结果的原始变量被定义为指针的原始变量。

• 使用一个复合参考元件表达式（composite reference component expression）。 在每种情况下，结果的原始变量都被定义为原始引用的原始变量。

  • 给定具有向量vector存储类型的引用，附加单字母向量访问短语会产生对向量的命名组件的引用。见§ 6.6.1.3 来自向量引用的分量引用。

  • 给定具有向量vector存储类型的引用，附加数组索引访问短语会产生对向量的索引组件的引用。见§ 6.6.1.3 来自向量引用的分量引用。

  • 给定具有矩阵matrix存储类型的引用，附加数组索引访问短语会产生对矩阵的索引列向量的引用。见§ 6.6.2 矩阵访问表达式。

  • 给定具有数组array存储类型的引用，附加数组索引访问短语会产生对数组索引元素的引用。见§ 6.6.3 数组访问表达式。

  • 给定具有结构structure 存储类型的引用，附加成员访问短语会产生对结构的命名成员的引用。见§ 6.6.4 结构访问表达式。

在所有情况下，结果的访问模式access mode与原始引用的访问模式相同。

struct S {
    age: i32,
    weight: f32
}
var<private> person: S;
// 'person' 的使用表示对变量底层内存的引用，并且类型为 ref<private,S,read_write>。

fn f() {
    var uv: vec2<f32>;
    // 'uv' 的使用表示对变量底层内存的引用，并且类型为 ref<function,vec2<f32>,read_write>。
    // 计算赋值的左侧： 计算 'uv.x' 以产生引用：
    // 1. 首先计算 'uv'，产生对 'uv' 变量内存的引用。结果的类型为 ref<function,vec2<f32>,read_write>。
    // 2. 然后应用'.x’向量访问短语，产生对由上一步中的引用值指向的向量的第一个分量的内存的引用。
    //    结果的类型为 ref<function,f32,read_write>。
    //    计算赋值的右侧会产生 f32 值 1.0。 将 f32 值 1.0 存储到 uv.x 引用的存储内存位置。
    uv.x = 1.0;

    // 计算赋值的左侧： 计算 'uv[1]' 以产生引用：
    // 1. 首先计算 'uv'，产生对 'uv' 变量内存的引用。 结果的类型为 ref<function,vec2<f32>,read_write>。
    // 2. 然后应用'[1]'数组索引短语，产生对上一步引用的向量的第二个组件的存储的引用。
    //    结果的类型为 ref<function,f32,read_write>。
    //    计算赋值的右侧会产生 f32 值 2.0。
    //    将 f32 值 2.0 存储到 uv[1] 引用的存储内存位置。
    uv[1] = 2.0;

    var m: mat3x2<f32>;
    // 计算’m[2]'时：
    // 1. 首先计算’m'，产生对“m”变量内存的引用。结果的类型为 ref<function,mat3x2<f32>,read_write>。
    // 2. 然后应用 '[2]' 数组索引短语，产生对存储由上一步中的参考值指向的第三列向量。
    //    因此，'m[2]' 表达式的类型为 ref<function,vec2<f32>,read_write>。
    //    'let' 声明是针对 vec2<f32> 类型的，因此声明语句要求初始化程序的类型为 vec2<f32>。
    //    加载规则适用（因为没有其他类型规则可以适用），以及初始化器的计算产生 vec2<f32> 值，
    //    该值是在执行声明时从 'm[2]' 引用的内存位置加载的。
    let p_m_col2: vec2<f32> = m[2];

    var A: array<i32,5>;
    // 计算’A[4]'时
    // 1. 首先计算“A”，产生对“A”变量内存的引用。
    //    结果的类型为 ref<function,array<i32,5>,read_write>。
    // 2. 然后应用'[4]'数组索引短语，产生对由上一步中的引用值引用的数组的第五个元素的存储的引用。
    //    结果值的类型为 ref<function,i32,read_write>。 他让声明要求右手边是 i32 类型。
    //    加载规则适用（因为没有其他类型规则可以应用），并且初始化器的计算产生从执行声明时由 'A[4]' 引用的内存位置加载的 i32 值。
    let A_4_value: i32 = A[4];

    // 计算’person.weight’时：
    // 1. 首先计算’person'，产生对在模块范围内声明的“person”变量的内存的引用。
    //    结果的类型为 ref<private,S,read_write>。
    // 2. 然后应用'.weight’成员访问短语，产生对存储器的第二个成员的内存的引用，该存储器由上一步中的引用值引用。
    //    结果的类型为 ref<private,f32,read_write>。
    //    let 声明要求右侧的类型为 f32。
    //    加载规则适用（因为没有其他类型规则可以应用），并且初始化器的计算产生从执行声明时’person.weight’引用的内存位置加载的 f32 值。

    let person_weight: f32 = person.weight;
}

指针值以下列方式之一形成：

• 对引用使用寻址（address-of (unary &)）操作。

  • 结果的原始变量被定义为引用的原始变量。

• 如果函数形式参数（formal parameter）具有指针类型，则在运行时调用该函数时，形参的使用表示在调用函数（calling function）的调用位置（call site）处提供给相应操作数的指针值。

  • 形式参数（在运行时）的原始变量被定义为在调用位置处指针操作数的原始变量。

在所有情况下，结果的访问模式与原始指针的访问模式相同。

// 在私有地址空间中声明一个变量，用于存储 f32 值。
var<private> x: f32;

fn f() {
    // 在函数地址空间中声明一个变量，用于存储 i32 值。
    var y: i32;

    // 名称’x'解析为模块范围变量’x'，并且具有引用类型 ref<private,f32,read_write>。
    // 应用一元“&”运算符将引用转换为指针。
    // 访问方式与原变量的访问方式相同，所以完全指定的类型为ptr<private,f32,read_write>。
    // 但是read_write是函数地址空间的默认访问方式，所以这种情况下不需要拼写read_write
    let x_ptr: ptr<private,f32> = &x;

    // 名称’y'解析为函数范围变量’y'，并且具有引用类型 ref<private,i32,read_write>。
    // 应用一元“&”运算符将引用转换为指针。 访问模式默认为“read_write”。
    let y_ptr: ptr<function,i32> = &y;

    // 一个新变量，不同于在模块范围内声明的变量。
    var x: u32;

    // 此处，名称’x'解析为前一条语句中声明的函数作用域变量’x'，其类型为 ref<function,u32,read_write>。
    // 应用一元'&'运算符将引用转换为指针。 访问模式默认为’read_write'。
    let inner_x_ptr: ptr<function,u32> = &x;
}

4.5.5. 与其他语言中引用和指针的比较

本节是信息性的，是不规范的。

WGSL中的引用和指针相较于其他语言受更多的限制。具体为：

• 在 WGSL 中，不能将引用直接声明为另一个引用或变量的别名，无论是变量还是形式参数。

• 在 WGSL中，指针和引用不可存储（storable）。也就是说，WGSL 变量的内容可能不包含指针或引用。

• 在 WGSL中，函数必须不返回指针或引用。

• 在 WGSL 中，无法在整数值和指针值之间进行转换。

• 在 WGSL中，无法将指针值的类型强制更改为另一种指针类型。

  • 复合组件引用表达式是不同的：它获取对复合值的引用，并产生对复合值内的组件或元素之一的引用。这些在 WGSL 中被认为是不同的引用，即使它们在实现抽象级别的较低级别可能具有相同的机器地址。

• 在 WGSL中，无法将引用值的类型强制更改为另一种引用类型

• 在 WGSL中，无法更改指针或引用的访问模式。

  • 相比之下，C++ 会自动将非常量指针转换为 const 指针，并有一个 const_cast 将一个 const 值转换为一个非常量值。

• 在 WGSL中，无法从“堆”分配新的存储空间。

• 在 WGSL 中，无法显式销毁变量。 只有当变量超出范围时，WGSL 变量的内存才会变得不可访问。

注意：根据上述规则，不可能形成“悬空”指针，即不引用有效的（或“活动的”）原始变量的内存的指针。

4.6. 纹理和采样器类型

纹素（texel）是用作纹理的最小可独立访问元素的标量或向量。texel这个词是纹理元素的缩写。

纹理（texture）是支持对渲染有用的特殊操作的纹素集合。 在 WGSL 中，这些操作是通过纹理内置函数调用的。 有关完整列表，见§ 16.8 纹理内置函数。

WGSL 纹理对应于 WebGPU GPUTexture。

纹理可以是数组形式的，也可以是非数组的：

• 非数组纹理（non-arrayed texture）是一个纹素网格，每个纹素有单独的网格坐标。

• 数组纹理（arrayed texture）是纹素网格的齐次数组。在数组纹理中，每个纹素通过其唯一的数组索引和网格坐标组合来标识。

纹理具有以下特性：

纹素格式

每个纹素中的数据，见§ 4.6.1 纹素格式。

维度

网格坐标中的维数，以及坐标的解释方式。维度数为 1、2 或 3。大多数纹理使用笛卡尔坐标。立方体纹理有六个方形面，并使用三维坐标进行采样，该坐标解释为从原点到以原点为中心的立方体的方向向量。

大小

沿每个维度的网格坐标的范围

mip 级别计数

采样纹理的 mip 级别计数至少为 1，存储纹理的 mip 级别计数至少为 1。Mip level 0 包含纹理的全尺寸版本。每个连续的 mip 级别都包含前一个 mip 级别的过滤版本，大小为前一个 mip 级别的一半（在舍入范围内）。在对纹理进行采样时，使用显式或隐式计算的细节级别来选择从中读取纹素数据的 mip 级别。 然后通过过滤将它们组合起来以产生采样值。

数组化

纹理是否为数组化

数组大小（array size）

齐次网格的数量，如果纹理为数组化的

纹理的表示通常针对渲染操作进行优化。为了实现这一点，许多细节对程序员是隐藏的，包括无法直接用着色器语言表达的数据布局、数据类型和内部操作。

因此，着色器无法直接访问纹理变量中的纹素内存。相反，访问是通过一个不透明的句柄来调解的：

• 在着色器内：

  • 声明一个模块范围变量，其中存储类型（store type）是后面部分中描述的纹理类型之一。 该变量存储底层纹理内存的不透明句柄，并自动放置在句柄（handle）地址空间中。

  • 在函数内部，调用其中一个纹理内置函数，并提供纹理变量作为第一个参数。

• 在构建 WebGPU 管道时，纹理变量的存储类型和绑定必须与相应的绑定组布局条目兼容。

这样，纹理类型支持的操作集由接受该纹理类型作为第一个参数的纹理内置函数的可用性决定。

注意：着色器无法更改纹理变量存储的句柄。也就是说，该变量是只读的，即使它提供访问的底层纹理可能是可变的（例如，只写存储纹理）。

采样器是一个不透明的句柄，它控制如何从采样纹理访问 texels。

WGSL 采样器映射到 WebGPU GPUSampler。

Texel 访问是通过采样器的几个属性控制的：

寻址方式

控制如何解决纹理边界和越界坐标。 每个纹理维度的寻址模式可以独立设置。 参见 WebGPU GPUAddressMode。

过滤模式

控制访问哪些纹素以产生最终结果。 过滤可以使用最近的纹素或在多个纹素之间进行插值。 多种过滤模式可独立设置。 参见 WebGPU GPUFilterMode。

LOD 钳位

控制访问的详细信息的最小和最大级别。

比较

控制为 comparison sampler 所做的比较类型。 参见 WebGPU GPUCompareFunction。

最大各向异性

控制采样器使用的最大各向异性值。

采样器不能在 WGSL 程序中创建，它们的状态（例如上面列出的属性）在着色器中是不可变的，只能由 WebGPU API 设置。

如果过滤采样器（即任何采样器 使用插值过滤）与具有不可过滤格式的纹理一起使用。

注意：着色器无法更改采样器变量存储的句柄。

4.6.1. 纹素格式

在 WGSL 中，某些纹理类型由纹素格式参数化。

纹素格式（texel format）被归类为：

频道（channels）

每个频道包含一个标量。纹素格式最多有四个通道：r、g、b 和 a，通常对应于红色、绿色、蓝色和 alpha 通道的概念。

频道格式（channel format）

通道中的位数，以及如何解释这些位。

WGSL 中的每个纹素格式都对应一个同名的 WebGPU GPUTextureFormat。

WGSL 源代码中仅使用某些纹素格式。用于定义这些纹素格式的频道格式在频道格式（Channel Formats）表中列出。 最后一列指定从存储的频道位到着色器中使用的值的转换。这也称为频道传递函数（channel transfer function）或 CTF。

存储纹理的纹素格式（Texel Formats for Storage Textures）表中列出的纹素格式对应于支持 WebGPU STORAGE 使用的 WebGPU 纯色格式（WebGPU plain color formats）。这些纹素格式用于参数化在§ 4.6.5 存储纹理类型中定义的存储纹理类型。

当纹素格式没有所有四个频道时，则：

• 当读取纹素时：

  • 如果纹素格式没有绿色频道，则着色器值的第二个分量为 0。

  • 如果纹素格式没有蓝色频道，则着色器值的第三个分量为 0。

  • 如果纹素格式没有 Alpha 频道，则着色器值的第四个分量为 1。

• 写入纹素时，缺失频道的着色器值组件将被忽略。

下表中的最后一列使用了频道格式（channel formats）表中特定于格式的频道传输函数（channel transfer function）。

4.6.2. 采样纹理类型

texture_1d<type>

texture_2d<type>

texture_2d_array<type>

texture_3d<type>

texture_cube<type>

texture_cube_array<type>

• 类型必须为 f32, i32 或 u32

• 图像的参数化类型是采样转换后的类型。例如：你可以使用带有 8 位 unorm 组件的纹素图像，但是当您对它们进行采样时，您会得到 32 位浮点结果（或 vec-of-f32）。

4.6.3. 多重采样纹理类型

texture_multisampled_2d<type>

• 类型必须为 f32, i32 或 u32

4.6.4. 外部采样纹理类型

texture_external

texture_external 是一种类似于 texture_2d<f32> 的不透明 2d 浮点采样纹理类型，但可能具有不同的表示。可以使用 textureLoad 或 textureSampleLevel 读取它，它们不透明地处理这些不同的表示。

详见WebGPU § GPUExternalTexture。

4.6.5. 存储纹理类型

存储纹理（storage texture

• 只写存储纹理（A write-only storage texture）支持写入单个纹素，并自动将着色器值转换为存储的纹素值。

存储纹理类型必须由一种存储纹理的纹素格式（texel formats for storage textures）进行参数化。纹素格式决定了 § 4.6.1 纹素格式 中指定的转换函数。

对于只写存储纹理，转换函数的逆函数用于将着色器值转换为存储的纹素。

见§ 16.8 纹理内置函数。

TODO(dentro)：将转换的描述移动到实际进行读取的内置函数中。

texture_storage_1d<texel_format,access>

texture_storage_2d<texel_format,access>

texture_storage_2d_array<texel_format,access>

texture_storage_3d<texel_format,access>

• texel format 必须为存储纹素格式（storage-texel-formats）中指定的纹素类型之一

• access 必须为 write

4.6.6. 深度纹理类型

texture_depth_2d

texture_depth_2d_array

texture_depth_cube

texture_depth_cube_array

texture_depth_multisampled_2d

4.6.7. 采样器类型

采样器（sampler）通过执行以下组合来调节对采样纹理或深度纹理的访问：

• 坐标转换。

• 可选地修改 mip-level 选择。

• 对于采样纹理，可选择过滤检索到的纹素值。

• 对于深度纹理，确定应用于检索到的纹素的比较函数。

采样器在被 WebGPU API 创建时参数化。它们在 WGSL 程序中不能被修改。

采样器仅能通过texture builtin functions被使用。

sampler

sampler_comparison

4.6.8. 纹理类型语法

4.7. 类型别名

type Arr = array<i32, 5>;

type RTArr = array<vec4<f32>>;
const pi_approx: single = 3.1415;
    fn two_pi() -> single {
      return single(2) * pi_approx;
    }

4.8. 类型声明语法

当类型声明是标识符（identifier）时，表达式必须在标识符声明（declaration）的范围内，作为类型别名或结构类型。

identifier
  Allows to specify types created by the type command

bool

f32

i32

u32

vec2<f32>

array<f32, 4>

array<f32>

mat2x3<f32>

// Storage buffers
@group(0) @binding(0)
var<storage,read> buf1: Buffer;       // Can read, cannot write.
@group(0) @binding(0)
var<storage> buf2: Buffer;            // Can read, cannot write.
@group(0) @binding(1)
var<storage,read_write> buf3: Buffer; // Can both read and write.

// Uniform buffer. Always read-only, and has more restrictive layout rules.
struct ParamsTable {weight: f32}
@group(0) @binding(2)
var<uniform> params: ParamsTable;     // Can read, cannot write.

5. 变量和值声明

5.1. 值声明

WGSL 作者可以使用 value declaration 为不可变值声明名称，这些名称可以是：

• let declarations，或者

• override declarations，或者

• creation-time constant

值声明没有任何关联的存储。 也就是说，没有与声明关联的 memory locations。

5.1.1. let 声明

let 声明（let declaration）指定值的名称。一旦计算出 let 声明的值，它就是不可变的。当标识符（identifier） 使用resolves至一个let-声明，该标识符表示该值。

当声明一个 let 标识符时没有明确指定的类型，例如 let foo = 4，类型是从 equals 标记右侧的表达式自动推断的。 let 声明的类型总是 concrete。 当指定类型时，例如 let foo: i32 = 4，初始化表达式必须计算为该类型。

let-声明只能出现在函数定义中。

// 'blockSize' 表示 i32 值 1024。
let blockSize: i32 = 1024;

// 'row_size' 表示 u32 值 16u。类型被指定。
let row_size = 16u;

变量（variable）是对内存的命名引用，可以包含特定可存储（storable）类型的值。

5.1.2. override 声明

覆盖声明 指定 pipeline-overridable 常量值的名称。 pipeline-overridable 常量的值在管道创建时是固定的。

该值是由 WebGPU 管道创建方法指定的值（如果指定），否则是其初始化表达式的值。 当 identifier 将 resolves 用于覆盖声明时，标识符表示该值。 override-声明必须满足以下限制：

• 声明只能出现在 module scope。

• 声明必须至少具有声明类型、初始化表达式或两者之一。

• 声明的类型（如果存在）必须是 scalar。

• 初始化表达式，如果存在，必须：

  • 评估为 scalar 类型。

  • 如果存在，则评估为声明的类型。

    • 仅由 creation-time expressions 或所有标识符 resolve 到可覆盖常量、creation-time constants 或 creation-time functions 的表达式组成。这样的表达式称为override expressio/dfn>。

• 如果声明应用了 id，则文字操作数称为 管道常量 ID，并且必须是 0 到 65535 之间的整数值。

• 管道常量 ID 在 WGSL 程序中必须是唯一的：两个“覆盖”-声明不得使用相同的管道常量 ID。

• 应用程序可以在管道创建时为常量指定自己的值。 管道创建 API 接受从可覆盖常量到常量类型值的映射。 常量由 pipeline-overridable 常量标识符字符串标识，如果指定，它是 pipeline constant ID 的 base-10 表示，否则声明的 name 常数。

• 管道可覆盖常量如果其声明具有初始化表达式，则具有默认值。 如果没有，则必须在管道创建时提供一个值。

注意：覆盖表达式是 creation-time expressions 的超集。

@id(0)    override has_point_light: bool = true;  // Algorithmic control
@id(1200) override specular_param: f32 = 2.3;     // Numeric control
@id(1300) override gain: f32;                     // Must be overridden
          override width: f32 = 0.0;              // Specified at the API level using
                                                  // the name "width".
          override depth: f32;                    // Specified at the API level using
                                                  // the name "depth".
                                                  // Must be overridden.
          override height = 2 * depth;            // The default value
                                                  // (if not set at the API level),
                                                  // depends on another
                                                  // overridable constant.

5.1.3. 创建时间常数

creation-time constant 指定固定在 shader-creation time 的值的名称。 一旦声明了常量，它的值就是不可变的。 当 identifier 将 resolves 用于创建时间常数时，标识符表示该值。

当声明创建时常量时没有明确指定类型， 例如 const foo = 4，类型是从 equals 标记右侧的表达式自动推断的。 创建时间常量的类型必须是：

• 一个 constructible 类型，或

• abstract numeric type，或

• 一个 vector，或

• 一个 matrix

当指定类型时，例如 const foo : i32 = 4，初始化表达式必须计算为该类型。

注意：由于 AbstractInt 和 AbstractFloat 在 WGSL 源代码中无法拼写，因此命名值只能通过类型推断来使用它们。

可以在模块范围或函数范围声明创建时间常量。 必须使用初始化程序声明创建时间常量，并且仅由 creation-time expressions 组成。

const a = 4; // 值为 4 的 AbstractInt。
const b : i32 = 4; // i32 值为 4。
const c : u32 = 4; // u32，值为 4。
const d : f32 = 4; // f32 值为 4。
const e = vec3(a, a, a); // AbstractInt 的 vec3，值为 (4, 4, 4)。
const f = 4.0; // 值为 4 的 AbstractFloat。
const g = mat2x2(a, f, a, f); // 值为 ((2, 4), (2, 4)) 的 AbstractFloat 的 mat2x2。

5.2. var 声明

变量声明（variable declaration）：

• 指定变量名称。

• 指定address spaces, store type, 和 access mode。这些一同构成变量的引用类型（reference type）。

• 确保执行环境在指定的地址空间中为存储类型的值分配内存，对变量的生命周期（lifetime）支持给定的访问模式。

• 如果变量在私有（private）或函数（function）地址空间中，可选择地有一个初始化表达式。 如果存在，初始化表达式必须计算为变量的存储类型。

当标识符（identifier）对变量声明使用resolves时，标识符是表示变量内存的引用内存视图（memory view）的表达式，其类型是变量的引用类型。见§ 6.12 变量标识符表达式。

有关可以在何处声明特定存储类中的变量以及何时需要、可选或禁止存储类修饰的规则，见§ 5.3 模块作用域变量和§ 5.5 函数作用域变量和常量。

访问方式总是有默认的，除了storage存储类中的变量，不得写入WGSL源文本。见§ 4.5.1 访问模式默认值。

一个变量的生命周期（lifetime）为该变量存在时进行着色器处理的期间。模块范围（module scope）变量的生命周期为着色器阶段的全部执行时间。

对于函数作用域（function scope）变量，每个调用都有其自己的独立变量版本。变量的生存周期取决于其作用域：

• 当控制写入变量声明时，它就开始了。

• 它包括从变量作用域内调用的任何函数的整个执行过程。

• 它在控制离开变量的作用域时结束，而不是从变量的作用域内调用函数。

两个生命周期重叠的变量不会有重叠的内存（overlapping memory）。当一个变量的生命周期结束时，它的内存可能会被另一个变量使用。

当一个变量被创建时，它的内存包含一个初始值，如下所示：

• 对于在私有（private）或函数（function）地址空间中的变量：

  • 如果变量声明没有初始化表达式，存储类型为零值（zero value）。

  • 否则它是对初始化表达式在程序执行处计算求值后的结果。

• 对于 workgroup 地址空间中的变量：

  • 当商店类型为 constructible 时，商店类型的 zero value。

  • 否则，商店类型是一个可构造元素的数组，每个元素都被初始化为零值。

• 其他地址空间的变量是通过 draw command 或 dispatch command 中的绑定设置的 resources。

考虑如下WGSL片段：

var i: i32;         // 初始值为 0。不是推荐的方式。
loop {
  var twice: i32 = 2 * i;   // 重新计算每次迭代。
  i = i++;
  if i == 5 { break; }
}

考虑如下WGSL片段：

var x: f32 = 1.0;
let y = x * x + x + 1;

然而，期望浏览器或驱动程序优化此中间表示，从而消除冗余负载。

5.3. 模块作用域变量

在整个程序中，变量名都是in scope。 模块作用域（module scope）内的变量被限制如下：

• 变量不能在函数（function）地址空间中。

• 变量在private, workgroup, uniform, 或 storage 地址空间中：

  • 必须使用显式地址空间修饰声明。

  • 必须使用§ 4.4.6 地址空间中所描述的存储类型（store type）。

• 如果存储类型（store type）是纹理类型或采样器类型，则变量声明不得具有存储类修饰。存储类将始终是句柄（handle）。

uniform 地址空间中的变量是 uniform buffer 变量。 它的 store type 必须是 host-shareable constructible 类型，并且必须满足 地址空间布局约束。

storage 地址空间中的变量是 storage buffer 变量。 它的 store type 必须是 host-shareable 类型并且必须满足 地址空间布局约束。 变量可以用 read 或 read_write 访问模式声明； 默认值为 read。

如§ 9.3.2 资源接口中所述，统一缓冲区、存储缓冲区、纹理和采样器构成了着色器的资源接口（resource interface of a shader）。 此类变量使用组（group）和绑定（binding）修饰声明。

WGSL 定义了以下可应用于全局变量的属性：

• binding

• group

var<private> decibels: f32;
var<workgroup> worklist: array<i32,10>;

struct Params {
  specular: f32,
  count: i32
}
@group(0) @binding(2)
var<uniform> param: Params;    // A uniform buffer

// A storage buffer, for reading and writing
@group(0) @binding(2)
var<storage,read_write> pbuf: array<vec2<f32>>;

// Textures and samplers are always in "handle" space.
@group(0) @binding(2)
var filter_params: sampler;

5.4. 模块常量

出现在所有函数之外的 value declaration 声明了一个 module-scope 常量。

模块范围的常量必须是 override declarations 或 creation-time constants。

整个程序的名称是 in scope。

// The golden ratio.
const golden: f32 = 1.61803398875;

// The second unit vector for three dimensions, with inferred type.
const e2 = vec3<i32>(0,1,0);

当在依赖于常量值的控制流中使用变量或特性时，则认为该变量或特性被程序使用。 无论常量的值如何，无论该值是来自常量声明的值还是来自管线覆盖的值，都是如此。

5.5. 函数作用域变量和常量

在函数体的声明语句中声明的变量或常量在 function scope 中。 该名称可在其声明语句之后立即使用，直到用大括号分隔的语句列表的末尾立即包含该声明。

let-declared 常量必须是 constructible 类型， 或 pointer type。

对于在函数作用域中声明的变量：

• 变量总是在 function 地址空间中。

• 地址空间address space 装饰是可选的。

• 存储类型（store type）必须为可构造（constructible）类型。

• 当指定了初始化表达式时，可以从声明中省略存储类型。在这种情况下，存储类型是初始化表达式的计算结果的类型。

fn f() {
   var<function> count: u32;  // 函数地址空间中的变量。
   var delta: i32;            // 函数地址空间中的另一个变量。
   var sum: f32 = 0.0;        // 拥有初始化值的函数存储类变量。
   var pi = 3.14159;          // 从初始化值推断出 f32 存储类型。
   let unit: i32 = 1;         // 未使用地址空间的let-声明常量。
}

在 for 语句的第一个子句中声明的变量或常量可用于第二个和第三个子句以及 for 语句的主体中。

函数作用域变量的实例是动态上下文（dynamic context）。每个一些调用范围内的变量都有一个重叠的生命周期（lifetime），因此，具有不重叠的内存。生命周期不重叠的变量可能会重用之前变量的内存；但是，不能保证同一变量的新实例使用相同的内存。

5.6. 变量和值声明语法总结

6. 表达式

表达式指定值如何被计算。

6.1. 创建时间表达式

在 shader-creation time 计算的表达式称为 creation-time 表达式。 为了在着色器创建时计算表达式，表达式使用的所有 identifiers 必须 resolve 到 creation-time constants 或 creation-time functions。

创建时表达式的类型可以解析为包含 abstract numeric types 的类型。

示例：(42)分析如下：

• 术语 42 是 AbstractInt 值 42。

• 用括号括住该术语会产生一个新表达式 (42)，其类型为 AbstractInt，值为 42。

示例：-5分析如下：

• 术语“5”是 AbstractInt 值 5。

• 在该术语前面加上 '-' 会产生一个新的表达式 -5，它的类型为 AbstractInt，值为 -5。

示例：-2147483648分析如下：

• 术语“2147483648”是 AbstractInt 值 2147483648。 请注意，此值不适合 32 位有符号整数。

• 在该术语前面加上 '-' 会产生一个新表达式 -2147483648，即 AbstractInt 类型，值为 -2147483648。

示例：const minint = -2147483648;分析如下：

• 如上所述，-2147483648 的计算结果为 AbstractInt 值 -2147483648。

• A creation-time constant 允许初始化器是 abstract numeric type。

• 结果是 minint 被声明为 AbstractInt 值 -2147483648。

示例：let minint = -2147483648;分析如下：

• 如上所述，-2147483648 的计算结果为 AbstractInt 值 -2147483648。

• let 声明 要求初始化程序是 constructible。

• let 声明没有明确的类型，所以使用 overload resolution。 应用的重载候选使用从 AbstractInt 到 i32、u32 或 f32 的 feasible automatic conversions。 排名最低的是 i32，所以 AbstractInt -2147483648 值转换为 i32 值 -2147483648。

• 结果是 minint被声明为i32值-2147483648。

6.2. 文字表达式

6.3. 括号表达式

6.4. 类型构造函数表达式

类型构造函数表达式显式创建给定 constructible 类型的值。 构造函数表达式共有三种：

• § 6.4.1 从组件构造

• § 6.4.2 零值表达式

• § 6.4.3 转换表达式

6.4.1. 从组件构造

本节中定义的表达式通过以下方式创建 constructible 值：

• 复制相同类型的现有值（即恒等函数），或

• 从一个显式的组件列表中创建一个复合值。

这里给出的标量形式是多余的，但提供了标量 conversion expressions 的对称性， 并可用于增强可读性。

向量和矩阵形式从具有匹配分量类型的分量和子向量的各种组合构造向量和矩阵值。 有用于构造向量和矩阵的重载overloads ，它们指定目标类型的维度而无需指定组件类型； 组件类型是从构造函数参数推断出来的。

6.4.2. 零值表达式

每个可构造的（constructible） T 都有一个唯一的零值（zero value），在 WGSL 中编写为，类型后面跟一对空括号：T ()。

零值如下：

• bool() 为 false

• i32() 为 0

• u32() 为 0

• f32() 为 0.0

• f16() is 0.0

• 具有 N 个类型 T 元素的向量的零值为一个 N-元素向量，每个分量值为 T 的零值。

• T 类型的 C 列 R 行矩阵的零值是用 T 的零值填充的那些维度的矩阵。

• 具有可构造（constructible）元素类型 E 的 N-元素向量的零值为一个有 N 个元素的数组，数组的每个分量为 E 的零值。

• 可构造（constructible）类型 S 的零值为构造值 S 的零值成员。

注意：WGSL中原子类型（atomic types），runtime-sized数组，以及其他非可构造（constructible）类型没有零值表达式。

vec2<f32>()                 // 两个 f32 元素的零值向量。
vec2<f32>(0.0, 0.0)         // 相同的值，明确写入。

vec3<i32>()                 // 三个 i32 元素的零值向量。
vec3<i32>(0, 0, 0)          // 相同的值，明确写入。