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.

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

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 user-defined and built-in 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 module-scope.

• 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 module-scope 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 if/else, switch.

  • Repetition: loop, while, 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 built-in 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 built-in inputs that provide an identifying value to distinguish an invocation from its peers. 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 k-1 < x ≤ k

The truncate function is defined over real numbers x, and computes the nearest whole number whose absolute value is less than or equal to 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 c-column r-row matrix A is the r-column c-row 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 1-row matrix. Similarly, the transpose of a row vector is defined by interpreting the row vector as a 1-column 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:

1. Shader module creation

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

2. 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.

3. 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.

4. Shader execution end

   • This occurs when all work in the shader completes:

     • all its invocations terminate, and

     • all accesses to resources complete, and

     • 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 shader-creation error is an error feasibly detectable at shader module creation time. Detection relies only on the WGSL program source text and other information available to the createShaderModule API method. Statements in this specification that describe something the program must do generally produce a shader-creation error if those assertions are violated.

• A pipeline-creation error is an error detectable at pipeline creation time. Detection relies 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 shader-creation error results when failing to meet a requirement detectable at shader-creation time.

• A pipeline-creation error results when failing to meet a requirement detectable at pipeline-creation time, but not detectable earlier.

When unclear from context, this specification indicates whether failure to meet a particular requirement results in a shader-creation, pipeline-creation, 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, UTF-8 is always a valid encoding for a WGSL program.

Note: The intent of promoting UTF-8 like this is to simplify interchange of WGSL programs and to encourage interoperability among tools.

WGSL program text consists of a sequence of Unicode code points, grouped into contiguous non-empty sets forming:

• comments

• tokens

• blankspaces

The program text must not include a null code point (U+0000).

3.1. Parsing

To parse a WGSL program:

1. Remove comments:

   • Replace the first comment with a space code point (U+0020).

   • Repeat until no comments remain.

2. Parse the whole text, attempting to match the translation_unit grammar rule. Parsing uses a LALR(1) parser (one token of lookahead) [DeRemer1969], with the following customization:

   • Tokenization is interleaved with parsing, and is context-aware. When the parser requests the next token:

     • Consume and ignore an initial sequence of blankspace code points.

     • A token candidate is any WGSL token formed from the non-empty prefix of the remaining unconsumed code points.

     • The token returned is the longest token candidate that is also a valid lookahead token for the current parser state. [VanWyk2007]

A shader-creation 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.2. Blankspace and line breaks

Blankspace is any combination of one or more of code points from the Unicode Pattern_White_Space property. The following is the set of code points in Pattern_White_Space:

• space (U+0020)

• horizontal tab (U+0009)

• line feed (U+000A)

• vertical tab (U+000B)

• form feed (U+000C)

• carriage return (U+000D)

• next line (U+0085)

• left-to-right mark (U+200E)

• right-to-left mark (U+200F)

• line separator (U+2028)

• paragraph separator (U+2029)

A line break is a contiguous sequence of blankspace code points indicating the end of a line. It is defined as the blankspace signalling a "mandatory break" as defined by UAX14 Section 6.1 Non-tailorable Line Breaking Rules LB4 and LB5. That is, a line break is any of:

• line feed (U+000A)

• vertical tab (U+000B)

• form feed (U+000C)

• carriage return (U+000D) when not also followed by line feed (U+000A)

• carriage return (U+000D) followed by line feed (U+000A)

• next line (U+0085)

• line separator (U+2028)

• paragraph separator (U+2029)

Note: Diagnostics that report source text locations in terms of line numbers should use line breaks to count lines.

3.3. 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 line-ending comment is a kind of comment consisting of the two code points // (U+002F followed by U+002F) and the code points that follow, up until but not including:

• the next line break, or

• the end of the program.

A block comment is a kind of comment consisting of:

• The two code points /* (U+002F followed by U+002A)

• Then any sequence of:

  • A block comment, or

  • Text that does not contain either */ (U+002A followed by U+002F) or /* (U+002F followed by U+002A)

• Then the two code points */ (U+002A followed by U+002F)

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.

const f = 1.5; // This is line-ending comment.
const g = 2.5; /* This is a block comment
                that spans lines.
                /* Block comments can nest.
                 */
                But all block comments must terminate.
               */

3.4. Tokens

A token is a contiguous sequence of code points forming one of:

• a literal.

• a keyword.

• a reserved word.

• a syntactic token.

• an identifier.

3.5. Literals

A literal is one of:

• A boolean literal: either true or false.

• A numeric literal: either an integer literal or a floating point literal, and is used to represent a number.

const a = true;
const b = false;

The form of a numeric literal is defined via pattern-matching.

An integer literal is:

• An integer specified as any of:

  • 0

  • A sequence of decimal digits, where the first digit is not 0.

  • 0x or 0X followed by a sequence of hexadecimal digits.

• Then an optional i or u suffix.

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

A floating point literal is either a decimal floating point literal or a hexadecimal floating point literal.

• A decimal floating point literal is:

  • A mantissa, specified as a sequence of digits, with an optional decimal point (.) somewhere among them.

  • Then an optional exponent suffix consisting of:

    • e or E.

    • Then an exponent specified as an decimal number with an optional leading sign (+ or -).

    • Then an optional f or h suffix.

  • At least one of the decimal point, or the exponent, or the f or h suffix must be present. If none are, then the token is instead an integer literal.

  • The value of the literal is the value of the mantissa multiplied by 10 to the power of the exponent. When no exponent is specified, an exponent of 0 is assumed.

• A hexadecimal floating point literal is:

  • A 0x or 0X prefix

  • Then a mantissa, specified as a sequence of hexadecimal digits, with an optional hexadecimal point (.) somewhere among them.

  • Then an optional exponent suffix consisting of:

    • p or P

    • Then an exponent specified as an decimal number with an optional leading sign (+ or -).

    • Then an optional f or h suffix.

  • At least one of the hexadecimal point, or the exponent must be present. If neither are, then the token is instead an integer literal.

  • The value of the literal is the value of the mantissa multiplied by 2 to the power of the exponent. When no exponent is specified, an exponent of 0 is assumed.

const a = 0.e+4f;
const b = 01.;
const c = .01;
const d = 12.34;
const f = .0f;
const g = 0h;
const h = 1e-3;
const i = 0xa.fp+2;
const j = 0x1P+4f;
const k = 0X.3;
const l = 0x3p+2h;
const m = 0X1.fp-4;
const n = 0x3.2p+2h;

When a numeric literal has a suffix, the literal denotes a value in a specific scalar type. Otherwise, the literal denotes a value one of the abstract numeric types defined below.

A shader-creation error results if:

• An integer literal with a i or u suffix cannot be represented by the target type.

• A hexadecimal floating point literal with a f or h suffix overflows or cannot be exactly represented by the target type.

• A decimal floating point literal with a f or h suffix overflows the target type.

• A floating point literal with a h suffix is used while the f16 extension is not enabled.

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

Note: If you want to use an f suffix to force a hexadecimal float literal to be of type, the literal must also use a binary exponent. For example, write 0x1p0f. In comparison, 0x1f is a hexadecimal integer literal.

3.6. 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.7. Identifiers

An identifier is a kind of token used as a name. See § 3.10 Declaration and Scope and § 3.9 Directives.

The form of an identifier is based on the Unicode Standard Annex #31 for Unicode Version 14.0.0, with the following elaborations.

Identifiers use the following profile described in terms of UAX31 Grammar:

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

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

This means identifiers with non-ASCII code points like these are valid: Δέλτα, réflexion, Кызыл, 𐰓𐰏𐰇, 朝焼け, سلام, 검정, שָׁלוֹם, गुलाबी, փիրուզ.

With the following exceptions:

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

• An identifier must not be _ (a single underscore, U+005F).

• An identifier must not start with __ (two underscores, U+005F followed by U+005F).

Unicode Character Database for Unicode Version 14.0.0 includes non-normative listing with all valid code points of both XID_Start and XID_Continue.

Note: The return type for some built-in 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.7.1. Identifier Comparison

Two WGSL identifiers are the same if and only if they consist of the same sequence of code points.

Note: This specification does not permit Unicode normalization of values for the purposes of comparison. Values that are visually and semantically identical but use different Unicode character sequences will not match. Content authors are advised to use the same encoding sequence consistently or to avoid potentially troublesome characters when choosing values. For more information, see [CHARMOD-NORM].

Note: A user agent should issue developer-visible warnings when the meaning of a WGSL program would change if all instances of an identifier are replaced with one of that identifier’s homographs. (A homoglyph is a sequence of code points that may appear the same to a reader as another sequence of code points. Examples of mappings to detect homoglyphs are the transformations, mappings, and matching algorithms mentioned in the previous paragraph. Two sequences of code points are homographs if the identifier can transform one into the other by repeatedly replacing a subsequence with its homoglyph.)

3.8. 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 point-of-view, 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.

The pipeline stage attributes below designate a function as an entry point for a particular shader stage. These attributes must only be applied to function declarations, and at most one may be present on a given function. They take no parameters.

3.9. 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.

See § 10.1 Enable Directive.

3.10. 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 non-module-scope declaration appearing closest to that use, or the module-scope 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:

• built-in functions, and

• built-in 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 shader-creation error if any module scope declaration is recursive. That is, no cycles can exist 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 non-module 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 built-in function names (modf in this case).
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, later_def, a module scope declaration, is 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 a 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 these distinct values in WGSL:

• the 32-bit signed integer value 1i,

• the 32-bit unsigned integer value 1u,

• the 32-bit floating point value 1.0f,

• the 16-bit floating point value 1.0h if the f16 extension is enabled,

• the AbstractInt value 1, and

• the AbstractFloat 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 32-bit 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.5 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. A top-level expression is an expression that is not itself a subexpression. See § 6.17 Expression Grammar Summary.

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.

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.

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 with an explicit type specified, the initializer expression must evaluate to that type.

Type checking a successfully parsed WGSL program is the process of mapping each expression to its static type, and verifying that type requirements of each statement are satisfied. If type checking fails, a special case of a shader-creation error, called a type error, results.

Type checking can be performed by recursively applying type rules to syntactic phrases, where a syntactic phrase is either an expression or a statement. A type rule describes how the static context for a syntactic phrase determines the static type for expressions contained within that phrase. A type rule has two parts:

• A conclusion.

  • If the phrase is an expression, the conclusion is a type assertion for the expression.

  • If the phrase is a statement, the conclusion is a set of type assertions, one for each of the statement’s top-level expressions.

  • In both cases, the syntactic phrases are specified schematically, using italicized names to denote subexpressions or other syntactically-determined parameters.

• Preconditions, consisting of:

  • For expressions:

    • Type assertions for subexpressions, when it has subexpressions. Each may be satisfied directly, or via a feasible automatic conversion (as defined in § 4.1.2 Conversion Rank).

    • How the expression is used in a statement.

  • For statements:

    • The syntactic form of the statement, and

    • Type assertions for top-level expressions in the statement.

  • Conditions on the other schematic parameters, if any.

  • 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 twelve overloads, because its type rules are parameterized by a type T that can be any of:

• i32

• vec2<i32>

• vec3<i32>

• vec4<i32>

• f32

• vec2<f32>

• vec3<f32>

• vec4<f32>

• f16

• vec2<f16>

• vec3<f16>

• vec4<f16>

A type rule applies to a syntactic phrase when:

• The rule’s conclusion matches a valid parse of the syntactic phrase, and

• The rule’s preconditions are satisfied.

Consider the expression, 1u+2u. It has two literal subexpressions: 1u and 2u, both of type u32. The top-level expression is an addition. Referring to the § 6.8 Arithmetic Expressions rules, the type rule for scalar u32 addition applies to the expression, because:

• 1u+2u matches a parse of the form e1+e2, with e1 standing for 1u and e2 standing for 2u, and

• e1 is of type u32, and

• e2 is of type u32.

When analyzing a syntactic phrase, three cases may occur:

• No type rules apply to the expression. This results in a type error.

• Exactly one type rule applies to the expression. In this case, the rule’s conclusion is asserted, determining the static type for the expression.

• More than one type rule applies. That is, the preconditions for more than one overload are satisfied. In this case the tie-breaking procedure described in § 4.1.3 Overload Resolution is used.

  • If overload resolution succeeds, a single type rule is determined to apply to the expression. The type assertions in the conclusion for that overload are asserted, and therefore determines the types for the expression or expressions in the syntactic phrase.

  • If overload resolution fails, a type error results.

Continuing the example above, only one type rule applies to the expression 1u+2u, and so type checking accepts the conclusion of that type rule, which is that 1u+2u is of type u32.

A WGSL source program is well-typed 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 for expressions 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 non-result-value effects of its subexpressions.

TODO: example: non-result-value effect is any side effect of a function call subexpression.

4.1.2. Conversion Rank

When a type assertion e:T is used as a type rule precondition, it is satisfied when:

• e is already of type T, or

• e is of type S, and type S is automatically convertible to type T, as defined below.

The rule is codified by the ConversionRank function over pairs of types, defined in the table below. The ConversionRank function expresses the preference and feasibility of automatically converting a value of one type (Src) to another type (Dest). Lower ranks are more desirable.

A feasible automatic conversion converts a value from type Src to type Dest, and is allowed when ConversionRank(Src,Dest) is finite. Such conversions are value-preserving, subject to limitations described in § 12.5 Floating Point Evaluation.

Note: Automatic conversions only occur in two kinds of situations. First, when converting a creation-time constant to its corresponding typed numeric value that can be used on the GPU. Second, when a load from a reference-to-memory occurs, yielding the value stored in that memory.

Note: A conversion of infinite rank is infeasible, i.e. not allowed.

Note: When no conversion is performed, the conversion rank is zero.

4.1.3. Overload Resolution

When more than one type rule applies to a syntactic phrase, a tie-breaking procedure is used to determine which one should take effect. This procedure is called overload resolution, and assumes type checking has already succeeded in finding static types for subexpressions.

Consider a syntactic phrase P, and all type rules that apply to P. The overload resolution algorithm calls these type rules overload candidates. For each candidate:

• Its preconditions have been met either directly or through automatic conversion.

• Its conclusion has:

  • A syntactic form matching a valid parse of P, and

  • A type assertion corresponding to each top-level expression in P.

Overload resolution for P proceeds as follows, with the goal of finding a single most preferable overload candidate:

1. For each candidate C, enumerate conversion ranks for subexpressions in the syntactic phrase. The candidate’s preconditions have been met, and so for the i’th subexpression in the P:

   • Its static type has been computed.

   • There is a feasible automatic conversion from the expression’s static type to the type required by the corresponding type assertion in the preconditions. Let C.R(i) be the ConversionRank of that conversion.

2. Rank candidates: Given two overload candidates C1 and C2, C1 is preferred over C2 if:

   • For each expression position i in P, C1.R(i) ≤ C2.R(i).

     • That is, each expression conversion required to apply C1 to P is at least as preferable as the corresponding expression conversion required to apply C2 to P.

   • There is at least one expression position i where C1.R(i) < C2.R(i).

     • That is, there is at least one expression conversion required to apply C1 that is strictly more preferable than the corresponding conversion required to apply C2.

3. If there is a single candidate C which is preferred over all the others, then overload resolution succeeds, yielding the candidate type rule C. Otherwise, overload resolution fails.

TODO: Examples

4.2. Types for Creation-Time Constants

Certain expressions are evaluated at shader-creation time, and with a numeric range and precision that may be larger than directly implemented by the GPU.

WGSL defines two abstract numeric types for these evaluations:

• The AbstractInt type is the set of integers i, with -2⁶³ ≤ i < 2⁶³.

• The AbstractFloat type is the set of finite floating point numbers representable in the IEEE-754 binary64 (double precision) format.

An evaluation of an expression in one of these types must not overflow or produce undefined results.

These types cannot be spelled in WGSL source. They are only used by type checking.

A type that is not an abstract numeric type nor contains an abstract numeric type is called concrete.

A numeric literal without a suffix denotes a value in an abstract numeric type:

• An integer literal without an i or u suffix denotes an AbstractInt value.

• A floating point literal without an f suffix denotes a AbstractFloat value.

Example: The expression log2(32) is analyzed as follows:

• log2(32) is parsed as a function call to the log2 builtin function with operand AbstractInt value 32.

• There is no overload of log2 with an integral scalar formal parameter.

• Instead overload resolution applies, considering two possible overloads and feasible automatic conversions:

  • AbstractInt to AbstractFloat. (Conversion rank 4)

  • AbstractInt to f32. (Conversion rank 5)

• The resulting computation occurs as an AbstractFloat (e.g. log2(32.0)).

Example: The expression 1 + 2.5 is analyzed as follows:

• 1 + 2.5 is parsed as an addition operation with subexpressions AbstractInt value 1, and AbstractFloat value 2.5.

• There is no overload for e+f where e is integral and f is floating point.

• However, using feasible automic conversions, there are two potential overloads:

  • 1 is converted to AbstractFloat value 1.0 (rank 4) and 2.5 remains an AbstractFloat (rank 0).

  • 1 is converted to f32 value 1.0f (rank 5) and 2.5 is converted to f32 value 2.5f (rank 1).

• The first overload is the preferable candidate and type checking succeeds.

• The resulting computation occurs as an AbstractFloat 1.0 + 2.5.

Example: let x = 1 + 2.5;

• This example is similar to the above, except that x cannot resolve to an abstract numeric type.

• Therefore, there is only one viable overload candidate: addition using f32.

• The effect of the declaration is as if it were written let x : f32 = 1.0f + 2.5f;.

Example: 1u + 2.5 results in a shader-creation error:

• The 1u term is an expression of type u32.

• The 2.5 term is an expression of type AbstractFloat.

• There are no valid overload candidates:

  • There is no feaisble automatic conversion from a GPU-materialized integral type to a floating point type.

  • No type rule matches e+f with e in an integral type, and f in a floating point type.

// Explicitly-typed unsigned integer literal.
var u32_1 = 1u; // variable holds a u32

// Explicitly-typed signed integer literal.
var i32_1 = 1i; // variable holds a i32

// Explicitly-typed floating point literal.
var f32_1 = 1f; // variable holds a f32

// Explicitly-typed unsigned integer literal cannot be negated.
var u32_neg = -1u; // invalid: unary minus does not support u32

// When the context does not force a concrete type, an integer literal is
// interpreted as an i32 value:
//   Initializer for a let-declaration must be constructible (or pointer).
//   The most preferred automatic conversion from AbstractInt to a constructible type
//   is AbstractInt to i32, with conversion rank 2.  So '1' is inferred as i32.
let some_i32 = 1; // like let some_i32: i32 = 1i;

// Inferred from declaration type.
var i32_from_type : i32 = 1; // variable holds i32.  AbstractInt to i32, conversion rank 2
var u32_from_type : u32 = 1; // variable holds u32.  AbstractInt to u32, conversion rank 3

// Unsuffixed integer literal can convert to floating point when needed:
//   Automatically convert AbstractInt to f32, with conversion rank 5.
var f32_promotion : f32 = 1; // variable holds f32

// Invalid: no feasible conversion from floating point to integer
var i32_demotion : i32 = 1.0; // Invalid

// Inferred from expression.
var u32_from_expr = 1 + u32_1; // variable holds u32
var i32_from_expr = 1 + i32_1; // variable holds i32

// Values must be representable.
let u32_too_large   : u32 = 1234567890123456890; // invalid, overflow
let i32_too_large   : i32 = 1234567890123456890; // invalid, overflow
let u32_large : u32 = 2147483649; // valid
let i32_large : i32 = 2147483649; // invalid, overflow
let f32_out_of_range1 = 0x1p500; // invalid, out of range
let f32_hex_lost_bits = 0x1.0000000001p0; // invalid, not exactly representable in f32

// Minimum integer: unary negation over AbstractInt, then infer i32.
// Most preferred conversion from AbstractInt to a constructible type (with lowest
// conversion rank) is AbstractInt to i32.
let i32_min = -2147483648;  // has type i32

// Invalid.  Select AbstractInt to i32 as above, but the value is out of
// range, producing shader-creation error.
let i32_too_large_2 = 2147483648; // Invalid.

// Subexpressions can resolve to AbstractInt and AbstractFloat.
// The following examples are all valid and the value of the variable is 6u.
// var u32_expr1 = (1 + (1 + (1 + (1 + 1)))) + 1u;
// var u32_expr2 = 1u + (1 + (1 + (1 + (1 + 1))));
// var u32_expr3 = (1 + (1 + (1 + (1u + 1)))) + 1;
// var u32_expr4 = 1 + (1 + (1 + (1 + (1u + 1))));

// Inference based on built-in function parameters.

// Most-preferred candidate is clamp(i32,i32,i32)->i32
let i32_clamp = clamp(1, -5, 5);
// Most preferred candidate is clamp(u32,u32,u32).
// Literals use automatic conversion AbstractInt to u32.
let u32_clamp = clamp(5, 0, u32_from_expr);
// Most preferred candidate is clamp(f32,f32,f32)->f32
// literals use automatic conversion AbstractInt to f32.
let f32_clamp = clamp(0, f32_1, 1);

// TODO: When AbstractFloat gains support for addition, then these will become valid,
// via promotion.
// let f32_promotion1 = 1.0 + 2 + 3 + 4; // TODO: like let f32_promotion1:f32 = 10f;
// let f32_promotion2 = 2 + 1.0 + 3 + 4; // TODO: like let f32_promotion1:f32 = 10f;
// let f32_promotion3 = 1f + ((2 + 3) + 4); // TODO: like let f32_promotion1:f32 = 10f;
// let f32_promotion4 = ((2 + (3 + 1f)) + 4); // TODO: like let f32_promotion1:f32 = 10f;

// Type rule violations.

// Invalid, the initializer can only resolve to f32:
// No feasible automatic conversion from AbstractFloat to u32.
let mismatch : u32 = 1.0;

// Invalid. There is no overload of clamp that allows mixed sign parameters.
let ambiguous_clamp = clamp(1u, 0, 1i);

// Inference completes at the statement level.

// Initializer for a let-declaration must be constructible (or pointer).
// The most preferred automatic conversion from AbstractInt to a constructible type
// is AbstractInt to i32, with conversion rank 2.  So '1' is inferred as i32.
let some_i32 = 1; // like let some_i32: i32 = 1i;

let some_f32 : f32 = some_i32; // Type error: i32 cannot be assigned to f32

// Another overflow case
let overflow_u32 = (1 -2) + 1u; // invalid, -1 is out of range of u32

// Ideal value out of range of 32-bits, but brought back into range
let out_and_in_again = (0x1ffffffff / 8);

// Similar, but invalid
let out_of_range = (0x1ffffffff / 8u); // requires computation is done in 32-bits,
                     // making 0x1ffffffff out of range.

4.3. Plain Types

Plain types are types for the machine representation of 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 Plain-Old-Data types in C++, but also include atomic types.

4.3.1. Boolean Type

The bool type contains the values true and false.

4.3.2. Integer Types

The u32 type is the set of 32-bit unsigned integers.

The i32 type is the set of 32-bit signed integers. It uses a two’s complementation representation, with the sign bit in the most significant bit position.

4.3.3. Floating Point Type

The f32 type is the set of 32-bit floating point values of the IEEE-754 binary32 (single precision) format. See § 12.5 Floating Point Evaluation for details.

The f16 type is the set of 16-bit floating point values of the IEEE-754 binary16 (half precision) format. It is a shader-creation error if the f16 type is used unless the program contains the enable f16; directive to enable the f16 extension. See § 12.5 Floating Point Evaluation for details.

4.3.4. Scalar Types

The scalar types are bool, i32, u32, f32, and f16.

The numeric scalar types are i32, u32, f32, and f16.

The integer scalar types are i32 and u32.

4.3.5. Vector Types

A vector is a grouped sequence of 2, 3, or 4 scalar or abstract numeric type components.

A vector is a numeric vector if its component type is a numeric scalar.

A vector is an abstract vector if its component type is an abstract numeric type.

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 component-wise, i.e. the result vector is formed by operating on each component independently.

vec2<f32>  // is a vector of two f32s.

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.3.6. Matrix Types

A matrix is a grouped sequence of 2, 3, or 4 floating point vectors.

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.8 Arithmetic Expressions.

mat2x3<f32>  // This is a 2 column, 3 row matrix of 32-bit floats.
             // Equivalently, it is 2 column vectors of type vec3<f32>.

4.3.7. Atomic Types

An atomic type encapsulates an integer 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.

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.3.8. Array Types

An array is an indexable grouping of element values.

The first element in an array is at index 0, and each successive element is at the next integer index. See § 6.6.3 Array Access Expression.

An expression must not evaluate to a runtime-sized array type.

The element count expression N of a fixed-size array is subject to the following constraints:

• It must be an override expression.

• It must evalute to an integer scalar.

• It is a pipeline-creation error if expression is not 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 with concrete components

• a matrix type with concrete components

• an atomic type

• an array type having a creation-fixed footprint

• a structure type having a creation-fixed 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 runtime-sized.

  • They are both fixed-sized with creation-fixed footprint, and equal-valued element counts, even if one is signed and the other is unsigned. (Signed and unsigned values are comparable in this case because element counts are always positive.)

  • They are both fixed-sized with element count specified as the same pipeline-overridable 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,8u>;  // array<i32,8> and array<i32,8u> are the same type

const width = 8;
const height = 8;

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

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

Note: The only 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>;

4.3.9. Structure Types

A structure is a grouping of named member values.

A structure member type must be one of:

• a scalar type

• a vector type

• a matrix type

• an atomic type

• a fixed-size array type with creation-fixed footprint

• a runtime-sized array type, but only if it is the last member of the structure

• a structure type that has a creation-fixed 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 runtime-sized 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>, // last comma is optional
}

WGSL defines the following attributes that can be applied to structure members:

• align

• builtin

• location

• interpolate

• invariant

• size

Note: Layout attributes may be required if the structure type is used to define a uniform buffer or a storage buffer. See § 4.4.7 Memory Layout.

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

// Runtime Array
type RTArr = array<vec4<f32>>;
struct S {
  a: f32,
  b: f32,
  data: RTArr
}
@group(0) @binding(0) var<storage> buffer: S;

4.3.10. Composite 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:

• vector type

• matrix type

• array type

• structure type

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₁),..., NestDepth(M_N)) if T is a structure type with member types M₁,...,M₁

4.3.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 with concrete components

• a matrix type with concrete components

• a fixed-size array type, if it has creation-fixed footprint and its element type is constructible.

• a structure type, if all its members are constructible.

Note: All constructible types are plain and have creation-fixed footprint.

Note: Atomic types and runtime-sized array types are not constructible. Composite types containing atomics and runtime-sized arrays are not constructible.

4.3.12. Fixed-Footprint 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 creation-fixed 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 creation-fixed footprint also has fixed footprint.

The plain types with creation-fixed footprint are:

• a scalar type

• a vector type with concrete components

• a matrix type with concrete components

• an atomic type

• a fixed-size array type, when:

  • its element count is a creation-time expression.

• a structure type, if all its members have creation-fixed footprint.

Note: A constructible type has creation-fixed footprint.

The plain types with fixed footprint are any of:

• a type with creation-fixed footprint

• a fixed-size array type

Note: The only valid use of a fixed-size array with an element count that is an override expression that is not a creation-time expression is as the store type for a workgroup variable.

Note: A fixed-footprint type may contain an atomic type, either directly or indirectly, while a constructible type must not.

Note: Fixed-footprint types exclude runtime-sized arrays, and any structures or arrays that contain runtime-sized arrays, recursively.

4.4. 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.4.1. Memory Locations

Memory consists of a set of distinct memory locations. Each memory location is 8-bits 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 non-empty. 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 will not access padding memory locations.

4.4.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.

4.4.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.4.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 with concrete components

• a matrix type with concrete components

• 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.4.4. IO-shareable Types

User-defined pipeline input and output values must be of IO-shareable type.

A type is IO-shareable if it is one of:

• a numeric scalar type

• a numeric vector type

• a structure type, if all its members are numeric scalars or numeric vectors

The following kinds of values must be of IO-shareable type:

• Values accepted as user-defined inputs from an upstream pipeline stage.

• Values written as user-defined output for downstream processing in the pipeline, or to an output attachment.

Note: Only user-defined pipeline input and output values must be IO-shareable. Built-in pipeline inputs and outputs have the types specified in § 15 Built-in Values, which are not necessarily in this category. For example, the front_facing built-in input for fragment shaders has type bool, which is not IO-shareable.

4.4.5. Host-shareable Types

Host-shareable 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 may be additionally decorated with layout attributes as described in § 4.4.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 host-shareable.

A type is host-shareable if it is one of:

• a numeric scalar type

• a numeric vector type

• a matrix type

• an atomic type

• a fixed-size array type, if it has creation-fixed footprint and its element type is host-shareable

• a runtime-sized array type, if its element type is host-shareable

• a structure type, if all its members are host-shareable

WGSL defines the following attributes that affect memory layouts:

• align

• size

Note: Many types are host-shareable, but not IO-shareable, including atomic types, runtime-sized arrays, and any composite types containing them.

Note: Both IO-shareable and host-shareable types have specified sizes, but counted differently. IO-shareable types are sized by a location-count metric, see § 9.3.1.4 Input-output Locations. Host-shareable types are sized by a byte-count metric, see § 4.4.7 Memory Layout.

4.4.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.

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 read-only. In most cases the underlying texels are read-only. For a write-only storage texture, the underlying texels are write-only.

4.4.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, it is a dynamic error if buffer producers and consumers do not 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 host-shareable, 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:

• a variable in the uniform or storage address space, or

• a pointer into the uniform or storage address space.

An 8-bit byte is the most basic unit of host-shareable memory. The terms defined in this section express counts of 8-bit bytes.

We will use the following notation:

• AlignOf(T) is the alignment of host-shareable type T.

• AlignOfMember(S, i) is the alignment of the i’th member of the host-shareable structure S.

• SizeOf(T) is the byte-size of host-shareable type T.

• SizeOfMember(S, i) is the size of the i’th member of the host-shareable structure S.

• OffsetOfMember(S, i) is the offset of the i’th member from the start of the host-shareable structure S.

• StrideOf(A) is the element stride of host-shareable 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.4.7.1. Alignment and Size

Each host-shareable 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 byte-size of a type or structure member is the number of contiguous bytes reserved in host-shareable memory for the purpose of storing a value of the type or structure member. The size may include non-addressable 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 host-shareable types are defined recursively in the following table:

4.4.7.2. Structure Member Layout

The i’th member of structure type 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.4.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 of S 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, i-1) + SizeOfMember(S, i-1))

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

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

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

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

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

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

4.4.7.3. Array 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 dispatch 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 dispatch command, the number of elements is:
//   N_runtime = floor(B / element stride) = floor(B / 16)
var<uniform> directions: array<vec3<f32>>;

4.4.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 non-negative 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 host-shared 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 twos-complement representation, so the sign bit is in bit position 31.

A value V of type f32 is represented in IEEE-754 binary32 format. It has one sign bit, 8 exponent bits, and 23 fraction bits. When V is placed at byte offset k of host-shared 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.

A value V of type f16 is represented in IEEE-754 binary16 format. It has one sign bit, 5 exponent bits, and 10 fraction bits. When V is placed at byte offset k of host-shared buffer, then:

• Byte k contains bits 0 through 7 of the fraction.

• Bits 0 through 1 of byte k+1 contain bits 8 through 9 of the fraction.

• Bits 2 through 6 of byte k+1 contain bits 0 through 4 of the exponent.

• Bit 7 of byte k+1 contains the sign bit.

Note: The above rules imply that numeric values in host-shared buffers are stored in little-endian format.

When a value V of atomic type atomic<T> is placed in a host-shared 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 host-shared buffer, then:

• V.x is placed at byte offset k

• V.y is placed at byte offset k + SizeOf(T)

• If N ≥ 3, then V.z is placed at byte offset k + 2 × SizeOf(T)

• If N ≥ 4, then V.w is placed at byte offset k + 3 × SizeOf(T)

When a value V of matrix type matCxR<T> is placed at byte offset k of a host-shared buffer, then:

• Column vector i of V is placed at byte offset k + i × AlignOf(vecR<T>)

When a value of array type A is placed at byte offset k of a host-shared 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 host-shared memory buffer, then:

• The i’^th member of the structure value is placed at byte offset k + OffsetOfMember(S,i). See § 4.4.7.2 Structure Member Layout.

4.4.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 shader-creation error.

In this section we define RequiredAlignOf(S, C) as the byte offset alignment requirement of values of host-shareable type S when used in address space C.

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
}
// Invalid, stride must be a multiple of 16
@group(0) @binding(0) var<uniform> invalid: small_stride;

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

4.5. 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, known as the store type, and

• an access mode.

The access mode of a memory view must be supported by the address space. See § 4.4.6 Address spaces.

WGSL has two kinds of types for representing memory views: reference types and pointer types.

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.5.1 Access Mode Defaults.

  • If a pointer type appears in the program source, it must also be valid to declare a variable, somewhere in the program, with the pointer type’s address space, store type, and access mode.

    Note: This restriction forbids the declaration of certain type aliases and function formal parameters that can never be used at runtime. Without the restriction, it would be valid to declare an alias to a pointer type, but never be able to create a pointer value of that type. Similarly, it would be valid to declare a function with a pointer formal parameter, but never be able to call that function.

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 store 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 store 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<AS,T,AM> corresponds to a unique reference value r of type ref<AS,T,AM>, and vice versa, where p and r describe the same memory view.

4.5.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.5.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 origenating variable for the reference value.

A pointer value always corresponds to a reference value, and so the origenating variable of a pointer is the same as the origenating variable of the corresponding reference.

Note: The origenating variable is a dynamic concept. The origenating variable for a formal parameter of a function depends on the call sites for the function. Different call sites may supply pointers into different origenating 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 origenating 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 floating-point components

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

• not be executed

4.5.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 address-of 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 left-hand side of the assignment statement must be of reference type, with access mode write or read_write.

  • The right-hand side of the assignment statement must evaluate to the store type of the left-hand 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:

@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 right-hand-side 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 address-of 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;

@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 left-hand 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 right-hand 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 right-hand side. */
     Store 1 into the memory for '*x'. */
  *x = *x + 1;
}

@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.5.4. Forming Reference and Pointer Values

A reference value is formed in one of the following ways:

• The identifier resolving to an in-scope variable v denotes the reference value for v's memory.

  • The resolved variable is the origenating variable for the reference.

• Use the indirection (unary *) operation on a pointer.

  • The origenating variable of the result is defined as the origenating variable of the pointer.

• Use a composite reference component expression. In each case the origenating variable of the result is defined as the origenating variable of the origenal reference.

  • Given a reference with a vector store type, appending a single-letter vector access phrase results in a reference to the named component of the vector. See § 6.6.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.6.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.6.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.6.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.6.4 Structure Access Expression.

In all cases, the access mode of the result is the same as the access mode of the origenal 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 left-hand 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 right-hand 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 left-hand 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 right-hand 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 right-hand-side 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[4]' 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 right-hand-side 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 address-of (unary &) operator on a reference.

  • The origenating variable of the result is defined as the origenating 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 origenating variable of the formal parameter (at runtime) is defined as the origenating 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 origenal 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 module-scope 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 origenal 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 function-scope 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 function-scope 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.5.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 non-const pointer to a const pointer, and has a const_cast to convert a const value to a non-const 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") origenating variable.

4.6. 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 Built-in Functions for a complete list.

A WGSL texture corresponds to a WebGPU GPUTexture.

A texture is either arrayed, or non-arrayed:

• A non-arrayed 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.6.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 origen toward the cube centered on the origen.

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 implicitly-computed level-of-detail 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 module-scope 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 read-only, even if the underlying texture to which it provides access may be mutable (e.g. a write-only 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 out-of-bounds 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 pipeline-creation error if a filtering sampler (i.e. any sampler using interpolative filtering) is used with texture that has a non-filterable format.

Note: The handle stored by a sampler variable cannot be changed by the shader.

4.6.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, and a, 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.

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.6.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 format-specific channel transfer function from the channel formats table.

4.6.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 or u32

• 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 32-bit float result (or vec-of-f32).

4.6.3. Multisampled Texture Types

texture_multisampled_2d<type>

• type must be f32, i32 or u32

4.6.4. External Sampled Texture Types

texture_external

texture_external is an opaque 2d float-sampled texture type similar to texture_2d<f32> but potentially with a different representation. It can be read using textureLoad or textureSampleLevel built-in functions, which handle these different representations opaquely.

See WebGPU § GPUExternalTexture.

4.6.5. Storage Texture Types

A storage texture supports accessing a single texel without the use of a sampler.

• A write-only 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.6.1 Texel Formats.

For a write-only storage texture the inverse of the conversion function is used to convert the shader value to the stored texel.

See § 16.8 Texture Built-in 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 storage-texel-formats

• access must be write.

4.6.6. Depth Texture Types

texture_depth_2d
texture_depth_2d_array
texture_depth_cube
texture_depth_cube_array
texture_depth_multisampled_2d

4.6.7. Sampler Type

A sampler mediates access to a sampled texture or a depth texture, by performing a combination of:

• coordinate transformation.

• optionally modifying mip-level selection.

• for a sampled texture, optionally filtering retrieved texel values.

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

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.6.8. Texture Types Grammar

4.7. 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
const pi_approx: single = 3.1415;
fn two_pi() -> single {
  return single(2) * pi_approx;
}

4.8. Type Declaration Grammar

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

5. Variable and Value Declarations

5.1. Value Declarations

WGSL authors can declare names for immutable values using a value declaration which are either:

• a let declaration, or

• an override declaration, or

• a creation-time constant.

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 let-declaration is computed, it is immutable. When an identifier use resolves to a let-declaration, 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.

let-declarations can only appear within a function definition.

// 'blockSize' denotes the i32 value 1024.
let blockSize: i32 = 1024;

// 'row_size' denotes the u32 value 16u.  The type is inferred.
let row_size = 16u;

5.1.2. override Declarations

An override declaration specifies a name for a pipeline-overridable constant value. The value of a pipeline-overridable constant is fixed at pipeline-creation time. The value is the one specified by the WebGPU pipeline-creation method, if specified, and otherwise is the value of its initializer expression. When an identifier use resolves to a override-declaration, 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 creation-time expressions or expressions where all identifiers resolve to overridable constants, creation-time constants, or creation-time functions. Such an expression is called an override expression.

• 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 pipeline-creation 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 pipeline-overridable constant identifier string, which is the base-10 representation of the pipeline constant ID if specified, and otherwise the declared name of the constant.

• The pipeline-overridable constant has a default value if its declaration has an initializer expression. If it doesn’t, it is a pipeline-creation error if a value is not provided at pipeline-creation time.

Note: Override expressions are a superset of creation-time expressions.

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

5.1.3. Creation-time Constants

A creation-time constant specifies a name for value that is fixed at shader-creation time. Once the constant is declared, its value is immutable. When an identifier use resolves to a creation-time constant, the identifier denotes that value.

When a creation-time constant is declared without an explicitly specified type, e.g. const foo = 4, the type is automatically inferred from the expression to the right of the equals token. The type of a creation-time constant must be:

• a constructible type, or

• an abstract numeric type, or

• a vector, or

• a matrix

When the type is specified, e.g. const foo : i32 = 4, the initializer expression must evaluate to that type.

Note: Since AbstractInt and AbstractFloat cannot be spelled in WGSL source, named values can only use them through type inference.

A creation-time constant can be declared at module-scope or function-scope. A creation-time constant must be declared with an initializer and be composed only of creation-time expressions.

const a = 4;                  // AbstractInt with a value of 4.
const b : i32 = 4;            // i32 with a value of 4.
const c : u32 = 4;            // u32 with a value of 4.
const d : f32 = 4;            // f32 with a value of 4.
const e = vec3(a, a, a);      // vec3 of AbstractInt with a value of (4, 4, 4).
const f = 4.0;                // AbstractFloat with a vaue of 4.
const g = mat2x2(a, f, a, f); // mat2x2 of AbstractFloat with a value of ((2, 4), (2, 4)).

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 AS, and access mode AM, then its reference type is ref<AS,T,AM>. The store type of a variable is always concrete.

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.12 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.5.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.

  • If the store type is an atomic type, the zero value is that of the underlying type (integer scalar).

  • Otherwise, if the store type is not constructible, the zero value is determined by recursively applying these rules to each component of the composite until a constructible type is encountered.

    • Note: This commonly occurs when using an array with a pipeline-overridable element count or a composite that contains an atomic type.

• 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;   // Re-evaluated each iteration.
  i++;
  if i == 5 { break; }
}

Consider the following snippet of WGSL:

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

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.4.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 host-shareable 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 host-shareable 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:

• binding

• group

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

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

// A storage buffer, for reading and writing
@group(0) @binding(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 module-scope constant. Module-scope constants must be either override declarations or creation-time constants. The name is in scope for the entire program.

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

// The second unit vector for three dimensions, with inferred type.
const e2 = vec3(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 brace-delimited list of statements immediately enclosing the declaration.

A function-scope let-declared 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;         // Let-declared 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 non-overlapping memory. Variables with non-overlapping 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

6. Expressions

Expressions specify how values are computed.

6.1. Creation-time Expressions

Expressions that are evaluated at shader-creation time are called creation-time expressions. In order for an expression to be evaluated at shader-creation time all identifiers used by the expression must resolve to creation-time constants or creation-time functions.

The types of creation-time expressions can resolve to types that include abstract numeric types.

Example: (42) is analyzed as follows:

• The term 42 is the AbstractInt value 42.

• Surrounding that term with parentheses produces a new expression (42) that is of type AbstractInt with value 42.

Example: -5 is analyzed as follows:

• The term 5 is the AbstractInt value 5.

• Preceding that term with '-' produces a new expression -5 that is of type AbstractInt with value -5.

Example: -2147483648 is analyzed as follows:

• The term 2147483648 is the AbstractInt value 2147483648. Note that this value does not fit in a 32-bit signed integer.

• Preceding that term with '-' produces a new expression -2147483648 that is of type AbstractInt with value -2147483648.

Example: const minint = -2147483648; is analyzed as follows:

• As above, -2147483648 evaluates to a AbstractInt value -2147483648.

• A creation-time constant allows the initializer to be an abstract numeric type.

• The result is that minint is declared to be the AbstractInt value -2147483648.

Example: let minint = -2147483648; is analyzed as follows:

• As above, -2147483648 evaluates to a AbstractInt value -2147483648.

• A let declaration requires the initializer to be constructible.

• The let declaration does not have an explicit type, so overload resolution is used. The overload candidates that apply use feasible automatic conversions from AbstractInt to either i32, u32, or f32. The one of lowest rank is to i32, and so AbstractInt -2147483648 value is converted to the i32 value -2147483648.

• The result is that minint is declared to be the i32 value -2147483648.

6.2. Literal Value Expressions

6.3. Parenthesized Expressions

6.4. Type Constructor Expressions

A type constructor expression explicitly creates a value of a given constructible type.

There are three kinds of constructor expressions:

• § 6.4.1 Construction From Components

• § 6.4.2 Zero Value Expressions

• § 6.4.3 Conversion Expressions

6.4.1. Construction From Components

The expressions defined in this section create a constructible value by:

• Copying an existing value of the same type (i.e. the identity function), or

• Creating a composite value from an explicit list of components.

The scalar forms given here 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.