1. Introduction
WebGPU Shading Language (WGSL) is the shader language for [WebGPU]. That is, an application using the WebGPU API uses WGSL to express the programs, known as shaders, that run on the GPU.
1.1. Technical Overview
WebGPU issues a unit of work to the GPU in the form of a GPU command. WGSL is concerned with two kinds of GPU commands:

a draw command executes a render pipeline in the context of inputs, outputs, and attached resources.

a dispatch command executes a compute pipeline in the context of inputs and attached resources.
Both kinds of pipelines use shaders written in WGSL.
A shader is the portion of a WGSL program that executes a shader stage in a pipeline. A shader comprises:

An entry point function.

The transitive closure of all called functions, starting with the entry point. This set includes both userdefined and builtin functions. (For a more rigorous definition, see "functions in a shader stage".)

The set of variables and constants statically accessed by all those functions.

The set of types used to define or analyze all those functions, variables, and constants.
When executing a shader stage, the implementation:

Computes the values of constants declared at modulescope.

Binds resources to variables in the shader’s resource interface, making the contents of those resources available to the shader during execution.

Allocates memory for other modulescope variables, and populates that memory with the specified initial values.

Populates the formal parameters of the entry point, if they exist, with the stage’s pipeline inputs.

Connects the entry point return value, if one exists, to the stage’s pipeline outputs.

Then it invokes the entry point.
A WGSL program is organized into:

Functions, which specify execution behaviour.

Statements, which are declarations or units of executable behaviour.

Literals, which are text representations for pure mathematical values.

Constants, each providing a name for a value computed at a specific time.

Variables, each providing a name for memory holding a value.

Expressions, each of which combines a set of values to produce a result value.

Types, each of which describes:

A set of values.

Constraints on supported expressions.

The semantics of those expressions.

WGSL is an imperative language: behaviour is specified as a sequence of statements to execute. Statements:

Declare constants or variables

Modify the contents of variables

Modify execution order using structured programming constructs:

Selective execution: if/else/else if, switch

Repetition: loop, for

Escaping a nested execution construct: break, continue

Refactoring: function call and return

Discard (fragment shaders only): terminating the invocation and throwing away the output


Evaluate expressions to compute values as part of the above behaviours.
WGSL is statically typed: each value computed by a particular expression is in a specific type, determined only by examining the program source.
WGSL has types to describe booleans, numbers, vectors, matrices, and aggregations of these in the form of arrays and structures. Additional types describe memory.
WGSL does not have implicit conversions or promotions between numeric or boolean types. Converting a value from one numeric or boolean type to another requires an explicit conversion, construction, or reinterpretation of bits. This also applies to vector types.
WGSL has texture and sampler types. Together with their associated builtin functions, these support functionality commonly used for graphics rendering, and commonly provided by GPUs.
The work of a shader stage is partitioned into one or more invocations, each of which executes the entry point, but under slightly different conditions. Invocations in a shader stage share access to certain variables:

All invocations in the stage share the resources in the shader interface.

In a compute shader, invocations in the same workgroup share variables in the workgroup address space. Invocations in different workgroups do not share those variables.
However, the invocations act on different sets of pipeline inputs, including builtin inputs that provide an identifying value to distinguish an invocation from its peers. Also, each invocation has its own independent memory space in the form of variables in the private and function address spaces.
Invocations within a shader stage execute concurrently, and may often execute in parallel. The shader author is responsible for ensuring the dynamic behaviour of the invocations in a shader stage:

Meet the uniformity requirements of certain primitive operations, including texture sampling and control barriers.

Coordinate potentially conflicting accesses to shared variables, to avoid race conditions.
WGSL sometimes permits several possible behaviours for a given feature. This is a portability hazard, as different implementations may exhibit the different behaviours. The design of WGSL aims to minimize such cases, but is constrained by feasibility, and goals for achieving high performance across a broad range of devices.
1.2. Notation
The floor expression is defined over real numbers x:

⌊x⌋ = k, where k is the unique integer such that k ≤ x < k+1
The ceiling expression is defined over real numbers x:

⌈x⌉ = k, where k is the unique integer such that k1 < x ≤ k
The truncate function is defined over real numbers x:

truncate(x) = ⌊x⌋ if x ≥ 0, and ⌈x⌉ if x < 0.
The roundUp function is defined for positive integers k and n as:

roundUp(k, n) = ⌈n ÷ k⌉ × k
The transpose of an ncolumn mrow matrix A is the mcolumn nrow matrix A^{T} formed by copying the rows of A as the columns of A^{T}:

transpose(A) = A^{T}

transpose(A)_{i,j} = A_{j,i}
The transpose of a column vector is defined by interpreting the column vector as a 1row matrix. Similarly, the transpose of a row vector is defined by interpreting the row vector as a 1column matrix.
2. Shader Lifecycle
There are four key events in the lifecycle of a WGSL program and the shaders it may contain. The first two correspond to the WebGPU API methods used to prepare a WGSL program for execution. The last two are the start and end of execution of a shader.
The events are:

Shader module creation

This occurs when the WebGPU createShaderModule method is called. The source text for a WGSL program is provided at this time.


Pipeline creation

This occurs when the WebGPU createComputePipeline method or the WebGPU createRenderPipeline method is invoked. These methods use one or more previously created shader modules, together with other configuration information.


Shader execution start

This occurs when a draw or dispatch command is issued to the GPU, begins executing the pipeline, and invokes the shader stage entry point function.



This occurs when all work in the shader completes:

all its invocations terminate

and all accesses to resources complete

outputs, if any, are passed to downstream pipeline stages.


The events are ordered due to:

data dependencies: shader execution requires a pipeline, and a pipeline requires a shader module.

causality: the shader must start executing before it can finish executing.
2.1. Processing Errors
A WebGPU implementation may fail to process a shader for two reasons:

A program error occurs if the shader does not satisfy the requirements of the WGSL or WebGPU specifications.

An uncategorized error may occur even when all WGSL and WebGPU requirements have been satisfied. Possible causes include:

The shaders are too complex, exceeding the capabilities of the implementation, but in a way not easily captured by prescribed limits. Simplifying the shaders may work around the issue.

A defect in the WebGPU implementation.

A processing error may occur during three phases in the shader lifecycle:

A shadercreation error is an error feasibly detectable at shader module creation time. Detection must rely only on the WGSL program source text and other information available to the
createShaderModule
API method. 
A pipelinecreation error is an error detectable at pipeline creation time. Detection must rely only on the WGSL program source text and other information available to the particular pipeline creation API method.

A dynamic error is an error occurring during shader execution. These errors may or may not be detectable.
Note: For example, a race condition may not be detectable.
Each requirement will be checked at the earliest opportunity. That is:

A shadercreation error results when failing to meet a requirement detectable at shadercreation time.

A pipelinecreation error results when failing to meet a requirement detectable at pipelinecreation time, but not detectable earlier.
When unclear from context, this specification indicates whether failure to meet a particular requirement results in a shadercreation, pipelinecreation, or dynamic error.
The WebGPU specification describes the consequences of each kind of error.
3. Textual Structure
A WGSL program is text. This specification does not prescribe a particular encoding for that text. However, UTF8 is always a valid encoding for a WGSL program.
Note: The intent of promoting UTF8 like this is to simplify interchange of WGSL programs and to encourage interoperability among tools.
WGSL program text consists of a sequence of characters, grouped into contiguous nonempty sets forming:
The program text must not include a null character.
Blankspace is any combination of one or more of the following characters:

space

horizontal tab

linefeed

vertical tab

formfeed

carriage return
To parse a WGSL program:

Remove comments:

Replace the first comment with a space character.

Repeat until no comments remain.


Scanning from beginning to end, group the remaining characters into tokens and blankspace in greedy fashion:

The next group is formed from the longest nonempty prefix of the remaining ungrouped characters, that is either:

a valid token, or

blankspace


Repeat until no ungrouped characters remain.


Discard the blankspace, leaving only tokens.

Parse the token sequence, attempting to match the translation_unit grammar rule.
A shadercreation error results if:

the entire source text cannot be converted into a finite sequence of valid tokens, or

the translation_unit grammar rule does not match the entire token sequence.
3.1. Comments
A comment is a span of text that does not influence the validity or meaning of a WGSL program, except that a comment can separate tokens. Shader authors can use comments to document their programs.
A lineending comment is a kind of comment consisting
of the two characters //
and the characters that follow,
up until but not including:

the next blankspace character other than a space or a horizontal tab, or

the end of the program.
A block comment is a kind of comment consisting of:

The two characters
/*

Then any sequence of:

A block comment, or

Text that does not contain either
*/
or/*


Then the two characters
*/
Note: Block comments can be nested. Since a block comment requires matching start and end text sequences, and allows arbitrary nesting, a block comment cannot be recognized with a regular expression. This is a consequence of the Pumping Lemma for Regular Languages.
let f = 1.5; // This is lineending comment. let g = 2.5; /* This is a block comment that spans lines. /* Block comments can nest. */ But all block comments must terminate. */
3.2. Tokens
A token is a contiguous sequence of characters forming one of:

a literal.

a keyword.

an identifier.
3.3. Literals
A literal is one of:
The form of a numeric literal is defined via patternmatching:
 /((?[09]*\.[09]+?[09]+\.[09]*)((eE)(\+)?[09]+)?f?)(?[09]+(eE)(\+)?[09]+f?)/
 /?0[xX]((([09afAF]*\.[09afAF]+[09afAF]+\.[09afAF]*)((pP)(\+)?[09]+f?)?)([09afAF]+(pP)(\+)?[09]+f?))/
 /?0[xX][09afAF]+0?[19][09]*/
 /0[xX][09afAF]+u0u[19][09]*u/
Note: literals are parsed greedily. This means that for statements like a 5
this will not parse as a
minus
5
but instead as a
5
which
may be unexpected. A space must be inserted after the 
if the first
expression is desired.
TODO(dneto): Describe how numeric literal tokens map to idealized values, and then to typed values.
3.4. Keywords
A keyword is a token which always refers to a predefined language concept. See § 14.1 Keyword Summary for the list of WGSL keywords.
3.5. Identifiers
An identifier is a kind of token used as a name. See § 3.8 Declaration and Scope and § 3.7 Directives.
The form of an identifier is defined via patternmatching, except that an identifier:

must not have the same spelling as a keyword or as a reserved word.

must not be
_
(a single underscore) 
must not start with two underscores.
 /([azAZ_][09azAZ][09azAZ_]*)([azAZ][09azAZ_]*)/
Note: The return type for some builtin functions are structure types whose name cannot be used WGSL source.
Those structure types are described as if they were predeclared with a name starting with two underscores.
The result value can be saved into newly declared let
or var
using type inferencing, or immediately have one of its members
immediately extracted by name. See example usages in the description of frexp
and modf
.
3.6. Attributes
An attribute modifies an object or type. WGSL provides a unified syntax for applying attributes. Attributes are used for a variety of purposes such as specifying the interface with the API. Generally speaking, from the language’s pointofview, attributes can be ignored for the purposes of type and semantic checking.
An attribute must not be specified more than once per object or type.
 attr ident paren_left ( literal_or_ident comma ) * literal_or_ident comma ? paren_right
Attribute  Valid Values  Description 

align
 positive i32 literal 
Must only be applied to a member of a structure type.
Must be a power of 2, and must satisfy the requiredalignment for the member type: If 
binding
 nonnegative i32 literal 
Must only be applied to a resource variable.
Specifies the binding number of the resource in a bind group. See § 9.3.2 Resource Interface. 
builtin
 identifier name for a builtin value 
Must only be applied to an entry point function parameter, entry point
return type, or member of a structure.
Declares a builtin value. See § 15 Builtin Values. 
group
 nonnegative i32 literal 
Must only be applied to a resource variable.
Specifies the binding group of the resource. See § 9.3.2 Resource Interface. 
id
 nonnegative i32 literal 
Must only be applied to an override declaration of scalar type.
Specifies a numeric identifier as an alternate name for a pipelineoverridable constant. 
interpolate

One or two parameters.
The first parameter must be an interpolation type. The second parameter, if present, must specify the interpolation sampling. 
Must only be applied to a declaration that is decorated with a location attribute.
Specifies how the userdefined IO must be interpolated. The attribute is only significant on userdefined vertex outputs and fragment inputs. See § 9.3.1.3 Interpolation. 
invariant
 None 
Must only be applied to the position builtin value.
When applied to the Note: this attribute maps to the 
location
 nonnegative i32 literal 
Must only be applied to an entry point function parameter, entry point
return type, or member of a structure type.
Must only be applied to declarations of numeric scalar or numeric
vector type.
Must not be used with the compute shader stage.
Specifies a part of the userdefined IO of an entry point. See § 9.3.1.4 Inputoutput Locations. 
size
 positive i32 literal 
Must only be applied to a member of a structure type.
The number of bytes reserved in the struct for this member. This number must be at least the bytesize of the type of the member: If 
stage
 compute , vertex , or fragment

Must only be applied to a function declaration.
Declares an entry point by specifying its pipeline stage. 
workgroup_size

One, two or three parameters.
Each parameter is either a literal or modulescope constant. All parameters must be of the same type, either i32 or u32. 
Must be applied to a compute shader entry point function.
Must not be applied to any other object.
Specifies the x, y, and z dimensions of the workgroup grid for the compute shader. The first parameter specifies the x dimension. The second parameter, if provided, specifies the y dimension, otherwise is assumed to be 1. The third parameter, if provided, specifies the z dimension, otherwise is assumed to be 1. Each dimension must be at least 1 and at most an upper bound specified by the WebGPU API. 
3.7. Directives
A directive is a token sequence which modifies how a WGSL program is processed by a WebGPU implementation.
Directives are optional. If present, all directives must appear before any declarations.
3.8. Declaration and Scope
A declaration associates an identifier with one of the following kinds of objects:

a type

a value

a variable

a function

a formal parameter
In other words, a declaration introduces a name for an object.
The scope of a declaration is the set of program locations where a use of the declared identifier potentially denotes its associated object. We say the identifier is in scope (of the declaration) at those source locations.
When an identifier is used, it must be in scope for some declaration, or as part of a directive. When an identifier is used in scope of one or more declarations for that name, the identifier will denote the object of the nonmodulescope declaration appearing closest to that use, or the modulescope declaration if no other declaration is in scope. We say the identifier use resolves to that declaration.
Where a declaration appears determines its scope. Generally, the scope is a span of text beginning immediately after the end of the declaration. Declarations at module scope are the exception, described below.
A declaration must not introduce a name when that identifier is already in scope with the same end of scope as another instance of that name.
Certain objects are provided by the WebGPU implementation, and are treated as if they have been declared by every WGSL program. We say such objects are predeclared. Their scope is the entire WGSL program. Examples of predeclared objects are:

builtin functions, and

builtin types.
A declaration is at module scope if the declaration appears outside the text of any other declaration. Module scope declarations are in scope for the entire program. That is, a declaration at module scope may be referenced by source text that follows or precedes that declaration.
It is a shadercreation error if any module scope declaration is recursive. That is, there must be no cycles among the declarations:
Consider the directed graph where:
Each node corresponds to a declaration D.
There is an edge from declaration D to declaration T when the definition for D mentions an identifier which resolves to T.
This graph must not have a cycle.
Note: The function body is part of the function declaration, thus functions must not be recursive, either directly or indirectly.
Note: Use of a nonmodule scope identifier must follow the declaration of that identifier in the text. This is not true, however, for module scope declarations, which may be referenced out of order in the text.
Note: Only a function declaration can contain other declarations.
// Invalid, cannot reuse builtin function names. var<private> modf: f32 = 0.0; // Valid, foo_1 is in scope for the entire program. var<private> foo: f32 = 0.0; // foo_1 // Valid, bar_1 is in scope for the entire program. var<private> bar: u32 = 0u; // bar_1 // Valid, my_func_1 is in scope for the entire program. // Valid, foo_2 is in scope until the end of the function. fn my_func(foo: f32) { // my_func_1, foo_2 // Any reference to 'foo' resolves to the function parameter. // Invalid, the scope of foo_2 ends at the of the function. var foo: f32; // foo_3 // Valid, bar_2 is in scope until the end of the function. var bar: u32; // bar_2 // References to 'bar' resolve to bar_2 { // Valid, bar_3 is in scope until the end of the compound statement. var bar: u32; // bar_3 // References to 'bar' resolve to bar_3 // Invalid, bar_4 has the same end scope as bar_3. var bar: i32; // bar_4 // Valid, i_1 is in scope until the end of the for loop for ( var i: i32 = 0; i < 10; i++ ) { // i_1 // Invalid, i_2 has the same end scope as i_1. var i: i32 = 1; // i_2. } } // Invalid, bar_5 has the same end scope as bar_2. var bar: u32; // bar_5 // Valid, module scope declarations are in scope for the entire program. var early_use : i32 = later_def; } // Invalid, bar_6 has the same scope as bar_1. var<private> bar: u32 = 1u; // bar_6 // Invalid, my_func_2 has the same end scope as my_func_1. fn my_func() { } // my_func_2 // Valid, my_foo_1 is in scope for the entire program. fn my_foo( //my_foo_1 // Valid, my_foo_2 is in scope until the end of the function. my_foo: i32 // my_foo_2 ) { } var<private> later_def : i32 = 1;
4. Types
Programs calculate values.
In WGSL, a type is set of values, and each value belongs to exactly one type. A value’s type determines the syntax and semantics of operations that can be performed on that value.
For example, the mathematical number 1 corresponds to three distinct values in WGSL:

the 32bit signed integer value
1
, 
the 32bit unsigned integer value
1u
, and 
the 32bit floating point value
1.0
.
WGSL treats these as different because their machine representation and operations differ.
A type is either predeclared, or created in WGSL source via a declaration.
We distinguish between the concept of a type and the syntax in WGSL to denote that type. In many cases the spelling of a type in this specification is the same as its WGSL syntax. For example:

the set of 32bit unsigned integer values is spelled
u32
in this specification, and also in a WGSL program. 
the spelling is different for structure types, or types containing structures.
Some WGSL types are only used for analyzing a source program and for determining the program’s runtime behaviour. This specification will describe such types, but they do not appear in WGSL source text.
Note: WGSL reference types are not written in WGSL programs. See § 4.4 Memory View Types.
4.1. Type Checking
A WGSL value is computed by evaluating an expression.
An expression is a segment of source text
parsed as one of the WGSL grammar rules whose name ends with "_expression
".
An expression E can contain subexpressions which are expressions properly contained
in the outer expression E.
The particular value produced by an expression evaluation depends on:

static context: the source text surrounding the expression, and

dynamic context: the state of the invocation evaluating the expression, and the execution context in which the invocation is running.
The values that may result from evaluating a particular expression will always belong to a specific WGSL type, known as the static type of the expression. The rules of WGSL are designed so that the static type of an expression depends only on the expression’s static context.
Statements often use expressions, and may place requirements on the static types of those expressions. For example:

The condition expression of an
if
statement must be of type bool. 
In a
let
declaration, the initializer must evaluate to the declared type of the constant.
Type checking a successfully parsed WGSL program is the process of mapping each expression to its static type, and determining if the type requirements of each statement are satisfied.
A type assertion is a mapping from some WGSL source expression to a WGSL type. The notation
e : T
is a type assertion meaning T is the static type of WGSL expression e.
Note: A type assertion is a statement of fact about the text of a program. It is not a runtime check.
Finding static types for expressions can be performed by recursively applying type rules. A type rule has two parts:

A conclusion, stated as a type assertion for an expression. The expression in the type assertion is specified schematically, using italicized names to denote subexpressions or other syntacticallydetermined parameters.

Preconditions, consisting of:

Type assertions for subexpressions, when there are subexpressions.

Conditions on the other schematic parameters, if any.

How the expression is used in a statement.

Optionally, other static context.

Each distinct type parameterization for a type rule is called an overload.
For example, unary negation (an expression of the form 
e)
has eight overloads, because its type rules are parameterized by a type T that can be any of:
A type rule applies to an expression when:

The rule’s conclusion matches a valid parse of the expression, and

The rule’s preconditions are satisfied.
TODO: write an example such as 1+2
, or 3  a
, where a
is inscope of a let declaration with i32
type.
The type rules are designed so that if parsing succeeds, at most one type rule will apply to each expression. If a type rule applies to an expression, then the conclusion is asserted, and therefore determines the static type of the expression.
A WGSL source program is welltyped when:

The static type can be determined for each expression in the program by applying the type rules, and

The type requirements for each statement are satisfied.
Otherwise there is a type error and the source program is not a valid WGSL program.
WGSL is a statically typed language because type checking a WGSL program will either succeed or discover a type error, while only having to inspect the program source text.
4.1.1. Type Rule Tables
The WGSL type rules are organized into type rule tables, with one row per type rule.
The semantics of an expression is the effect of evaluating that expression, and is primarily the production of a result value. The Description column of the type rule that applies to an expression will specify the expression’s semantics. The semantics usually depends on the values of the type rule parameters, including the assumed values of any subexpressions. Sometimes the semantics of an expression includes effects other than producing a result value, such as the nonresultvalue effects of its subexpressions.
TODO: example: nonresultvalue effect is any side effect of a function call subexpression.
4.2. Plain Types
Plain types are the types for representing boolean values, numbers, vectors, matrices, or aggregations of such values.
A plain type is either a scalar type, an atomic type, or a composite type.
Note: Plain types in WGSL are similar to PlainOldData types in C++, but also include atomic types.
4.2.1. Boolean Type
The bool type contains the values true
and false
.
4.2.2. Integer Types
The u32 type is the set of 32bit unsigned integers.
The i32 type is the set of 32bit signed integers. It uses a two’s complementation representation, with the sign bit in the most significant bit position.
4.2.3. Floating Point Type
The f32 type is the set of 32bit floating point values of the IEEE754 binary32 (single precision) format. See § 12.5 Floating Point Evaluation for details.
4.2.4. Scalar Types
The scalar types are bool, i32, u32, and f32.
The numeric scalar types are i32, u32, and f32.
The integer scalar types are i32 and u32.
4.2.5. Vector Types
A vector is a grouped sequence of 2, 3, or 4 scalar components.
Type  Description 

vecN<T>  Vector of N components of type T. N must be in {2, 3, 4} and T must be one of the scalar types. We say T is the component type of the vector. 
A vector is a numeric vector if its component type is a numeric scalar.
Key use cases of a vector include:

to express both a direction and a magnitude.

to express a position in space.

to express a color in some color space. For example, the components could be intensities of red, green, and blue, while the fourth component could be an alpha (opacity) value.
Many operations on vectors act componentwise, i.e. the result vector is formed by operating on each component independently.
let x :vec3 < f32 > = a + b ; // a and b are vec3<f32> // x[0] = a[0] + b[0] // x[1] = a[1] + b[1] // x[2] = a[2] + b[2]
4.2.6. Matrix Types
A matrix is a grouped sequence of 2, 3, or 4 floating point vectors.
Type  Description 

matNxM<f32>  Matrix of N columns and M rows, where N and M are both in {2, 3, 4}. Equivalently, it can be viewed as N column vectors of type vecM<f32>. 
The key use case for a matrix is to embody a linear transformation. In this interpretation, the vectors of a matrix are treated as column vectors.
The product operator (*
) is used to either:

scale the transformation by a scalar magnitude.

apply the transformation to a vector.

combine the transformation with another matrix.
See § 6.9 Arithmetic Expressions.
mat2x3 < f32 > // This is a 2 column, 3 row matrix of 32bit floats. // Equivalently, it is 2 column vectors of type vec3<f32>.
4.2.7. Atomic Types
An atomic type encapsulates a scalar type such that:

atomic objects provide certain guarantees to concurrent observers, and

the only valid operations on atomic objects are the atomic builtin functions.
Type  Description 

atomic<T>  Atomic of type T. T must be either u32 or i32. 
An expression must not evaluate to an atomic type.
Atomic types may only be instantiated by variables in the workgroup address space or by storage buffer variables with a read_write access mode.
The memory scope of operations on the type is determined by the address space it is instantiated in.
Atomic types in the workgroup address space have a memory
scope of Workgroup
, while those in the storage address space have a memory scope of QueueFamily
.
An atomic modification is any operation on an atomic object which sets the content of the object. The operation counts as a modification even if the new value is the same as the object’s existing value.
In WGSL, atomic modifications are mutually ordered, for each object. That is, during execution of a shader stage, for each atomic object A, all agents observe the same order of modification operations applied to A. The ordering for distinct atomic objects may not be related in any way; no causality is implied. Note that variables in workgroup space are shared within a workgroup, but are not shared between different workgroups.
4.2.8. Array Types
An array is an indexable grouping of element values.
Type  Description 

array<E,N>  A fixedsize array with N elements of type E. N is called the element count of the array. 
array<E>  A runtimesized array of elements of type E.
These may only appear in specific contexts. 
The first element in an array is at index 0, and each successive element is at the next integer index. See § 6.7.3 Array Access Expression.
An expression must not evaluate to a runtimesized array type.
The element count expression N of a fixedsize array must:

be a literal, or the name of a modulescope constant (possibly pipelineoverridable), and

evaluate to an integer scalar with value greater than zero.
Note: The element count value is fully determined at pipeline creation time.
An array element type must be one of:

a scalar type

a vector type

a matrix type

an atomic type

an array type having a creationfixed footprint

a structure type having a creationfixed footprint.
Note: The element type must be a plain type.
Two array types are the same if and only if all of the following are true:

They have the same element type.

Their element count specifications match, i.e. one of the following is true:

They are both runtimesized.

They are both fixedsized with creationfixed footprint, and equalvalued element counts, even if one is signed and the other is unsigned. (Signed and unsigned values are comparable in this case because element counts must be greater than zero.)

They are both fixedsized with element count specified as the same pipelineoverridable constant.

// 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 , 8 u > ; // array<i32,8> and array<i32,8u> are the same type let width = 8 ; let 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 > ;
Note: The valid use of an array sized by an overridable constant is as the store type of a variable in workgroup space.
override blockSize = 16; var<workgroup> odds: array<i32,blockSize>; var<workgroup> evens: array<i32,blockSize>; // An invalid example, because the overridable element count may only occur // at the outer level. // var<workgroup> both: array<array<i32,blockSize>,2>; // An invalid example, because the overridable element count is only // valid for workgroup variables. // var<private> bad_address_space: array<i32,blockSize>;
 array less_than type_decl ( comma element_count_expression ) ? greater_than
4.2.9. Structure Types
A structure is a grouping of named member values.
Type  Description 

struct<T_{1},...,T_{N}>  An ordered tuple of N members of types T_{1} through T_{N}, with N being an integer greater than 0. A structure type declaration specifies an identifier name for each member. Two members of the same structure type must not have the same name. 
A structure member type must be one of:

a scalar type

a vector type

a matrix type

an atomic type

a fixedsize array type with creationfixed footprint

a runtimesized array type, but only if it is the last member of the structure

a structure type that has a creationfixed footprint
Note: Each member type must be a plain type.
Some consequences of the restrictions structure member and array element types are:

A pointer, texture, or sampler must not appear in any level of nesting within an array or structure.

When a runtimesized array is part of a larger type, it may only appear as the last element of a structure, which itself cannot be part of an enclosing array or structure.
// A structure with four members. struct Data { a :i32 ; b :vec2 < f32 > ; c :array < i32 , 10 > ; d :array < f32 > ; }
WGSL defines the following attributes that can be applied to structure members:
Note: Layout attributes may be required if the structure type is used to define a uniform buffer or a storage buffer. See § 4.3.7 Memory Layout.
// Runtime Array type RTArr = array<vec4<f32>>; struct S { a: f32; b: f32; data: RTArr; } @group(0) @binding(0) var<storage> buffer: S;
4.2.10. Composite Types
A type is composite if it has internal structure expressed as a composition of other types. The internal parts do not overlap, and are called components.
The composite types are:
For a composite type T, the nesting depth of T, written NestDepth(T) is:

1 for a vector type

2 for a matrix type

1 + NestDepth(E) for an array type with element type E

1 + max(NestDepth(M_{1}),..., NestDepth(M_{N})) if T is a structure type with member types M_{1},...,M_{1}
4.2.11. Constructible Types
Many kinds of values can be created, loaded, stored, passed into functions, and returned from functions. We call these constructible.
A type is constructible if it is one of:

a scalar type

a vector type

a matrix type

a fixedsize array type, if it has creationfixed footprint and its element type is constructible.

a structure type, if all its members are constructible.
Note: All constructible types are plain and have creationfixed footprint.
Note: Atomic types and runtimesized array types are not constructible. Composite types containing atomics and runtimesized arrays are not constructible.
4.2.12. FixedFootprint Types
The memory footprint of a variable is the number of memory locations used to store the contents of the variable. The memory footprint of a variable depends on its store type and becomes finalized at some point in the shader lifecycle. Most variables are sized very early, at shader creation time. Some variables may be sized later, at pipeline creation time, and others as late as the start of shader execution.
A plain type has a creationfixed footprint if its size is fully determined at shader creation time.
A plain type has a fixed footprint if its size is fully determined at pipeline creation time.
Note: Pipeline creation depends on shader creation, so a type with creationfixed footprint also has fixed footprint.
The plain types with creationfixed footprint are:

a scalar type

a vector type

a matrix type

an atomic type

a fixedsize array type, when:

its element count is a literal, or the name of a modulescope let declaration


a structure type, if all its members have creationfixed footprint.
Note: A constructible type has creationfixed footprint.
The plain types with fixed footprint are any of:

a type with creationfixed footprint

a fixedsize array type, where its element count is a pipelineoverridable constant.
Note: The only valid use of a fixedsize array with an element count that is a pipelineoverridable constant is as the store type for a workgroup variable.
Note: A fixedfootprint type may contain an atomic type, either directly or indirectly, while a constructible type must not.
Note: Fixedfootprint types exclude runtimesized arrays, and any structures or arrays that contain runtimesized arrays, recursively.
4.3. Memory
In WGSL, a value of storable type may be stored in memory, for later retrieval. This section describes the structure of memory, and how WGSL types are used to describe the contents of memory.
4.3.1. Memory Locations
Memory consists of a set of distinct memory locations. Each memory location is 8bits in size. An operation affecting memory interacts with a set of one or more memory locations.
Two sets of memory locations overlap if the intersection of their sets of memory locations is nonempty. Each variable declaration has a set of memory locations that does not overlap with the sets of memory locations of any other variable declaration. Memory operations on structures and arrays may access padding between elements, but must not access padding at the end of the structure or array.
4.3.2. Memory Access Mode
A memory access is an operation that acts on memory locations.

A read access observes the contents of memory locations.

A write access sets the contents of memory locations.
A single operation can read, write, or both read and write.
Particular memory locations may support only certain kinds of accesses, expressed as the memory’s access mode:
 read

Supports read accesses, but not writes.
 write

Supports write accesses, but not reads.
 read_write

Supports both read and write accesses.
 'read'
 'write'
 'read_write'
4.3.3. Storable Types
The value contained in a variable must be of a storable type. A storable type may have an explicit representation defined by WGSL, as described in § 4.3.7.4 Internal Layout of Values, or it may be opaque, such as for textures and samplers.
A type is storable if it is one of:

a scalar type

a vector type

a matrix type

an atomic type

an array type

a structure type

a texture type

a sampler type
Note: That is, the storable types are the plain types, texture types, and sampler types.
4.3.4. IOshareable Types
Pipeline input and output values must be of IOshareable type.
A type is IOshareable if it is one of:

a scalar type

a numeric vector type

a structure type, if all its members are scalars or numeric vectors
The following kinds of values must be of IOshareable type:

Values read from or written to builtin values.

Values accepted as inputs from an upstream pipeline stage.

Values written as output for downstream processing in the pipeline, or to an output attachment.
Note: Only builtin pipeline inputs may have a boolean type. A user input or output data attribute must not be of bool type or contain a bool type. See § 9.3.1 Pipeline Input and Output Interface.
4.3.5. Hostshareable Types
Hostshareable types are used to describe the contents of buffers which are shared between the host and the GPU, or copied between host and GPU without format translation. When used for this purpose, the type must be additionally decorated with layout attributes as described in § 4.3.7 Memory Layout. We will see in § 5.3 Module Scope Variables that the store type of uniform buffer and storage buffer variables must be hostshareable.
A type is hostshareable if it is one of:

a numeric scalar type

a numeric vector type

a matrix type

an atomic type

a fixedsize array type, if it has creationfixed footprint and its element type is hostshareable

a runtimesized array type, if its element type is hostshareable

a structure type, if all its members are hostshareable
WGSL defines the following attributes that affect memory layouts:
Note: An IOshareable type T is hostshareable if T is not bool and does not contain bool. Many types are hostshareable, but not IOshareable, including atomic types, runtimesized arrays, and any composite types containing them.
Note: Both IOshareable and hostshareable types have concrete sizes, but counted differently. IOshareable types are sized by a locationcount metric, see § 9.3.1.4 Inputoutput Locations. Hostshareable types are sized by a bytecount metric, see § 4.3.7 Memory Layout.
4.3.6. Address spaces
Memory locations are partitioned into address spaces. Each address space has unique properties determining mutability, visibility, the values it may contain, and how to use variables with it.
Address space  Sharing among invocations  Supported access modes  Variable scope  Restrictions on stored values  Notes 

function  Same invocation only  read_write  Function scope  Constructible type  
private  Same invocation only  read_write  Module scope  Constructible type  
workgroup  Invocations in the same compute shader workgroup  read_write  Module scope  Plain type with fixed footprint  The element count of an outermost array may be a pipelineoverridable constant. 
uniform  Invocations in the same shader stage  read  Module scope  Constructible hostshareable types  For uniform buffer variables 
storage  Invocations in the same shader stage  read_write, read (default)  Module scope  Hostshareable  For storage buffer variables 
handle  Invocations in the same shader stage  read  Module scope  Sampler types or texture types  For sampler and texture variables. 
Note: The token handle
is reserved: it is never used in a WGSL program.
Note: A texture variable holds an opaque handle which is used to access the underlying grid of texels. The handle itself is always readonly. In most cases the underlying texels are readonly. For a writeonly storage texture, the underlying texels are writeonly.
4.3.7. Memory Layout
Uniform buffer and storage buffer variables are used to share bulk data organized as a sequence of bytes in memory. Buffers are shared between the CPU and the GPU, or between different shader stages in a pipeline, or between different pipelines.
Because buffer data are shared without reformatting or translation, buffer producers and consumers must agree on the memory layout, which is the description of how the bytes in a buffer are organized into typed WGSL values.
The store type of a buffer variable must be hostshareable, with fully elaborated memory layout, as described below.
Each buffer variable must be declared in either the uniform or storage address spaces.
The memory layout of a type is significant only when evaluating an expression with:
An 8bit byte is the most basic unit of hostshareable memory. The terms defined in this section express counts of 8bit bytes.
We will use the following notation:

AlignOf(T) is the alignment of hostshareable type T.

AlignOfMember(S, i) is the alignment of the i’th member of the hostshareable structure S.

SizeOf(T) is the bytesize of hostshareable type T.

SizeOfMember(S, i) is the size of the i’th member of the hostshareable structure S.

OffsetOfMember(S, i) is the offset of the i’th member from the start of the hostshareable structure S.

StrideOf(A) is the element stride of hostshareable array type A, defined as the number of bytes from the start of one array element to the start of the next element. It equals the size of the array’s element type, rounded up to the alignment of the element type:
StrideOf(array<E, N>) = roundUp(AlignOf(E), SizeOf(E))
StrideOf(array<E>) = roundUp(AlignOf(E), SizeOf(E))
4.3.7.1. Alignment and Size
Each hostshareable data type T has an alignment and size.
The alignment of a type is a constraint on where values of that type may be placed in memory, expressed as an integer: a type’s alignment must evenly divide the byte address of the starting memory location of a value of that type. Alignments enable use of more efficient hardware instructions for accessing the values, or satisfy more restrictive hardware requirements on certain address spaces. (See address space layout constraints).
Note: Each alignment value is always a power of two, by construction.
The bytesize of a type or structure member is the number of contiguous bytes reserved in hostshareable memory for the purpose of storing a value of the type or structure member. The size may include nonaddressable padding at the end of the type. Consequently, loads and stores of a value might access fewer memory locations than the value’s size.
Alignment and size for hostshareable types are defined recursively in the following table:
Hostshareable type T  AlignOf(T)  SizeOf(T) 

i32, u32, or f32  4  4 
atomic<T>  4  4 
vec2<T>  8  8 
vec3<T>  16  12 
vec4<T>  16  16 
matNxM (colmajor) (General form)  AlignOf(vecM)  SizeOf(array<vecM, N>) 
mat2x2<f32>  8  16 
mat3x2<f32>  8  24 
mat4x2<f32>  8  32 
mat2x3<f32>  16  32 
mat3x3<f32>  16  48 
mat4x3<f32>  16  64 
mat2x4<f32>  16  32 
mat3x4<f32>  16  48 
mat4x4<f32>  16  64 
struct S with members M_{1}...M_{N}  max(AlignOfMember(S,1), ... , AlignOfMember(S,N))  roundUp(AlignOf(S), justPastLastMember) where justPastLastMember = OffsetOfMember(S,N) + SizeOfMember(S,N) 
array<E, N>  AlignOf(E)  N × roundUp(AlignOf(E), SizeOf(E)) 
array<E>  AlignOf(E)  N_{runtime} × roundUp(AlignOf(E),SizeOf(E)) where N_{runtime} is the runtimedetermined number of elements of T 
4.3.7.2. Structure Member Layout
The i’th member of structure S has a size and alignment, denoted by SizeOfMember(S, i) and AlignOfMember(S, i), respectively. The member sizes and alignments are used to calculate each member’s byte offset from the start of the structure, as described in § 4.3.7.4 Internal Layout of Values.
SizeOfMember(S, i) is k if the i’th member of S has attribute size(k). Otherwise, it is SizeOf(T) where T is the type of the member.
AlignOfMember(S, i) is k if the i’th member has attribute align(k). Otherwise, it is AlignOf(T) where T is the type of the member.
If a structure member is decorated with the size attribute, the value must be at least as large as the size of the member’s type:
SizeOfMember(S, i) ≥ SizeOf(T)
Where T is the type of the i’th member of S.
The first structure member always has a zero byte offset from the start of the structure:
OffsetOfMember(S, 1) = 0
Each subsequent member is placed at the lowest offset that satisfies the member type alignment, and which avoids overlap with the previous member. For each member index i > 1:
OffsetOfMember(S, i) = roundUp(AlignOfMember(S, i ), OffsetOfMember(S, i1) + SizeOfMember(S, i1))
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.3.7.3. Array Layout Examples
// 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 // If B is the effective buffer binding size for the binding on the // draw or dispach command, the number of elements is: // 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) // If B is the effective buffer binding size for the binding on the // draw or dispach command, the number of elements is: // N_runtime = floor(B / element stride) = floor(B / 16) var < uniform > directions :array < vec3 < f32 >> ;
4.3.7.4. Internal Layout of Values
This section describes how the internals of a value are placed in the byte locations of a buffer, given an assumed placement of the overall value. These layouts depend on the value’s type, and the align and size attributes on structure members.
The buffer byte offset at which a value is placed must satisfy the type alignment requirement: If a value of type T is placed at buffer offset k, then k = c × AlignOf(T), for some nonnegative integer c.
The data will appear identically regardless of the address space.
When a value V of type u32 or i32 is placed at byte offset k of a hostshared buffer, then:

Byte k contains bits 0 through 7 of V

Byte k+1 contains bits 8 through 15 of V

Byte k+2 contains bits 16 through 23 of V

Byte k+3 contains bits 24 through 31 of V
Note: Recall that i32 uses twoscomplement representation, so the sign bit is in bit position 31.
A value V of type f32 is represented in IEEE754 binary32 format. It has one sign bit, 8 exponent bits, and 23 fraction bits. When V is placed at byte offset k of hostshared buffer, then:

Byte k contains bits 0 through 7 of the fraction.

Byte k+1 contains bits 8 through 15 of the fraction.

Bits 0 through 6 of byte k+2 contain bits 16 through 22 of the fraction.

Bit 7 of byte k+2 contains bit 0 of the exponent.

Bits 0 through 6 of byte k+3 contain bits 1 through 7 of the exponent.

Bit 7 of byte k+3 contains the sign bit.
Note: The above rules imply that numeric values in hostshared buffers are stored in littleendian format.
When a value V of atomic type atomic
<T> is placed in a hostshared buffer,
it has the same internal layout as a value of the underlying type T.
When a value V of vector type vecN<T> is placed at byte offset k of a hostshared buffer, then:

V.x is placed at byte offset k

V.y is placed at byte offset k+4

If N ≥ 3, then V.z is placed at byte offset k+8

If N ≥ 4, then V.w is placed at byte offset k+12
When a matrix value M is placed at byte offset k of a hostshared memory buffer, then:

If M has 2 rows, then:

Column vector i of M is placed at byte offset k + 8 × i


If M has 3 or 4 rows, then:

Column vector i of M is placed at byte offset k + 16 × i

When a value of array type A is placed at byte offset k of a hostshared memory buffer, then:

Element i of the array is placed at byte offset k + i × StrideOf(A)
When a value of structure type S is placed at byte offset k of a hostshared memory buffer, then:

The i’^{th} member of the structure value is placed at byte offset k + OffsetOfMember(S,i). See § 4.3.7.2 Structure Member Layout.
4.3.7.5. Address Space Layout Constraints
The storage and uniform address spaces have different buffer layout constraints which are described in this section.
All structure and array types directly or indirectly referenced by a variable must obey the constraints of the variable’s address space. Violations of an address space constraint results in a shadercreation error.
In this section we define RequiredAlignOf(S, C) as the byte offset alignment requirement of values of hostshareable type S when used in address space C.
Hostshareable type S  RequiredAlignOf(S, storage)  RequiredAlignOf(S, uniform) 

i32, u32, or f32  AlignOf(S)  AlignOf(S) 
atomic<T>  AlignOf(S)  AlignOf(S) 
vecN<T>  AlignOf(S)  AlignOf(S) 
matNxM<f32>  AlignOf(S)  AlignOf(S) 
array<T, N>  AlignOf(S)  roundUp(16, AlignOf(S)) 
array<T>  AlignOf(S)  roundUp(16, AlignOf(S)) 
struct S  AlignOf(S)  roundUp(16, AlignOf(S)) 
Structure members of type T must have a byte offset from the start of the structure that is a multiple of the RequiredAlignOf(T, C) for the address space C:
OffsetOfMember(S, M) = k × RequiredAlignOf(T, C)
Where k is a positive integer and M is a member of structure S with type T
Arrays of element type T must have an element stride that is a multiple of the RequiredAlignOf(T, C) for the address space C:
StrideOf(array<T, N>) = k × RequiredAlignOf(T, C)
StrideOf(array<T>) = k × RequiredAlignOf(T, C)
Where k is a positive integer
Note: RequiredAlignOf(T, C) does not impose any additional restrictions on the values permitted for an align decoration, nor does it affect the rules of AlignOf(T). Data is laid out with the rules defined in previous sections and then the resulting layout is validated against the RequiredAlignOf(T, C) rules.
The uniform address space also requires that:

Array elements are aligned to 16 byte boundaries. That is, StrideOf(array<T,N>) = 16 × k’ for some positive integer k’.

If a structure member itself has a structure type
S
, then the number of bytes between the start of that member and the start of any following member must be at least roundUp(16, SizeOf(S)).
Note: The following examples show how to use align and size attributes on structure members to satisfy layout requirements for uniform buffers. In particular, these techniques can be used mechanically transform a GLSL buffer with std140 layout to WGSL.
struct S { x :f32 ; } struct Invalid { a :S ; b :f32 ; // invalid: offset between a and b is 4 bytes, but must be at least 16 } @ group ( 0 ) @ binding ( 0 ) var < uniform > invalid :Invalid ; struct Valid { a :S ; @ align ( 16 ) b :f32 ; // valid: offset between a and b is 16 bytes } @ 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.4. Memory View Types
In addition to calculating with plain values, a WGSL program will also often read values from memory or write values to memory, via memory access operations. Each memory access is performed via a memory view.
A memory view comprises:

a set of memory locations in a particular address space,

an interpretation of the contents of those locations as a WGSL type, and

an access mode.
The access mode of a memory view must be supported by the address space. See § 4.3.6 Address spaces.
WGSL has two kinds of types for representing memory views: reference types and pointer types.
Constraint  Type  Description 

S is a address space, T is a storable type, A is an access mode  ref<S,T,A>  The reference type identified with the set of memory views for memory locations in S holding values of type T,
supporting memory accesses described by mode A. In this context T is known as the store type. Reference types are not written in WGSL program source; instead they are used to analyze a WGSL program. 
S is a address space, T is a storable type, A is an access mode  ptr<S,T,A>  The pointer type identified with the set of memory views for memory locations in S holding values of type T,
supporting memory accesses described by mode A. In this context T is known as the pointee type. Pointer types may appear in WGSL program source. 
When analyzing a WGSL program, reference and pointer types are fully parameterized by an address space, a storable type, and an access mode. In code examples in this specification, the comments show this fully parameterized form.
However, in WGSL source text:

Reference types must not appear.

Pointer types may appear. A pointer type is spelled with parameterization by:

address space,

store type, and

sometimes by access mode, as specified in § 4.4.1 Access Mode Defaults.

fn my_function ( /* 'ptr<function,i32,read_write>' is the type of a pointer value that references memory for keeping an 'i32' value, using memory locations in the 'function' address space. Here 'i32' is the pointee type. The implied access mode is 'read_write'. See below for access mode defaults. */ ptr_int :ptr < function , i32 > , // 'ptr<private,array<f32,50>,read_write>' is the type of a pointer value that // refers to memory for keeping an array of 50 elements of type 'f32', using // memory locations in the 'private' address space. // Here the pointee type is 'array<f32,50>'. // The implied access mode is 'read_write'. See below for access mode defaults. ptr_array :ptr < private , array < f32 , 50 >> ) { }
Reference types and pointer types are both sets of memory views: a particular memory view is associated with a unique reference value and also a unique pointer value:
Each pointer value p of type ptr<S,T,A> corresponds to a unique reference value r of type ref<S,T,A>, and vice versa, where p and r describe the same memory view.
4.4.1. Access Mode Defaults
The access mode for a memory view is often determined by context:

The storage address space supports both read and read_write access modes.

Each other address space supports only one access mode, as described in the address space table.
When writing a variable declaration or a pointer type in WGSL source:

For the storage address space, the access mode is optional, and defaults to read.

For other address spaces, the access mode must not be written.
4.4.2. Originating Variable
In WGSL a reference value always corresponds to the memory view for some or all of the memory locations for some variable. This defines the originating variable for the reference value.
A pointer value always corresponds to a reference value, and so the originating variable of a pointer is the same as the originating variable of the corresponding reference.
Note: The originating variable is a dynamic concept. The originating variable for a formal parameter of a function depends on the call sites for the function. Different call sites may supply pointers into different originating variables.
If a reference or pointer access is out of bounds, an invalid memory reference is produced. Loads from an invalid reference return one of:

a value from any memory location(s) of the WebGPU buffer bound to the originating variable

the zero value for store type of the reference

if the loaded value is a vector, the value (0, 0, 0, x), where x is:

0, 1, or the maximum positive value for integer components

0.0 or 1.0 for floatingpoint components


store the value to any memory location(s) of the WebGPU buffer bound to the originating variable

not be executed
4.4.3. Use Cases for References and Pointers
References and pointers are distinguished by how they are used:

The type of a variable is a reference type.

The addressof operation (unary
&
) converts a reference value to its corresponding pointer value. 
The indirection operation (unary
*
) converts a pointer value to its corresponding reference value. 
A let declaration can be of pointer type, but not of reference type.

A formal parameter can be of pointer type, but not of reference type.

A simple assignment statement performs a write access to update the contents of memory via a reference, where:

The lefthand side of the assignment statement must be of reference type, with access mode write or read_write.

The righthand side of the assignment statement must evaluate to the store type of the lefthand side.


The Load Rule: Inside a function, a reference is automatically dereferenced (read from) to satisfy type rules:

In a function, when a reference expression r with store type T is used in a statement or an expression, where

r has an access mode of read or read_write, and

The only potentially matching type rules require r to have a value of type T, then

That type rule requirement is considered to have been met, and

The result of evaluating r in that context is the value (of type T) stored in the memory locations referenced by r at the time of evaluation. That is, a read access is performed to produce the result value.

Defining references in this way enables simple idiomatic use of variables:
@ stage ( compute ) fn main () { // 'i' has reference type ref<function,i32,read_write> // The memory locations for 'i' store the i32 value 0. var i :i32 = 0 ; // 'i + 1' can only match a type rule where the 'i' subexpression is of type i32. // So the expression 'i + 1' has type i32, and at evaluation, the 'i' subexpression // evaluates to the i32 value stored in the memory locations for 'i' at the time // of evaluation. let one :i32 = i + 1 ; // Update the value in the locations referenced by 'i' so they hold the value 2. i = one + 1 ; // Update the value in the locations referenced by 'i' so they hold the value 5. // The evaluation of the righthandside occurs before the assignment takes effect. i = i + 3 ; }
var < private > age :i32 ; fn get_age () >i32 { // The type of the expression in the return statement must be 'i32' since it // must match the declared return type of the function. // The 'age' expression is of type ref<private,i32,read_write>. // Apply the Load Rule, since the store type of the reference matches the // required type of the expression, and no other type rule applies. // The evaluation of 'age' in this context is the i32 value loaded from the // memory locations referenced by 'age' at the time the return statement is // executed. return age ; } fn caller () { age = 21 ; // The copy_age constant will get the i32 value 21. let copy_age :i32 = get_age (); }
Defining pointers in this way enables two key use cases:

Using a let declaration with pointer type, to form a short name for part of the contents of a variable.

Using a formal parameter of a function to refer to the memory of a variable that is accessible to the calling function.

The call to such a function must supply a pointer value for that operand. This often requires using an addressof operation (unary
&
) to get a pointer to the variable’s contents.

Note: The following examples use WGSL features explained later in this specification.
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 ; @ stage ( compute ) fn main () { // Form a pointer to a specific Particle in storage memory. 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 > ) { /* Update the locations for 'x' to contain the next higher integer value, (or to wrap around to the largest negative i32 value). On the lefthand side, unary '*' converts the pointer to a reference that can then be assigned to. It has a read_write access mode, by default. /* On the righthand side:  Unary '*' converts the pointer to a reference, with a read_write access mode.  The only matching type rule is for addition (+) and requires '*x' to have type i32, which is the store type for '*x'. So the Load Rule applies and '*x' evaluates to the value stored in the memory for '*x' at the time of evaluation, which is the i32 value for 0.  Add 1 to 0, to produce a final value of 1 for the righthand side. */ Store 1 into the memory for '*x'. */ * x = * x + 1 ; } @ stage ( compute ) fn main () { var i :i32 = 0 ; // Modify the contents of 'i' so it will contain 1. // Use unary '&' to get a pointer value for 'i'. // This is a clear signal that the called function has access to the memory // for 'i', and may modify it. add_one ( & i ); let one :i32 = i ; // 'one' has value 1. }
4.4.4. Forming Reference and Pointer Values
A reference value is formed in one of the following ways:

The identifier resolving to an inscope variable v denotes the reference value for v's memory.

The resolved variable is the originating variable for the reference.


Use the indirection (unary
*
) operation on a pointer.
The originating variable of the result is defined as the originating variable of the pointer.


Use a composite reference component expression. In each case the originating variable of the result is defined as the originating variable of the original reference.

Given a reference with a vector store type, appending a singleletter vector access phrase results in a reference to the named component of the vector. See § 6.7.1.3 Component Reference from Vector Reference.

Given a reference with a vector store type, appending an array index access phrase results in a reference to the indexed component of the vector. See § 6.7.1.3 Component Reference from Vector Reference.

Given a reference with a matrix store type, appending an array index access phrase results in a reference to the indexed column vector of the matrix. See § 6.7.2 Matrix Access Expression.

Given a reference with an array store type, appending an array index access phrase results in a reference to the indexed element of the array. See § 6.7.3 Array Access Expression.

Given a reference with a structure store type, appending a member access phrase results in a reference to the named member of the structure. See § 6.7.4 Structure Access Expression.

In all cases, the access mode of the result is the same as the access mode of the original reference.
struct S { age :i32 ; weight :f32 ; } var < private > person :S ; // Uses of 'person' denote the reference to the memory underlying the variable, // and will have type ref<private,S,read_write>. fn f () { var uv :vec2 < f32 > ; // Uses of 'uv' denote the reference to the memory underlying the variable, // and will have type ref<function,vec2<f32>,read_write>. // Evaluate the lefthand side of the assignment: // Evaluate 'uv.x' to yield a reference: // 1. First evaluate 'uv', yielding a reference to the memory for // the 'uv' variable. The result has type ref<function,vec2<f32>,read_write>. // 2. Then apply the '.x' vector access phrase, yielding a reference to // the memory for the first component of the vector pointed at by the // reference value from the previous step. // The result has type ref<function,f32,read_write>. // Evaluating the righthand side of the assignment yields the f32 value 1.0. // Store the f32 value 1.0 into the storage memory locations referenced by uv.x. uv . x = 1.0 ; // Evaluate the lefthand side of the assignment: // Evaluate 'uv[1]' to yield a reference: // 1. First evaluate 'uv', yielding a reference to the memory for // the 'uv' variable. The result has type ref<function,vec2<f32>,read_write>. // 2. Then apply the '[1]' array index phrase, yielding a reference to // the memory for second component of the vector referenced from // the previous step. The result has type ref<function,f32,read_write>. // Evaluating the righthand side of the assignment yields the f32 value 2.0. // Store the f32 value 2.0 into the storage memory locations referenced by uv[1]. uv [ 1 ] = 2.0 ; var m :mat3x2 < f32 > ; // When evaluating 'm[2]': // 1. First evaluate 'm', yielding a reference to the memory for // the 'm' variable. The result has type ref<function,mat3x2<f32>,read_write>. // 2. Then apply the '[2]' array index phrase, yielding a reference to // the memory for the third column vector pointed at by the reference // value from the previous step. // Therefore the 'm[2]' expression has type ref<function,vec2<f32>,read_write>. // The 'let' declaration is for type vec2<f32>, so the declaration // statement requires the initializer to be of type vec2<f32>. // The Load Rule applies (because no other type rule can apply), and // the evaluation of the initializer yields the vec2<f32> value loaded // from the memory locations referenced by 'm[2]' at the time the declaration // is executed. let p_m_col2 :vec2 < f32 > = m [ 2 ]; var A :array < i32 , 5 > ; // When evaluating 'A[4]' // 1. First evaluate 'A', yielding a reference to the memory for // the 'A' variable. The result has type ref<function,array<i32,5>,read_write>. // 2. Then apply the '[4]' array index phrase, yielding a reference to // the memory for the fifth element of the array referenced by // the reference value from the previous step. // The result value has type ref<function,i32,read_write>. // The let declaration requires the righthandside to be of type i32. // The Load Rule applies (because no other type rule can apply), and // the evaluation of the initializer yields the i32 value loaded from // the memory locations referenced by 'A[5]' at the time the declaration // is executed. let A_4_value :i32 = A [ 4 ]; // When evaluating 'person.weight' // 1. First evaluate 'person', yielding a reference to the memory for // the 'person' variable declared at module scope. // The result has type ref<private,S,read_write>. // 2. Then apply the '.weight' member access phrase, yielding a reference to // the memory for the second member of the memory referenced by // the reference value from the previous step. // The result has type ref<private,f32,read_write>. // The let declaration requires the righthandside to be of type f32. // The Load Rule applies (because no other type rule can apply), and // the evaluation of the initializer yields the f32 value loaded from // the memory locations referenced by 'person.weight' at the time the // declaration is executed. let person_weight :f32 = person . weight ; }
A pointer value is formed in one of the following ways:

Use the addressof (unary '&') operator on a reference.

The originating variable of the result is defined as the originating variable of the reference.


If a function formal parameter has pointer type, then when the function is invoked at runtime the uses of the formal parameter denote the pointer value provided to the corresponding operand at the call site in the calling function.

The originating variable of the formal parameter (at runtime) is defined as the originating variable of the pointer operand at the call site.

In all cases, the access mode of the result is the same as the access mode of the original pointer.
// Declare a variable in the private address space, for storing an f32 value. var < private > x :f32 ; fn f () { // Declare a variable in the function address space, for storing an i32 value. var y :i32 ; // The name 'x' resolves to the modulescope variable 'x', // and has reference type ref<private,f32,read_write>. // Applying the unary '&' operator converts the reference to a pointer. // The access mode is the same as the access mode of the original variable, so // the fully specified type is ptr<private,f32,read_write>. But read_write // is the default access mode for function address space, so read_write does not // have to be spelled in this case let x_ptr :ptr < private , f32 > = & x ; // The name 'y' resolves to the functionscope variable 'y', // and has reference type ref<private,i32,read_write>. // Applying the unary '&' operator converts the reference to a pointer. // The access mode defaults to 'read_write'. let y_ptr :ptr < function , i32 > = & y ; // A new variable, distinct from the variable declared at module scope. var x :u32 ; // Here, the name 'x' resolves to the functionscope variable 'x' declared in // the previous statement, and has type ref<function,u32,read_write>. // Applying the unary '&' operator converts the reference to a pointer. // The access mode defaults to 'read_write'. let inner_x_ptr :ptr < function , u32 > = & x ; }
4.4.5. Comparison with References and Pointers in Other Languages
This section is informative, not normative.
References and pointers in WGSL are more restricted than in other languages. In particular:

In WGSL a reference can’t directly be declared as an alias to another reference or variable, either as a variable or as a formal parameter.

In WGSL pointers and references are not storable. That is, the content of a WGSL variable may not contain a pointer or a reference.

In WGSL a function must not return a pointer or reference.

In WGSL there is no way to convert between integer values and pointer values.

In WGSL there is no way to forcibly change the type of a pointer value into another pointer type.

A composite component reference expression is different: it takes a reference to a composite value and yields a reference to one of the components or elements inside the composite value. These are considered different references in WGSL, even though they may have the same machine address at a lower level of implementation abstraction.


In WGSL there is no way to forcibly change the type of a reference value into another reference type.

In WGSL there is no way to change the access mode of a pointer or reference.

By comparison, C++ automatically converts a nonconst pointer to a const pointer, and has a
const_cast
to convert a const value to a nonconst value.


In WGSL there is no way to allocate new memory from a "heap".

In WGSL there is no way to explicitly destroy a variable. The memory for a WGSL variable becomes inaccessible only when the variable goes out of scope.
Note: From the above rules, it is not possible to form a "dangling" pointer, i.e. a pointer that does not reference the memory for a valid (or "live") originating variable.
4.5. Texture and Sampler Types
A texel is a scalar or vector used as the smallest independently accessible element of a texture. The word texel is short for texture element.
A texture is a collection of texels supporting special operations useful for rendering. In WGSL, those operations are invoked via texture builtin functions. See § 16.8 Texture Builtin Functions for a complete list.
A WGSL texture corresponds to a WebGPU GPUTexture.
A texture is either arrayed, or nonarrayed:

A nonarrayed texture is a grid of texels. Each texel has a unique grid coordinate.

An arrayed texture is a homogeneous array of grids of texels. In an arrayed texture, each texel is identified with its unique combination of array index and grid coordinate.
A texture has the following features:
 texel format

The data in each texel. See § 4.5.1 Texel Formats
 dimensionality

The number of dimensions in the grid coordinates, and how the coordinates are interpreted. The number of dimensions is 1, 2, or 3. Most textures use cartesian coordinates. Cube textures have six square faces, and are sampled with a three dimensional coordinate interpreted as a direction vector from the origin toward the cube centered on the origin.
 size

The extent of grid coordinates along each dimension
 mip level count

The mip level count is at least 1 for sampled textures, and equal to 1 for storage textures.
Mip level 0 contains a full size version of the texture. Each successive mip level contains a filtered version of the previous mip level at half the size (within rounding) of the previous mip level.
When sampling a texture, an explicit or implicitlycomputed levelofdetail is used to select the mip levels from which to read texel data. These are then combined via filtering to produce the sampled value.  arrayed

whether the texture is arrayed
 array size

the number of homogeneous grids, if the texture is arrayed
A texture’s representation is typically optimized for rendering operations. To achieve this, many details are hidden from the programmer, including data layouts, data types, and internal operations that cannot be expressed directly in the shader language.
As a consequence, a shader does not have direct access to the texel memory within a texture variable. Instead, access is mediated through an opaque handle:

Within the shader:

Declare a modulescope variable where the store type is one of the texture types described in later sections. The variable stores an opaque handle to the underlying texture memory, and is automatically placed in the handle address space.

Inside a function, call one of the texture builtin functions, and provide the texture variable or function parameter as the builtin function’s first parameter.


When constructing the WebGPU pipeline, the texture variable’s store type and binding must be compatible with the corresponding bind group layout entry.
In this way, the set of supported operations for a texture type is determined by the availability of texture builtin functions accepting that texture type as the first parameter.
Note: The handle stored by a texture variable cannot be changed by the shader. That is, the variable is readonly, even if the underlying texture to which it provides access may be mutable (e.g. a writeonly storage texture).
A sampler is an opaque handle that controls how texels are accessed from a sampled texture.
A WGSL sampler maps to a WebGPU GPUSampler.
Texel access is controlled via several properties of the sampler:
 addressing mode

Controls how texture boundaries and outofbounds coordinates are resolved. The addressing mode for each texture dimension can be set independently. See WebGPU GPUAddressMode.
 filter mode

Controls which texels are accessed to produce the final result. Filtering can either use the nearest texel or interpolate between multiple texels. Multiple filter modes can be set independently. See WebGPU GPUFilterMode.
 LOD clamp

Controls the min and max levels of details that are accessed.
 comparison

Controls the type of comparison done for comparison sampler. See WebGPU GPUCompareFunction.
 max anisotropy

Controls the maximum anisotropy value used by the sampler.
Samplers cannot be created in WGSL programs and their state (e.g. the properties listed above) are immutable within a shader and can only be set by the WebGPU API.
It is a pipelinecreation error if a filtering sampler (i.e. any sampler using interpolative filtering) is used with texture that has a nonfilterable format.
Note: The handle stored by a sampler variable cannot be changed by the shader.
4.5.1. Texel Formats
In WGSL, certain texture types are parameterized by texel format.
A texel format is characterized by:
 channels

Each channel contains a scalar. A texel format has up to four channels:
r
,g
,b
, anda
, normally corresponding to the concepts of red, green, blue, and alpha channels.  channel format

The number of bits in the channel, and how those bits are interpreted.
Each texel format in WGSL corresponds to a WebGPU GPUTextureFormat with the same name.
Only certain texel formats are used in WGSL source code. The channel formats used to define those texel formats are listed in the Channel Formats table. The last column specifies the conversion from the stored channel bits to the value used in the shader. This is also known as the channel transfer function, or CTF.
Channel format  Number of stored bits  Interpretation of stored bits  Shader type  Shader value (Channel Transfer Function) 

8unorm  8  unsigned integer v ∈ {0,...,255}  f32  v ÷ 255 
8snorm  8  signed integer v ∈ {128,...,127}  f32  max(1, v ÷ 127) 
8uint  8  unsigned integer v ∈ {0,...,255}  u32  v 
8sint  8  signed integer v ∈ {128,...,127}  i32  v 
16uint  16  unsigned integer v ∈ {0,...,65535}  u32  v 
16sint  16  signed integer v ∈ {32768,...,32767}  i32  v 
16float  16  IEEE754 binary16 16bit floating point value v, with 1 sign bit, 5 exponent bits, 10 mantissa bits  f32  v 
32uint  32  32bit unsigned integer value v  u32  v 
32sint  32  32bit signed integer value v  i32  v 
32float  32  IEEE754 binary32 32bit floating point value v  f32  v 
The texel formats listed in the Texel Formats for Storage Textures table correspond to the WebGPU plain color formats which support the WebGPU STORAGE usage. These texel formats are used to parameterize the storage texture types defined in § 4.5.5 Storage Texture Types.
When the texel format does not have all four channels, then:

When reading the texel:

If the texel format has no green channel, then the second component of the shader value is 0.

If the texel format has no blue channel, then the third component of the shader value is 0.

If the texel format has no alpha channel, then the fourth component of the shader value is 1.


When writing the texel, shader value components for missing channels are ignored.
The last column in the table below uses the formatspecific channel transfer function from the channel formats table.
Texel format  Channel format  Channels in memory order  Corresponding shader value 

rgba8unorm  8unorm  r, g, b, a  vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) 
rgba8snorm  8snorm  r, g, b, a  vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) 
rgba8uint  8uint  r, g, b, a  vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a)) 
rgba8sint  8sint  r, g, b, a  vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a)) 
rgba16uint  16uint  r, g, b, a  vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a)) 
rgba16sint  16sint  r, g, b, a  vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a)) 
rgba16float  16float  r, g, b, a  vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) 
r32uint  32uint  r  vec4<u32>(CTF(r), 0u, 0u, 1u) 
r32sint  32sint  r  vec4<i32>(CTF(r), 0, 0, 1) 
r32float  32float  r  vec4<f32>(CTF(r), 0.0, 0.0, 1.0) 
rg32uint  32uint  r, g  vec4<u32>(CTF(r), CTF(g), 0.0, 1.0) 
rg32sint  32sint  r, g  vec4<i32>(CTF(r), CTF(g), 0.0, 1.0) 
rg32float  32float  r, g  vec4<f32>(CTF(r), CTF(g), 0.0, 1.0) 
rgba32uint  32uint  r, g, b, a  vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a)) 
rgba32sint  32sint  r, g, b, a  vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a)) 
rgba32float  32float  r, g, b, a  vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) 
4.5.2. Sampled Texture Types
texture_1d<type>
texture_2d<type>
texture_2d_array<type>
texture_3d<type>
texture_cube<type>
texture_cube_array<type>

type must be
f32
,i32
oru32

The parameterized type for the images is the type after conversion from sampling. E.g. you can have an image with texels with 8bit unorm components, but when you sample them you get a 32bit float result (or vecoff32).
4.5.3. Multisampled Texture Types
texture_multisampled_2d<type>

type must be
f32
,i32
oru32
4.5.4. External Sampled Texture Types
texture_external
texture_external
is an opaque 2d floatsampled texture type similar to texture_2d<f32>
but potentially with a different representation.
It can be read using textureLoad
or textureSampleLevel
,
which handle these different representations opaquely.
See WebGPU § GPUExternalTexture.
4.5.5. Storage Texture Types
A storage texture supports accessing a single texel without the use of a sampler.

A writeonly storage texture supports writing a single texel, with automatic conversion of the shader value to a stored texel value.
A storage texture type must be parameterized by one of the texel formats for storage textures. The texel format determines the conversion function as specified in § 4.5.1 Texel Formats.
For a writeonly storage texture the inverse of the conversion function is used to convert the shader value to the stored texel.
See § 16.8 Texture Builtin Functions.
TODO(dneto): Move description of the conversion to the builtin function that actually does the reading.
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
must be one of the texel types specified in storagetexelformats 
access
must be write.
4.5.6. Depth Texture Types
texture_depth_2d
texture_depth_2d_array
texture_depth_cube
texture_depth_cube_array
texture_depth_multisampled_2d
4.5.7. Sampler Type
A sampler mediates access to a sampled texture or a depth texture, by performing a combination of:

coordinate transformation.

optionally modifying miplevel selection.

for a sampled texture, optionally filtering retrieved texel values.

for a depth texture, determining the comparison function applied to the retrieved texel.
Type  Description 

sampler  Sampler. Mediates access to a sampled texture. 
sampler_comparison  Comparison sampler. Mediates access to a depth texture. 
Samplers are parameterized when created in the WebGPU API. They cannot be modified by a WGSL program.
Samplers can only be used by the texture builtin functions.
sampler sampler_comparison
4.5.8. Texture Types Grammar
 sampled_texture_type less_than type_decl greater_than
 multisampled_texture_type less_than type_decl greater_than
 storage_texture_type less_than texel_format comma access_mode greater_than
 sampler
 'rgba8unorm'
 'rgba8snorm'
 'rgba8uint'
 'rgba8sint'
 'rgba16uint'
 'rgba16sint'
 'rgba16float'
 'r32uint'
 'r32sint'
 'r32float'
 'rg32uint'
 'rg32sint'
 'rg32float'
 'rgba32uint'
 'rgba32sint'
 'rgba32float'
4.6. Type Aliases
A type alias declares a new name for an existing type. The declaration must appear at module scope, and its scope is the entire program.
type Arr = array<i32, 5>; type RTArr = array<vec4<f32>>; type single = f32; // Declare an alias for f32 let pi_approx: single = 3.1415; fn two_pi() > single { return single(2) * pi_approx; }
4.7. Type Declaration Grammar
 ident
 bool
 float32
 int32
 uint32
 vec2 less_than type_decl greater_than
 vec3 less_than type_decl greater_than
 vec4 less_than type_decl greater_than
 pointer less_than address_space comma type_decl ( comma access_mode ) ? greater_than
 mat2x2 less_than type_decl greater_than
 mat2x3 less_than type_decl greater_than
 mat2x4 less_than type_decl greater_than
 mat3x2 less_than type_decl greater_than
 mat3x3 less_than type_decl greater_than
 mat3x4 less_than type_decl greater_than
 mat4x2 less_than type_decl greater_than
 mat4x3 less_than type_decl greater_than
 mat4x4 less_than type_decl greater_than
When the type declaration is an identifier, then the expression must be in scope of a declaration of the identifier as a type alias or structure type.
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 readonly, and has more restrictive layout rules. struct ParamsTable {weight: f32;} @group(0) @binding(2) var<uniform> params: ParamsTable; // Can read, cannot write.
5. Variable and Value Declarations
5.1. Value Declarations
WGSL authors can declare names for immutable values using one of the following methods:
Value declarations do not have any associated storage. That is, there are no memory locations associated with the declaration.
5.1.1. let
Declarations
A let declaration specifies a name for a value. Once the value for a letdeclaration is computed, it is immutable. When an identifier use resolves to a letdeclaration, the identifier denotes that value.
When a let
identifier is declared without an explicitly specified type,
e.g. let foo = 4
, the type is automatically inferred from the expression to the right of the equals token (=
).
When the type is specified, e.g let foo: i32 = 4
, the initializer expression must evaluate to that type.
Some rules about let
declarations depend on where the declaration appears.
See § 5.4 Module Constants and § 5.5 Function Scope Variables and Constants.
// 'blockSize' denotes the i32 value 1024. let blockSize :i32 = 1024 ; // 'row_size' denotes the u32 value 16u. The type is inferred. let row_size = 16 u ;
5.1.2. override
Declarations
An override declaration specifies a name for a pipelineoverridable constant value.
The value of a pipelineoverridable constant is fixed at
pipelinecreation time.
The value is the one specified by the WebGPU pipelinecreation method, if
specified, and otherwise is the value of its initializer expression.
When an identifier use resolves to a overridedeclaration, the identifier denotes that value. override
declarations must meet the following restrictions:

The declaration must only occur at module scope.

The declaration must have at least one of a declared type, an initializer expression, or both.

The declared type, if present, must be a scalar.

The initializer expression, if present, must:

evaluate to a scalar type.

evaluate to the declared type if it is present.

be composed only const_expressions or expressions where all identifiers resolve to overridable constants.


If the declaration has the id applied, the literal operand is known as the pipeline constant ID, and must be an integer value between 0 and 65535.

Pipeline constant IDs must be unique within the WGSL program: Two
override
declarations must not use the same pipeline constant ID. 
The application can specify its own value for the constant at pipelinecreation time. The pipeline creation API accepts a mapping from overridable constant to a value of the constant’s type. The constant is identified by a pipelineoverridable constant identifier string, which is the base10 representation of the pipeline constant ID if specified, and otherwise the declared name of the constant.

The pipelineoverridable constant has a default value if its declaration has an initializer expression. If it doesn’t, a value must be provided at pipelinecreation time.
@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.2. var
Declarations
A variable is a named reference to memory that can contain a value of a particular storable type.
Two types are associated with a variable: its store type (the type of value that may be placed in the referenced memory) and its reference type (the type of the variable itself). If a variable has store type T, address space S, and access mode A, then its reference type is ref<S,T,A>.
A variable declaration:

Specifies the variable’s name.

Specifies the address space, store type, and access mode. Together these comprise the variable’s reference type.

Ensures the execution environment allocates memory for a value of the store type, in the specified address space, supporting the given access mode, for the lifetime of the variable.

Optionally has an initializer expression, if the variable is in the private or function address spaces. If present, the initializer expression must evaluate to the variable’s store type.
When an identifier use resolves to a variable declaration, the identifier is an expression denoting the reference memory view for the variable’s memory, and its type is the variable’s reference type. See § 6.13 Variable Identifier Expression.
See § 5.3 Module Scope Variables and § 5.5 Function Scope Variables and Constants for rules about where a variable in a particular address space can be declared, and when the address space decoration is required, optional, or forbidden.
The access mode always has a default, and except for variables in the storage address space, must not be written in WGSL source text. See § 4.4.1 Access Mode Defaults.
The lifetime of a variable is the period during shader execution for which the variable exists. The lifetime of a module scope variable is the entire execution of the shader stage.
For a function scope variable, each invocation has its own independent version of the variable. The lifetime of the variable is determined by its scope:

It begins when control enters the variable’s declaration.

It includes the entire execution of any function called from within the variable’s scope.

It ends when control leaves the variable’s scope, other than calling a function from within the variable’s scope.
Two variables with overlapping lifetimes will not have overlapping memory. When a variable’s lifetime ends, its memory may be used for another variable.
When a variable is created, its memory contains an initial value as follows:

For variables in the private or function address spaces:

The zero value for the store type, if the variable declaration has no initializer.

Otherwise, it is the result of evaluating the initializer expression at that point in the program execution.


For variables in the workgroup address space:

When the store type is constructible, the zero value for the store type.

Otherwise, the store type is an array of construcible elements, and each element is initialized to its zero value.


Variables in other address spaces are resources set by bindings in the draw command or dispatch command.
Consider the following snippet of WGSL:
var i :i32 ; // Initial value is 0. Not recommended style. loop { var twice :i32 = 2 * i ; // Reevaluated each iteration. i ++ ; if i == 5 { break ; } }
i
will take on values 0, 1, 2, 3, 4, 5, and variable twice
will take on values 0, 2, 4, 6, 8.
Consider the following snippet of WGSL:
Becausex
is a variable, all accesses to it turn into load and store operations.
However, it is expected that either the browser or the driver optimizes this intermediate representation
such that the redundant loads are eliminated.
5.3. Module Scope Variables
A variable declared outside all functions is at module scope. The variable name is in scope for the entire program.
Variables at module scope are restricted as follows:

The variable must not be in the function address space.

A variable in the private, workgroup, uniform, or storage address spaces:

Must be declared with an explicit address space decoration.

Must use a store type as described in § 4.3.6 Address spaces.


If the store type is a texture type or a sampler type, then the variable declaration must not have an address space decoration. The address space will always be handle.
A variable in the uniform address space is a uniform buffer variable. Its store type must be a hostshareable constructible type, and must satisfy address space layout constraints.
A variable in the storage address space is a storage buffer variable. Its store type must be a hostshareable type and must satisfy address space layout constraints. The variable may be declared with a read or read_write access mode; the default is read.
As described in § 9.3.2 Resource Interface, uniform buffers, storage buffers, textures, and samplers form the resource interface of a shader. Such variables are declared with group and binding decorations.
WGSL defines the following attributes that can be applied to global variables:
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(0) var<storage,read_write> pbuf: array<vec2<f32>>; // Textures and samplers are always in "handle" space. @group(0) @binding(1) var filter_params: sampler;
5.4. Module Constants
A value declaration appearing outside all functions declares a modulescope constant. The name is in scope for the entire program.
A modulescope letdeclared constant must be of constructible type.
An initializer expression must be present for a modulescope let
declaration
and the name denotes the value of that expression.
TODO: define creationtime constant that allows const_expression and module scope let declarations.
A pipelineoverridable constant must be one of the scalar types. An initializer expression is optional for pipelineoverridable constants.
// The golden ratio. let golden: f32 = 1.61803398875; // The second unit vector for three dimensions, with inferred type. let e2 = vec3<i32>(0,1,0);
When a variable or feature is used within control flow that depends on the value of a constant, then that variable or feature is considered to be used by the program. This is true regardless of the value of the constant, whether that value is the one from the constant’s declaration or from a pipeline override.
5.5. Function Scope Variables and Constants
A variable or constant declared in a declaration statement in a function body is in function scope. The name is available for use immediately after its declaration statement, and until the end of the bracedelimited list of statements immediately enclosing the declaration.
A functionscope letdeclared constant must be of constructible type, or of pointer type.
For a variable declared in function scope:

The variable is always in the function address space.

The address space decoration is optional.

The store type must be a constructible type.

When an initializer is specified, the store type may be omitted from the declaration. In this case the store type is the type of the result of evaluating the initializer.
fn f () { var < function > count :u32 ; // A variable in function address space. var delta :i32 ; // Another variable in the function address space. var sum :f32 = 0.0 ; // A function address space variable with initializer. var pi = 3.14159 ; // Infer the f32 store type from the initializer. let unit :i32 = 1 ; // Letdeclared constants don’t use an address space. }
A variable or constant declared in the first clause of a for
statement is available for use in the second
and third clauses and in the body of the for
statement.
An instance of a function scope variable is a dynamic context. Each variable that is in scope for some invocation has an overlapping lifetime and, therefore, has nonoverlapping memory. Variables with nonoverlapping lifetimes may reuse the memory of previous variables; however, new instances of the same variable are not guaranteed to use the same memory.
5.6. Variable and Value Declaration Grammar Summary
 variable_decl equal expression
 let ( ident  variable_ident_decl ) equal expression
 var variable_qualifier ? ( ident  variable_ident_decl )
 less_than address_space ( comma access_mode ) ? greater_than
 attribute * variable_decl ( equal const_expression ) ?
 let ( ident  variable_ident_decl ) equal const_expression
 attribute * override ( ident  variable_ident_decl ) ( equal expression ) ?
 type_decl paren_left ( ( const_expression comma ) * const_expression comma ? ) ? paren_right
6. Expressions
Expressions specify how values are computed.
6.1. Literal Expressions
Precondition  Conclusion  Notes 

true : bool
 true boolean value.
 
false : bool
 false boolean value.
 
INT_LITERAL: i32  Signed integer value.  
UINT_LITERAL: u32  Unsigned integer value.  
FLOAT_LITERAL: f32  Floatingpoint value. 
6.2. Parenthesized Expressions
Precondition  Conclusion  Description 

e : T  ( e ) : T
 Evaluates to e. Use parentheses to isolate an expression from the surrounding text. 
6.3. Type Constructor Expressions
Type constructor expressions explicitly create a value of a given constructible type. Type constructor expressions explicitly create a value of a given type.
The scalar forms are redundant, but provide symmetry with scalar conversion expressions, and can be used to enhance readability.
The vector and matrix forms construct vector and matrix values from various combinations of components and subvectors with matching component types. There are overloads for constructing vectors and matrices that specify the dimensions of the target type without having to specify the component type; the component type is inferred from the constructor arguments.
See also § 6.4 Zero Value Expressions and § 6.5 Conversion Expressions.
Precondition  Conclusion  

e: bool  bool(e) : bool
 Identity. 
e: i32  i32(e) : i32
 Identity. 
e: u32  u32(e) : u32
 Identity. 
e: f32  f32(e) : f32
 Identity. 
Precondition  Conclusion  Notes 

e: T  vec N< T>( e) : vecN<T>
 Evaluates e once. Results in the Ncomponent vector where each component has the value of e. 
vec N( e) : vecN<T>
 
e1: T e2: T  vec2<T>(e1,e2) : vec2<T>
 
vec2(e1,e2) : vec2<T>
 
e: vec2<T>  vec2<T>(e) : vec2<T>
 Identity. The result is e. 
vec2(e) : vec2<T>
 
e1: T e2: T e3: T  vec3<T>(e1,e2,e3) : vec3<T>
 
vec3(e1,e2,e3) : vec3<T>
 
e1: T e2: vec2<T>  vec3<T>(e1,e2) : vec3<T>vec3<T>(e2,e1) : vec3<T>
 
vec3(e1,e2) : vec3<T>vec3(e2,e1) : vec3<T>
 
e: vec3<T>  vec3<T>(e) : vec3<T>
 Identity. The result is e. 
vec3(e) : vec3<T>
 
e1: T e2: T e3: T e4: T  vec4<T>(e1,e2,e3,e4) : vec4<T>
 
vec4(e1,e2,e3,e4) : vec4<T>
 
e1: T e2: T e3: vec2<T>  vec4<T>(e1,e2,e3) : vec4<T>vec4<T>(e1,e3,e2) : vec4<T>vec4<T>(e3,e1,e2) : vec4<T>
 
vec4(e1,e2,e3) : vec4<T>vec4(e1,e3,e2) : vec4<T>vec4(e3,e1,e2) : vec4<T>
 
e1: vec2<T> e2: vec2<T>  vec4<T>(e1,e2) : vec4<T>
 
vec4(e1,e2) : vec4<T>
 
e1: T e2: vec3<T>  vec4<T>(e1,e2) : vec4<T>vec4<T>(e2,e1) : vec4<T>
 
vec4(e1,e2) : vec4<T>vec4(e2,e1) : vec4<T>
 
e: vec4<T>  vec4<T>(e) : vec4<T>
 Identity. The result is e. 
vec4(e) : vec4<T>

Precondition  Conclusion  Notes 

e1: f32 ... eN: f32  mat2x2<f32>(e1,e2,e3,e4) : mat2x2<f32>mat3x2<f32>(e1,...,e6) : mat3x2<f32>mat2x3<f32>(e1,...,e6) : mat2x3<f32>mat4x2<f32>(e1,...,e8) : mat4x2<f32>mat2x4<f32>(e1,...,e8) : mat2x4<f32>mat3x3<f32>(e1,...,e9) : mat3x3<f32>mat4x3<f32>(e1,...,e12) : mat4x3<f32>mat3x4<f32>(e1,...,e12) : mat3x4<f32>mat4x4<f32>(e1,...,e16) : mat4x4<f32>
 Columnmajor construction by elements. 
mat2x2(e1,e2,e3,e4) : mat2x2<f32>mat3x2(e1,...,e6) : mat3x2<f32>mat2x3(e1,...,e6) : mat2x3<f32>mat4x2(e1,...,e8) : mat4x2<f32>mat2x4(e1,...,e8) : mat2x4<f32>mat3x3(e1,...,e9) : mat3x3<f32>mat4x3(e1,...,e12) : mat4x3<f32>mat3x4(e1,...,e12) : mat3x4<f32>mat4x4(e1,...,e16) : mat4x4<f32>
 
e1: vec2<f32> e2: vec2<f32> e3: vec2<f32> e4: vec2<f32>  mat2x2<f32>(e1,e2) : mat2x2<f32>mat3x2<f32>(e1,e2,e3) : mat3x2<f32>mat4x2<f32>(e1,e2,e3,e4) : mat4x2<f32>
 Column by column construction. 
mat2x2(e1,e2) : mat2x2<f32>mat3x2(e1,e2,e3) : mat3x2<f32>mat4x2(e1,e2,e3,e4) : mat4x2<f32>
 
e1: vec3<f32> e2: vec3<f32> e3: vec3<f32> e4: vec3<f32>  mat2x3<f32>(e1,e2) : mat2x3<f32>mat3x3<f32>(e1,e2,e3) : mat3x3<f32>mat4x3<f32>(e1,e2,e3,e4) : mat4x3<f32>
 Column by column construction. 
mat2x3(e1,e2) : mat2x3<f32>mat3x3(e1,e2,e3) : mat3x3<f32>mat4x3(e1,e2,e3,e4) : mat4x3<f32>
 
e1: vec4<f32> e2: vec4<f32> e3: vec4<f32> e4: vec4<f32>  mat2x4<f32>(e1,e2) : mat2x4<f32>mat3x4<f32>(e1,e2,e3) : mat3x4<f32>mat4x4<f32>(e1,e2,e3,e4) : mat4x4<f32>
 Column by column construction. 
mat2x4(e1,e2) : mat2x4<f32>mat3x4(e1,e2,e3) : mat3x4<f32>mat4x4(e1,e2,e3,e4) : mat4x4<f32>

Precondition  Conclusion  Notes 

e1: T ... eN: T, T is constructible  array< T,N>( e1,...,eN) : array<T,N>

Construction of an array from elements.
Note: array<T,N> is constructible because its element count is equal to the number of arguments to the constructor, and hence fully determined at shadercreation time. 
Precondition  Conclusion  Notes 

e1: T1 ... eN: TN, S is a constructible structure type with members having types T1 ... TN. The expression is in the scope of declaration of S.  S( e1,...,eN) : S
 Construction of a structure from members. 
6.4. Zero Value Expressions
Each constructible T has a unique zero value written in WGSL as the type followed by an empty pair of parentheses: T ()
.
The zero values are as follows:

bool()
isfalse

i32()
is 0 
u32()
is 0 
f32()
is 0.0 
The zero value for an Ncomponent vector of type T is the Ncomponent vector of the zero value for T.

The zero value for an Ncolumn Mrow matrix of
f32
is the matrix of those dimensions filled with 0.0 entries. 
The zero value for a constructible Nelement array with element type E is an array of N elements of the zero value for E.

The zero value for a constructible structure type S is the structure value S with zerovalued members.
Note: WGSL does not have zero expression for atomic types, runtimesized arrays, or other types that are not constructible.
Precondition  Conclusion  Notes 

bool() : bool
 false Zero value  
i32() : i32
 0 Zero value  
u32() : u32
 0u Zero value  
f32() : f32
 0.0 Zero value 
Precondition  Conclusion  Notes 

vec2<T>() : vec2<T>
 Zero value  
vec3<T>() : vec3<T>
 Zero value  
vec4<T>() : vec4<T>
 Zero value 
vec2 < f32 > () // The zerovalued vector of two f32 components. vec2 < f32 > ( 0.0 , 0.0 ) // The same value, written explicitly. vec3 < i32 > () // The zerovalued vector of three i32 components. vec3 < i32 > ( 0 , 0 , 0 ) // The same value, written explicitly.
Precondition  Conclusion  Notes 

mat2x2<f32>() : mat2x2<f32>mat3x2<f32>() : mat3x2<f32>mat4x2<f32>() : mat4x2<f32>
 Zero value  
mat2x3<f32>() : mat2x3<f32>mat3x3<f32>() : mat3x3<f32>mat4x3<f32>() : mat4x3<f32>
 Zero value  
mat2x4<f32>() : mat2x4<f32>mat3x4<f32>() : mat3x4<f32>mat4x4<f32>() : mat4x4<f32>
 Zero value 
Precondition  Conclusion  Notes 

T is a constructible  array< T,N>() : array<T,N>
 Zerovalued array 
array < bool , 2 > () // The zerovalued array of two booleans. array < bool , 2 > ( false , false ) // The same value, written explicitly.
Precondition  Conclusion  Notes 

S is a constructible structure type. The expression is in the scope of declaration of S.  S() : S
 Zerovalued structure: a structure of type S where each member is the zero value for its member type. 
struct Student { grade :i32 ; GPA :f32 ; attendance :array < bool , 4 > ; } fn func () { var s :Student ; // The zero value for Student s = Student (); // The same value, written explicitly. s = Student ( 0 , 0.0 , array < bool , 4 > ( false , false , false , false )); // The same value, written with zerovalued members. s = Student ( i32 (), f32 (), array < bool , 4 > ()); }
6.5. Conversion Expressions
WGSL does not implicitly convert or promote a numeric or boolean value to another type. Instead use conversion expressions as defined in the tables below.
For details on conversion to and from floating point types, see § 12.5.2 Floating Point Conversion.
See also § 6.3 Type Constructor Expressions.
Precondition  Conclusion  Notes 

e: u32  bool( e) : bool
 Coercion to boolean. The result is false if e is 0, and true otherwise. 
e: i32  bool( e) : bool
 Coercion to boolean. The result is false if e is 0, and true otherwise. 
e: f32  bool( e) : bool
 Coercion to boolean. The result is false if e is 0.0 or 0.0, and true otherwise. In particular NaN and infinity values map to true. 
e: bool  i32( e) : i32
 Conversion of a boolean value to a signed integer The result is 1 if e is true and 0 otherwise. 
e: u32  i32( e) : i32
 Reinterpretation of bits. The result is the unique value in i32 that is equal to (e mod 2^{32}). 
e: f32  i32( e) : i32
 Value conversion, rounding toward zero. 
e: bool  u32( e) : u32
 Conversion of a boolean value to an unsigned integer The result is 1u if e is true and 0u otherwise. 
e: i32  u32( e) : u32
 Reinterpretation of bits. The result is the unique value in u32 that is equal to (e mod 2^{32}). 
e: f32  u32( e) : u32
 Value conversion, rounding toward zero. 
e: bool  f32( e) : f32
 Conversion of a boolean value to floating point The result is 1.0 if e is true and 0.0 otherwise. 
e: i32  f32( e) : f32
 Value conversion, including invalid cases. 
e: u32  f32( e) : f32
 Value conversion, including invalid cases. 
Details of conversion to and from floating point are explained in § 12.5.2 Floating Point Conversion.
Precondition  Conclusion  Notes 

e: vecN<u32>  vec N<bool >( e) : vecN<bool>
 Componentwise coercion of a unsigned integer vector to a boolean vector. 
e: vecN<i32>  vec N<bool >( e) : vecN<bool>
 Componentwise coercion of a signed integer vector to a boolean vector. 
e: vecN<f32>  vec N<bool >( e) : vecN<bool>
 Componentwise coercion of a floating point vector to a boolean vector. 
e: vecN<bool>  vec N<i32 >( e) : vecN<i32>
 Componentwise conversion of a boolean vector to signed. Component i of the result is i32( e[ i])

e: vecN<u32>  vec N<i32 >( e) : vecN<i32>
 Componentwise reinterpretation of bits. Component i of the result is i32( e[ i])

e: vecN<f32>  vec N<i32 >( e) : vecN<i32>
 Componentwise value conversion to signed integer, including invalid cases. 
e: vecN<bool>  vec N<u32 >( e) : vecN<u32>
 Componentwise conversion of a boolean vector to unsigned. Component i of the result is u32( e[ i])

e: vecN<i32>  vec N<u32 >( e) : vecN<u32>
 Componentwise reinterpretation of bits. 
e: vecN<f32>  vec N<u32 >( e) : vecN<u32>
 Componentwise value conversion to unsigned integer, including invalid cases. 
e: vecN<bool>  vec N<u32 >( e) : vecN<u32>
 Componentwise conversion of a boolean vector to floating point. Component i of the result is f32( e[ i])

e: vecN<i32>  vec N<f32 >( e) : vecN<f32>
 Componentwise value conversion to floating point, including invalid cases. 
e: vecN<u32>  vec N<f32 >( e) : vecN<f32>
 Componentwise value conversion to floating point, including invalid cases. 
6.6. Reinterpretation of Representation Expressions
A bitcast
expression is used to reinterpet the bit representation of a
value in one type as a value in another type.
Precondition  Conclusion  Notes 

e: T T is a numeric scalar or numeric vector type  bitcast<T>(e): T  Identity transform. Componentwise when T is a vector. The result is e. 
e: T1 T1 is a numeric scalar or numeric vector type T2 is not T1 and is a numeric scalar type if T1 is a scalar, or a numeric vector type if T1 is a vector  bitcast<T2>(e): T2  Reinterpretation of bits as T2. Componentwise when T1 is a vector. The result is the reinterpretation of the bits in e as a T2 value. 
6.7. Composite Value Decomposition Expressions
6.7.1. Vector Access Expression
Accessing components of a vector can be done either using array subscripting (e.g. a[2]
) or using a sequence of convenience names, each mapping to a component of the source vector.
 The colour set of convenience names:
r
,g
,b
,a
for vector components 0, 1, 2, and 3 respectively.  The dimensional set of convenience names:
x
,y
,z
,w
for vector components 0, 1, 2, and 3, respectively.
The convenience names are accessed using the .
notation. (e.g. color.bgra
).
NOTE: the convenience letterings can not be mixed. (i.e. you can not use rybw
).
Using a convenience letter, or array subscript, which accesses a component past the end of the vector is an error.
The convenience letterings can be applied in any order, including duplicating letters as needed. You can provide 1 to 4 letters when extracting components from a vector. Providing more then 4 letters is an error.
The result type depends on the number of letters provided. Assuming a vec4<f32>
Accessor  Result type 

r  f32

rg  vec2<f32>

rgb  vec3<f32>

rgba  vec4<f32>

var a :vec3 < f32 > = vec3 < f32 > ( 1. , 2. , 3. ); var b :f32 = a . y ; // b = 2.0 var c :vec2 < f32 > = a . bb ; // c = (3.0, 3.0) var d :vec3 < f32 > = a . zyx ; // d = (3.0, 2.0, 1.0) var e :f32 = a [ 1 ]; // e = 2.0
6.7.1.1. Vector Single Component Selection
Precondition  Conclusion  Description 

e: vecN<T>  e.x : Te .r : T
 Select the first component of e 
e: vecN<T>  e.y : Te .g : T
 Select the second component of e 
e: vecN<T> N is 3 or 4  e.z : Te .b : T
 Select the third component of e 
e: vec4<T>  e.w : Te .a : T
 Select the fourth component of e 
e: vecN<T> i: i32 or u32  e[i]: T  Select the i’^{th} component of vector The first component is at index i=0. If i is outside the range [0,N1], then any valid value for T may be returned. 
6.7.1.2. Vector Multiple Component Selection
Precondition  Conclusion  Description 

e: vecN<T> I is the letter x , y , z , or w J is the letter x , y , z , or w  e. IJ: vec2<T>  Computes the twocomponent vector with first component e.I, and second component e.J. Letter z is valid only when N is 3 or 4.Letter w is valid only when N is 4.

e: vecN<T> I is the letter r , g , b , or a J is the letter r , g , b , or a  e. IJ: vec2<T>  Computes the twocomponent vector with first component e.I, and second component e.J. Letter b is valid only when N is 3 or 4.Letter a is valid only when N is 4.

e: vecN<T> I is the letter x , y , z , or w J is the letter x , y , z , or w K is the letter x , y , z , or w  e. IJK: vec3<T>  Computes the threecomponent vector with first component e.I, second component e.J, and third component e.K. Letter z is valid only when N is 3 or 4.Letter w is valid only when N is 4.

e: vecN<T> I is the letter r , g , b , or a J is the letter r , g , b , or a K is the letter r , g , b , or a  e. IJK: vec3<T>  Computes the threecomponent vector with first component e.I, second component e.J, and third component e.K. Letter b is only valid when N is 3 or 4.Letter a is only valid when N is 4.

e: vecN<T> I is the letter x , y , z , or w J is the letter x , y , z , or w K is the letter x , y , z , or w L is the letter x , y , z , or w  e. IJKL: vec4<T>  Computes the fourcomponent vector with first component e.I, second component e.J, third component e.K, and fourth component e.L. Letter z is valid only when N is 3 or 4.Letter w is valid only when N is 4.

e: vecN<T> I is the letter r , g , b , or a J is the letter r , g , b , or a K is the letter r , g , b , or a L is the letter r , g , b , or a  e. IJKL: vec4<T>  Computes the fourcomponent vector with first component e.I, second component e.J, third component e.K, and fourth component e.L. Letter b is only valid when N is 3 or 4.Letter a is only valid when N is 4.

6.7.1.3. Component Reference from Vector Reference
A write access to component of a vector may access all of the memory locations associated with that vector.
Note: This means accesses to different components of a vector by different invocations must be synchronized if at least one access is a write access. See § 16.12 Synchronization Builtin Functions.
Precondition  Conclusion  Description 

r: ref<S,vecN<T>>  r.x : ref<S,T>r .r : ref<S,T>  Compute a reference to the first component of the vector referenced by the reference r. The originating variable of the resulting reference is the same as the originating variable of r. 
r: ref<S,vecN<T>>  r.y : ref<S,T>r .g : ref<S,T>  Compute a reference to the second component of the vector referenced by the reference r. The originating variable of the resulting reference is the same as the originating variable of r. 
r: ref<S,vecN<T>> N is 3 or 4  r.z : ref<S,T>r .b : ref<S,T>  Compute a reference to the third component of the vector referenced by the reference r. The originating variable of the resulting reference is the same as the originating variable of r. 
r: ref<S,vec4<T>>  r.w : ref<S,T>r .a : ref<S,T>  Compute a reference to the fourth component of the vector referenced by the reference r. The originating variable of the resulting reference is the same as the originating variable of r. 
r: ref<S,vecN<T>> i: i32 or u32  r[i] : ref<S,T> 
Compute a reference to the i’^{th} component of the vector
referenced by the reference r.
If i is outside the range [0,N1], then the expression evaluates to invalid memory reference. The originating variable of the resulting reference is the same as the originating variable of r. 
6.7.2. Matrix Access Expression
Precondition  Conclusion  Description 

e: matNxM<T> i: i32 or u32  e[i]: vecM<T> 
The result is the i’^{th} column vector of e.
If i is outside the range [0,N1], then any valid value for vecM<T> may be returned. 
Precondition  Conclusion  Description 

r: ref<S,matNxM<T>> i: i32 or u32  r[i] : ref<vecM<S,T>> 
Compute a reference to the i’^{th} column vector of the
matrix referenced by the reference r.
If i is outside the range [0,N1], then the expression evaluates to invalid memory reference. The originating variable of the resulting reference is the same as the originating variable of r. 
6.7.3. Array Access Expression
Precondition  Conclusion  Description 

e: array<T,N> i: i32 or u32  e[i] : T 
The result is the value of the i’^{th} element of the array value e.
If i is outside the range [0,N1], then any valid value for T may be returned. 
Precondition  Conclusion  Description 

r: ref<S,array<T,N>> i: i32 or u32  r[i] : ref<S,T> 
Compute a reference to the i’^{th} element of the array
referenced by the reference r.
If i is outside the range [0,N1], then the expression evaluates to an invalid memory reference. The originating variable of the resulting reference is the same as the originating variable of r. 
r: ref<S,array<T>> i: i32 or u32  r[i] : ref<S,T> 
Compute a reference to the i’^{th} element of the
runtimesized array referenced by the reference r.
If at runtime the array has N elements, and i is outside the range [0,N1], then the expression evaluates to an invalid memory reference. The originating variable of the resulting reference is the same as the originating variable of r. 
6.7.4. Structure Access Expression
Precondition  Conclusion  Description 

S is a structure type M is the identifier name of a member of S, having type T e: S  e.M: T  The result is the value of the member with name M from the structure value e. 
Precondition  Conclusion  Description 

S is a structure type M is the name of a member of S, having type T r: ref<S,S>  r.M: ref<S,T>  Given a reference to a structure, the result is a reference to the structure member with identifier name M. The originating variable of the resulting reference is the same as the originating variable of r. 
6.8. Logical Expressions
Precondition  Conclusion  Notes 

e: bool T is bool or vecN<bool>  ! e: T
 Logical negation.
The result is true when e is false and false when e is true . Componentwise when T is a vector.

Precondition  Conclusion  Notes 

e1: bool e2: bool  e1  e2: bool
 Shortcircuiting "or". Yields true if either e1 or e2 are true;
evaluates e2 only if e1 is false.

e1: bool e2: bool  e1 && e2: bool
 Shortcircuiting "and". Yields true if both e1 and e2 are true;
evaluates e2 only if e1 is true.

e1: T e2: T T is bool or vecN<bool>  e1  e2: T
 Logical "or". Componentwise when T is a vector. Evaluates both e1 and e2. 
e1: T e2: T T is bool or vecN<bool>  e1 & e2: T
 Logical "and". Componentwise when T is a vector. Evaluates both e1 and e2. 
6.9. Arithmetic Expressions
Precondition  Conclusion  Notes 

e: T T is i32 or vecN<i32>   e: T
 Signed integer negation. Componentwise when T is a vector. If e evaluates to the largest negative value, then the result is e. 
e: T T is f32 or vecN<f32>   e: T
 Floating point negation. Componentwise when T is a vector. 
Precondition  Conclusion  Notes 

e1 : T e2 : T T is i32, u32, vecN<i32>, or vecN<u32>  e1 + e2 : T
 Integer addition, modulo 2^{32}. Componentwise when T is a vector. 
e1 : T e2 : T T is f32 or vecN<f32>  e1 + e2 : T
 Floating point addition. Componentwise when T is a vector. 
e1 : T e2 : T T is i32, u32, vecN<i32>, or vecN<u32>  e1  e2 : T
 Integer subtraction, modulo 2^{32}. Componentwise when T is a vector. 
e1 : T e2 : T T is f32 or vecN<f32>  e1  e2 : T
 Floating point subtraction. Componentwise when T is a vector. 
e1 : T e2 : T T is i32, u32, vecN<i32>, or vecN<u32>  e1 * e2 : T
 Integer multiplication, modulo 2^{32}. Componentwise when T is a vector. 
e1 : T e2 : T T is f32 or vecN<f32>  e1 * e2 : T
 Floating point multiplication. Componentwise when T is a vector. 
e1 : T e2 : T T is i32 or vecN<i32>  e1 / e2 : T

Signed integer division. Componentwise when T is a vector.
In the scalar case, evaluates to:
Note: The need to ensure truncation behaviour may require an implementation to perform more operations than when computing an unsigned division. Use unsigned division when both operands are known to have the same sign. 
e1 : T e2 : T T is u32 or vecN<u32>  e1 / e2 : T

Unsigned integer division. Componentwise when T is a vector.
In the scalar case, evaluates to:

e1 : T e2 : T T is f32 or vecN<f32>  e1 / e2 : T
 Floating point division. Componentwise when T is a vector. 
e1 : T e2 : T T is i32 or vecN<i32>  e1 % e2 : T

Signed integer remainder. Componentwise when T is a vector.
In the scalar case, evaluates e1 and e2 once, and evaluates to:
Note: When nonzero, the result has the same sign as e1. Note: The need to ensure consistent behaviour may require an implementation to perform more operations than when computing an unsigned remainder. 
e1 : T e2 : T T is u32 or vecN<u32>  e1 % e2 : T

Unsigned integer remainder. Componentwise when T is a vector.
In the scalar case, evaluates to:

e1 : T e2 : T T is f32 or vecN<f32>  e1 % e2 : T
 Floating point remainder, where sign of nonzero result matches sign of e1. Componentwise when T is a vector. Result equal to: e1  e2 * trunc(e1 / e2) 
Preconditions  Conclusions  Semantics 

S is one of f32, i32, u32 V is vecN<S> es: S ev: V  ev + es: V
 ev + V(es)

es + ev: V
 V(es) + ev
 
ev  es: V
 ev  V(es)
 
es  ev: V
 V(es)  ev
 
ev * es: V
 ev * V(es)
 
es * ev: V
 V(es) * ev
 
ev / es: V
 ev / V(es)
 
es / ev: V
 V(es) / ev
 
ev % es: V
 ev % V(es)
 
es % ev: V
 V(es) % ev

Preconditions  Conclusions  Semantics 

e1, e2: matMxN<f32>  e1 + e2: matMxN<f32>  Matrix addition: column i of the result is e1[i] + e2[i] 
e1  e2: matMxN<f32>
 Matrix subtraction: column i of the result is e1[i]  e2[i]  
m: matMxN<f32> s: f32  m * s: matMxN<f32>  Componentwise scaling: (m * s)[i][j] is m[i][j] * s

s * m: matMxN<f32>  Componentwise scaling: (s * m)[i][j] is m[i][j] * s
 
m: matMxN<f32> v: vecM<f32>  m * v: vecN<f32>  Linear algebra matrixcolumnvector product:
Component i of the result is dot (m[i],v)

m: matMxN<f32> v: vecN<f32>  v * m: vecM<f32>  Linear algebra rowvectormatrix product: transpose(transpose(m) * transpose(v))

e1: matKxN<f32> e2: matMxK<f32>  e1 * e2: matMxN<f32>  Linear algebra matrix product. 
6.10. Comparison Expressions
Precondtion  Conclusion  Notes 

e1: T e2: T T is bool or vecN<bool>  e1 == e2: T
 Equality. Componentwise when T is a vector. 
e1: T e2: T T is bool or vecN<bool>  e1 != e2: T
 Inequality. Componentwise when T is a vector. 
e1: TI e2: TI TI is i32, u32, vecN<i32>, or vecN<u32> TB is bool if TI is scalar, or vecN<bool> if TI is a vector  e1 == e2: TB
 Equality. Componentwise when TI is a vector. 
e1: TI e2: TI TI is i32, u32, vecN<i32>, or vecN<u32> TB is bool if TI is scalar, or vecN<bool> if TI is a vector  e1 != e2: TB
 Inequality. Componentwise when TI is a vector. 
e1: TI e2: TI TI is i32 or vecN<i32> TB is bool if TI is scalar, or vecN<bool> if TI is a vector  e1 < e2: TB
 Less than. Componentwise when TI is a vector. 
e1: TI e2: TI TI is i32 or vecN<i32> TB is bool if TI is scalar, or vecN<bool> if TI is a vector  e1 <= e2: TB
 Less than or equal. Componentwise when TI is a vector. 
e1: TI e2: TI TI is i32 or vecN<i32> TB is bool if TI is scalar, or vecN<bool> if TI is a vector  e1 > e2: TB
 Greater than. Componentwise when TI is a vector. 
e1: TI e2: TI TI is i32 or vecN<i32> TB is bool if TI is scalar, or vecN<bool> if TI is a vector  e1 >= e2: TB
 Greater than or equal. Componentwise when TI is a vector. 
e1: TI e2: TI TI is u32 or vecN<u32> TB is bool if TI is scalar, or vecN<bool> if TI is a vector  e1 < e2: TB
 Less than. Componentwise when TI is a vector. 
e1: TI e2: TI TI is u32 or vecN<u32> TB is bool if TI is scalar, or vecN<bool> if TI is a vector  e1 <= e2: TB
 Less than or equal. Componentwise when TI is a vector. 
e1: TI e2: TI TI is u32 or vecN<u32> TB is bool if TI is scalar, or vecN<bool> if TI is a vector  e1 > e2: TB
 Greater than. Componentwise when TI is a vector. 
e1: TI e2: TI TI is u32 or vecN<u32> TB is bool if TI is scalar, or vecN<bool> if TI is a vector  e1 >= e2: TB
 Greater than or equal. Componentwise when TI is vector. 
e1: TF e2: TF TF is f32 or vecN<f32> TB is bool if TF is scalar, or vecN<bool> if TF is a vector.  e1 == e2: TB  Equality. Componentwise when TF is a vector. 
e1: TF e2: TF TF is f32 or vecN<f32> TB is bool if TF is scalar, or vecN<bool> if TF is a vector.  e1 != e2: TB  Inequality. Componentwise when TF is a vector. 
e1: TF e2: TF TF is f32 or vecN<f32> TB is bool if TF is scalar, or vecN<bool> if TF is a vector.  e1 < e2: TB  Less than. Componentwise when TF is a vector. 
e1: TF e2: TF TF is f32 or vecN<f32> TB is bool if TF is scalar, or vecN<bool> if TF is a vector.  e1 <= e2: TB  Less than or equal. Componentwise when TF is a vector. 
e1: TF e2: TF TF is f32 or vecN<f32> TB is bool if TF is scalar, or vecN<bool> if TF is a vector.  e1 > e2: TB  Greater than. Componentwise when TF is a vector. 
e1: TF e2: TF TF is f32 or vecN<f32> TB is bool if TF is scalar, or vecN<bool> if TF is a vector.  e1 >= e2: TB  Greater than or equal. Componentwise when TF is a vector. 
6.11. Bit Expressions
Precondition  Conclusion  Notes 

e: T T is i32, u32, vecN<i32>, or vecN<u32>  ~ e : T
 Bitwise complement on e. Each bit in the result is the opposite of the corresponding bit in e. Componentwise when T is a vector. 
Precondition  Conclusion  Notes 

e1: T e2: T T is i32, u32, vecN<i32>, or vecN<u32>  e1  e2 : T
 Bitwiseor. Componentwise when T is a vector. 
e1: T e2: T T is i32, u32, vecN<i32>, or vecN<u32>  e1 & e2 : T
 Bitwiseand. Componentwise when T is a vector. 
e1: T e2: T T is i32, u32, vecN<i32>, or vecN<u32>  e1 ^ e2 : T
 Bitwiseexclusiveor. Componentwise when T is a vector. 
Precondition  Conclusion  Notes 

e1: T e2: TS T is i32, u32, vecN<i32>, or vecN<u32> TS is u32 if e1 is a scalar, or vecN<u32>.  e1 << e2: T
 Logical shift left: Shift e1 left, inserting zero bits at the least significant positions, and discarding the most significant bits. The number of bits to shift is the value of e2 modulo the bit width of e1. Componentwise when T is a vector. 
e1: T e2: T T is u32 or vecN<u32>  e1 >> e2: T  Logical shift right: Shift e1 right, inserting zero bits at the most significant positions, and discarding the least significant bits. The number of bits to shift is the value of e2 modulo the bit width of e1. Componentwise when T is a vector. 
e1: T e2: TS T is i32 or vecN<i32> TS is u32 if e1 is a scalar, or vecN<u32>.  e1 >> e2: T  Arithmetic shift right: Shift e1 right, copying the sign bit of e1 into the most significant positions, and discarding the least significant bits. The number of bits to shift is the value of e2 modulo the bit width of e1. Componentwise when T is a vector. 
6.12. Function Call Expression
A function call expression executes a function call where the called function has a return type. If the called function does not return a value, a function call statement should be used instead. See § 7.5 Function Call Statement.
6.13. Variable Identifier Expression
Precondition  Conclusion  Description 

v is an identifier resolving to an inscope variable declared in address space S with store type T  v: ref<S,T>  Result is a reference to the memory for the named variable v. 
6.14. Formal Parameter Expression
Precondition  Conclusion  Description 

a is an identifier resolving to an inscope formal parameter declaration with type T  a: T  Result is the value supplied for the corresponding function call operand at the call site invoking this instance of the function. 
6.15. AddressOf Expression
The addressof operator converts a reference to its corresponding pointer.
Precondition  Conclusion  Description 

r: ref<S,T,A>  & r: ptr<S,T,A>

Result is the pointer value corresponding to the
same memory view as the reference value r.
If r is an invalid memory reference, then the resulting pointer is also an invalid memory reference. It is a shadercreation error if S is the handle address space. It is a shadercreation error if r is a reference to a vector component. 
6.16. Indirection Expression
The indirection operator converts a pointer to its corresponding reference.
Precondition  Conclusion  Description 

p: ptr<S,T,A>  * p: ref<S,T,A>

Result is the reference value corresponding to the
same memory view as the pointer value p.
If p is an invalid memory reference, then the resulting reference is also an invalid memory reference. 
6.17. Constant Identifier Expression
Precondition  Conclusion  Description 

c is an identifier resolving to an inscope override declaration with type T  c: T  If pipeline creation specified a value for the constant ID,
then the result is that value.
This value may be different for different pipeline instances. Otherwise, the result is the value computed for the initializer expression. Pipelineoverridable constants appear at modulescope, so evaluation occurs before the shader begins execution. Note: Pipeline creation fails if no initial value was specified in the API call and the let declaration has no initializer expression.

c is an identifier resolving to
an inscope let declaration with type T
 c: T  Result is the value computed for the initializer expression. For a let declaration at module scope, evaluation occurs before the shader begins execution.For a let declaration inside a function, evaluation occurs each time control reaches
the declaration. 
6.18. Expression Grammar Summary
 ident argument_expression_list ?
 paren_left ( ( expression comma ) * expression comma ? ) ? paren_right
 bracket_left expression bracket_right postfix_expression ?
 ( star  and ) * core_lhs_expression postfix_expression ?
 ident
 multiplicative_expression star unary_expression
 shift_expression less_than shift_expression
 shift_expression greater_than shift_expression
 shift_expression less_than_equal shift_expression
 shift_expression greater_than_equal shift_expression
 short_circuit_and_expression and_and relational_expression
 short_circuit_or_expression or_or relational_expression
 short_circuit_and_expression and_and relational_expression
 binary_and_expression and unary_expression
7. Statements
Statements are program fragments that control its execution. Statements are generally executed in sequential order; however, control flow statements may cause a program to execute in nonsequential order.
7.1. Compound Statement
A compound statement is a braceenclosed sequence of zero or more statements. When a declaration is one of those statements, its identifier is in scope from the start of the next statement until the end of the compound statement.
7.2. Assignment Statement
An assignment evaluates an expression, and optionally stores it in memory (thus updating the contents of a variable).
 lhs_expression ( equal  compound_assignment_operator ) expression
The text to the left of the operator token is the lefthand side, and the expression to the right of the operator token is the righthand side.
7.2.1. Simple Assignment
An assignment is a simple assignment when the lefthand side is an expression, and the operator is the equal token. In this case the value of the righthand side is written to the memory referenced by the lefthand side.
Precondition  Statement  Description 

r: ref<S,T,A>, A is write or read_write e: T, T is a constructible type, S is a writable address space  r = e  Evaluates e, evaluates r, then writes the value computed for e into the memory locations referenced by r. Note: if the reference is an invalid memory reference, the write may not execute, or may write to a different memory location than expected. 
In the simplest case, the left hand side is the name of a variable. See § 4.4.4 Forming Reference and Pointer Values for other cases.
struct S { age :i32 ; weight :f32 ; } var < private > person :S ; fn f () { var a :i32 = 20 ; a = 30 ; // Replace the contents of 'a' with 30. person . age = 31 ; // Write 31 into the age field of the person variable. var uv :vec2 < f32 > ; uv . y = 1.25 ; // Place 1.25 into the second component of uv. let uv_x_ptr :ptr < function , f32 > = & uv . x ; * uv_x_ptr = 2.5 ; // Place 2.5 into the first component of uv. var friend :S ; // Copy the contents of the 'person' variable into the 'friend' variable. friend = person ; }
7.2.2. Phony Assignment
An assignment is a phony assignment when the lefthand side is an underscore token. In this case the righthand side is evaluated, and then ignored.
Precondition  Statement  Description 

e: T, T is constructible, a pointer type, a texture type, or a sampler type  _ = e 
Evaluates e.
Note: The resulting value is not stored.
The 
A phonyassignment is useful for:

Calling a function that returns a value, but clearly expressing that the resulting value is not needed.

Statically accessing a variable, thus establishing it as a part of the shader’s resource interface.
Note: A buffer variable’s store type may not be constructible, e.g. it contains an atomic type, or a runtimesized array. In these cases, use a pointer to the variable’s contents instead.
var < private > counter :i32 ; fn increment_and_yield_previous () >i32 { let previous = counter ; counter = counter + 1 ; return previous ; } fn user () { // Increment the counter, but don’t use the result. _ = increment_and_yield_previous (); }
struct BufferContents { counter :atomic < u32 > ; data :array < vec4 < f32 >> ; } @ group ( 0 ) @ binding ( 0 ) var < storage > buf :BufferContents ; @ group ( 0 ) @ binding ( 1 ) var t :texture_2d < f32 > ; @ group ( 0 ) @ binding ( 2 ) var s :sampler ; @ stage ( fragment ) fn shade_it () >@ location ( 0 ) vec4 < f32 > { // Declare that buf, t, and s are part of the shader interface, without // using them for anything. _ = & buf ; _ = t ; _ = s ; return vec4 < f32 > (); }
7.2.3. Compound assignment
An assignment is a compound assignment when the lefthand side is an expression, and the operator is one of the compound_assignment_operators.
 or_equal
The type requirements, semantics, and behavior of each statement is defined as if the compound assignment expands as in the following table, except that the reference expression e1 is evaluated only once.
Statement  Expansion 

e1 += e2  e1 = e1 + (e2) 
e1 = e2  e1 = e1  (e2) 
e1 *= e2  e1 = e1 * (e2) 
e1 /= e2  e1 = e1 / (e2) 
e1 %= e2  e1 = e1 % (e2) 
e1 &= e2  e1 = e1 & (e2) 
e1 = e2  e1 = e1  (e2) 
e1 ^= e2  e1 = e1 ^ (e2) 
e1 >>= e2  e1 = e1 >> (e2) 
e1 <<= e2  e1 = e1 << (e2) 
Note: The syntax does not allow a compound assignment to also be a phony assignment.
Note: Even though the reference e1 is evaluated once, its underlying memory is accessed twice: first a read access gets the old value, and then a write access stores the updated value.
var<private> next_item: i32 = 0; fn advance_item() > i32 { next_item += 1; // Adds 1 to next_item. return next_item  1; } fn bump_item() { var data: array<f32,10>; next_item = 0; // Adds 5.0 to data[0], calling advance_item() only once. data[advance_item()] += 5.0; // next_item will be 1 here. } fn precedence_example() { var value = 1; // The righthand side of a compound assignment is its own expression. value *= 2 + 3; // Same as value = value * (2 + 3); // 'value' now holds 5. }
Note: A compound assignment can rewritten as different WGSL code that uses a simple assignment instead. The idea is to use a pointer to hold the result of evaluating the reference once.
For example,
when e1 is not a reference to a component inside a vector, then e1+=
e2 can be rewritten as {let p = &(
e1); *p = *p + (
e2);}
,
where the identifier p
is chosen to be different from all other identifiers in the program.
When e1 is a reference to a component inside a vector, the above technique
needs to be modified because WGSL does not allow taking the address in that case.
For example, if ev is a reference to a vector, the statement ev[
c] +=
e2 can be rewritten as {let p = &(
ev); let c0 =
c; (*p)[c0] = (*p)[c0] + (
e2);}
, where
identifiers c0
and p
are chosen to be different from all other identifiers in the program.
7.3. Increment and Decrement Statements
An increment statement adds 1 to the contents of a variable. A decrement statement subtracts 1 from the contents of a variable.
The expression must evaluate to a reference with an integer scalar store type and read_write access mode.
Precondition  Statement  Description 

r : ref<SC,T,read_write>, T is integer scalar  r++
 Adds 1 to the contents of memory referenced by r. Same as r += T(1) 
r : ref<SC,T,read_write>, T is integer scalar  r
 Subtracts 1 from the contents of memory referenced by r. Same as r = T(1) 
fn f () { var a :i32 = 20 ; a ++ ; // Now a contains 21 a  ; // Now a contains 20 }
7.4. Control Flow
Control flow statements may cause the program to execute in nonsequential order.
7.4.1. If Statement
 if expression compound_statement ( else else_statement ) ?
An if statement conditionally executes at most one compound statement based on the evaluation of the condition expressions.
The if
statements in WGSL use an if/else if/else structure, that contains a single required if
clause, zero or more else if
clauses and a single optional else
clause.
Each of the expressions for the if
and else if
clause conditions must be scalar boolean expressions.
An if
statement is executed as follows:

The condition associated with the
if
clause is evaluated. If the result istrue
, control transfers to the first compound statement (immediately after the parenthesized condition expression). 
Otherwise, the condition of the next
else if
clause in textual order (if one exists) is evaluated and, if the result istrue
, control transfers to the associated compound statement.
This behavior is repeated for all
else if
clauses until one of the conditions evaluates totrue
.


If no condition evaluates to
true
, then control transfers to the compound statement associated with theelse
clause (if it exists).
7.4.2. Switch Statement
 case case_selectors colon brace_left case_body ? brace_right
 const_literal ( comma const_literal ) * comma ?
A switch statement transfers control to one of a set of case clauses, or to the default
clause,
depending on the evaluation of a selector expression.
The selector expression must be of a scalar integer type.
If the selector value equals a value in a case selector list, then control is transferred to
the body of that case clause.
If the selector value does not equal any of the case selector values, then control is
transferred to the default
clause.
Each switch statement must have exactly one default clause.
The case selector values must have the same type as the result of evaluating the selector expression.
A literal value must not appear more than once in the case selectors for a switch statement.
Note: The value of the literal is what matters, not the spelling.
For example 0
, 00
, and 0x0000
all denote the zero value.
When control reaches the end of a case body, control normally transfers to the first statement
after the switch statement.
Alternately, executing a fallthrough statement
transfers control to the body of the next case clause or
default clause, whichever appears next in the switch body.
A fallthrough
statement must not appear as the last statement in the last clause of a switch.
When a declaration appears in a case body, its identifier is in scope from
the start of the next statement until the end of the case body.
Note: Identifiers declared in a case body are not in scope of case bodies
which are reachable via a fallthrough
statement.
7.4.3. Loop Statement
 loop brace_left statement * continuing_statement ? brace_right
A loop statement repeatedly executes a loop body; the loop body is specified as a compound statement. Each execution of the loop body is called an iteration.
The identifier of a declaration in a loop is in scope from the start of the next statement until the end of the loop body. The declaration is executed each time it is reached, so each new iteration creates a new instance of the variable or constant, and reinitializes it.
This repetition can be interrupted by a break, return, or discard statement.
Optionally, the last statement in the loop body may be a continuing statement.
Note: The loop statement is one of the biggest differences from other shader languages.
This design directly expresses loop idioms commonly found in compiled code. In particular, placing the loop update statements at the end of the loop body allows them to naturally use values defined in the loop body.
let a: i32 = 2; var i: i32 = 0; // <1> loop { if i >= 4 { break; } a = a * 2; i++; }
 <1> The initialization is listed before the loop.
int a = 2; let int step = 1; for (int i = 0; i < 4; i += step) { if i % 2 == 0 continue; a *= 2; }
var a: i32 = 2; var i: i32 = 0; loop { if i >= 4 { break; } let step: i32 = 1; i = i + step; if i % 2 == 0 { continue; } a = a * 2; }
var a: i32 = 2; var i: i32 = 0; loop { if i >= 4 { break; } let step: i32 = 1; if i % 2 == 0 { continue; } a = a * 2; continuing { // <2> i = i + step; } }
 <2> The continue construct is placed at the end of the
loop
7.4.4. For Statement
 for_init ? semicolon expression ? semicolon for_update ?
The for statement takes the form for (initializer; condition; update_part) { body }
and is syntactic sugar on top of a loop statement with the same body
.
Additionally:

If
initializer
is nonempty, it is executed inside an additional scope before the first iteration. The scope of a declaration in the initializer extends to the end of the loop body. 
If
condition
is nonempty, it is checked at the beginning of the loop body and if unsatisfied then a § 7.4.6 Break Statement is executed. 
If
update_part
is nonempty, it becomes a continuing statement at the end of the loop body.
The initializer
of a for loop is executed once prior to executing the loop.
When a declaration appears in the initializer, its identifier is in scope until the end of the body
.
Unlike declarations in the body
, the declaration is not reinitialized each iteration.
The condition
, body
and update_part
execute in that order to form a loop iteration.
The body
is a special form of compound statement.
The identifier of a declaration in the body
is in scope from the start of
the next statement until the end of the body
.
The declaration is executed each time it is reached, so each new iteration
creates a new instance of the variable or constant, and reinitializes it.
for(var i: i32 = 0; i < 4; i++) { if a == 0 { continue; } a = a + 2; }
Converts to:
{ // Introduce new scope for loop variable i var i: i32 = 0; var a: i32 = 0; loop { if !(i < 4) { break; } if a == 0 { continue; } a = a + 2; continuing { i++; } } }
7.4.5. While Statement
The while statement is a kind of loop parameterized by a condition. At the start of each loop iteration, a boolean condition is evaluated. If the condition is false, the while loop ends execution. Otherwise, the rest of the iteration is executed.
A while loop can be viewed as syntactic sugar over either a loop or for statement. The following statement forms are equivalent:

while
condition{
body_statements}

loop { if !
condition{break;}
body_statements}

for (;
condition;) {
body_statements}
7.4.6. Break Statement
 break
A break statement transfers control to the first statement
after the body of the nearestenclosing loop or switch statement.
A break
statement must only be used within loop, for, and switch statements.
When a break
statement is placed such that it would exit from a loop’s continuing statement,
then:

The
break
statement must appear as either:
The only statement in the
if
clause of anif
statement that has:
no
else
clause or an emptyelse
clause 
no
else if
clauses


The only statement in the
else
clause of anif
statement that has an emptyif
clause and noelse if
clauses.


That
if
statement must appear last in thecontinuing
clause.
var a: i32 = 2; var i: i32 = 0; loop { let step: i32 = 1; if i % 2 == 0 { continue; } a = a * 2; continuing { i = i + step; if i >= 4 { break; } } }
var a: i32 = 2; var i: i32 = 0; loop { let step: i32 = 1; if i % 2 == 0 { continue; } a = a * 2; continuing { i = i + step; if i < 4 {} else { break; } } }
var a: i32 = 2; var i: i32 = 0; loop { let step: i32 = 1; if i % 2 == 0 { continue; } a = a * 2; continuing { i = i + step; break; // Invalid: too early if i < 4 { i++; } else { break; } // Invalid: if is too complex, and too early if i >= 4 { break; } else { i++; } // Invalid: if is too complex } }
7.4.7. Continue Statement
 continue
A continue statement transfers control in the nearestenclosing loop:

forward to the continuing statement at the end of the body of that loop, if it exists.

otherwise backward to the first statement in the loop body, starting the next iteration.
A continue
statement must only be used in a loop or for statement.
A continue
statement must not be placed such that it would transfer
control to an enclosing continuing statement.
(It is a forward branch when branching to a continuing
statement.)
A continue
statement must not be placed such that it would transfer
control past a declaration used in the targeted continuing statement.
var i: i32 = 0; loop { if i >= 4 { break; } if i % 2 == 0 { continue; } // <3> let step: i32 = 2; continuing { i = i + step; } }
 <3> The
continue
is invalid because it bypasses the declaration ofstep
used in thecontinuing
construct
7.4.8. Continuing Statement
A continuing statement specifies a compound statement to be executed at the end of a loop iteration. The construct is optional.
The compound statement must not contain a return at any compound statement nesting level.
The compound statement must not contain a discard at any compound statement nesting level nor through function calls. See § 7.7 Statements Behavior Analysis for a more formal description of this rule.
7.4.9. Return Statement
 return expression ?
A return statement ends execution of the current function. If the function is an entry point, then the current shader invocation is terminated. Otherwise, evaluation continues with the next expression or statement after the evaluation of the call site of the current function invocation.
If the function does not have a return type, then the return statement is optional. If the return statement is provided for such a function, it must not supply a value. Otherwise the expression must be present, and is called the return value. In this case the call site of this function invocation evaluates to the return value. The type of the return value must match the return type of the function.
7.4.10. Discard Statement
A discard statement immediately ends execution of a fragment shader invocation and throws away the fragment.
The discard
statement must only be used in a fragment shader stage.
More precisely, executing a discard
statement will:

immediately terminate the current invocation, and

prevent evaluation and generation of a return value for the entry point, and

prevent the current fragment from being processed downstream in the GPURenderPipeline.
Only statements
executed prior to the discard
statement will have observable effects.
Note: A discard
statement may be executed by any function in a fragment stage and the effect is the same:
immediate termination of the invocation.
After a discard
statement is executed, control flow is nonuniform for the
duration of the entry point.
var<private> will_emit_color: bool = false; fn discard_if_shallow(pos: vec4<f32>) { if pos.z < 0.001 { // If this is executed, then the will_emit_color flag will // never be set to true. discard; } will_emit_color = true; } @stage(fragment) fn main(@builtin(position) coord_in: vec4<f32>) > @location(0) vec4<f32> { discard_if_shallow(coord_in); // Set the flag and emit red, but only if the helper function // did not execute the discard statement. will_emit_color = true; return vec4<f32>(1.0, 0.0, 0.0, 1.0); }
7.5. Function Call Statement
A function call statement executes a function call.
Note: If the function returns a value, that value is ignored.
7.6. Statements Grammar Summary
 func_call_statement semicolon
 variable_statement semicolon
 continue_statement semicolon
 assignment_statement semicolon
7.7. Statements Behavior Analysis
7.7.1. Rules
Some statements affecting controlflow are only valid in some contexts.
For example, fallthrough
is invalid outside of a switch, and continue
is invalid outside of a loop.
Additionally, the uniformity analysis (see § 12.2 Uniformity) needs to know when control flow can exit a statement in multiple different ways.
Both goals are achieved by a system for summarizing execution behaviors of statements and expressions. Behavior analysis maps each statement and expression to the set of possible ways execution proceeds after evaluation of the statement or expression completes. As with type analysis for values and expressions, behavior analysis proceeds bottom up: first determine behaviors for certain basic statements, and then determine behavior for higher level constructs by applying combining rules.
A behavior is a set, whose elements may be:

Return

Discard

Break

Continue

Fallthrough

Next
Each of those correspond to a way to exit a compound statement: either through a keyword, or by falling to the next statement ("Next").
We note "s: B" to say that s respects the rules regarding behaviors, and has behavior B.
For each function:

Its body must be a valid statement by these rules.

If the function has a return type, the behavior of its body must be one of {Return} or {Return, Discard}.

Otherwise, the behavior of its body must be a subset of {Next, Return, Discard}.
We assign a behavior to each function: it is its body’s behavior (treating the body as a regular statement), with any "Return" replaced by "Next". As a consequence of the rules above, a function behavior is always one of {}, {Next}, {Discard}, or {Next, Discard}.
Similarly, we assign a behavior to each expression, since expressions can include function calls, which can discard. Like functions, expression behaviors are always one of {}, {Next}, {Discard}, or {Next, Discard}.
Note: There is currently no valid program with an expression that does not have Next in its behavior. The reason is that only functions without a return type can have such a behavior, and there is no compound expression in which such a function can be called.
Statement  Preconditions  Resulting behavior 

empty statement  {Next}  
{s}  s: B  B 
s1 s2
Note: s1 often ends in a semicolon.  s1: B1 Next in B1 s2: B2  (B1∖{Next}) ∪ B2 
var x:T;  {Next}  
let x = e;  e: B  B 
var x = e;  e: B  B 
x = e;  x: B1 e: B2 x is not _
 B1 ∪ B2 
_ = e;  e: B  B 
f(e1, ..., en);  e1: B1 ... en: Bn f has behavior B  B ∪ ((B1 ∪ ... ∪ Bn)∖{Next}) 
return;  {Return}  
return e;  e: B  (B∖{Next}) ∪ {Return} 
discard;  {Discard}  
break;  {Break}  
continue;  {Continue}  
fallthrough;  {Fallthrough}  
if e s1 else s2  e: B s1: B1 s2: B2  (B∖{Next}) ∪ B1 ∪ B2 
loop {s1 continuing {s2}}  s1: B1 s2: B2 None of {Continue, Return, Discard} are in B2 Break is not in (B1 ∪ B2)  (B1 ∪ B2)∖{Continue, Next} 
s1: B1 s2: B2 None of {Continue, Return, Discard} are in B2 Break is in (B1 ∪ B2)  (B1 ∪ B2 ∪ {Next})∖{Break, Continue}  
switch e {case c1: s1 ... case cn: sn}  e: B s1: B1 ... sn: Bn Fallthrough is not in Bn Break is not in (B1 ∪ ... ∪ Bn)  ((B∖{Next}) ∪ B1 ∪ ... ∪ Bn)∖{Fallthrough} 
e: B s1: B1 ... sn: Bn Fallthrough is not in Bn Break is in (B1 ∪ ... ∪ Bn)  (B ∪ B1 ∪ ... ∪ Bn ∪ {Next})∖{Break, Fallthrough} 
Note: The empty statement case occurs when a loop
has an empty body, or when a for
loop lacks an initialization or update statement.
For the purpose of this analysis:

for
loops get desugared (see § 7.4.4 For Statement) 
while
loops get desugared (see § 7.4.5 While Statement) 
loop {s}
is treated asloop {s continuing {}}

if
statements without anelse
branch are treated as if they had an empty else branch (which adds Next to their behavior) 
if
statements withelse if
branches are treated as if they were nested simpleif/else
statements 
a switch_body starting with
default
behaves just like a switch_body starting withcase _:
Expression  Preconditions  Resulting behavior 

f(e1, ..., en)  e1: B1 ... en: Bn f has behavior B  B ∪ ((B1 ∪ ... ∪ Bn)∖{Next}) 
Any literal  {Next}  
Any variable reference  {Next}  
e1[e2]  e1: B1 e2: B2  B1 ∪ B2 
e.field  e: B  B 
e1  e2  e1: B1 e2: B2  B1 ∪ B2 
e1 && e2  e1: B1 e2: B2  B1 ∪ B2 
Each builtin function has a behavior of {Next}. And each operator application not listed in the table above has the same behavior as if it were a function call with the same operands and with a function’s behavior of {Next}.
A shadercreation error results if behavior analysis fails:

Behavior analysis must be able to determine a nonempty behavior for each statement, expression, and function.

The function behaviors must satisfy the rules given above.

The behaviors of compute and vertex entry points must not contain Discard.
7.7.2. Notes
This section is informative, nonnormative.
Here is the full list of ways that these rules can cause a program to be rejected (this is just restating information already listed above):

The body of a function (treated as a regular statement) has a behavior not included in {Next, Return, Discard}.

The body of a function with a return type has a behavior which is neither {Return} nor {Return, Discard}

A statement without Next in its behavior is not the last in a sequence of statements

The behavior of a continuing block contains any of Continue, Return, or Discard

The behavior of the last case of a switch contains Fallthrough

The behavior of a compute or vertex entry point function contains Discard

Some obviously infinite loops have an empty behaviour set, and are therefore invalid.
This analysis can be run in linear time, by analyzing the callgraph bottomup (since the behavior of a function call can depend on the function’s code).
7.7.3. Examples
Here are some examples showing this analysis in action:
fn simple () >i32 { var a :i32 ; return 0 ; // Behavior: {Return} a = 1 ; // Error: by the rule for sequences of statements, the previous statement should have a behavior that include "Next" return 2 ; }
fn nested () >i32 { var a :i32 ; { // The start of a compound statement. a = 2 ; // Behavior: {Next} return 1 ; // Behavior: {Return} } // The compound statement as a whole has behavior {Return} a = 1 ; // Error: for the same reason as in the previous example, the previous statement is missing "Next" in its behavior return 2 ; }
fn if_example () { var a :i32 = 0 ; loop { if a == 5 { break ; // Behavior: {Break} } // Behavior of the whole if compound statement: {Break, Next}, as the if has an implicit empty else a = a + 1 ; // valid, as the previous statement had "Next" in its behavior } }
fn if_example () { var a :i32 = 0 ; loop { if a == 5 { break ; // Behavior: {Break} } else { continue ; // Behavior: {Continue} } // Behavior of the whole if compound statement: {Break, Continue} a = a + 1 ; // Error: the previous statement is missing "Next" in its behavior } }
fn if_example () { var a :i32 = 0 ; loop { // if e1 s1 else if e2 s2 else s3 // is identical to // if e1 else { if e2 s2 else s3 } if a == 5 { break ; // Behavior: {Break} } else if a == 42 { continue ; // Behavior: {Continue} } else { return ; // Behavior {Return} } // Behavior of the whole if compound statement {Break, Continue, Return} } // Behavior of the whole loop compound statement {Next, Return} } // Behavior of the whole function {Next}
fn switch_example () { var a :i32 = 0 ; switch a { default :{ break ; // Behavior: {Break} } } // Behavior: {Next}, as switch replaces Break by Next a = 5 ; // valid, as the previous statement had Next in its behavior }
fn invalid_infinite_loop () { loop { } // Behavior: { }. Invalid because it’s empty. }
fn conditional_continue () { var a :i32 ; loop { if a == 5 { break ; } // Behavior: {Break, Next} if a % 2 == 1 { // valid, as the previous statement has Next in its behavior continue ; // Behavior: {Continue} } // Behavior: {Continue, Next} a = a * 2 ; // valid, as the previous statement has Next in its behavior continuing { // valid as the continuing statement has behavior {Next} which does not include any of {Break, Continue, Discard, Return} a = a + 1 ; } } // The loop as a whole has behavior {Next}, as it absorbs "Continue" and "Next", then replaces "Break" with "Next" }
fn redundant_continue_with_continuing () { var a :i32 ; loop { if a == 5 { break ; } continue ; // Valid. This is redundant, branching to the next statement. continuing { a = a + 1 ; } } }
fn continue_end_of_loop_body () { for ( var i :i32 = 0 ; i < 5 ; i ++ ) { continue ; // Valid. This is redundant, branching to the end of the loop body. } // Behavior: {Next}, as loops absorb "Continue", and "for" loops always add "Next" }
for
loops desugar to loop
with a conditional break. As shown in a previous example, the conditional break has behavior {Break, Next}, which leads to adding "Next" to the loop’s behavior.
fn always_discard () { discard ; } // The whole function has behavior {Discard} fn code_after_discard () { var a :i32 ; always_discard (); // Behavior: {Discard} a = a + 1 ; // Error: the previous statement is missing "Next" in its behavior }
fn sometimes_discard ( a :i32 ) { if a { discard ; // Behavior: {Discard} } // Behavior: {Next, Discard} } // The whole function has behavior {Next, Discard} fn code_after_discard () { var a :i32 ; a = 42 ; sometimes_discard ( a ); // Behavior: {Next, Discard} a = a + 1 ; // Valid } // The whole function has behavior {Next, Discard}
fn missing_return () >i32 { var a :i32 = 0 ; if a == 42 { return a ; // Behavior: {Return} } // Behavior: {Next, Return} } // Error: Next is invalid in the body of a function with a return type
fn continue_out_of_loop () { var a :i32 = 0 ; if a > 0 { continue ; // Behavior: {Continue} } // Behavior: {Next, Continue} } // Error: Continue is invalid in the body of a function
8. Functions
A function performs computational work when invoked.
A function is invoked in one of the following ways:

By evaluating a function call expression. See § 6.12 Function Call Expression.

By executing a function call statement. See § 7.5 Function Call Statement.

An entry point function is invoked by the WebGPU implementation to perform the work of a shader stage in a pipeline. See § 9 Entry Points
There are two kinds of functions:

A builtin function is provided by the WGSL implementation, and is always available to a WGSL program. See § 16 Builtin Functions.

A userdefined function is declared in a WGSL program.
8.1. Declaring a Userdefined Function
A function declaration creates a userdefined function, by specifying:

An optional set of attributes.

The name of the function.

The formal parameter list: an ordered sequence of zero or more formal parameter declarations, separated by commas, and surrounded by parentheses.

An optional, possibly decorated, return type.

The function body. This is the set of statements to be executed when the function is called.
A function declaration must only occur at module scope. A function name is in scope for the entire program.
A formal parameter declaration specifies an identifier name and a type for a value that must be provided when invoking the function. A formal parameter may have attributes. See § 8.2 Function Calls. The identifier is in scope until the end of the function. Two formal parameters for a given function must not have the same name.
The return type, if specified, must be constructible.
 fn ident paren_left param_list ? paren_right ( arrow attribute * type_decl ) ?
WGSL defines the following attributes that can be applied to function declarations:
WGSL defines the following attributes that can be applied to function parameters and return types:
// Declare the add_two function. // It has two formal paramters, i and b. // It has a return type of i32. // It has a body with a return statement. fn add_two(i: i32, b: f32) > i32 { return i + 2; // A formal parameter is available for use in the body. } // A compute shader entry point function, 'main'. // It has no specified return type. // It invokes the ordinary_two function, and captures // the resulting value in the named value 'two'. @stage(compute) fn main() { let six: i32 = add_two(4, 5.0); }
8.2. Function Calls
A function call is a statement or expression which invokes a function.
The function containing the function call is the calling function, or caller. The function being invoked is the called function, or callee.
The function call:

Names the called function, and

Provides a parenthesized, commaseparated list of argument value expressions.
The function call must supply the same number of argument values as there are formal parameters in the called function. Each argument value must evaluate to the same type as the corresponding formal parameter, by position.
In summary, when calling a function:

Execution of the calling function is suspended.

The called function executes until it returns.

Execution of the calling function resumes.
A called function returns as follows:

A builtin function returns when its work has completed.

A userdefined function with a return type returns when it executes a return statement.

A userdefined function with no return type returns when it executes a return statement, or when execution reaches the end of its function body.
In detail, when a function call is executed the following steps occur:

Function call argument values are evaluated. The relative order of evaluation is lefttoright.

Execution of the calling function is suspended. All function scope variables and constants maintain their current values.

If the called function is userdefined, memory is allocated for each function scope variable in the called function.

Initialization occurs as described in § 5 Variable and Value Declarations.


Values for the formal parameters of the called function are determined by matching the function call argument values by position. For example, in the body of the called function the first formal parameter will denote the value of the first argument at the call site.

If the called function is userdefined, control is transferred to the first statement in its body.

The called function is executed, until it returns.

Control is transferred back to the calling function, and the called function’s execution is unsuspended. If the called function returns a value, that value is supplied for the value of the function call expression.
Note: The current function will not resume execution if the called function or any descendent called function executes a discard statement.
The location of a function call is referred to as a call site. Call sites are a dynamic context. As such, the same textual location may represent multiple call sites.
8.3. Restrictions on functions

A vertex shader must return the
position
builtin output value. See § 15 Builtin Values. 
An entry point must never be the target of a function call.

If a function has a return type, it must be a constructible type.

A function parameter must one the following types:

a constructible type

a pointer type

a texture type

a sampler type


Each function call argument must evaluate to the type of the corresponding function parameter.

In particular, an argument that is a pointer must agree with the formal parameter on address space, pointee type, and access mode.


For userdefined functions, a parameter of pointer type must be in one of the following address spaces:

For builtin functions, a parameter of pointer type must be in one of the following address spaces:

Each argument of pointer type to a userdefined function must be one of:

An addressof expression of a variable identifier expression

A function parameter


WGSL assumes no aliasing is present between any combination of function parameters and variables. As such, a function parameter of pointer type must not be used to read or write to any memory locations of its originating variable that are also written via:

Another function parameter in the same function

A statement or expression in the function using the originating variable directly

Note: The aliasing restriction applies to memory location written by function calls in the function.
Note: Recursion is disallowed because cycles are not permitted among any kinds of declarations.
Revisit aliasing rules for clarity.
9. Entry Points
An entry point is a userdefined function that performs the work for a particular shader stage.
9.1. Shader Stages
WebGPU issues work to the GPU in the form of draw or dispatch commands. These commands execute a pipeline in the context of a set of inputs, outputs, and attached resources.
A pipeline describes the work to be performed on the GPU, as a sequence of stages, some of which are programmable. In WebGPU, a pipeline is created before scheduling a draw or dispatch command for execution. There are two kinds of pipelines: GPUComputePipeline, and GPURenderPipeline.
A dispatch command uses a GPUComputePipeline to run a compute shader stage over a logical grid of points with a controllable amount of parallelism, while reading and possibly updating buffer and image resources.
A draw command uses a GPURenderPipeline to run a multistage process with two programmable stages among other fixedfunction stages:

A vertex shader stage maps input attributes for a single vertex into output attributes for the vertex.

Fixedfunction stages map vertices into graphic primitives (such as triangles) which are then rasterized to produce fragments.

A fragment shader stage processes each fragment, possibly producing a fragment output.

Fixedfunction stages consume a fragment output, possibly updating external state such as color attachments and depth and stencil buffers.
The WebGPU specification describes pipelines in greater detail.
WGSL defines three shader stages, corresponding to the programmable parts of pipelines:

compute

vertex

fragment
Each shader stage has its own set of features and constraints, described elsewhere.
9.2. Entry Point Declaration
To create an entry point, declare a userdefined function with a stage attribute.
When configuring a pipeline in the WebGPU API,
the entry point’s function name maps to the entryPoint
attribute of the WebGPU § GPUProgrammableStage object.
The entry point’s formal parameters form the stage’s pipeline inputs. The entry point’s return type, if specified, forms the stage’s pipeline output. Each input and output must be an entry point IO type.
Note: Compute entry points never have a return type.
@stage(vertex) fn vert_main() > @builtin(position) vec4<f32> { return vec4<f32>(0.0, 0.0, 0.0, 1.0); } @stage(fragment) fn frag_main(@builtin(position) coord_in: vec4<f32>) > @location(0) vec4<f32> { return vec4<f32>(coord_in.x, coord_in.y, 0.0, 1.0); } @stage(compute) fn comp_main() { }
The set of functions in a shader stage is the union of:

The entry point function for the stage.

The targets of function calls from within the body of a function in the shader stage, whether or not that call is executed.
The union is applied repeatedly until it stabilizes. It will stabilize in a finite number of steps.
9.2.1. Function Attributes for Entry Points
WGSL defines the following attributes that can be applied to entry point declarations:
Can we query upper bounds on workgroup size dimensions? Is it independent of the shader, or a property to be queried after creating the shader module?
@stage(compute) @workgroup_size(8,4,1) fn sorter() { } @stage(compute) @workgroup_size(8u) fn reverser() { } // Using an pipelineoverridable constant. @id(42) override block_width = 12u; @stage(compute) @workgroup_size(block_width) fn shuffler() { } // Error: workgroup_size must be specified on compute shader @stage(compute) fn bad_shader() { }
9.3. Shader Interface
The shader interface is the set of objects through which the shader accesses data external to the shader stage, either for reading or writing. The interface includes:

Pipeline inputs and outputs

Buffer resources

Texture resources

Sampler resources
These objects are represented by modulescope variables in certain address spaces.
When an identifier used in a function declaration resolves to a modulescope variable,
then we say the variable is statically accessed by the function.
Static access of a let
declared constant is defined similarly.
Note that being statically accessed is independent of whether an execution of the shader
will actually evaluate the expression referring to the variable,
or even execute the statement that may enclose the expression.
More precisely, the interface of a shader stage consists of:

all parameters of the entry point

the result value of the entry point

all module scope variables that are statically accessed by functions in the shader stage, and which are in address spaces uniform, storage, or handle.
9.3.1. Pipeline Input and Output Interface
The entry point IO types include the following:

Builtin values. See § 9.3.1.1 Builtin Inputs and Outputs.

Userdefined IO. See § 9.3.1.2 Userdefined Inputs and Outputs

Structures containing only builtin values and userdefined IO. The structure must not contain a nested structure.
A pipeline input is data provided to the shader stage from upstream in the pipeline. A pipeline input is denoted by the arguments of the entry point.
A pipeline output is data the shader provides for further processing downstream in the pipeline. A pipeline output is denoted by the return type of the entry point.
Each pipeline input or output is one of:

A builtin value. See § 9.3.1.1 Builtin Inputs and Outputs.

A userdefined value. See § 9.3.1.2 Userdefined Inputs and Outputs.
9.3.1.1. Builtin Inputs and Outputs
A builtin input value provides access to systemgenerated control information. The set of builtin inputs are listed in § 15 Builtin Values.
A builtin input for stage S with name X and type T_{X} is accessed via a formal parameter to an entry point for shader stage S, in one of two ways:

The parameter has attribute
builtin(
X)
and is of type T_{X}. 
The parameter has structure type, where one of the structure members has attribute
builtin(
X)
and is of type T_{X}.
Conversely, when a parameter or member of a parameter for an entry point has a builtin
attribute,
the corresponding builtin must be an input for the entry point’s shader stage.
A builtin output value is used by the shader to convey control information to later processing steps in the pipeline. The set of builtin outputs are listed in § 15 Builtin Values.
A builtin output for stage S with name Y and type T_{Y} is set via the return value for an entry point for shader stage S, in one of two ways:

The entry point return type has attribute
builtin(
Y)
and is of type T_{Y}. 
The entry point return type has structure type, where one of the structure members has attribute
builtin(
Y)
and is of type T_{Y}.
Conversely, when the return type or member of a return type for an entry point has a builtin
attribute,
the corresponding builtin must be an output for the entry point’s shader stage.
Note: The position
builtin is both an output of a vertex shader, and an input to the fragement shader.
9.3.1.2. Userdefined Inputs and Outputs
Userdefined data can be passed as input to the start of a pipeline, passed between stages of a pipeline or output from the end of a pipeline. Userdefined IO must not be passed to compute shader entry points. Userdefined IO must be of numeric scalar or numeric vector type, or of a structure type whose members are numeric scalars or vectors. All userdefined IO must be assigned locations (See § 9.3.1.4 Inputoutput Locations).
9.3.1.3. Interpolation
Authors can control how userdefined IO data is interpolated through the use of the interpolate attribute. WGSL offers two aspects of interpolation to control: the type of interpolation, and the sampling of the interpolation.
The interpolation type must be one of:

perspective
 Values are interpolated in a perspective correct manner. 
linear
 Values are interpolated in a linear, nonperspective correct manner. 
flat
 Values are not interpolated. Interpolation sampling is not used withflat
interpolation.
The interpolation sampling must be one of:

center
 Interpolation is performed at the center of the pixel. 
centroid
 Interpolation is performed at a point that lies within all the samples covered by the fragment within the current primitive. This value is the same for all samples in the primitive. 
sample
 Interpolation is performed per sample. The fragment shader is invoked once per sample when this attribute is applied.
For userdefined IO of scalar or vector floatingpoint type:

If the interpolation attribute is not specified, then
@interpolate(perspective, center)
is assumed. 
If the interpolation attribute is specified with an interpolation type:

If the interpolation type is
flat
, then interpolation sampling must not be specified. 
If the interpolation type is
perspective
orlinear
, then:
Any interpolation sampling is valid.

If interpolation sampling is not specified,
center
is assumed.


Userdefined IO of scalar or vector integer type must always be specified as @interpolate(flat)
.
Interpolation attributes must match between vertex outputs and fragment inputs with the same location assignment within the same pipeline.
9.3.1.4. Inputoutput Locations
Each location can store a value up to 16 bytes in size. The byte size of a type is defined using the SizeOf column in § 4.3.7.1 Alignment and Size. For example, a fourcomponent vector of floatingpoint values occupies a single location.
Locations are specified via the location attribute.
Every userdefined input and output must have a fully specified set of locations. Each structure member in the entry point IO must be one of either a builtin value (see § 9.3.1.1 Builtin Inputs and Outputs), or assigned a location.
Locations must not overlap within each of the following sets:

Members within a structure type. This applies to any structure, not just those used in pipeline inputs or outputs.

An entry point’s pipeline inputs, i.e. locations for its formal parameters, or for the members of its formal parameters of structure type.
Note: Location numbering is distinct between inputs and outputs: Location numbers for an entry point’s pipeline inputs do not conflict with location numbers for the entry point’s pipeline outputs.
Note: No additional rule is required to prevent location overlap within an entry point’s outputs. When the output is a structure, the first rule above prevents overlap. Otherwise, the output is a scalar or a vector, and can have only a single location assigned to it.
Note: The number of available locations for an entry point is defined by the WebGPU API.
Userdefined IO can be mixed with builtin values in the same structure. For example,
// Mixed builtins and userdefined inputs. struct MyInputs { @location(0) x: vec4<f32>; @builtin(front_facing) y: bool; @location(1) @interpolate(flat) z: u32; } struct MyOutputs { @builtin(frag_depth) x: f32; @location(0) y: vec4<f32>; } @stage(fragment) fn fragShader(in1: MyInputs) > MyOutputs { // ... }
struct A { @location(0) x: f32; // Invalid, x and y cannot share a location. @location(0) y: f32; } struct B { @location(0) x: f32; } struct C { // Invalid, structures with userdefined IO cannot be nested. b: B; } struct D { x: vec4<f32>; } @stage(fragment) // Invalid, location cannot be applied to a structure type. fn fragShader1(@location(0) in1: D) { // ... } @stage(fragment) // Invalid, in1 and in2 cannot share a location. fn fragShader2(@location(0) in1: f32, @location(0) in2: f32) { // ... } @stage(fragment) // Invalid, location cannot be applied to a structure. fn fragShader3(@location(0) in1: vec4<f32>) > @location(0) D { // ... }
9.3.2. Resource Interface
A resource is an object, other than a pipeline input or output, which provides access to data external to a shader stage. Resources are shared by all invocations of the shader.
There are four kinds of resources:

textures

samplers
The resource interface of a shader is the set of modulescope resource variables statically accessed by functions in the shader stage.
Each resource variable must be declared with both group and binding attributes. Together with the shader’s stage, these identify the binding address of the resource on the shader’s pipeline. See WebGPU § GPUPipelineLayout.
Bindings must not alias within a shader stage: two different variables in the resource interface of a given shader must not have the same group and binding values, when considered as a pair of values.
9.3.3. Resource Layout Compatibility
WebGPU requires that a shader’s resource interface match the layout of the pipeline using the shader.
Each WGSL variable in a resource interface must be bound to a WebGPU resource with a compatible resource type and binding type, where compatibility is defined by the following table.
WGSL resource  WebGPU Resource type  WebGPU Binding type  

uniform buffer  GPUBufferBinding  GPUBufferBindingType  uniform 
storage buffer with read_write access  storage  
storage buffer with read access  readonlystorage  
sampler  GPUSampler  GPUSamplerBindingType  filtering 
nonfiltering  
sampler_comparison  comparison  
sampled texture  GPUTextureView  GPUTextureSampleType  float 
unfilterablefloat  
sint  
uint  
depth  
writeonly storage texture  GPUTextureView  GPUStorageTextureAccess  writeonly 
See the WebGPU API specification for interface validation requirements.
10. Language Extensions
The WGSL language is expected to evolve over time.
An extension is a named grouping for a coherent set of modifications to a particular version of the WGSL specification, consisting of any combination of:

Addition of new concepts and behaviours via new syntax, including:

declarations, statements, attributes, and builtin functions.


Removal of restrictions in the current specification or in previously published extensions.

Syntax for reducing the set of permissible behaviours.

Syntax for limiting the features available to a part of the program.

A description of how the extension interacts with the existing specification, and optionally with other extensions.
Hypothetically, extensions could be used to:

Add numeric scalar types, such as 16bit integers.

Add syntax to constrain floating point rounding mode.

Add syntax to signal that a shader does not use atomic types.

Add new kinds of statements.

Add new builtin functions.

Add constraints on how shader invocations execute.

Add new shader stages.
10.1. Enable Directive
An enable directive indicates that the functionality
described by a particular named extension may be used in the source text after the directive itself.
That is, language functionality described by the extension may be used in any
source text after the enable
directive.
The directive must not appear inside the text of any declaration. (If it were a declaration, it would be at module scope.)
The directive uses an identifier to name the extension, but does not create a scope for the identifier. Use of the identifier by the directive does not conflict with the use of that identifier as the name in any declaration.
Note: The grammar rule includes the terminating semicolon token,
ensuring the additional functionality is usable only after that semicolon.
Therefore any WGSL implementation can parse the entire enable
directive.
When an implementation encounters an enable directive for an unsupported extension,
the implementation can issue a clear diagnostic.
// Enable a hypothetical IEEE754 binary16 floating point extension. enable f16; enable f16; // A redundant enable directive is ok. // Enable a hypothetical extension adding syntax for controlling // the rounding mode on f16 arithmetic. enable rounding_mode_f16; // Assuming the f16 extension enables use of the f16 type: //  as function return value //  as the type for let declaration //  as a type constructor, with an i32 argument //  as operands to the division operator: / fn halve_it(x: f16) > f16 { let two: f16 = f16(2); return x / two; } @round_to_even_f16 // Attribute enabled by the rounding_mode_f16 extension fn triple_it(x: f16) > f16 { return x * f16(3); // Uses roundtoeven. }
11. WGSL Program
A WGSL program is a sequence of optional directives followed by module scope declarations.
 global_directive * global_decl *
 global_variable_decl semicolon
11.1. Limits
A program must satisfy the following limits:
Limit  Maximum value 

Number of members in a structure type  16383 
Nesting depth of a composite type  255 
Number of parameters for a function  255 
Number of case selector values in a switch statement  16383 
12. Execution
§ 1.1 Technical Overview describes how a shader is invoked and partitioned into invocations. This section describes further constraints on how invocations execute, individually and collectively.
12.1. Program Order Within an Invocation
Each statement in a WGSL program may be executed zero or more times during execution. For a given invocation, each execution of a given statement represents a unique dynamic statement instance.
When a statement includes an expression, the statement’s semantics determines:

Whether the expression is evaluated as part of statement execution.

The relative ordering of evaluation between independent expressions in the statement.
Expression nesting defines data dependencies which must be satisfied to
complete evaluation.
That is, a nested expression must be evaluated before the enclosing expression
can be evaluated.
The order of evaluation for operands of an expression is lefttoright in
WGSL.
For example, foo() + bar()
must evaluate foo()
before bar()
.
See § 6 Expressions.
Statements in a WGSL program are executed in control flow order. See § 7 Statements and § 8.2 Function Calls.
12.2. Uniformity
12.2.1. Terminology and concepts
The following definitions are merely informative, trying to give an intuition for what the analysis in the next subsection is computing. The analysis is what actually defines these concepts, and when a program is valid or breaks the uniformity rules.
For a given group of invocations:

control flow is said to be uniform at a point in the program if all invocations execute in lockstep at that program point.

an expression is said to be uniform if it is in uniform control flow, and all invocations compute the same value for it.

a local variable is said to be uniform, if at every point where it is live, all invocations have the same value in that variable.
12.2.2. Uniformity analysis overview
Some functions (e.g. barriers and derivatives) are only safe to call in uniform control flow. In this section we specify an analysis that verifies that these functions are only called in such a context.

Sound (meaning that it rejects every program that would break the uniformity requirements of builtins)

Linear time complexity (in the number of tokens in the program)

Refactoring a piece of code into a function, or inlining a function, cannot make a shader invalid if it was valid before the transformation

If the analysis refuses a program, it provides a straightforward chain of implications that can be used by the user agent to craft a good error message
The analysis analyzes each function, verifying that there is a context where it is safe to call this function. It rejects the program as invalid if there is no such context.
At the same time, it computes metadata about the function to help analyze its callers in turn. This means that the call graph must first be built, and functions must be analyzed from the leaves upwards, i.e. from functions that call no function outside the standard library toward the entry point. This way, whenever a function is analyzed, the metadata for all of its callees has already been computed. There is no risk of being trapped in a cycle, as recurrence is forbidden in the language.
Note: another way of saying the same thing is that we do a topological sort of functions ordered by the "is a (possibly indirect) callee of" partial order, and analyze them in that order.
12.2.3. Analyzing the uniformity requirements of a function
Each function is analyzed in two phases.
The first phase walks over the syntax of the function, building a directed graph along the way based on the rules in the following subsections. The second phase explores that graph, resulting in either rejecting the program, or computing the constraints on calling this function.

The control flow at a specific program point is required to be uniform

The value of an expression is required to be uniform

A variable is required to be uniform
An edge can be understood as an implication from the statement corresponding to its source node to the statement corresponding to its target node.
To express that something must always be uniform (e.g. the control flow at the call site of a derivative), we add an edge from RequiredToBeUniform to the corresponding node. One way to understand this, is that RequiredToBeUniform corresponds to the proposition True, so that RequiredToBeUniform > X is the same as saying that X is true.
Reciprocally, to express that we cannot ensure the uniformity of something (e.g. a variable which holds the thread id), we add an edge from the corresponding node to MayBeNonUniform. One way to understand this, is that MayBeNonUniform corresponds to the proposition False, so that X > MayBeNonUniform is the same as saying that X is false.
A consequence of this interpretation is that every node reachable from RequiredToBeUniform corresponds to something which must be uniform for the program to be valid, and every node from which cannotBeUniform is reachable corresponds to something whose uniformity we cannot guarantee. It follows that we have a uniformity violation (and thus reject the program) if there is any path from RequiredToBeUniform to MayBeNonUniform.
In more detail, here is what we compute for each function:

A tag about the controlflow in which the function can be called: one of {AlwaysRequiredToBeUniform, RequiredToBeUniformForControlFlowAfter}

A tag for each parameter of the function: one of {AlwaysRequiredToBeUniform, RequiredToBeUniformForControlFlowAfter, RequiredToBeUniformForReturnValue, NoRestriction}

A tag for the function as a whole: one of {ControlFlowMayBeNonUniformAfter, ReturnValueMayBeNonUniform, NoRestriction}
Here is the algorithm for computing this data for a given function:

Create nodes called "RequiredToBeUniform", "MayBeNonUniform", "CF_start", "CF_return", and if the function is nonvoid a node called "Value_return"

Create one node for each parameter of the function which we’ll call "arg_i"

Walk over the syntax of the function, adding nodes and edges to the graph following the rules of the next sections (§ 12.2.4 Uniformity rules for statements, § 12.2.5 Uniformity rules for function calls, § 12.2.6 Uniformity rules for expressions), using CF_start as the starting controlflow for the function’s body.

Look at which nodes are reachable from "RequiredToBeUniform"

If this set includes the node "MayBeNonUniform", then reject the program.

If this set includes "CF_start", then the controlflow in which the function can be called is AlwaysRequiredToBeUniform.

Otherwise it is RequiredToBeUniformForControlFlowAfter

For each "arg_i" in this set, the corresponding parameter is AlwaysRequiredToBeUniform

Remove from the graph all nodes that have been visited


Look at which nodes are reachable from "CF_return"

If this set includes "MayBeNonUniform", then the function is ControlFlowMayBeNonUniformAfter

For each "arg_i" in this set, the corresponding parameter is RequiredToBeUniformForControlFlowAfter

Remove from the graph all nodes that have been visited


If "Value_return" exists, look at which nodes are reachable from it

If this set includes "MayBeNonUniform", then the function is ReturnValueMayBeNonUniform

For each "arg_i" in this set, the corresponding parameter is RequiredToBeUniformForReturnValue


If the function has not been assigned a tag, then it is NoRestriction.

For each parameter, if it has not been assigned a tag, then it is NoRestriction.
Note: The entire graph can be destroyed at this point. The tags listed above are all that we need to remember to analyze callers of this function.
12.2.4. Uniformity rules for statements
The rules for analyzing statements take as argument both the statement itself and the node corresponding to control flow at the beginning of it (which we’ll note "CF" below) and return both of the following:

A node corresponding to control flow at the exit of it

A set of new nodes and edges to add to the graph
In the table below, (CF1, S) => CF2
means "run the analysis on S starting with control flow CF1, apply the required changes to the graph, and name the resulting control flow CF2".
Similarly, (CF1, E) => (CF2, V)
means "run the analysis on expression E, starting with control flow CF1, apply the required changes to the graph, and name the resulting control flow node CF2 and the resulting value node V" (see next section for the analysis of expressions).
We have a similar set of rules for expressions in leftvalue positions, that we denote by LValue: (CF, E) => (CF, L)
. Instead of computing the node which corresponds to the uniformity of the value, it computes the node which corresponds to the uniformity of the variable we are addressing.
When several edges have to be created we use X > {Y, Z}
as a shorthand for X > Y, X > Z
.
Statement  New nodes  Recursive analyses  Resulting control flow node  New edges 

{s}  (CF, s) => CF'  CF'  
s1 s2
Note: s1 often ends in a semicolon.  (CF, s1) => CF1 (CF1, s2) => CF2  CF2  
if e s1 else s2 , with behavior {Next}  (CF, e) => (CF', V) (V, s1) => CF1 (V, s2) => CF2  CF  
if e s1 else s2 , with another behavior  CFend  CFend  CFend > {CF1, CF2}  
loop {s1 continuing {s2}} , with behavior {Next}  CF'  (CF', s1) => CF1 (CF1, s2) => CF2  CF  CF' > {CF2, CF} 
loop {s1 continuing {s2}} , with another behavior  CF'  
loop {s1}, with behavior {Next}  CF'  (CF', s1) => CF1  CF  CF' > {CF1, CF} 
loop {s1}, with another behavior  CF'  
switch e case _: s_1 .. case _: s_n , with behavior {Next}  (CF, e) => (CF', V) (V, s_1) => CF_1 ... if s_(n1) may fallthrough, (CF_(n1), s_n) => CF_n else (V, s_n) => CF_n  CF  
switch e case _: s_1 .. case _: s_n , with another behavior  CFend  CFend  CFend > {CF_1, ..., CF_n}  
var x: T;  CF  
break;  
continue;  
fallthrough;  
discard;  
return;  CF  CF_return > CF  
return e;  (CF, e) => (CF', V)  CF'  CF_return > CF' Value_return > V  
e2 = e1;  (CF, e1) => (CF1, V1) LValue: (CF1, e2) => (CF2, L2)  CF2  L2 > V1  
_ = e  (CF, e) => (CF', V)  CF'  
let x = e;  (CF, e) => (CF', V)  CF'  
var x = e; 
Analysis of for and while loops follows from their respective desugaring translations to loop statements.
In switch, a default
block is treated exactly like a case block with regards to uniformity.
Note: If the set of behaviors (see § 7.7 Statements Behavior Analysis) for an if, switch, or loop statement is {Next}, this means that we either did not diverge within the statement, or we reconverged, so we pick the node corresponding to control flow at the start of the statement as the node corresponding to control flow at the exit of the statement.
12.2.5. Uniformity rules for function calls
The most complex rule is for function calls:

For each argument, apply the corresponding expression rule, with the control flow at the exit of the previous argument (using the control flow at the beginning of the function call for the first argument). Name the corresponding value nodes "arg_i" and the corresponding control flow nodes "CF_i"

Create two new nodes, named "Result" and "CF_after"

If the controlflow at the start of the function was tagged AlwaysRequiredToBeUniform, then add an edge from RequiredToBeUniform to the last CF_i

Otherwise add an edge from CF_after to the last CF_i

If the function was tagged ControlFlowMayBeNonUniformAfter, then add an edge from CF_after to MayBeNonUniform

Else if it was tagged ReturnValueMayBeNonUniform, then add an edge from Result to MayBeNonUniform

Add an edge from Result to CF_after

For each argument i:

If the corresponding parameter was tagged AlwaysRequiredToBeUniform, then add an edge from RequiredToBeUniform to arg_i

Else if it was tagged RequiredToBeUniformForControlFlowAfter, then add an edge from CF_after to arg_i

Else if it was tagged RequiredToBeUniformForReturnValue, then add an edge from Result to arg_i

Note: Notice that this rule only requires adding a number of edges bounded by 3 + the number of parameters of the functions, independently of how complex the implementation of the function might be. This is key to the linear complexity of the overall algorithm.
Most builtin functions have tags of:

for the controlflow in which they can be called: RequiredToBeUniformForControlFlowAfter

for the function as a whole: NoRestriction

for each parameter: RequiredToBeUniformForReturnValue
Here is the list of exceptions:

All functions in § 16.12 Synchronization Builtin Functions are AlwaysRequiredToBeUniform

All functions in § 16.7 Derivative Builtin Functions, § 16.8.8 textureSample, § 16.8.9 textureSampleBias, and § 16.8.10 textureSampleCompare are AlwaysRequiredToBeUniform and ReturnValueMayBeNonUniform
12.2.6. Uniformity rules for expressions
The rules for analyzing expressions take as argument both the expression itself and the node corresponding to control flow at the beginning of it (which we’ll note "CF" below) and return the following:

A node corresponding to control flow at the exit of it

A node corresponding to its value

A set of new nodes and edges to add to the graph
Expression  New nodes  Recursive analyses  Resulting control flow node, value node  New edges 

e1  e2  (CF, e1) => (CF1, V1) (V1, e2) => (CF2, V2)  CF2, V2  
e1 && e2  
Literal  CF, CF  
reference to functionscope variable, letdeclaration, or nonbuiltin parameter "x"  Result  X is the node corresponding to "x"  CF, Result  Result > {CF, X} 
reference to uniform builtin value "x"  CF, CF  
reference to nonuniform builtin value "x"  CF, MayBeNonUniform  
reference to readonly modulescope variable "x"  CF, CF  
reference to nonreadonly modulescope variable "x"  CF, MayBeNonUniform  
op e, where "op" is a unary operator  (CF, e) => (CF', V)  CF', V  
e.field  
e1 op e2, where op is a nonshortcircuiting binary operator  Result  (CF, e1) => (CF1, V1) (CF1, e2) => (CF2, V2)  CF2, Result  Result > {V1, V2} 
e1[e2] 
The following builtin input variables are considered uniform:

workgroup_id

num_workgroups
All other ones (see § 15 Builtin Values) are considered nonuniform.
Expression  New nodes  Recursive analyses  Resulting control flow node, variable node  New edges 

reference to functionscope variable, letdeclaration, or parameter "x"  X is the node corresponding to "x"  CF, X  
reference to modulescope variable "x"  CF, MayBeNonUniform  
e.field  LValue: (CF, e) => (CF1, L1)  CF1, L1  
e1[e2]  LValue: (CF, e1) => (CF1, L1) (CF1, e2) => (CF2, V2)  CF2, L1  L1 > V2 
12.2.7. Annotating the uniformity of every point in the controlflow
This entire subsection is nonnormative.
If implementers want to provide developers with a diagnostic mode that shows for each point in the controlflow of the entire shader whether it is uniform or not (and thus whether it would be valid to call a function that requires uniformity there), we suggest the following:

Run the (mandatory, normative) analysis described in the previous subsections, keeping the graph for every function.

Reverse all edges in all of those graphs

Go through each function, starting with the entry point and never visiting a function before having visited all of its callers:

Add an edge from MayBeNonUniform to every argument that was nonuniform in at least one caller

Add an edge from MayBeNonUniform to CF_start if the function was called in nonuniform controlflow in at least one caller

Look at which nodes are reachable from MayBeNonUniform. Every node visited is an expression or point in the controlflow whose uniformity cannot be proven by the analysis

Any node which is not visited by these reachability analyses can be proven to be uniform by the analysis (and so it would be safe to call a derivative or similar function there).
Note: The bottomup analysis is still required, as it lets us know what edges to add to the graphs when encountering calls.
12.3. Compute Shaders and Workgroups
A workgroup is a set of invocations which concurrently execute a compute shader stage entry point, and share access to shader variables in the workgroup address space.
The workgroup grid for a compute shader is the set of points with integer coordinates (i,j,k) with:

0 ≤ i < workgroup_size_x

0 ≤ j < workgroup_size_y

0 ≤ k < workgroup_size_z
where (workgroup_size_x, workgroup_size_y, workgroup_size_z) is the value specified for the workgroup_size attribute of the entry point.
There is exactly one invocation in a workgroup for each point in the workgroup grid.
An invocation’s local invocation ID is the coordinate triple for the invocation’s corresponding workgroup grid point.
When an invocation has local invocation ID (i,j,k), then its local invocation index is
i + (j * workgroup_size_x) + (k * workgroup_size_x * workgroup_size_y)
Note that if a workgroup has W invocations, then each invocation I the workgroup has a unique local invocation index L(I) such that 0 ≤ L(I) < W, and that entire range is covered.
A compute shader begins execution when a WebGPU implementation removes a dispatch command from a queue and begins the specified work on the GPU. The dispatch command specifies a dispatch size, which is an integer triple (group_count_x, group_count_y, group_count_z) indicating the number of workgroups to be executed, as described in the following.
The compute shader grid for a particular dispatch is the set of points with integer coordinates (CSi,CSj,CSk) with:

0 ≤ CSi < workgroup_size_x × group_count_x

0 ≤ CSj < workgroup_size_y × group_count_y

0 ≤ CSk < workgroup_size_z × group_count_z
where workgroup_size_x, workgroup_size_y, and workgroup_size_z are as above for the compute shader entry point.
The work to be performed by a compute shader dispatch is to execute exactly one invocation of the entry point for each point in the compute shader grid.
An invocation’s global invocation ID is the coordinate triple for the invocation’s corresponding compute shader grid point.
The invocations are organized into workgroups, so that each invocation (CSi, CSj, CSk) is identified with the workgroup grid point
( CSi mod workgroup_size_x , CSj mod workgroup_size_y , CSk mod workgroup_size_z )
in workgroup ID
( ⌊ CSi ÷ workgroup_size_x ⌋, ⌊ CSj ÷ workgroup_size_y ⌋, ⌊ CSk ÷ workgroup_size_z ⌋).
WebGPU provides no guarantees about:

Whether invocations from different workgroups execute concurrently. That is, you cannot assume more than one workgroup executes at a time.

Whether, once invocations from a workgroup begin executing, that other workgroups are blocked from execution. That is, you cannot assume that only one workgroup executes at a time. While a workgroup is executing, the implementation may choose to concurrently execute other workgroups as well, or other queued but unblocked work.

Whether invocations from one particular workgroup begin executing before the invocations of another workgroup. That is, you cannot assume that workgroups are launched in a particular order.
12.4. Collective Operations
12.4.1. Barriers
A barrier is a synchronization builtin function that orders memory operations in a program. A control barrier is executed by all invocations in the same workgroup as if it were executed concurrently. As such, control barriers must only be executed in uniform control flow in a compute shader.
12.4.2. Derivatives
A partial derivative is the rate of change of a value along an axis.
Fragment shader invocations operating on neighbouring fragments (in screenspace coordinates) collaborate to compute approximate partial derivatives. These neighbouring fragments are referred to as a quad.
Partial derivatives of the fragment coordinate are computed implicitly as part of operation of the following builtin functions:
For these, the derivatives help determine the mip levels of texels to be sampled, or in the case of textureSampleCompare
, sampled and compared against a reference value.
Partial derivatives of invocationspecified values are computed by the builtin functions described in § 16.7 Derivative Builtin Functions:

dpdx, dpdxCoarse, and dpdxFine compute partial derivatives along the x axis.

dpdy, dpdyCoarse, and dpdyFine compute partial derivatives along the y axis.

fwidth, fwidthCoarse, and fwidthFine compute the Manhattan metric over the associated x and y partial derivatives.
Because neighbouring invocations must collaborate to compute derivatives, these functions must only be invoked in uniform control flow in a fragment shader.
12.5. Floating Point Evaluation
WGSL follows the IEEE754 standard for floating point computation with the following exceptions:

No floating point exceptions are generated.

Signaling NaNs may not be generated. Any signaling NaN may be converted to a quiet NaN.

Implementations may assume that NaNs, infinities are not present

Note: This means some functions (e.g.
min
andmax
) may not return the expected result due to optimizations about the presence of NaNs and infinities.


Implementations may ignore the sign of a zero. That is, a zero with a positive sign may behave like a zero a with a negative sign, and vice versa.

No rounding mode is specified.

Implementations may flush denormalized value on the input and/or output of any operation listed in § 12.5.1 Floating Point Accuracy.

Other operations are required to preserve denormalized numbers.


The accuracy of operations is given in § 12.5.1 Floating Point Accuracy.
12.5.1. Floating Point Accuracy

x, when x is in T,

Otherwise:

the smallest value in T greater than x, or

the largest value in T less than x.

That is, the result may be rounded up or down: WGSL does not specify a rounding mode.
Note: Floating point types include positive and negative infinity, so the correctly rounded result may be finite or infinite.
The units in the last place, ULP, for a floating point
number x
is the minimum distance between two nonequal floating point numbers a
and b
such that a
≤ x
≤ b
(i.e. ulp(x) =
min
_{a,b}b  a
).
In the following tables, the accuracy of an operation is provided among five possibilities:

Correct result (for nonfloating point return values)

A relative error bound expressed as ULP

A function that the accuracy is inherited from. That is, the accuracy is equal to implementing the operation in terms of the derived function.

An absolute error bound
For any accuracy values specified over a range, the accuracy is undefined for results outside that range.
If an allowable return value for any operation is greater in magnitude than the largest representable finite floatingpoint value, then that operation may additionally return either the infinity with the same sign or the largest finite value with the same sign.
Expression  Accuracy  

x + y
 Correctly rounded  
x  y
 Correctly rounded  
x * y
 Correctly rounded  
x / y
 2.5 ULP for y in the range [2^{126}, 2^{126}]
 
x % y
 Derived from x  y * trunc(x/y)
 
x
 Correctly rounded  
x == y
 Correct result  
x != y
 Correct result  
x < y
 Correct result  
x <= y
 Correct result  
x > y
 Correct result  
x >= y
 Correct result 
Builtin Function  Accuracy 

abs(x)
 Correctly rounded 
acos(x)
 Inherited from atan2(sqrt(1.0  x * x), x)

asin(x)
 Inherited from atan2(x, sqrt(1.0  x * x))

atan(x)
 4096 ULP 
atan2(y, x)
 4096 ULP 
ceil(x)
 Correctly rounded 
clamp(x,low,high)
 Correctly rounded 
cos(x)
 Absolute error ≤ 2^{11} inside the range of [π, π] 
cosh(x)
 Inherited from (exp(x)  exp(x)) * 0.5

cross(x, y)
 Inherited from (x[i] * y[j]  x[j] * y[i])

degrees(x)
 Inherited from x * 57.295779513082322865

distance(x, y)
 Inherited from length(x  y)

exp(x)
 3 + 2 * x ULP 
exp2(x)
 3 + 2 * x ULP 
faceForward(x, y, z)
 Inherited from select(x, x, dot(z, y) < 0.0)

floor(x)
 Correctly rounded 
fma(x, y, z)
 Inherited from x * y + z

fract(x)
 Correctly rounded 
frexp(x)
 Correctly rounded 
inverseSqrt(x)
 2 ULP 
ldexp(x, y)
 Correctly rounded 
length(x)
 Inherited from sqrt(dot(x, x))

log(x)
 3 ULP outside the range [0.5, 2.0]. Absolute error < 2^{21} inside the range [0.5, 2.0] 
log2(x)
 3 ULP outside the range [0.5, 2.0]. Absolute error < 2^{21} inside the range [0.5, 2.0] 
max(x, y)
 Correctly rounded 
min(x, y)
 Correctly rounded 
mix(x, y, z)
 Inherited from x * (1.0  z) + y * z

modf(x)
 Correctly rounded 
normalize(x)
 Inherited from x / length(x)

pow(x, y)
 Inherited from exp2(y * log2(x))

radians(x)
 Inherited from x * 0.017453292519943295474

reflect(x, y)
 Inherited from x  2.0 * dot(x, y) * y

refract(x, y, z)
 Inherited from z * x  (z * dot(y, x) + sqrt(k)) * y ,where k = 1.0  z * z * (1.0  dot(y, x) * dot(y, x)) If k < 0.0 the result is precisely 0.0

round(x)
 Correctly rounded 
sign(x)
 Correctly rounded 
sin(x)
 Absolute error ≤ 2^{11} inside the range [π, π] 
sinh(x)
 Inherited from (exp(x)  exp(x)) * 0.5

smoothStep(low, high, x)
 Inherited from t * t * (3.0  2.0 * t) ,where t = clamp((x  low) / (high  low), 0.0, 1.0)

sqrt(x)
 Inherited from 1.0 / inverseSqrt(x)

step(edge, x)
 Correctly rounded 
tan(x)
 Inherited from sin(x) / cos(x)

tanh(x)
 Inherited from sinh(x) / cosh(x)

trunc(x)
 Correctly rounded 
Reassociation is the reordering of operations in an expression such that the answer is the same if computed exactly. For example:

(a + b) + c
reassociates toa + (b + c)

(a  b) + c
reassociates to(a + c)  b

(a * b) / c
reassociates to(a / c) * b
However, the result may not be the same when computed in floating point. The reassociated result may be inaccurate due to approximation, or may trigger an overflow or NaN when computing intermediate results.
An implementation may reassociate operations.
An implementation may fuse operations if the transformed expression is at least as accurate as the original formulation. For example, some fused multiplyadd implementations can be more accurate than performing a multiply followed by an addition.
12.5.2. Floating Point Conversion
In this section, a floating point type may be any of:

The f32 type in WGSL.

A hypothetical type corresponding to a binary format defined by the IEEE754 floating point standard.
Note: The binary16 format is referenced in this way.
Note: Recall that the f32 WGSL type corresponds to the IEEE754 binary32 format.
When converting a floating point scalar value to an integral type:

If the original value is exactly representable in the destination type, then the result is that value.

Otherwise, the original value is rounded toward zero.

If the rounded value is exactly representable in the destination type, the result is that value.

Otherwise, the result is the value in the destination type that is closest to the rounded value.

Note: In other words, floating point to integer conversion rounds toward zero, then saturates.
Note: The result in the overflow case may not yield the value with the maximum magnitude in the target type, because
that value may not be exactly representable in the original floating point type.
For example, the maximum value in u32 is 4294967295, but 4294967295.0 is not exactly representable in f32
.
For any real number x with 4294967040 ≤ x ≤ 4294967295,
the f32 value nearest to x is either larger than 429467295 or rounds down to 4294967040.
Therefore the maximum u32 value resulting from a floating point conversion is 4294967040u.
When converting a value to a floating point type:

If the original value is exactly representable in the destination type, then the result is that value.

If the original value is zero and of integral type, then the resulting value has a zero sign bit.


Otherwise, the original value is not exactly representable.

If the original value is different from but lies between two adjacent values representable in the destination type, then the result is one of those two values. WGSL does not specify whether the larger or smaller representable value is chosen, and different instances of such a conversion may choose differently.

Otherwise, if the original value lies outside the range of the destination type.

This does not occur when the original types is one of i32 or u32 and the destination type is f32.

This does not occur when the source type is a floating point type with fewer exponent and mantissa bits.

If the source type is a floating point type with more mantissa bits than the destination type, then:

The extra mantissa bits of the source value may be discarded (treated as if they are 0).

If the resulting value is the maximum normal value of the destination type, then that is the result.


Otherwise the result is the infinity value with the same sign as the source value.



Otherwise, if the original value is a NaN for the source type, then the result is a NaN in the destination type.

NOTE: An integer value may lie between two adjacent representable floating point values. In particular, the f32 type uses 23 explicit fractional bits. Additionally, when the floating point value is in the normal range (the exponent is neither extreme value), then the mantissa is the set of fractional bits together with an extra 1bit at the most significant position at bit position 23. Then, for example, integers 2^{28} and 1+2^{28} both map to the same floating point value: the difference in the least significant 1 bit is not representable by the floating point format. This kind of collision occurs for pairs of adjacent integers with a magnitude of at least 2^{25}.
Check behaviour of the f32 to f16 conversion for numbers just beyond the max normal f16 values. I’ve written what an NVIDIA GPU does. See https://github.com/google/amber/pull/918 for an executable test case.
13. Memory Model
In general, WGSL follows the Vulkan Memory Model. The remainder of this section describes how WGSL programs map to the Vulkan Memory Model.
Note: The Vulkan Memory Model is a textual version of a formal Alloy model.
13.1. Memory Operation
In WGSL, a read access is equivalent to a memory read operation in the Vulkan Memory Model. A WGSL, a write access is equivalent to a memory write operation in the Vulkan Memory Model.
A read access occurs when an invocation executes one of the following:

An evaluation of the Load Rule

Any texture builtin function except:

Any atomic builtin function except atomicStore
A write access occurs when an invocation executes one of the following:

An assignment statement

A textureStore builtin function

Any atomic builtin function except atomicLoad

atomicCompareExchangeWeak only performs a write if the
exchanged
member of the returned result istrue

Atomic readmodifywrite builtin functions perform a single memory operation that is both a read access and a write access.
Read and write accesses do not occur under any other circumstances. Read and write accesses are collectively known as memory operations in the Vulkan Memory Model.
A memory operation accesses exactly the set of locations associated with the particular memory view used in the operation. For example, a memory read that accesses a u32 from a struct containing multiple members, only reads the memory locations associated with that u32 member.
struct S { a : f32; b : u32; c : f32; } @group(0) @binding(0) var<storage> v : S; fn foo() { let x = v.b; // Does not access memory locations for v.a or v.c. }
13.2. Memory Model Reference
Each modulescope variable in WGSL forms a unique memory model reference for the lifetime of a given entry point. Each functionscope variable in WGSL forms a unique memory model reference for the lifetime of the variable.
13.3. Scoped Operations
When an invocation performs a scoped operation, it will affect one or two sets of invocations. These sets are the memory scope and the execution scope. The memory scope specifies the set of invocations that will see any updates to memory contents affected by the operation. For synchronization builtin functions, this also means that all affected memory operations program ordered before the function are visible to affected operations program ordered after the function. The execution scope specifies the set of invocations which may participate in an operation (see § 12.4 Collective Operations).
Atomic builtin functions map to atomic operations whose memory scope is:

Workgroup
if the atomic pointer is in the workgroup address space 
QueueFamily
if the atomic pointer is in the storage address space
Synchronization builtin functions map to control
barriers whose execution and memory scopes are Workgroup
.
Implicit and explicit derivatives have an implicit quad execution scope.
Note: If the Vulkan memory model is not enabled in generated shaders, Device
scope should be used instead of QueueFamily
.
13.4. Memory Semantics
All Atomic builtin functions use Relaxed
memory semantics and, thus, no address space
semantics.
workgroupBarrier uses AcquireRelease
memory semantics and WorkgroupMemory
semantics. storageBarrier uses AcquireRelease
memory semantics and UniformMemory
semantics.
Note: A combined workgroupBarrier
and storageBarrier
uses AcquireRelease
ordering semantics and both WorkgroupMemory
and UniformMemory
memory
semantics.
Note: No atomic or synchronization builtin functions use MakeAvailable
or MakeVisible
semantics.
13.5. Private vs Nonprivate
All nonatomic read accesses in the storage or workgroup address spaces are considered nonprivate and correspond to read operations with NonPrivatePointer  MakePointerVisible
memory operands with the Workgroup
scope.
All nonatomic write accesses in the storage or workgroup address spaces are considered nonprivate and correspond to write operations
with NonPrivatePointer  MakePointerAvailable
memory operands with the Workgroup
scope.
https://github.com/gpuweb/gpuweb/issues/1621
14. Keyword and Token Summary
14.1. Keyword Summary
14.1.1. Typedefining Keywords
 'array'
 'atomic'
 'bool'
 'f32'
 'i32'
 'mat2x2'
 'mat2x3'
 'mat2x4'
 'mat3x2'
 'mat3x3'
 'mat3x4'
 'mat4x2'
 'mat4x3'
 'mat4x4'
 'override'
 'ptr'
 'sampler'
 'sampler_comparison'
 'struct'
 'texture_1d'
 'texture_2d'
 'texture_2d_array'
 'texture_3d'
 'texture_cube'
 'texture_cube_array'
 'texture_multisampled_2d'
 'texture_storage_1d'
 'texture_storage_2d'
 'texture_storage_2d_array'
 'texture_storage_3d'
 'texture_depth_2d'
 'texture_depth_2d_array'
 'texture_depth_cube'
 'texture_depth_cube_array'
 'texture_depth_multisampled_2d'
 'u32'
 'vec2'
 'vec3'
 'vec4'
14.1.2. Other Keywords
 'bitcast'
 'break'
 'case'
 'continue'
 'continuing'
 'default'
 'discard'
 'else'
 'enable'
 'fallthrough'
 'false'
 'fn'
 'for'
 'function'
 'if'
 'let'
 'loop'
 'private'
 'return'
 'storage'
 'switch'
 'true'
 'type'
 'uniform'
 'var'
 'while'
 'workgroup'
14.2. Reserved Words
A reserved word is a token which is reserved for future use. A WGSL program must not contain a reserved word.
The following are reserved words:
 'asm'
 'bf16'
 'const'
 'demote'
 'demote_to_helper'
 'do'
 'enum'
 'f16'
 'f64'
 'handle'
 'i8'
 'i16'
 'i64'
 'mat'
 'premerge'
 'regardless'
 'std'
 'typedef'
 'u8'
 'u16'
 'u64'
 'unless'
 'using'
 'vec'
 'void'
 'wgsl'
14.3. Syntactic Tokens
A syntactic token is a sequence of special characters, used:

to spell an expression operator, or

as punctuation: to group, sequence, or separate other grammar elements.
 '&'
 '&&'
 '>'
 '@'
 '/'
 '!'
 '['
 ']'
 '{'
 '}'
 ':'
 ','
 '='
 '=='
 '!='
 '>'
 '>='
 '>>'
 '<'
 '<='
 '<<'
 '%'
 ''
 ''
 '.'
 '+'
 '++'
 ''
 ''
 '('
 ')'
 ';'
 '*'
 '~'
 '_'
 '^'
 '+='
 '='
 '*='
 '/='
 '%='
 '&='
 '='
 '^='
 '>>='
 '<<='
15. Builtin Values
The following table lists the available builtin input values and builtin output values.
See § 9.3.1.1 Builtin Inputs and Outputs for how to declare a builtin value.
Name  Stage  Input or Output  Type  Description 

vertex_index
 vertex  input  u32 
Index of the current vertex within the current APIlevel draw command,
independent of draw instancing.
For a nonindexed draw, the first vertex has an index equal to the For an indexed draw, the index is equal to the index buffer entry for
vertex, plus the 
instance_index
 vertex  input  u32 
Instance index of the current vertex within the current APIlevel draw command.
The first instance has an index equal to the 
position
 vertex  output  vec4<f32>  Output position of the current vertex, using homogeneous coordinates. After homogeneous normalization (where each of the x, y, and z components are divided by the w component), the position is in the WebGPU normalized device coordinate space. See WebGPU § Coordinate Systems. 
position
 fragment  input  vec4<f32>  Framebuffer position of the current fragment, using normalized homogeneous coordinates. (The x, y, and z components have already been scaled such that w is now 1.) See WebGPU § Coordinate Systems. 
front_facing
 fragment  input  bool  True when the current fragment is on a frontfacing primitive. False otherwise. See WebGPU § Frontfacing. 
frag_depth
 fragment  output  f32  Updated depth of the fragment, in the viewport depth range. See WebGPU § Coordinate Systems. 
local_invocation_id
 compute  input  vec3<u32>  The current invocation’s local invocation ID, i.e. its position in the workgroup grid. 
local_invocation_index
 compute  input  u32  The current invocation’s local invocation index, a linearized index of the invocation’s position within the workgroup grid. 
global_invocation_id
 compute  input  vec3<u32>  The current invocation’s global invocation ID, i.e. its position in the compute shader grid. 
workgroup_id
 compute  input  vec3<u32>  The current invocation’s workgroup ID, i.e. the position of the workgroup in the workgroup grid. 
num_workgroups
 compute  input  vec3<u32>  The dispatch size, vec<u32>(group_count_x, group_count_y, group_count_z) , of the compute shader dispatched by the API.

sample_index
 fragment  input  u32  Sample index for the current fragment.
The value is least 0 and at most sampleCount 1, where sampleCount is the number of MSAA samples specified for the GPU render pipeline. See WebGPU § GPURenderPipeline. 
sample_mask
 fragment  input  u32  Sample coverage mask for the current fragment.
It contains a bitmask indicating which samples in this fragment are covered
by the primitive being rendered. See WebGPU § Sample Masking. 
sample_mask
 fragment  output  u32  Sample coverage mask control for the current fragment.
The last value written to this variable becomes the shaderoutput mask.
Zero bits in the written value will cause corresponding samples in
the color attachments to be discarded. See WebGPU § Sample Masking. 
struct VertexOutput { @builtin(position) my_pos: vec4<f32>; } @stage(vertex) fn vs_main( @builtin(vertex_index) my_index: u32, @builtin(instance_index) my_inst_index: u32, ) > VertexOutput {} struct FragmentOutput { @builtin(frag_depth) depth: f32; @builtin(sample_mask) mask_out: u32; } @stage(fragment) fn fs_main( @builtin(front_facing) is_front: bool, @builtin(position) coord: vec4<f32>, @builtin(sample_index) my_sample_index: u32, @builtin(sample_mask) mask_in: u32, ) > FragmentOutput {} @stage(compute) fn cs_main( @builtin(local_invocation_id) local_id: vec3<u32>, @builtin(local_invocation_index) local_index: u32, @builtin(global_invocation_id) global_id: vec3<u32>, ) {}
16. Builtin Functions
Certain functions are predeclared, provided by the implementation, and therefore always available for use in a WGSL program. These are called builtin functions.
A builtin function is a family of functions, all with the same name, but distinguished by the number, order, and types of their formal parameters. Each of these distinct function variations is an overload.
Note: Each userdefined function only has one overload.
Each overload is described below via:

Type parameterizations, if any.

The builtin function name, a parenthesized list of formal parameters, and optionally a return type.

The behaviour of this overload of the function.
Since a builtin function is always in scope, it is an error to attempt to redefine one or to use the name of a builtin function as an identifier for any other kind of declaration.
When calling a builtin function, all arguments to the function are evaluated before function evaluation begins. See § 8.2 Function Calls.
16.1. Logical Builtin Functions
Parameterization  Overload  Description 

all (e: vecN<bool>) > bool
 Returns true if each component of e is true.  
e: bool  all( e) : bool
 Returns e. 
any (e: vecN<bool>) > bool
 Returns true if any component of e is true.  
e: bool  any( e) : bool
 Returns e. 
T is scalar or vector  select (f: T, t: T, cond: bool) > T
 Returns t when cond is true, and f otherwise. 
T is scalar  select (f: vecN<T>, t: vecN<T>, cond: vecN<bool>) > vecN<T>
 Componentwise selection. Result component i is evaluated
as select( f[ i], t[ i], cond[ i]) .

16.2. Array Builtin Functions
Parameterization  Overload  Description 

arrayLength (e: ptr<storage,array<T>> ) > u32
 Returns the number of elements in the runtimesized array. 
16.3. Float Builtin Functions
Parameterization  Overload  Description 

T is f32 or vecN<f32>  abs( e: T ) > T
 Returns the absolute value of e (e.g. e with a positive sign bit). Componentwise when T is a vector. 
T is f32 or vecN<f32>  acos( e: T ) > T
 Returns the arc cosine of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  asin( e: T ) > T
 Returns the arc sine of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  atan( e: T ) > T
 Returns the arc tangent of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  atan2( e1: T , e2: T ) > T
 Returns the arc tangent of e1 over e2. Componentwise when T is a vector. 
T is f32 or vecN<f32>  ceil( e: T ) > T
 Returns the ceiling of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  clamp( e: T , low: T , high: T) > T
 Returns min(max( e, low), high) . Componentwise when T is a vector.

T is f32 or vecN<f32>  cos( e: T ) > T
 Returns the cosine of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  cosh( e: T ) > T
 Returns the hyperbolic cosine of e. Componentwise when T is a vector 
T is f32  cross( e1: vec3<T> , e2: vec3<T>) > vec3<T>
 Returns the cross product of e1 and e2. 
T is f32 or vecN<f32>  degrees( e1: T ) > T
 Converts radians to degrees, approximating e1 × 180 ÷ π. Componentwise when T is a vector 
T is f32 or vecN<f32>  distance( e1: T , e2: T ) > f32
 Returns the distance between e1 and e2 (e.g. length( e1 e2) ).

T is f32 or vecN<f32>  exp( e1: T ) > T
 Returns the natural exponentiation of e1 (e.g. e ^{e1}). Componentwise when T is a vector.

T is f32 or vecN<f32>  exp2( e: T ) > T
 Returns 2 raised to the power e (e.g. 2 ^{e}). Componentwise when T is a vector.

T is vecN<f32>  faceForward( e1: T , e2: T , e3: T ) > T
 Returns e1 if dot( e2, e3) is negative, and  e1 otherwise.

T is f32 or vecN<f32>  floor( e: T ) > T
 Returns the floor of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  fma( e1: T , e2: T , e3: T ) > T
 Returns e1 * e2 + e3. Componentwise when T is a vector.

T is f32 or vecN<f32>  fract( e: T ) > T
 Returns the fractional part of e, computed as e  floor( e) .Componentwise when T is a vector. 
T is f32  frexp( e: T) > __frexp_result 
Splits e into a significand and exponent of the form significand * 2 ^{exponent}.
Returns the __frexp_result builtin structure, defined as if as follows:
The magnitude of the significand is in the range of [0.5, 1.0) or 0. Note: A value cannot be explicitly declared with the type 
T is vecN<f32>  frexp( e: T) > __frexp_result_vec N 
Splits the components of e into a significand and exponent of the form significand * 2 ^{exponent}.
Returns the __frexp_result_vec N builtin structure, defined as if as follows:
The magnitude of each component of the significand is in the range of [0.5, 1.0) or 0. Note: A value cannot be explicitly declared with the type 
T is f32 or vecN<f32>  inverseSqrt( e: T ) > T
 Returns the reciprocal of sqrt( e) . Componentwise when T is a vector.

T is f32 or vecN<f32> I is i32 or vecN<i32>, where I is a scalar if T is a scalar, or a vector when T is a vector  ldexp( e1: T , e2: I ) > T
 Returns e1 * 2 ^{e2}. Componentwise when T is a vector.

T is f32 or vecN<f32>  length( e: T ) > f32
 Returns the length of e (e.g. abs( e) if T is a scalar, or sqrt( e[0] ^{2} + e[1] ^{2} + ...) if T is a vector).

T is f32 or vecN<f32>  log( e: T ) > T
 Returns the natural logarithm of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  log2( e: T ) > T
 Returns the base2 logarithm of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  max( e1: T , e2: T ) > T
 Returns e2 if e1 is less than e2, and e1 otherwise. If one operand is a NaN, the other is returned. If both operands are NaNs, a NaN is returned. Componentwise when T is a vector. 
T is f32 or vecN<f32>  min( e1: T , e2: T ) > T
 Returns e2 if e2 is less than e1, and e1 otherwise. If one operand is a NaN, the other is returned. If both operands are NaNs, a NaN is returned. Componentwise when T is a vector. 
T is f32 or vecN<f32>  mix( e1: T , e2: T , e3: T) > T
 Returns the linear blend of e1 and e2 (e.g. e1*(1 e3)+ e2* e3). Componentwise when T is a vector. 
T is vecN<f32>  mix( e1: T , e2: T , e3: f32 ) > T
 Returns the componentwise linear blend of e1 and e2,
using scalar blending factor e3 for each component. Same as mix( e1, e2, T( e3)) .

T is f32  modf( e: T) > __modf_result 
Splits e into fractional and whole number parts.
Returns the __modf_result builtin structure, defined as if as follows:
Note: A value cannot be explicitly declared with the type 
T is vecN<f32>  modf( e: T) > __modf_result_vec N 
Splits the components of e into fractional and whole number parts.
Returns the __modf_result_vec N builtin structure, defined as if as follows:
Note: A value cannot be explicitly declared with the type 
T is f32  normalize( e: vecN<T> ) > vecN<T>
 Returns a unit vector in the same direction as e. 
T is f32 or vecN<f32>  pow( e1: T , e2: T ) > T
 Returns e1 raised to the power e2. Componentwise when T is a vector. 
T is f32 or vecN<f32>  quantizeToF16( e: T ) > T

Quantizes a 32bit floating point value e as if e were converted to a IEEE 754 binary16 value,
and then converted back to a IEEE 754 binary32 value. See § 12.5.2 Floating Point Conversion. Componentwise when T is a vector. Note: The vec2<f32> case is the same as 
T is f32 or vecN<f32>  radians( e1: T ) > T
 Converts degrees to radians, approximating e1 × π ÷ 180. Componentwise when T is a vector 
T is vecN<f32>  reflect( e1: T , e2: T ) > T
 For the incident vector e1 and surface orientation e2, returns the reflection direction e12*dot( e2, e1)* e2.

T is vecN<f32> I is f32  refract( e1: T , e2: T , e3: I ) > T
 For the incident vector e1 and surface normal e2, and the ratio of indices of refraction e3,
let k = 1.0  e3* e3* (1.0  dot( e2, e1) * dot( e2, e1)) . If k < 0.0 , returns the
refraction vector 0.0, otherwise return the refraction vector e3* e1 ( e3* dot( e2, e1) + sqrt(k)) * e2.

T is f32 or vecN<f32>  round( e: T ) > T
 Result is the integer k nearest to e, as a floating point value. When e lies halfway between integers k and k+1, the result is k when k is even, and k+1 when k is odd. Componentwise when T is a vector. 
T is f32 or vecN<f32>  sign( e: T ) > T
 Returns the sign of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  sin( e: T ) > T
 Returns the sine of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  sinh( e: T ) > T
 Returns the hyperbolic sine of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  smoothStep( low: T , high: T , x: T ) > T

Returns the smooth Hermite interpolation between 0 and 1. Componentwise when T is a vector.
For scalar T, the result is t * t * (3.0  2.0 * t), where t = clamp((x  low) / (high  low), 0.0, 1.0). 
T is f32 or vecN<f32>  sqrt( e: T ) > T
 Returns the square root of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  step( edge: T , x: T ) > T
 Returns 1.0 if edge ≤ x, and 0.0 otherwise. Componentwise when T is a vector. 
T is f32 or vecN<f32>  tan( e: T ) > T
 Returns the tangent of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  tanh( e: T ) > T
 Returns the hyperbolic tangent of e. Componentwise when T is a vector. 
T is f32 or vecN<f32>  trunc( e: T ) > T
 Returns the nearest whole number whose absolute value is less than or equal to e. Componentwise when T is a vector. 
16.4. Integer Builtin Functions
Parameterization  Overload  Description 

T is i32 or vecN<i32>  abs (e: T ) > T
 The absolute value of e. Componentwise when T is a vector. If e evaluates to the largest negative value, then the result is e. 
T is u32 or vecN<u32>  abs (e: T ) > T
 Result is e. This is provided for symmetry with abs for signed integers. Componentwise when T is a vector.

T is u32 or vecN<u32>  clamp( e: T , low: T, high: T) > T
 Returns min(max( e, low), high) . Componentwise when T is a vector.

T is i32 or vecN<i32>  clamp( e: T , low: T, high: T) > T
 Returns min(max( e, low), high) . Componentwise when T is a vector.

T is i32, u32, vecN<i32>, or vecN<u32>  countLeadingZeros( e: T ) > T
 The number of consecutive 0 bits starting from the most significant bit
of e, when T is a scalar type. Componentwise when T is a vector. Also known as "clz" in some languages. 
T is i32, u32, vecN<i32>, or vecN<u32>  countOneBits( e: T ) > T
 The number of 1 bits in the representation of e. Also known as "population count". Componentwise when T is a vector. 
T is i32, u32, vecN<i32>, or vecN<u32>  countTrailingZeros( e: T ) > T
 The number of consecutive 0 bits starting from the least significant bit
of e, when T is a scalar type. Componentwise when T is a vector. Also known as "ctz" in some languages. 
T is i32 or vecN<i32>  firstLeadingBit( e: T ) > T

For scalar T, the result is:
Note: Since signed integers use twoscomplement representation, the sign bit appears in the most significant bit position. Componentwise when T is a vector. 
T is u32 or vecN<u32>  firstLeadingBit( e: T ) > T

For scalar T, the result is:

T is i32, u32, vecN<i32>, or vecN<u32>  firstTrailingBit( e: T ) > T

For scalar T, the result is:

T is i32 or vecN<i32>  extractBits( e : T, offset : u32, count : u32) > T

Reads bits from an integer, with sign extension.
When T is a scalar type, then:

T is u32 or vecN<u32>  extractBits( e : T, offset : u32, count : u32) > T

Reads bits from an integer, without sign extension.
When T is a scalar type, then:

T is i32, u32, vecN<i32>, or vecN<u32>  insertBits( e : T, newbits : T, offset : u32, count : u32) > T

Sets bits in an integer.
When T is a scalar type, then:

T is u32 or vecN<u32>  max( e1: T , e2: T) > T
 Returns e2 if e1 is less than e2, and e1 otherwise. Componentwise when T is a vector. 
T is i32 or vecN<i32>  max( e1: T , e2: T) > T
 Returns e2 if e1 is less than e2, and e1 otherwise. Componentwise when T is a vector. 
T is u32 or vecN<u32>  min( e1: T , e2: T) > T
 Returns e1 if e1 is less than e2, and e2 otherwise. Componentwise when T is a vector. 
T is i32 or vecN<i32>  min( e1: T , e2: T) > T
 Returns e1 if e1 is less than e2, and e2 otherwise. Componentwise when T is a vector. 
T is i32, u32, vecN<i32>, or vecN<u32>  reverseBits( e: T ) > T
 Reverses the bits in e: The bit at position k of the result equals the
bit at position 31k of e. Componentwise when T is a vector. 
16.5. Matrix Builtin Functions
Parameterization  Overload  Description 

T is f32  determinant( e: matNxN<T> ) > T
 Returns the determinant of e. 
T is f32  transpose( e: matMxN<T> ) > matNxM<T>
 Returns the transpose of e. 
16.6. Vector Builtin Functions
Parameterization  Overload  Description 

T is f32  dot( e1: vecN< T>, e2: vecN< T>) > T
 Returns the dot product of e1 and e2. 
T is i32  dot( e1: vecN< T>, e2: vecN< T>) > T
 Returns the dot product of e1 and e2. 
T is u32  dot( e1: vecN< T>, e2: vecN< T>) > T
 Returns the dot product of e1 and e2. 
16.7. Derivative Builtin Functions
See § 12.4.2 Derivatives.
These functions:

Must only be used in a fragment shader stage.

Must only be invoked in uniform control flow.
Parameterization  Overload  Description 

T is f32 or vecN<f32>  dpdx( e: T) > T
 Partial derivative of e with respect to window x coordinates.
The result is the same as either dpdxFine( e) or dpdxCoarse( e) .

dpdxCoarse( e: T) > T
 Returns the partial derivative of e with respect to window x coordinates using local differences.
This may result in fewer unique positions that dpdxFine( e) .
 
dpdxFine( e: T) > T
 Returns the partial derivative of e with respect to window x coordinates.  
dpdy( e: T) > T
 Partial derivative of e with respect to window y coordinates.
The result is the same as either dpdyFine( e) or dpdyCoarse( e) .
 
dpdyCoarse( e: T) > T
 Returns the partial derivative of e with respect to window y coordinates using local differences.
This may result in fewer unique positions that dpdyFine( e) .
 
dpdyFine( e: T) > T
 Returns the partial derivative of e with respect to window y coordinates.  
fwidth( e: T) > T
 Returns abs(dpdx( e)) + abs(dpdy( e)) .
 
fwidthCoarse( e: T) > T
 Returns abs(dpdxCoarse( e)) + abs(dpdyCoarse( e)) .
 
fwidthFine( e: T) > T
 Returns abs(dpdxFine( e)) + abs(dpdyFine( e)) .

16.8. Texture Builtin Functions
In this section, texture types are shown with the following parameters:

T, a sampled type.

F, a texel format.

A, an access mode.
Parameter values must be valid for the respective texture types.
16.8.1. textureDimensions
Returns the dimensions of a texture, or texture’s mip level in texels.
textureDimensions ( t :texture_1d < T > ) >i32 textureDimensions ( t :texture_1d < T > , level :i32 ) >i32 textureDimensions ( t :texture_2d < T > ) >vec2 < i32 > textureDimensions ( t :texture_2d < T > , level :i32 ) >vec2 < i32 > textureDimensions ( t :texture_2d_array < T > ) >vec2 < i32 > textureDimensions ( t :texture_2d_array < T > , level :i32 ) >vec2 < i32 > textureDimensions ( t :texture_3d < T > ) >vec3 < i32 > textureDimensions ( t :texture_3d < T > , level :i32 ) >vec3 < i32 > textureDimensions ( t :texture_cube < T > ) >vec2 < i32 > textureDimensions ( t :texture_cube < T > , level :i32 ) >vec2 < i32 > textureDimensions ( t :texture_cube_array < T > ) >vec2 < i32 > textureDimensions ( t :texture_cube_array < T > , level :i32 ) >vec2 < i32 > textureDimensions ( t :texture_multisampled_2d < T > ) >vec2 < i32 > textureDimensions ( t :texture_depth_2d ) >vec2 < i32 > textureDimensions ( t :texture_depth_2d , level :i32 ) >vec2 < i32 > textureDimensions ( t :texture_depth_2d_array ) >vec2 < i32 > textureDimensions ( t :texture_depth_2d_array , level :i32 ) >vec2 < i32 > textureDimensions ( t :texture_depth_cube ) >vec2 < i32 > textureDimensions ( t :texture_depth_cube , level :i32 ) >vec2 < i32 > textureDimensions ( t :texture_depth_cube_array ) >vec2 < i32 > textureDimensions ( t :texture_depth_cube_array , level :i32 ) >vec2 < i32 > textureDimensions ( t :texture_depth_multisampled_2d ) >vec2 < i32 > textureDimensions ( t :texture_storage_1d < F , A > ) >i32 textureDimensions ( t :texture_storage_2d < F , A > ) >vec2 < i32 > textureDimensions ( t :texture_storage_2d_array < F , A > ) >vec2 < i32 > textureDimensions ( t :texture_storage_3d < F , A > ) >vec3 < i32 > textureDimensions ( t :texture_external ) >vec2 < i32 >
Parameters:
t
 The sampled, multisampled, depth, storage, or external texture. 
level
 The mip level, with level 0 containing a full size version of the texture. If omitted, the dimensions of level 0 are returned. 
Returns:
The dimensions of the texture in texels.
For textures based on cubes, the results are the dimensions of each face of the cube. Cube faces are square, so the x and y components of the result are equal.
If level
is outside the range [0, textureNumLevels(t))
then any valid value
for the return type may be returned.
16.8.2. textureGather
A texture gather operation reads from a 2D, 2D array, cube, or cube array texture, computing a fourcomponent vector as follows:

Find the four texels that would be used in a sampling operation with linear filtering, from mip level 0:

Use the specified coordinate, array index (when present), and offset (when present).

The texels are adjacent, forming a square, when considering their texture space coordinates (u,v).

Selected texels at the texture edge, cube face edge, or cube corners are handled as in ordinary texture sampling.


For each texel, read one channel and convert it into a scalar value.

For nondepth textures, a zerobased
component
parameter specifies the channel to use.
If the texture format supports the specified channel, i.e. has more than
component
channels:
Yield scalar value
v[component]
when the texel value isv
.


Otherwise:

Yield 0.0 when
component
is 1 or 2. 
Yield 1.0 when
component
is 3 (the alpha channel).



For depth textures, yield the texel value. (Depth textures only have one channel.)


Yield the fourcomponent vector, arranging scalars produced by the previous step into components according to the relative coordinates of the texels, as follows:

Result component Relative texel coordinate x (u_{min},v_{max}) y (u_{max},v_{max}) z (u_{max},v_{min}) w (u_{min},v_{min})

TODO: The four texels are the "sample footprint" that should be described by the WebGPU spec. https://github.com/gpuweb/gpuweb/issues/2343
textureGather ( component :i32 , t :texture_2d < T > , s :sampler , coords :vec2 < f32 > ) >vec4 < T > textureGather ( component :i32 , t :texture_2d < T > , s :sampler , coords :vec2 < f32 > , offset :vec2 < i32 > ) >vec4 < T > textureGather ( component :i32 , t :texture_2d_array < T > , s :sampler , coords :vec2 < f32 > , array_index :i32 ) >vec4 < T > textureGather ( component :i32 , t :texture_2d_array < T > , s :sampler , coords :vec2 < f32 > , array_index :i32 , offset :vec2 < i32 > ) >vec4 < T > textureGather ( component :i32 , t :texture_cube < T > , s :sampler , coords :vec3 < f32 > ) >vec4 < T > textureGather ( component :i32 , t :texture_cube_array < T > , s :sampler , coords :vec3 < f32 > , array_index :i32 ) >vec4 < T > textureGather ( t :texture_depth_2d , s :sampler , coords :vec2 < f32 > ) >vec4 < f32 > textureGather ( t :texture_depth_2d , s :sampler , coords :vec2 < f32 > , offset :vec2 < i32 > ) >vec4 < f32 > textureGather ( t :texture_depth_2d_array , s :sampler , coords :vec2 < f32 > , array_index :i32 ) >vec4 < f32 > textureGather ( t :texture_depth_2d_array , s :sampler , coords :vec2 < f32 > , array_index :i32 , offset :vec2 < i32 > ) >vec4 < f32 > textureGather ( t :texture_depth_cube , s :sampler , coords :vec3 < f32 > ) >vec4 < f32 > textureGather ( t :texture_depth_cube_array , s :sampler , coords :vec3 < f32 > , array_index :i32 ) >vec4 < f32 >
Parameters:
component

Only applies to nondepth textures. The index of the channel to read from the selected texels. When provided, the component expression must be either:

t
 The sampled or depth texture to read from. 
s
 The sampler type. 
coords
 The texture coordinates. 
array_index
 The 0based texture array index. 
offset

The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes. The offset expression must be either:
offset component must be at least 8 and at most 7 . Values outside
of this range will result in a shadercreation error.

Returns:
A four component vector with components extracted from the specified channel from the selected texels, as described above.
@group(0) @binding(0) var t: texture_2d<f32>; @group(0) @binding(1) var dt: texture_depth_2d; @group(0) @binding(2) var s: sampler; fn gather_x_components(c: vec2<f32>) > vec4<f32> { return textureGather(0,t,s,c); } fn gather_y_components(c: vec2<f32>) > vec4<f32> { return textureGather(1,t,s,c); } fn gather_z_components(c: vec2<f32>) > vec4<f32> { return textureGather(2,t,s,c); } fn gather_depth_components(c: vec2<f32>) > vec4<f32> { return textureGather(dt,s,c); }
16.8.3. textureGatherCompare
A texture gather compare operation performs a depth comparison on four texels in a depth texture and collects the results into a single vector, as follows:

Find the four texels that would be used in a depth sampling operation with linear filtering, from mip level 0:

Use the specified coordinate, array index (when present), and offset (when present).

The texels are adjacent, forming a square, when considering their texture space coordinates (u,v).

Selected texels at the texture edge, cube face edge, or cube corners are handled as in ordinary texture sampling.


For each texel, perform a comparison against the depth reference value, yielding a 0.0 or 1.0 value, as controlled by the comparison sampler parameters.

Yield the fourcomponent vector where the components are the comparison results with the texels with relative texel coordinates as follows:

Result component Relative texel coordinate x (u_{min},v_{max}) y (u_{max},v_{max}) z (u_{max},v_{min}) w (u_{min},v_{min})

Parameters:textureGatherCompare ( t :texture_depth_2d , s :sampler_comparison , coords :vec2 < f32 > , depth_ref :f32 ) >vec4 < f32 > textureGatherCompare ( t :texture_depth_2d , s :sampler_comparison , coords :vec2 < f32 > , depth_ref :f32 , offset :vec2 < i32 > ) >vec4 < f32 > textureGatherCompare ( t :texture_depth_2d_array , s :sampler_comparison , coords :vec2 < f32 > , array_index :i32 , depth_ref :f32 ) >vec4 < f32 > textureGatherCompare ( t :texture_depth_2d_array , s :sampler_comparison , coords :vec2 < f32 > , array_index :i32 , depth_ref :f32 , offset :vec2 < i32 > ) >vec4 < f32 > textureGatherCompare ( t :texture_depth_cube , s :sampler_comparison , coords :vec3 < f32 > , depth_ref :f32 ) >vec4 < f32 > textureGatherCompare ( t :texture_depth_cube_array , s :sampler_comparison , coords :vec3 < f32 > , array_index :i32 , depth_ref :f32 ) >vec4 < f32 >
t
 The depth texture to read from. 
s
 The sampler comparison. 
coords
 The texture coordinates. 
array_index
 The 0based texture array index. 
depth_ref
 The reference value to compare the sampled depth value against. 
offset

The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes. The offset expression must be either:
offset component must be at least 8 and at most 7 . Values outside
of this range will result in a shadercreation error.

Returns:
A four component vector with comparison result for the selected texels, as described above.
@group(0) @binding(0) var dt: texture_depth_2d; @group(0) @binding(1) var s: sampler; fn gather_depth_compare(c: vec2<f32>, depth_ref: f32) > vec4<f32> { return textureGatherCompare(dt,s,c,depth_ref); }
16.8.4. textureLoad
Reads a single texel from a texture without sampling or filtering.
textureLoad ( t :texture_1d < T > , coords :i32 , level :i32 ) >vec4 < T > textureLoad ( t :texture_2d < T > , coords :vec2 < i32 > , level :i32 ) >vec4 < T > textureLoad ( t :texture_2d_array < T > , coords :vec2 < i32 > , array_index :i32 , level :i32 ) >vec4 < T > textureLoad ( t :texture_3d < T > , coords :vec3 < i32 > , level :i32 ) >vec4 < T > textureLoad ( t :texture_multisampled_2d < T > , coords :vec2 < i32 > , sample_index :i32 ) >vec4 < T > textureLoad ( t :texture_depth_2d , coords :vec2 < i32 > , level :i32 ) >f32 textureLoad ( t :texture_depth_2d_array , coords :vec2 < i32 > , array_index :i32 , level :i32 ) >f32 textureLoad ( t :texture_depth_multisampled_2d , coords :vec2 < i32 > , sample_index :i32 ) >f32 textureLoad ( t :texture_external , coords :vec2 < i32 > ) >vec4 < f32 >
Parameters:
t
 The sampled, multisampled, depth, or external texture. 
coords
 The 0based texel coordinate. 
array_index
 The 0based texture array index. 
level
 The mip level, with level 0 containing a full size version of the texture. 
sample_index
 The 0based sample index of the multisampled texture. 
Returns:
The unfiltered texel data.
An out of bounds access occurs if:

any element of
coords
is outside the range[0, textureDimensions(t, level))
for the corresponding element, or 
array_index
is outside the range[0, textureNumLayers(t))
, or 
level
is outside the range[0, textureNumLevels(t))
If an out of bounds access occurs, the builtin function returns one of:

The data for some texel within bounds of the texture

A vector (0,0,0,0) or (0,0,0,1) of the appropriate type for nondepth textures

0.0 for depth textures
16.8.5. textureNumLayers
Returns the number of layers (elements) of an array texture.
textureNumLayers ( t :texture_2d_array < T > ) >i32 textureNumLayers ( t :texture_cube_array < T > ) >i32 textureNumLayers ( t :texture_depth_2d_array ) >i32 textureNumLayers ( t :texture_depth_cube_array ) >i32 textureNumLayers ( t :texture_storage_2d_array < F , A > ) >i32
Parameters:
t
 The sampled, multisampled, depth or storage array texture. 
Returns:
If the number of layers (elements) of the array texture.
16.8.6. textureNumLevels
Returns the number of mip levels of a texture.
textureNumLevels ( t :texture_1d < T > ) >i32 textureNumLevels ( t :texture_2d < T > ) >i32 textureNumLevels ( t :texture_2d_array < T > ) >i32 textureNumLevels ( t :texture_3d < T > ) >i32 textureNumLevels ( t :texture_cube < T > ) >i32 textureNumLevels ( t :texture_cube_array < T > ) >i32 textureNumLevels ( t :texture_depth_2d ) >i32 textureNumLevels ( t :texture_depth_2d_array ) >i32 textureNumLevels ( t :texture_depth_cube ) >i32 textureNumLevels ( t :texture_depth_cube_array ) >i32
Parameters:
t
 The sampled or depth texture. 
Returns:
If the number of mip levels for the texture.
16.8.7. textureNumSamples
Returns the number samples per texel in a multisampled texture.
textureNumSamples ( t :texture_multisampled_2d < T > ) >i32 textureNumSamples ( t :texture_depth_multisampled_2d ) >i32
Parameters:
t
 The multisampled texture. 
Returns:
If the number of samples per texel in the multisampled texture.
16.8.8. textureSample
Samples a texture.
Must only be used in a fragment shader stage. Must only be invoked in uniform control flow.
textureSample ( t :texture_1d < f32 > , s :sampler , coords :f32 ) >vec4 < f32 > textureSample ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > ) >vec4 < f32 > textureSample ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , offset :vec2 < i32 > ) >vec4 < f32 > textureSample ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 ) >vec4 < f32 > textureSample ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , offset :vec2 < i32 > ) >vec4 < f32 > textureSample ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > ) >vec4 < f32 > textureSample ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , offset :vec3 < i32 > ) >vec4 < f32 > textureSample ( t :texture_cube < f32 > , s :sampler , coords :vec3 < f32 > ) >vec4 < f32 > textureSample ( t :texture_cube_array < f32 > , s :sampler , coords :vec3 < f32 > , array_index :i32 ) >vec4 < f32 > textureSample ( t :texture_depth_2d , s :sampler , coords :vec2 < f32 > ) >f32 textureSample ( t :texture_depth_2d , s :sampler , coords :vec2 < f32 > , offset :vec2 < i32 > ) >f32 textureSample ( t :texture_depth_2d_array , s :sampler , coords :vec2 < f32 > , array_index :i32 ) >f32 textureSample ( t :texture_depth_2d_array , s :sampler , coords :vec2 < f32 > , array_index :i32 , offset :vec2 < i32 > ) >f32 textureSample ( t :texture_depth_cube , s :sampler , coords :vec3 < f32 > ) >f32 textureSample ( t :texture_depth_cube_array , s :sampler , coords :vec3 < f32 > , array_index :i32 ) >f32
Parameters:
t
 The sampled, depth, or external texture to sample. 
s
 The sampler type. 
coords
 The texture coordinates used for sampling. 
array_index
 The 0based texture array index to sample. 
offset

The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes. The offset expression must be either:
offset component must be at least 8 and at most 7 . Values outside
of this range will result in a shadercreation error.

Returns:
The sampled value.
16.8.9. textureSampleBias
Samples a texture with a bias to the mip level.
Must only be used in a fragment shader stage. Must only be invoked in uniform control flow.
textureSampleBias ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , bias :f32 ) >vec4 < f32 > textureSampleBias ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , bias :f32 , offset :vec2 < i32 > ) >vec4 < f32 > textureSampleBias ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , bias :f32 ) >vec4 < f32 > textureSampleBias ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , bias :f32 , offset :vec2 < i32 > ) >vec4 < f32 > textureSampleBias ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , bias :f32 ) >vec4 < f32 > textureSampleBias ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , bias :f32 , offset :vec3 < i32 > ) >vec4 < f32 > textureSampleBias ( t :texture_cube < f32 > , s :sampler , coords :vec3 < f32 > , bias :f32 ) >vec4 < f32 > textureSampleBias ( t :texture_cube_array < f32 > , s :sampler , coords :vec3 < f32 > , array_index :i32 , bias :f32 ) >vec4 < f32 >
Parameters:
t
 The texture to sample. 
s
 The sampler type. 
coords
 The texture coordinates used for sampling. 
array_index
 The 0based texture array index to sample. 
bias
 The bias to apply to the mip level before sampling. bias must be between 16.0 and 15.99 .

offset

The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes. The offset expression must be either:
offset component must be at least 8 and at most 7 . Values outside
of this range will result in a shadercreation error.

Returns:
The sampled value.
16.8.10. textureSampleCompare
Samples a depth texture and compares the sampled depth values against a reference value.
Must only be used in a fragment shader stage. Must only be invoked in uniform control flow.
textureSampleCompare ( t :texture_depth_2d , s :sampler_comparison , coords :vec2 < f32 > , depth_ref :f32 ) >f32 textureSampleCompare ( t :texture_depth_2d , s :sampler_comparison , coords :vec2 < f32 > , depth_ref :f32 , offset :vec2 < i32 > ) >f32 textureSampleCompare ( t :texture_depth_2d_array , s :sampler_comparison , coords :vec2 < f32 > , array_index :i32 , depth_ref :f32 ) >f32 textureSampleCompare ( t :texture_depth_2d_array , s :sampler_comparison , coords :vec2 < f32 > , array_index :i32 , depth_ref :f32 , offset :vec2 < i32 > ) >f32 textureSampleCompare ( t :texture_depth_cube , s :sampler_comparison , coords :vec3 < f32 > , depth_ref :f32 ) >f32 textureSampleCompare ( t :texture_depth_cube_array , s :sampler_comparison , coords :vec3 < f32 > , array_index :i32 , depth_ref :f32 ) >f32
Parameters:
t
 The depth texture to sample. 
s
 The sampler comparision type. 
coords
 The texture coordinates used for sampling. 
array_index
 The 0based texture array index to sample. 
depth_ref
 The reference value to compare the sampled depth value against. 
offset

The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes. The offset expression must be either:
offset component must be at least 8 and at most 7 . Values outside
of this range will result in a shadercreation error.

Returns:
A value in the range [0.0..1.0]
.
Each sampled texel is compared against the reference value using the comparision
operator defined by the sampler_comparison
, resulting in either a 0
or 1
value for each texel.
If the sampler uses bilinear filtering then the returned value is the filtered average of these values, otherwise the comparision result of a single texel is returned.
16.8.11. textureSampleCompareLevel
Samples a depth texture and compares the sampled depth values against a reference value.
textureSampleCompareLevel ( t :texture_depth_2d , s :sampler_comparison , coords :vec2 < f32 > , depth_ref :f32 ) >f32 textureSampleCompareLevel ( t :texture_depth_2d , s :sampler_comparison , coords :vec2 < f32 > , depth_ref :f32 , offset :vec2 < i32 > ) >f32 textureSampleCompareLevel ( t :texture_depth_2d_array , s :sampler_comparison , coords :vec2 < f32 > , array_index :i32 , depth_ref :f32 ) >f32 textureSampleCompareLevel ( t :texture_depth_2d_array , s :sampler_comparison , coords :vec2 < f32 > , array_index :i32 , depth_ref :f32 , offset :vec2 < i32 > ) >f32 textureSampleCompareLevel ( t :texture_depth_cube , s :sampler_comparison , coords :vec3 < f32 > , depth_ref :f32 ) >f32 textureSampleCompareLevel ( t :texture_depth_cube_array , s :sampler_comparison , coords :vec3 < f32 > , array_index :i32 , depth_ref :f32 ) >f32
The textureSampleCompareLevel
function is the same as textureSampleCompare
, except that:

textureSampleCompareLevel
always samples texels from mip level 0.
The function does not compute derivatives.

There is no requirement for
textureSampleCompareLevel
to be invoked in uniform control flow.


textureSampleCompareLevel
may be invoked in any shader stage.
16.8.12. textureSampleGrad
Samples a texture using explicit gradients.
textureSampleGrad ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , ddx :vec2 < f32 > , ddy :vec2 < f32 > ) >vec4 < f32 > textureSampleGrad ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , ddx :vec2 < f32 > , ddy :vec2 < f32 > , offset :vec2 < i32 > ) >vec4 < f32 > textureSampleGrad ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , ddx :vec2 < f32 > , ddy :vec2 < f32 > ) >vec4 < f32 > textureSampleGrad ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , ddx :vec2 < f32 > , ddy :vec2 < f32 > , offset :vec2 < i32 > ) >vec4 < f32 > textureSampleGrad ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , ddx :vec3 < f32 > , ddy :vec3 < f32 > ) >vec4 < f32 > textureSampleGrad ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , ddx :vec3 < f32 > , ddy :vec3 < f32 > , offset :vec3 < i32 > ) >vec4 < f32 > textureSampleGrad ( t :texture_cube < f32 > , s :sampler , coords :vec3 < f32 > , ddx :vec3 < f32 > , ddy :vec3 < f32 > ) >vec4 < f32 > textureSampleGrad ( t :texture_cube_array < f32 > , s :sampler , coords :vec3 < f32 > , array_index :i32 , ddx :vec3 < f32 > , ddy :vec3 < f32 > ) >vec4 < f32 >
Parameters:
t
 The texture to sample. 
s
 The sampler type. 
coords
 The texture coordinates used for sampling. 
array_index
 The 0based texture array index to sample. 
ddx
 The x direction derivative vector used to compute the sampling locations. 
ddy
 The y direction derivative vector used to compute the sampling locations. 
offset

The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes. The offset expression must be either:
offset component must be at least 8 and at most 7 . Values outside
of this range will result in a shadercreation error.

Returns:
The sampled value.
16.8.13. textureSampleLevel
Samples a texture using an explicit mip level, or at mip level 0.
textureSampleLevel ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , level :f32 ) >vec4 < f32 > textureSampleLevel ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , level :f32 , offset :vec2 < i32 > ) >vec4 < f32 > textureSampleLevel ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , level :f32 ) >vec4 < f32 > textureSampleLevel ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , level :f32 , offset :vec2 < i32 > ) >vec4 < f32 > textureSampleLevel ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , level :f32 ) >vec4 < f32 > textureSampleLevel ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , level :f32 , offset :vec3 < i32 > ) >vec4 < f32 > textureSampleLevel ( t :texture_cube < f32 > , s :sampler , coords :vec3 < f32 > , level :f32 ) >vec4 < f32 > textureSampleLevel ( t :texture_cube_array < f32 > , s :sampler , coords :vec3 < f32 > , array_index :i32 , level :f32 ) >vec4 < f32 > textureSampleLevel ( t :texture_depth_2d , s :sampler , coords :vec2 < f32 > , level :i32 ) >f32 textureSampleLevel ( t :texture_depth_2d , s :sampler , coords :vec2 < f32 > , level :i32 , offset :vec2 < i32 > ) >f32 textureSampleLevel ( t :texture_depth_2d_array , s :sampler , coords :vec2 < f32 > , array_index :i32 , level :i32 ) >f32 textureSampleLevel ( t :texture_depth_2d_array , s :sampler , coords :vec2 < f32 > , array_index :i32 , level :i32 , offset :vec2 < i32 > ) >f32 textureSampleLevel ( t :texture_depth_cube , s :sampler , coords :vec3 < f32 > , level :i32 ) >f32 textureSampleLevel ( t :texture_depth_cube_array , s :sampler , coords :vec3 < f32 > , array_index :i32 , level :i32 ) >f32 textureSampleLevel ( t :texture_external , s :sampler , coords :vec2 < f32 > ) >vec4 < f32 >
Parameters:
t
 The sampled or depth texture to sample. 
s
 The sampler type. 
coords
 The texture coordinates used for sampling. 
array_index
 The 0based texture array index to sample. 
level
 The mip level, with level 0 containing a full size version of the texture.
For the functions where level is a f32 , fractional values may interpolate
between two levels if the format is filterable according to the Texture Format Capabilities.When not specified, mip level 0 is sampled. 
offset

The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes. The offset expression must be either:
offset component must be at least 8 and at most 7 . Values outside
of this range will result in a shadercreation error.

Returns:
The sampled value.
16.8.14. textureStore
Writes a single texel to a texture.
textureStore ( t :texture_storage_1d < F , write > , coords :i32 , value :vec4 < T > ) textureStore ( t :texture_storage_2d < F , write > , coords :vec2 < i32 > , value :vec4 < T > ) textureStore ( t :texture_storage_2d_array < F , write > , coords :vec2 < i32 > , array_index :i32 , value :vec4 < T > ) textureStore ( t :texture_storage_3d < F , write > , coords :vec3 < i32 > , value :vec4 < T > )
The channel format T
depends on the storage texel format F
. See the texel format table for the mapping of texel
format to channel format.
Parameters:
t
 The writeonly storage texture. 
coords
 The 0based texel coordinate. 
array_index
 The 0based texture array index. 
value
 The new texel value. 
Note:
An outofbounds access occurs if:

any element of
coords
is outside the range[0, textureDimensions(t))
for the corresponding element, or 
array_index
is outside the range of[0, textureNumLayers(t))
If an outofbounds access occurs, the builtin function may do any of the following:

not be executed

store
value
to some in bounds texel
TODO:
TODO(dsinclair): Need gather operations
16.9. Atomic Builtin Functions
Atomic builtin functions can be used to read/write/readmodifywrite atomic objects. They are the only operations allowed on § 4.2.7 Atomic Types.
All atomic builtin functions use a relaxed
memory
ordering. This means synchronization and ordering guarantees only apply among
atomic operations acting on the same memory locations. No synchronization
or ordering guarantees apply between atomic and nonatomic memory accesses, or
between atomic accesses acting on different memory locations.
Atomic builtin functions must
not be used in a vertex shader stage.
The address space SC
of the atomic_ptr
parameter in all atomic builtin
functions must
be either storage or workgroup.
The access mode A
in all atomic builtin functions must be read_write.
16.9.1. Atomic Load
atomicLoad ( atomic_ptr :ptr < SC , atomic < T > , A > ) >T
Returns the atomically loaded the value pointed to by atomic_ptr
.
It does not modify the object.
16.9.2. Atomic Store
atomicStore ( atomic_ptr :ptr < SC , atomic < T > , A > , v :T )
Atomically stores the value v
in the atomic object pointed to by atomic_ptr
.
16.9.3. Atomic Readmodifywrite
Each function performs the following steps atomically:atomicAdd ( atomic_ptr :ptr < SC , atomic < T > , A > , v :T ) >T atomicSub ( atomic_ptr :ptr < SC , atomic < T > , A > , v :T ) >T atomicMax ( atomic_ptr :ptr < SC , atomic < T > , A > , v :T ) >T atomicMin ( atomic_ptr :ptr < SC , atomic < T > , A > , v :T ) >T atomicAnd ( atomic_ptr :ptr < SC , atomic < T > , A > , v :T ) >T atomicOr ( atomic_ptr :ptr < SC , atomic < T > , A > , v :T ) >T atomicXor ( atomic_ptr :ptr < SC , atomic < T > , A > , v :T ) >T

Load the original value pointed to by
atomic_ptr
. 
Obtains a new value by performing the operation (e.g. max) from the function name with the value v.

Store the new value using
atomic_ptr
.
Each function returns the original value stored in the atomic object.
atomicExchange ( atomic_ptr :ptr < SC , atomic < T > , A > , v :T ) >T
Atomically stores the value v
in the atomic object pointed to atomic_ptr
and returns the original value stored in the atomic object.
atomicCompareExchangeWeak ( atomic_ptr :ptr < SC , atomic < T > , A > , cmp :T , v :T ) >__atomic_compare_exchange_result < T > struct __atomic_compare_exchange_result < T > { old_value :T ; // old value stored in the atomic exchanged :bool ; // true if the exchange was done }
Note: A value cannot be explicitly declared with the type __atomic_compare_exchange_result
, but a value may infer the type.
Performs the following steps atomically:

Load the original value pointed to by
atomic_ptr
. 
Compare the original value to the value
v
using an equality operation. 
Store the value
v
only if
the result of the equality comparison wastrue
.
Returns a two member structure, where the first member, old_value
, is the
original value of the atomic object and the second member, exchanged
, is
whether or not the comparison succeeded.
Note: the equality comparison may spuriously fail on some implementations. That
is, the second component of the result vector may be false
even if the first
component of the result vector equals cmp
.
16.10. Data Packing Builtin Functions
Data packing builtin functions can be used to encode values using data formats that do not correspond directly to types in WGSL. This enables a program to write many densely packed values to memory, which can reduce a shader’s memory bandwidth demand.
Overload  Description 
pack4x8snorm (e: vec4<f32>) > u32
 Converts four normalized floating point values to 8bit signed integers, and then combines them
into one u32 value.Component e[i] of the input is converted to an 8bit twos complement integer value ⌊ 0.5 + 127 × min(1, max(1, e[i])) ⌋ which is then placed in bits 8 × i through 8 × i + 7 of the result. 
pack4x8unorm (e: vec4<f32>) > u32
 Converts four normalized floating point values to 8bit unsigned integers, and then combines them
into one u32 value.Component e[i] of the input is converted to an 8bit unsigned integer value ⌊ 0.5 + 255 × min(1, max(0, e[i])) ⌋ which is then placed in bits 8 × i through 8 × i + 7 of the result. 
pack2x16snorm (e: vec2<f32>) > u32
 Converts two normalized floating point values to 16bit signed integers, and then combines them
into one u32 value.Component e[i] of the input is converted to a 16bit twos complement integer value ⌊ 0.5 + 32767 × min(1, max(1, e[i])) ⌋ which is then placed in bits 16 × i through 16 × i + 15 of the result. 
pack2x16unorm (e: vec2<f32>) > u32
 Converts two normalized floating point values to 16bit unsigned integers, and then combines them
into one u32 value.Component e[i] of the input is converted to a 16bit unsigned integer value ⌊ 0.5 + 65535 × min(1, max(0, e[i])) ⌋ which is then placed in bits 16 × i through 16 × i + 15 of the result. 
pack2x16float (e: vec2<f32>) > u32
 Converts two floating point values to halfprecision floating point numbers, and then combines
them into one u32 value.Component e[i] of the input is converted to a IEEE754 binary16 value, which is then placed in bits 16 × i through 16 × i + 15 of the result. See § 12.5.2 Floating Point Conversion. 
16.11. Data Unpacking Builtin Functions
Data unpacking builtin functions can be used to decode values in data formats that do not correspond directly to types in WGSL. This enables a program to read many densely packed values from memory, which can reduce a shader’s memory bandwidth demand.
Overload  Description 
unpack4x8snorm (e: u32) > vec4<f32>
 Decomposes a 32bit value into four 8bit chunks, then reinterprets
each chunk as a signed normalized floating point value. Component i of the result is max(v ÷ 127, 1), where v is the interpretation of bits 8×i through 8×i+7 of e as a twoscomplement signed integer. 
unpack4x8unorm (e: u32) > vec4<f32>
 Decomposes a 32bit value into four 8bit chunks, then reinterprets
each chunk as an unsigned normalized floating point value. Component i of the result is v ÷ 255, where v is the interpretation of bits 8×i through 8×i+7 of e as an unsigned integer. 
unpack2x16snorm (e: u32) > vec2<f32>
 Decomposes a 32bit value into two 16bit chunks, then reinterprets
each chunk as a signed normalized floating point value. Component i of the result is max(v ÷ 32767, 1), where v is the interpretation of bits 16×i through 16×i+15 of e as a twoscomplement signed integer. 
unpack2x16unorm (e: u32) > vec2<f32>
 Decomposes a 32bit value into two 16bit chunks, then reinterprets
each chunk as an unsigned normalized floating point value. Component i of the result is v ÷ 65535, where v is the interpretation of bits 16×i through 16×i+15 of e as an unsigned integer. 
unpack2x16float (e: u32) > vec2<f32>
 Decomposes a 32bit value into two 16bit chunks, and reinterpets each chunk
as a floating point value. Component i of the result is the f32 representation of v, where v is the interpretation of bits 16×i through 16×i+15 of e as an IEEE754 binary16 value. See § 12.5.2 Floating Point Conversion. 
16.12. Synchronization Builtin Functions
WGSL provides the following synchronization functions:
storageBarrier () workgroupBarrier ()
All synchronization functions execute a control barrier with
Acquire/Release memory ordering.
That is, all synchronization functions, and affected memory and atomic
operations are ordered in program order relative to the
synchronization function.
Additionally, the affected memory and atomic operations programordered before
the synchronization function must be visible to all other threads in the
workgroup before any affected memory or atomic operation programordered after
the synchronization function is executed by a member of the workgroup.
All synchronization functions use the Workgroup
memory scope.
All synchronization functions have a Workgroup
execution scope.
All synchronization functions must only be used in the compute shader
stage.
storageBarrier affects memory and atomic operations in the storage address space.
workgroupBarrier affects memory and atomic operations in the workgroup address space.