---

9.3.1  Address Spaces and Alignment

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

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

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 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_factor*b.
• 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

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

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