跳转至

9.3 WGSL

WGSL

WGSL（WebGPU Shader Language）是WebGPU的着色器编程语言。它具有类似于C和JavaScript中的控制结构，但有一些变化和增加。它拥有类似于GLSL中的数据类型和大量内置函数。但是，正如我们在前几节中看到的，它的变量和函数声明方式有显著的不同。

这个相当技术性的章节涵盖了WGSL的语法和语义的主要方面。注意，这里没有涵盖处理纹理的语言部分；它们被推迟到下一部分。一些关于计算着色器的详细信息也被推迟到第9.6节。我假设您已经熟悉像C或JavaScript这样的语言，但如果您需要复习，请参见附录A。熟悉GLSL（第6.3节）也会很有帮助，但不是必需的。虽然我没有给出WGSL语言的完整规范，但我尽量覆盖了大部分重要特性。有关非常长的完整规范，请访问 https://www.w3.org/TR/WGSL/。

由于网络问题，上述网页链接可能无法成功解析。如果用户需要该网页的内容，请告知用户检查网页链接的合法性或稍后再试。如果用户不需要这个链接的内容解析，可以继续回答用户的问题。

WGSL is the shader programming language for WebGPU. It has control structures that are similar to those in C and JavaScript, with some changes and additions. And it has data types and a large set of built in functions that are similar to those in GLSL. But, as we have seen in previous sections, it has significantly different variable and function declarations.

This rather technical section covers major aspects of the syntax and semantics of WGSL. Note that the parts of the language that deal with textures are not covered here; they are postponed until the next section. And some details about working with compute shaders are postponed until Section 9.6. I will assume that you are already familiar with a language like C or JavaScript, but see Appendix A if you need a refresher. Familiarity with GLSL (Section 6.3) would also be useful, but not essential. While I do not give a complete specification of the WGSL language, I try to cover most of the important features. For the very long complete specification, see https://www.w3.org/TR/WGSL/.

9.3.1 地址空间和对齐

Address Spaces and Alignment

要避免在使用WGSL数据值时感到沮丧，您需要理解WGSL中两个不常见于其他编程语言的方面：地址空间和对齐。

GPU可访问的内存被划分为不同的地址空间，这些地址空间具有不同的可访问性规则，并且可能以不同的物理方式被访问。每个变量都存在于特定的地址空间中，而该地址空间是变量类型的一部分。例如，我们已经看到如何使用 var<uniform> 声明一个全局变量。该变量位于 uniform 地址空间，该空间保存的值通常来自程序的JavaScript端。这里是可用的地址空间列表：

  • function 地址空间 — function 地址空间用于函数中的局部变量和参数。它基本上是GPU中单个处理器的函数调用栈，存储在该处理器专用的本地内存中。局部变量可以使用 var<function> 声明，但 function 地址空间是局部变量的唯一选择，它们可以简单地使用 var 声明。
  • private 地址空间 — private 地址空间用于着色器程序中的全局变量，但每个GPU处理器拥有该变量的自己的副本，存储在该处理器专用的本地内存中。作为一个全局变量，使用 var<private> 声明的变量可以在着色器程序的任何函数中使用，但是给定副本的变量仅由同一着色器实例中的函数调用共享。
  • uniform 地址空间 — uniform 地址空间保存所有GPU处理器共享的全局变量。Uniform 变量是只读的。使用 var<uniform> 声明的变量不能包含变量的初始值，着色器也不能为变量分配新值。Uniform 变量中的值是来自绑定组的“资源”，并且每个 uniform 变量声明都必须有 @group 和 @binding 注释，用于指定资源的来源。
  • storage 地址空间 — 存储地址空间与 uniform 空间类似。存储变量需要 @group 和 @binding 注释，并且在着色器程序中不能分配初始值。存储变量默认情况下是只读的，但也可以是读写的。（具有读写访问权限的存储变量可以用于片段和计算着色器，但不能用于顶点着色器。）使用读写访问权限的存储变量使用 var<storage,read_write> 声明。
  • workgroup 地址空间 — 这个地址空间只能在计算着色器中使用，稍后将介绍。

Uniform 和 storage 变量的值来自绑定组。程序的 JavaScript 端使用缓冲区、绑定组和绑定组布局（9.1.3小节）提供它们的值。有一些特定要求：对于 uniform 变量，device.createBuffer() 中缓冲区的 usage 属性必须包含 GPUBufferUsage.UNIFORM，绑定组布局中的缓冲区的 type 属性必须设置为 "uniform"（这是默认值）。在绑定组本身，每个条目的 offset 属性必须是 256 的倍数。这是对齐规则的一个示例。例如，如果有两个 uniform 变量在着色器程序中：

@group(0) @binding(0) var<uniform> x : f32;
@group(0) @binding(1) var<uniform> y : f32;

如果使用一个缓冲区来保存两个变量，那么缓冲区必须至少有 300 字节，绑定组可能是这样的：

bindGroup = device.createBindGroup({
    layout: bindGroupLayout,
    entries: [{
        binding: 0,
        resource: {
            buffer: buffer, offset: 0, size: 4
        }
    },
    {
        binding: 1,
        resource: {
            buffer: buffer, offset: 256, size: 4
        }
    }]
});

对于存储变量，对齐规则是相同的。创建缓冲区时的用法必须包含 GPUBufferUsage.STORAGE。在绑定组布局中，类型必须是 "read-only-storage" 用于默认的只读存储变量，或者 "storage" 用于读写存储变量。

In addition to the alignment rule for uniform and storage bindings, GLSL has alignment rules for its data types. The alignment value for a data type can be 4, 8, or 16. An alignment is always a power of 2. (Alignment 2 is also possible for a 16-bit floating point type that can only be used if a language extension is enabled; 16-bit floats are not covered in this textbook.) If the alignment for a data type is N, then the memory address of any value of that type must be a multiple of N. When the value is part of a data structure, the offset of that value from the start of the data structure must be a multiple of N.

Ordinarily, you will only need to worry about alignment for data structures in the uniform or storage address space. But in that case, knowing the alignment is essential, since you have to take it into account on the JavaScript side when writing data to buffers.

The basic (scalar) data types in WGSL are 4-byte integers and floats, which have alignment 4. WGSL has vectors of 2, 3, and 4 scalar values, which have size 8, 12, and 16. The alignments for 2-vectors and 4-vectors are 8 and 16, as you might expect. But the size of a 3-vector is 12, which is not a legal alignment, so the alignment for 3-vectors is 16. That is, the address in memory of the first byte of a 3-vector must be a multiple of 16.

For an array data structure, the elements of the array must be aligned within the array. This means that in an array of 3-vectors, each element must start at a multiple of 16 bytes from the start of the array. Since a 3-vector such as a vec3f only occupies 12 bytes, four extra bytes of padding must be inserted after each element. No padding is needed in an array of 2-vectors or 4-vectors. So, an array of vec3f takes up just as much memory as an array of vec4f with the same number of elements. The alignment of an array type is equal to the alignment of its elements.

For structs, each element of the struct must satisfy the alignment rule for the data type of that element, which might require padding between some elements. The alignment for the struct itself is the maximum of the alignments of its elements. And the size of the struct must be a multiple of its alignment, which might require some padding at the end.

Let's look at an example that might appear in a shader program that does 3D graphics (see the next section). Some of the syntax has not been covered yet, but you should be able to follow it:

struct LightProperties {
    position : vec4f,      //  size 16,  offset  0
    color : vec3f,         //  size 12,  offset 16 bytes (4 floats)
    spotDirection: vec3f,  //  size 12,  offset 32 bytes (8 floats)
    spotCosineCutoff: f32, //  size  4,  offset 44 bytes (11 floats)
    spotExponent: f32,     //  size  4,  offset 48 bytes (12 floats)
    enabled : f32          //  size  4,  offset 52 bytes (13 floats)
}

@group(0) @binding(0) var<uniform> lights : array<LightProperties,4>

The first vec3f in the struct, color, ends with byte number 27, but the next vec3f, spotDirection, can't start at byte 28 because the alignment rule says that it must start at a multiple of 16. So, four bytes of padding are added. Then, spotDirection starts at byte number 32 and ends with byte number 43. The next element is the 32-bit float spotCosineCutoff, with alignment 4, and it can start at the next byte number, 44. Note that there is no padding after spotDirection. The alignment rule for vec3f does not say that every vec3f is followed by four bytes of padding. Alignment rules are restrictions on where a variable can start. (Yes, this example did trip me up the first time I tried it.)

The array in the example, lights, is an array of four structs of type LightProperties. The alignment for a LightProperties struct is 16 (the maximum of the alignments of its elements). The size, which must be a multiple of the alignment, is 64, with 8 bytes of padding at the end. The size of the array is therefore 256 bytes, or 64 32-bit floats. On the JavaScript side, data for the WGSL array could come from a Float32Array of length 64. When storing values into that Float32Array, you would have to be very careful to take the data alignments into account.

WGSL also has data types for matrices of floating point values. A matrix in WGSL is essentially an array of column vectors, and it follows the same alignment rules. In particular, a matrix with 3 rows is an array of vec3f, with four bytes of padding after each column. This will become important when we work with normal transformation metrics in 3D graphics.

To avoid a lot of frustration when working with WGSL data values, you will need to understand two aspects of WGSL that are not common in other programming languages: address spaces and alignment.

Memory that is accessible to a GPU is divided into address spaces, which have different accessibility rules and which might be physically accessed in different ways. Every variable lives in a particular address space, and that address space is part of the variable's type. For example, we have seen how a global variable can be declared using var<uniform>. That variable lives in the uniform address space, which holds values that generally come from the JavaScript side of the program. Here are the available address spaces:

  • function address space — The function address space is for local variables and parameters in functions. It is basically the function call stack for a single processor in the GPU, which is stored in the dedicated local memory for that processor. Local variables can be declared using var<function>, but the function address space is the only possibility for local variables, and they can declared using simply var.
  • private address space — The private address space is used for global variables in shader programs, but each GPU processor has its own copy of the variable, stored in the dedicated local memory for that processor. As a global variable, a variable declared using var<private> can be used in any function in the shader program, but a given copy of the variable is only shared by function calls in the same invocation of the shader.
  • uniform address space — The uniform address space holds global variables that are shared by all GPU processors. Uniform variables are read-only. A variable declaration using var<uniform> cannot include an initial value for the variable, and a shader cannot assign a new value to the variable. The values in a uniform variable are "resources" that come from a bind group, and every uniform variable declaration must have @group and @binding annotations that are used to specify the source of the resource.
  • storage address space — The storage address space is similar to the uniform space. Storage variables require @group and @binding annotations and cannot be assigned an initial value in the shader program. Storage variables by default are read-only, but read-write access is also possible. (A storage variable with read-write access can be used in fragment and compute shaders, but not in vertex shaders.) A storage variable with read-write access is declared using var<storage,read_write>.
  • workgroup address space — This address space can only be used in compute shaders and will be covered later.

Values for uniform and storage variables come from bind groups. The JavaScript side of the program provides their values using buffers, bind groups, and bind group layouts (Subsection 9.1.3). There are certain requirements: For a uniform variable, the usage property of the buffer in device.createBuffer() must include GPUBufferUsage.UNIFORM, and the buffer in the bind group layout must have its type property set to "uniform" (which is the default). In the bind group itself, the offset property for each entry must be a multiple of 256. This is an example of an alignment rule. For example, if there are two uniform variables in the shader program

@group(0) @binding(0) var<uniform> x : f32;
@group(0) @binding(1) var<uniform> y : f32;

and if one buffer is used to hold both variables, then the buffer must be at least 300 bytes and the bind group would be something like

bindGroup = device.createBindGroup({
layout: bindGroupLayout,
entries: [{
    binding: 0,
    resource: {
        buffer: buffer, offset: 0, size: 4
    }
},
{
    binding: 1,
    resource: {
        buffer: buffer, offset: 256, size: 4
    }
}]
});

For storage variables the alignment rule is the same. The usage when creating the buffer must include GPUBufferUsage.STORAGE. And the type in the bind group layout must be "read-only-storage" for the default read-only storage variables, or "storage" for read-write storage variables.

In addition to the alignment rule for uniform and storage bindings, GLSL has alignment rules for its data types. The alignment value for a data type can be 4, 8, or 16. An alignment is always a power of 2. (Alignment 2 is also possible for a 16-bit floating point type that can only be used if a language extension is enabled; 16-bit floats are not covered in this textbook.) If the alignment for a data type is N, then the memory address of any value of that type must be a multiple of N. When the value is part of a data structure, the offset of that value from the start of the data structure must be a multiple of N.

Ordinarily, you will only need to worry about alignment for data structures in the uniform or storage address space. But in that case, knowing the alignment is essential, since you have to take it into account on the JavaScript side when writing data to buffers.

The basic (scalar) data types in WGSL are 4-byte integers and floats, which have alignment 4. WGSL has vectors of 2, 3, and 4 scalar values, which have size 8, 12, and 16. The alignments for 2-vectors and 4-vectors are 8 and 16, as you might expect. But the size of a 3-vector is 12, which is not a legal alignment, so the alignment for 3-vectors is 16. That is, the address in memory of the first byte of a 3-vector must be a multiple of 16.

For an array data structure, the elements of the array must be aligned within the array. This means that in an array of 3-vectors, each element must start at a multiple of 16 bytes from the start of the array. Since a 3-vector such as a vec3f only occupies 12 bytes, four extra bytes of padding must be inserted after each element. No padding is needed in an array of 2-vectors or 4-vectors. So, an array of vec3f takes up just as much memory as an array of vec4f with the same number of elements. The alignment of an array type is equal to the alignment of its elements.

For structs, each element of the struct must satisfy the alignment rule for the data type of that element, which might require padding between some elements. The alignment for the struct itself is the maximum of the alignments of its elements. And the size of the struct must be a multiple of its alignment, which might require some padding at the end.

Let's look at an example that might appear in a shader program that does 3D graphics (see the next section). Some of the syntax has not been covered yet, but you should be able to follow it:

struct LightProperties {
    position : vec4f,      //  size 16,  offset  0
    color : vec3f,         //  size 12,  offset 16 bytes (4 floats)
    spotDirection: vec3f,  //  size 12,  offset 32 bytes (8 floats)
    spotCosineCutoff: f32, //  size  4,  offset 44 bytes (11 floats)
    spotExponent: f32,     //  size  4,  offset 48 bytes (12 floats)
    enabled : f32          //  size  4,  offset 52 bytes (13 floats)
}

@group(0) @binding(0) var<uniform> lights : array<LightProperties,4>

The first vec3f in the struct, color, ends with byte number 27, but the next vec3f, spotDirection, can't start at byte 28 because the alignment rule says that it must start at a multiple of 16. So, four bytes of padding are added. Then, spotDirection starts at byte number 32 and ends with byte number 43. The next element is the 32-bit float spotCosineCutoff, with alignment 4, and it can start at the next byte number, 44. Note that there is no padding after spotDirection. The alignment rule for vec3f does not say that every vec3f is followed by four bytes of padding. Alignment rules are restrictions on where a variable can start. (Yes, this example did trip me up the first time I tried it.)

The array in the example, lights, is an array of four structs of type LightProperties. The alignment for a LightProperties struct is 16 (the maximum of the alignments of its elements). The size, which must be a multiple of the alignment, is 64, with 8 bytes of padding at the end. The size of the array is therefore 256 bytes, or 64 32-bit floats. On the JavaScript side, data for the WGSL array could come from a Float32Array of length 64. When storing values into that Float32Array, you would have to be very careful to take the data alignments into account.

WGSL also has data types for matrices of floating point values. A matrix in WGSL is essentially an array of column vectors, and it follows the same alignment rules. In particular, a matrix with 3 rows is an array of vec3f, with four bytes of padding after each column. This will become important when we work with normal transformation metrics in 3D graphics.

9.3.2 数据类型

Data Types

WGSL中的基本或“标量”类型包括布尔类型bool，其值为true和false；32位无符号整型类型u32；32位有符号整型类型i32；以及32位浮点类型f32。特别要注意的是，没有8位、16位或64位的数值类型（尽管16位浮点类型f16作为语言扩展是可用的）。

布尔类型bool不是“宿主共享”的，这意味着bool类型的变量不能位于存储或uniform地址空间，也不能从JavaScript端获取其值。这也意味着任何包含bool的的数据结构都不能位于存储或uniform地址空间。

整型字面量可以以通常的十进制形式书写，或者以十六进制形式书写，以0x或0X开头。u32类型的整型字面量以"u"后缀书写，i32类型的以"i"后缀书写。一些例子：17i, 0u, 0xfadeu, 0X7Fi。没有后缀的整型字面量也是可能的；它们被认为是“抽象整数”。奇怪的是，抽象整数可以自动转换为u32、i32或f32，尽管WGSL不会在常规类型之间进行自动转换。因此，如果N是f32类型的变量，那么表达式N+2是合法的，抽象整数2被自动转换为f32。但是表达式N+2u是非法的，因为u32 2u不会自动转换为f32。抽象整数的主要目的似乎是使得能够以更自然的方式书写涉及常量的表达式。

浮点字面量包括小数点、指数或"f"后缀。带有"f"后缀的浮点字面量类型为f32。没有后缀的，它是“抽象浮点数”，可以自动转换为f32类型。例子包括：.0, 17.0, 42f, 0.03e+10f。（也有十六进制浮点字面量，但这里不涉及。）

WGSL具有具有2、3和4个元素的向量类型。向量中的元素可以是任何标量类型：bool、u32、i32或f32。向量类型有官方名称，如vec3<f32>表示包含三个f32值的向量，vec4<bool>表示包含四个bool的向量。但是数值向量的类型名称有更常用的“别名”：vec4f是vec4<f32>的别名，vec2i是vec2<i32>的别名，vec3u是vec3<u32>的别名。

向量类似于数组，可以使用数组表示法引用向量的元素。例如，如果V是vec4f，那么它的元素是V[0]、V[1]、V[2]和V[3]。元素也可以使用swizzlers表示为V.x、V.y、V.z和V.w。通过在点后面使用多个字母，可以构造由V的选定元素组成的向量。例如，V.yx是一个包含V的前两个元素并以相反顺序排列的vec4f，V.zzzz是由V的第三个元素的四份副本组成的vec4f。字母rgba也可以代替xyzw。（这与GLSL类似，6.3.1小节。）

WGSL还有矩阵类型，但只有用于浮点值的矩阵。有2x2、3x3和4x4矩阵的N-by-M类型，官方名称如mat3x2<f32>和mat4x4<f32>。但是这些类型也有更简单的别名，如mat3x2f和mat4x4f。

数组的元素以列主序存储：第一列的元素，然后是第二列的元素，依此类推。每一列是一个向量，列向量可以使用数组表示法访问。例如，如果M是mat4x4f，那么M[1]是M的第二列的vec4f，M[1][0]是该向量的第一个元素。

对于构建数据结构，WGSL有数组和结构体。具有元素类型T和长度N的数组的数据类型是array<T,N>。数组长度必须是常量。没有长度的数组类型也是可能的，但只在存储地址空间中。像通常一样引用数组元素；例如，A[i]。

结构体数据类型包含成员声明列表，可以是不同类型的。参见上面LightProperties类型的示例定义。成员可以是标量、向量、矩阵、数组或结构体。使用通常的点表示法访问成员。例如，如果L是LightProperties类型，那么L.color是L的颜色成员。我将指出，结构体的各个成员可以有注解。例如，

struct VertexOutput {
    @builtin(position) position: vec4f,
    @location(0) color : vec3f
}

WGSL有指针类型，但据我所知，它们只能用于函数定义中的形式参数类型。指针类型名称采用ptr<A,T>的形式，其中A是地址空间名称，T是类型；例如：ptr<function,i32>或ptr<private,array<f32,5>>。类型ptr<A,T>的指针只能指向地址空间A中类型T的值。

如果P是指针，那么*P是它指向的值。如果V是变量，那么&V是指向V的指针。指针类型可用于实现传递给函数的引用。例如，

fn array5sum( A : ptr<function,array<f32,5>> ) -> f32 {
    var sum = 0;
    for (var i = 0; i < 5; i++) {
        sum += (*A)[i];
    }
    return sum;
}

注意使用*A来命名A指向的数组。在(*A)[i]中使用括号是由优先级规则要求的。这个函数可以这样调用：array5sum(&Nums)，其中Nums是函数地址空间中的array<f32,5>类型的变量。（也就是说，Nums必须是局部变量。）

标量类型、向量、矩阵、数组和结构体都是可构造的。也就是说，可以从适当的值列表构造给定类型的值。符号看起来像函数调用，函数名称是类型的名称。这里有一些例子：

var a = u32(23.67f);           // a是23u
var b = f32(a);                // b是23.0f
var c = vec3f(1, 2, 3);        // 抽象整数1,2,3被转换为f32
var d = vec4f(c.xy, 0, 1);     // c.xy向vec4f贡献了两个值
var e = mat2x2f(1, 0, 0, 1);   // 构造2x2单位矩阵
var f = mat3x3f(c, c, c);      // f的每一列是vec3f c
var g = array<u32,4>(1,2,3,4); // 构造长度为4的数组
var h = MyStruct( 17u, 42f );  // MyStruct是由u32和f32组成的结构体
var i = vec4i(2);              // 与vec4i(2,2,2,2)相同；2被重复

The basic, or "scalar," types in WGSL include the boolean type, bool, with values true and false; the 32-bit unsigned integer type, u32; the 32-bit signed integer type, i32; and the 32-bit floating point type, f32. Note in particular that there are no 8-bit, 16-bit, or 64-bit numeric types (although the 16-bit floating point type, f16, is available as a language extension).

The bool type is not "host sharable," which means that a variable of type bool cannot be in the storage or uniform address space, and it can't get its value from the JavaScript side. This also means that any data structure that includes a bool cannot be in the storage or uniform address space.

Literals of integer type can be written in the usual decimal form, or in hexadecimal form with a leading 0x or 0X. An integer literal of type u32 is written with a "u" suffix, and one of type i32 with an "i" suffix. Some examples: 17i, 0u, 0xfadeu, 0X7Fi. Integer literals without suffixes are also possible; they are considered to be "abstract integers." Curiously, an abstract integer can be automatically converted into a u32, i32, or f32, even though WGSL will not do automatic conversions between the regular types. So, if N is a variable of type f32, then the expression N+2 is legal, with the abstract integer 2 being automatically converted into an f32. But the expression N+2u is illegal because the u32 2u is not automatically converted to f32. The main point of abstract integers seems to be to make it possible to write expressions involving constants in a more natural way.

Floating point literals include either a decimal point, or an exponent, or an "f" suffix." A floating point literal with an "f" suffix has type f32. Without the suffix, it is an "abstract float," which can be automatically converted to type f32. Examples include: .0, 17.0, 42f, 0.03e+10f. (There are also hexadecimal floating point literals, but they are not covered here.)

WGSL has vector types with 2, 3, and 4 elements. The elements in a vector can be any scalar type: bool, u32, i32, or f32. The vector types have official names like vec3<f32> for a vector of three f32 values and vec4<bool> for a vector of four bools. But the type names for numeric vectors have "aliases" that are more commonly used instead of the official names: vec4f is an alias for vec4<f32>, vec2i is an alias for vec2<i32>, and vec3u is an alias for vec3<u32>.

Vectors are similar to arrays, and the elements of a vector can be referred to using array notation. For example, if V is a vec4f, then its elements are V[0], V[1], V[2], and V[3]. The elements can also be referred to using swizzlers as V.x, V.y, V.z, and V.w. By using multiple letters after the dot, you can construct vectors made up of selected elements of V. For example, V.yx is a vec4f containing the first two elements of V in reversed order, and V.zzzz is a vec4f made up of four copies of the third element of V. The letters rgba can also be used instead of xyzw. (All this is similar to GLSL, Subsection 6.3.1.)

WGSL also has matrix types, but only for matrices of floating point values. There are types for N-by-M matrices for all a N and M equal to 2, 3, or 4, with official names like mat3x2<f32> and mat4x4<f32>. But again these types have simpler aliases like mat3x2f and mat4x4f.

The elements of an array are stored in column-major order: the elements of the first column, followed by the elements of the second column, and so on. Each column is a vector, and the column vectors can be accessed using array notation. For example, if M is a mat4x4f, then M[1] is the vec4f that is the second column of M, and M[1][0] is the first element of that vector.

For building data structures, WGSL has arrays and structs. The data type for an array with element type T and length N is array<T,N>. The array length must be a constant. Array types without a length are also possible, but only in the storage address space. Array elements are referred to as usual; for example, A[i].

A struct data type contains a list of member declarations, which can be of different types. See, for example, the definition of the LightProperties type, above. A member can be a scalar, a vector, a matrix, an array, or a struct. Members are accessed using the usual dot notation. For example, if L is of type LightProperties, then L.color is the color member of L. I will note that the individual members of a struct can have annotations. For example,

struct VertexOutput {
@builtin(position) position: vec4f,
@location(0) color : vec3f
}

WGSL has pointer types, but as far as I can tell, they can only be used for the types of formal parameters in function definitions. A pointer type name takes the form ptr<A,T>, where A is an address space name and T is a type; for example:ptr<function,i32> or ptr<private,array<f32,5>>. A pointer of type ptr<A,T> can only point to a value of type T in address space A.

If P is a pointer, then *P is the value that it points to. If V is a variable, then &V is a pointer to V. Pointer types can be used to implement pass-by-reference to a function. For example,

fn array5sum( A : ptr<function,array<f32,5>> ) -> f32 {
    var sum = 0;
    for (var i = 0; i < 5; i++) {
        sum += (*A)[i];
    }
    return sum;
}

Note the use of *A to name the array that A points to. The parentheses in (*A)[i] are required by precedence rules. This function could be called as array5sum(&Nums) where Nums is a variable of type array<f32,5> in the function address space. (That is, Nums must be a local variable.)

Scalar types, vectors, matrices, arrays, and structs are constructible. That is, a value of the given type can be constructed from an appropriate list of values. The notation looks like a function call, with the function name being the name of the type. Here are some examples:

var a = u32(23.67f);           // a is 23u
var b = f32(a);                // b is 23.0f
var c = vec3f(1, 2, 3);        // the abstract ints 1,2,3 are converted to f32
var d = vec4f(c.xy, 0, 1);     // c.xy contributes two values to the vec4f
var e = mat2x2f(1, 0, 0, 1);   // constructs the 2-by-2 identity matrix
var f = mat3x3f(c, c, c);      // each column of f is the vec3f c
var g = array<u32,4>(1,2,3,4); // construct an array of length 4
var h = MyStruct( 17u, 42f );  // MyStruct is a struct made of a u32 and an f32
var i = vec4i(2);              // Same as vec4i(2,2,2,2); the 2 is repeated

9.3.3 声明和注解

Declarations and Annotations

我们已经看到如何使用 var<A> 声明变量，其中 A 是地址空间。函数中的局部变量可以使用 var<function> 或者简单地使用 var 来声明。对于全局变量，需要一个地址空间 —— 私有的、uniform、存储的或工作组的（但是与纹理相关的全局变量遵循不同的规则）。

可以在声明中通过在变量名后跟一个冒号然后是类型名称来指定变量的类型。例如：

var<private> sum : f32;

在函数或私有地址空间中的变量声明可以包含变量的初始值。初始值可以是一个常量、一个变量或一个表达式。当声明中包含初始值时，通常不需要指定变量的类型，因为 GLSL 编译器可以从初始值确定类型。当使用抽象整数初始化变量且未指定类型时，类型被视为 i32。

在函数体中，可以使用 let 而不是 var 来声明一个标识符。结果是一个命名值而不是变量。let 声明必须包含初始值。初始化后值不能更改。声明中可以可选地包含类型，但通常不必要。不能指定地址空间。使用 let 明确表示你不期望值会改变，并使意外更改值变得不可能。

也可以使用 const 声明命名值，但是在 const 声明中的初始值必须是已知于编译时的常量。初始值可以作为表达式给出，只要表达式只包含常量。虽然 let 只能在函数中使用，const 声明可以用于任何地方。

一个声明只能声明一个标识符。所以像 "var a = 1, b = 2;" 这样的东西是不合法的。这适用于 const 和 let，以及 var。

我们已经看到，像 @location(0) 这样的注解可以用于变量声明、函数定义、函数形式参数和函数的返回类型。（WGSL 文档称它们为 "attributes"，但我更倾向于保留 "attribute" 这个术语用于顶点属性。）这本教科书只涵盖了最常见的注解。我们在前面的章节中遇到了其中一些，在讨论计算着色器时，稍后会看到更多。常见的注解包括：

  • group(N) 和 @binding(M)，N 和 M 是整数，用于 uniform 和 storage 地址空间的 var 声明中，指定资源的来源。关联是通过绑定组布局指定的。见 9.1.3小节。
  • @vertex、@fragment 和 @compute 用于函数定义，指定该函数可以用作顶点、片段或计算着色器的入口点函数。见 9.1.2小节。
  • @location(N)，N 是一个整数，可以用于顶点着色器和片段着色器入口点函数的输入和输出。它可以应用于它们的形式参数和返回类型，以及用于指定它们的正式参数和返回类型的结构体成员。含义取决于上下文。在顶点着色器入口点的输入上，它指定了顶点缓冲区中输入的来源 (9.1.6小节)。在片段着色器入口点函数的返回类型上，它指定了该输出的目标颜色附件 (9.1.3小节)。当用于顶点着色器输出或片段着色器输入时，它将顶点着色器的特定输出与片段着色器的相应输入关联起来 (9.1.6小节)。
  • @interpolate(flat) 可以应用于顶点着色器入口点函数的输出和片段着色器程序的相应输入。如果它应用于其中一个，那么必须同时应用于两者。通常，片段着色器输入的值是从正在绘制的三角形或线条的所有顶点处的顶点着色器输出值进行插值得到的。@interpolate(flat) 注解关闭了插值；相反，使用第一个顶点的值用于所有片段。这个注解对于整数或布尔值是必需的，也可以应用于浮点值。
  • @builtin(vertex_index) 和 @builtin(instance_index) 用于顶点着色器入口点函数的输入，以指定正在处理的顶点编号或实例编号。见 9.2.4小节。
  • @builtin(position) 当用于顶点着色器入口点函数的输出时，指定输出是顶点在裁剪坐标系中的 (x,y,z,w) 坐标。每个顶点着色器入口点函数都需要有一个带有此注解的输出。当用于片段着色器程序的输入时，它指定输入是正在处理的片段的插值位置，以视口坐标表示。（见 9.4.2小节 讨论 WebGPU 中的坐标系。）
  • @builtin(front_facing) 用于片段着色器程序的布尔类型输入。如果正在处理的片段是前向三角形的一部分，则该值将为 true。这在 3D 图形中进行双面照明时可能很有用 (7.2.4小节)。

We have seen how to declare variables using var<A>, where A is an address space. Local variables in functions can be declared using either var<function> or simply var. For global variables, an address space—private, uniform, storage, or workgroup—is required (but texture-related global variables follow a different rule).

The type of a variable can be specified in a declaration by following the variable name with a colon and then the name of the type. For example

var<private> sum : f32;

The declaration of a variable in the function or private address space can include an initial value for the variable. The initial value can be a constant, a variable, or an expression. When an initial value is included in the declaration, the type of the variable generally does not have to be specified because the GLSL compiler can determine the type from the initial value. When a variable is initialized using an abstract int, and no type is specified, the type is taken to be i32.

In a function body, an identifier can be declared using let instead of var. The result is a named value rather than a variable. A let declaration must include an initial value. The value cannot be changed after initialization. The declaration can optionally include a type, but it is usually not necessary. An address space cannot be specified. Using let makes it clear that you do not expect the value to change and makes it impossible to change the value accidentally.

Named values can also be declared using const, but the initial value in a const declaration must be a constant that is known at compile time. The initial value can be given as an expression, as long as the expression only contains constants. While let can only be used in functions, const declarations can be used anywhere.

A declaration can only declare one identifier. So something like "var a = 1, b = 2;" is not legal. This applies to const and let, as well as to var.

We have seen that annotations like @location(0) can be used on variable declarations, function definitions, function formal parameters, and the return type of a function. (The WGSL documentation calls them "attributes", but I prefer to save the term "attribute" for vertex attributes.) This textbook only covers the most common annotations. We encountered some of them in previous sections, and a few more will come up later when we discuss compute shaders. Common annotations include:

  • group(N) and @binding(M), where N and M are integers, are used on var declarations in the uniform and storage address spaces to specify the source of resource. The association is specified by a bind group layout. See Subsection 9.1.3.
  • @vertex, @fragment, and @compute are used on a function definition to specify that that function can be used as the entry point function for a vertex, fragment, or compute shader. See Subsection 9.1.2.
  • @location(N), where N is an integer, can be used on inputs and outputs of vertex shader and fragment shader entry point functions. It can be applied to their formal parameters and return types and to members of structs that are used to specify the type of their formal parameters and return types. The meaning depends on context. On an input to a vertex shader entry point, it specifies the source of the input in a vertex buffer (Subsection 9.1.6). On the return type of a fragment shader entry point function, it specifies the color attachment that is the destination of that output (Subsection 9.1.3.) And when used on a vertex shader output or a fragment shader input, it associates a particular output of the vertex shader with the corresponding input to the fragment shader (Subsection 9.1.6).
  • @interpolate(flat) can be applied to an output from the vertex shader entry point function and the corresponding input to the fragment shader program. If it is applied to one, it must be applied to both. Usually, the values for a fragment shader input are interpolated from the output values of the vertex shader at all vertices of the triangle or line that is being drawn. The @interpolate(flat) annotation turns off interpolation; instead, the value from the first vertex is used for all fragments. This annotation is required for values of integer or boolean type and can also be applied to floating point values.
  • @builtin(vertex_index) and @builtin(instance_index) are used on inputs to a vertex shader entry point function to specify the vertex number or instance number that is being processed. See Subsection 9.2.4.
  • @builtin(position) when used on an output from a vertex shader entry point function specifies that the output is the (x,y,z,w) coordinates of the vertex in the clip coordinate system. Every vertex shader entry point function is required to have an output with this annotation. When used on an input to a fragment shader program, it specifies that the input is the interpolated position of the fragment being processed, in viewport coordinates. (See Subsection 9.4.2 for a discussion of coordinate systems in WebGPU.)
  • @builtin(front_facing) is used on an input of type bool to a fragment shader program. The value will be true if the fragment that is being processed is part of a front facing triangle. This can be useful, for example, when doing two-sided lighting in 3D graphics (Subsection 7.2.4).

9.3.4 表达式和内置函数

Expressions and Built-in Functions

WGSL包含了所有熟悉的算术、逻辑、位和比较运算符：+, -, *, /, %, &&, ||, !, &, |, ~, ^, <<, >>, ==, !=, <, >, <=, >=。它没有条件运算符?:，但有一个等价的内置函数select(false_case, true_case, boolean)。注意，赋值（=, += 等）不是运算符；也就是说，A = B 是一个语句，而不是一个表达式，它不像在 C 或 JavaScript 中那样有值。

有趣的是，运算符在许多方面被扩展，可以与向量和矩阵一起使用，也可以与标量一起使用。例如，如果 A 是一个 n-by-m 矩阵，B 是一个 m-by-r 矩阵，那么 A*B 计算 A 和 B 的矩阵乘积。如果 V 是一个包含 m 个浮点数的向量，那么 A*V 是矩阵和向量的线性代数乘积得到的向量。

算术运算符也可以应用于两个相同数值类型的向量。操作是逐分量应用的。也就是说，

vec3f(2.0f, 3.0f, 7.0f) / vec3f(5.0f, 8.0f, 9.0f)

结果是 vec3f(2.0f/5.0f, 3.0f/8.0f, 7.0f/9.0f)。相同数值类型的数值向量也可以使用比较运算符进行组合。结果是相同长度的布尔向量。

更有趣的是，算术运算符可以应用于向量和标量。然后，操作应用于向量的每个分量：2+vec2f(5,12) 的结果是 vec2f(7,14)，vec4i(2,5,10,15)/2 的结果是 vec4i(1,2,5,7)。

当然，表达式也可以包括对内置函数和用户定义函数的调用。WGSL有许多内置函数。它有数学函数，如 abs、cos、atan、exp、log 和 sqrt（log 是自然对数）。（除了 abs，参数必须是浮点类型。参数可以是标量或向量。当它是向量时，函数是逐分量应用的：sqrt(vec2f(16.0,9.0)) 的结果是 vec2f(4.0,3.0)。）

有几个内置函数用于对向量执行线性代数运算，包括：length(v) 计算向量 v 的长度；normalize(v) 计算与 v 指向相同方向的单位向量；dot(v,w) 计算 v 和 w 的点积；cross(v,w) 计算两个 3 维向量的叉积；以及 distance(v,w) 计算 v 和 w 之间的距离。在所有情况下，这些函数只适用于浮点向量。有几个函数执行在计算机图形学中常见的操作：

  • clamp(value, min, max) 将值 clamp 限制在 min 到 max 的范围内，即如果 value 在 min 和 max 之间则返回 value，如果 value <= min 则返回 min，如果 value >= max 则返回 max。
  • mix(a, b, blend_factor) 返回 a 和 b 的加权平均值，即返回 (1-blend_factor)*a + blend_factor*b。
  • step(edge, x) 如果 x <= edge 返回 0，如果 x > edge 返回 1。
  • smoothstep(low_edge, high_edge, x) 如果 x < low_edge 返回 0，如果 x > high_edge 返回 1，随着 x 从 low_edge 增加到 high_edge，返回值从 0 平滑增加到 1。
  • reflect(L,N)，其中 L 和 N 是单位向量，计算由法向量 N 反射的向量 L。（见 4.1.4小节，只是那个部分的插图中的 L 指向从表面向光源的方向，但 reflect(L,N) 中的 L 指向从光源向表面的方向。）
  • refract(L,N,ior)，其中 L 和 N 是单位向量，ior 是折射率的比值，计算当光线从方向 L 击中具有法向量 N 的表面时的折射向量，该表面分隔了具有不同折射率的区域。

WGSL has all the familiar arithmetic, logical, bitwise, and comparison operators:+, -, *, /, %, &&, ||, !, &, |, ~, ^, <<, >>, ==, !=, <, >, <=, >=. It does not have the conditional ?: operator, but it has an equivalent built-in function, select(false_case,true_case,boolean). Note that assignment (=, +=, etc.) is not an operator; that is, A = B is a statement, not an expression, and it does not have a value like it would in C or JavaScript.

The interesting thing is that operators are extended in many ways to work with vectors and matrices as well as with scalars. For example, if A is an n-by-m matrix and B is an m-by-r matrix, then A*B computes the matrix product of A and B. And if V is a vector of m floats, then A*V is the vector that is the linear algebra product of the matrix and the vector.

The arithmetic operators can be applied to two vectors of the same numeric type. The operation is applied component-wise. That is,

vec3f(2.0f, 3.0f, 7.0f) / vec3f(5.0f, 8.0f, 9.0f)

is vec3f(2.0f/5.0f, 3.0f/8.0f, 7.0f/9.0f). Numeric vectors of the same numeric type can also be combined using a comparison operator. The result is a bool vector of the same length.

Even more interesting, the arithmetic operators can be applied to a vector and a scalar. The operation then applies to each component of the vector: 2+vec2f(5,12) is vec2f(7,14), and vec4i(2,5,10,15)/2 is vec4i(1,2,5,7).

Expressions, of course, can also include calls to functions, both built-in and user-defined. WGSL has many built-in functions. It has mathematical functions such as abs, cos, atan, exp, log, and sqrt. (log is the natural logarithm.) Except for abs, the parameter must be of floating point type. The parameter can be either a scalar or a vector. When it is a vector, the function is applied component-wise: sqrt(vec2f(16.0,9.0)) is vec2f(4.0,3.0).

There are several built-in functions for doing linear algebra operations on vectors, including: length(v) for the length of vector v; normalize(v) for a unit vector pointing in the same direction as v; dot(v,w) for the dot product of v and w; cross(v,w) for the cross product of two 3-vectors; and distance(v,w) for the distance between v and w. In all cases, these functions only work for vectors of floats. There are several functions that do operations that are common in computer graphics:

  • clamp(value, min, max) clamps value to the range min to max, that is, returns value if value is between min and max, returns min if value <= min, and returns max if value >= max. mix(a, b, blend_factor) returns the weighted average of a and b, that is, returns (1-blend_factor)a + blend_factorb.
  • step(edge, x) returns 0 if x <= edge and 1 if x > edge.
  • smoothstep(low_edge, high_edge, x) returns 0 if x < low_edge, returns 1 if x > high_edge, and the return value increases smoothly from 0 to 1 as x increases from low_edge to high_edge.
  • reflect(L,N), where L and N are unit vectors, computes the vector L reflected by a surface with normal vector N. (See Subsection 4.1.4, except that the L in the illustration in that section points from the surface towards the light source, but the L in reflect(L,N) points from the light source towards the surface.)
  • refract(L,N,ior), where L and N are unit vectors, and ior is the ratio of indices of refraction, computes the refracted vector when light from direction L hits a surface with normal vector N separating regions with different indices of refraction.

9.3.5 语句和控制

Statements and Control

WGSL 中的语句在很大程度上与 C 语言中的类似，但存在一些限制和扩展。

WGSL 中的基本语句包括赋值（使用 =）；复合赋值（使用 +=、*= 等）；递增（使用 ++ 如 x++）；递减（使用 --）；函数调用语句；返回语句；break；continue；以及 discard。递增和递减只能是后缀形式；即允许 x++，但不允许 ++x。并且 —— 像赋值语句一样 —— 递增和递减语句不是表达式；也就是说，它们没有值，不能作为更大表达式的一部分。discard 语句只能在片段着色器入口点函数中使用。它阻止片段着色器的输出被写入其目的地。

就控制结构而言，WGSL 中的 for 循环、while 循环和 if 语句的形式与 C、Java 和 JavaScript 中的形式相同，只不过大括号 {} 始终需要包围循环体和 if 语句内的语句，即使大括号只包围一个单独的语句。break 和 continue 可以在循环中像通常一样使用，但请注意，语句不能有标签，也没有带标签的 break 或 continue 语句。WGSL 中还有一种额外的循环语句形式：

loop {
    statements
}

这种循环通过 break 或 return 语句退出。它基本上与 "while(true)" 循环相同。

WGSL 中的 switch 语句与通常的形式有显著变化。可以合并情况（例如 case 1,2,3）。case 后的冒号是可选的。每个 case 中的代码必须用大括号包围。在没有 break 语句的情况下，一个 case 内的代码不会流入下一个 case，因此 break 语句在 case 中是可选的。然而，break 和 return 仍然可以用于提前结束一个 case。必须有一个 default 情况。switch 表达式必须是 i32 或 u32 类型，所有的 case 常量必须类型相同，或者是抽象整数。示例请参见 webgpu/indices_in_shader.html 中着色器源代码中的 switch 语句。

WGSL 没有异常的概念，也没有 try..catch 语句。

Statements in WGSL are in large part similar to those in C, but there are some restrictions and extensions.

Basic statements in WGSL include assignment (using =); compound assignment (using +=, *=, etc.); increment (using ++ as in x++); decrement (using --); function call statements; return statements; break; continue; and discard. Increment and decrement are postfix only; that is, x++ is allowed, but not ++x. And—like assignment statements—increment and decrement statements are not expressions; that is, they don't have a value and cannot be used as part of a larger expression. The discard statement can only be used in a fragment shader entry point function. It stops the output of the fragment shader from being written to its destination.

As for control structures, for loops, while loops, and if statements in WGSL have the same form as in C, Java, and JavaScript, except that braces, { and }, are always required around the body of a loop and around the statements inside an if statement, even if the braces enclose just a single statement. break and continue can be used in loops as usual, but note that statements cannot have labels and there is no labeled break or labeled continue statement. There is an additional looping statement in WGSL that takes the form

loop {
statements
}

This kind of loop is exited with a break or return statement. It is basically the same as a "while(true)" loop.

The switch statement in WGSL is significantly changed from its usual form. Cases can be combined (case 1,2,3). The colon after a case is optional. The code in each case must be enclosed in braces. There is no fallthrough from one case to the next in the absence of a break statement, so break statements are optional in cases. However, break and return can still be used to end a case early. A default case is required. The switch expression must be of type i32 or u32, and all of the case constants must either be of the same type, or be abstract integers. For an example, see the switch statement in the shader source code in webgpu/indices_in_shader.html.

WGSL does not have the concept of exceptions, and there is no try..catch statement.

9.3.6 函数定义

Function Definitions

我们在 第9.1节 和 第9.2节 中看到了函数定义的例子。这些部分中的所有示例都是带有 @vertex 或 @fragment 注解的着色器入口点函数。在着色器中定义额外的函数是可能的，然后可以按照通常的方式调用这些函数。但请注意，调用入口点函数是不合法的；它们只能由系统作为管线的一部分调用。

我会指出，管线的顶点着色器和片段着色器可以在不同的着色器模块中定义。此外，一个着色器模块可以包含任意数量的着色器入口点。管线使用的入口点函数在管线描述符中指定（见 9.1.3小节）。

使用 fn 后跟函数名，然后是形式参数列表，然后是可选的 -> 和返回类型，最后是函数体（必须用大括号括起来）来定义函数。除了入口点函数外，用户定义的函数可以从同一着色器模块的任何位置调用。

函数有一些限制。不允许递归，无论是直接的还是间接的。没有嵌套：函数定义不能在另一个函数定义内部。数组参数必须有指定的大小。参数的指针类型必须在函数或私有命名空间中。函数名不能被重载；也就是说，你不能有两个同名的函数，即使它们有不同的参数列表。（但一些内置函数是重载的。）同样，函数不能与全局变量同名。

为了结束这一节，这里有一些用户定义的函数：

fn invertedColor(color: vec4f) -> vec4f { // 返回反转的颜色
    return vec4f(1 - color.rgb, color.a);
}

fn grayify(color: ptr<function, vec4f>) { // 就地修改颜色
    let c = *color;
    let gray = c.r * 0.3 + c.g * 0.59 + c.b * 0.11;
    *color = vec4f(gray, gray, gray, c.a);
}

fn min10(A: array<f32, 10>) -> f32 { // 参数按值传递！
    var min = A[0];
    for (var i = 1; i < 10; i++) {
        if (A[i] < min) {
            min = A[i];
        }
    }
    return min;
}

fn simpleLighting(N: vec3f, L: vec3f, V: vec3f, diffuse: vec3f) -> vec3f {
    // N 是单位表面法向量。
    // L 是指向光源的单位向量。
    // V 是指向观察者的单位向量
    if (dot(N, L) <= 0) { // 表面照明的错误面
        return vec3f(0);   // 返回零向量（黑色）
    }
    var color = diffuse * dot(N, L);
    let R = -reflect(L, N);  // 反射光线；
    if (dot(R, V) > 0) { // 添加镜面照明
        // 镜面颜色是灰色，镜面指数是 10
        color += vec3f(0.5) * pow(dot(R, V), 10);
    }
    return color;
}

We have seen examples of function definitions in Section 9.1 and Section 9.2. All of the examples in those sections were shader entry point functions, annotated with @vertex or @fragment. It is possible to define additional functions in a shader, and those functions can then be called in the usual way. Note however that it is not legal to call an entry point function; they can only be called by the system as part of a pipeline.

I will remark that the vertex shader and the fragment shader for a pipeline can be defined in different shader modules. Also, a shader module can contain any number of shader entry points. The entry point functions to be used by a pipeline are specified in a pipeline descriptor (Subsection 9.1.3).

A function is defined using fn followed by the function name, then the formal parameter list, followed optionally by -> and the return type, and finally the function body, which must be enclosed in braces. A user-defined function, other than an entry point function, can be called from anywhere in the same shader module.

There are some restrictions on functions. Recursion, direct or indirect, is not allowed. There is no nesting: a function definition cannot be inside another function definition. Array parameters must have a specified size. Pointer types for parameters must be in the function or private namespace. Function names can't be overloaded; that is, you can't have two functions with the same name, even if they have different parameter lists. (But some of the built-in functions are overloaded.) Also, a function cannot have the same name as a global variable.

To finish this section, here are a few user-defined functions:

fn invertedColor( color : vec4f ) -> vec4f { // return the inverted color
return vec4f( 1 - color.rgb, color.a );
}

fn grayify( color : ptr<function,vec4f> ) { // modify color in place
    let c = *color;
    let gray = c.r * 0.3 + c.g * 0.59 + c.b * 0.11;
    *color = vec4f( gray, gray, gray, c.a );  
}

fn min10( A : array<f32,10> ) -> f32 { // parameter is passed by value!
    var min = A[0];
    for (var i = 1; i < 5; i++) {
    if ( A[i] < min ) {
        min = A[i];
    }
    }
    return min;
}

fn simpleLighting(N : vec3f, L : vec3f, V : vec3f, diffuse : vec3f) -> vec3f {
    // N is the unit surface normal vector.
    // L is the unit vector pointing towards the light.
    // V is the unit vector pointing towards viewer
    if ( dot(N,L) <= 0 ) { // wrong side of surface to be illuminated
        return vec3f(0);   // return the zero vector (black)
    }
    var color = diffuse * dot(N,L);
    let R = -reflect(L,N);  // reflected ray;
    if ( dot(R,V) > 0 ) { // add in specular lighting
        // specular color is gray, specular exponent is 10
    color += vec3f(0.5) * pow(dot(R,V), 10);
    }
    return color;
}