WebGPU Shading Language

Editor’s Draft,

More details about this document
This version:
https://gpuweb.github.io/gpuweb/wgsl/
Latest published version:
https://www.w3.org/TR/WGSL/
Feedback:
GitHub
Inline In Spec
Editors:
(Google)
(Apple Inc.)
Former Editor:
(Google)
Participate:
File an issue (open issues)
Tests:
WebGPU CTS shader/

Abstract

Shading language for WebGPU.

Status of this document

This specification was published by the GPU for the Web Community Group. It is not a W3C Standard nor is it on the W3C Standards Track. Please note that under the W3C Community Contributor License Agreement (CLA) there is a limited opt-out and other conditions apply. Learn more about W3C Community and Business Groups.

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:

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:

When executing a shader stage, the implementation:

A WGSL program is organized into:

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

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:

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:

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:

The ceiling expression is defined over real numbers x:

The truncate function is defined over real numbers x:

The roundUp function is defined for positive integers k and n as:

The transpose of an c-column r-row matrix A is the r-column c-row matrix AT formed by copying the rows of A as the columns of AT:

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

  3. Shader execution start

  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:

2.1. Processing Errors

A WebGPU implementation may fail to process a shader for two reasons:

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

Note: For example, a race condition may not be detectable.

Each requirement will be checked at the earliest opportunity. That is:

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:

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. Scanning from beginning to end, group the remaining code points into tokens and blankspace in greedy fashion:

    • The next group is formed from the longest non-empty prefix of the remaining ungrouped code points, that is either:

      • a valid token, or

      • blankspace

    • Repeat until no ungrouped code points remain.

  3. Discard the blankspace, leaving only tokens.

  4. Parse the token sequence, attempting to match the translation_unit grammar rule.

A shader-creation error results if:

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:

_blankspace :

| /[\u0020\u0009\u000a\u000b\u000c\u000d\u0085\u200e\u200f\u2028\u2029]/uy

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:

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:

A block comment is a kind of comment consisting of:

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.

EXAMPLE: Comments
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:

3.5. Literals

A literal is one of:

EXAMPLE: boolean literals
const a = true;
const b = false;
bool_literal :

| true

| false

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

An integer literal is:

EXAMPLE: integer literals
const a = 0x123;
const b = 0X123u;
const c = 1u;
const d = 123;
const e = 0;
const f = 0i;
const g = 0x3f;
int_literal :

| /(0[xX][0-9a-fA-F]+|0|[1-9][0-9]*)[iu]?/

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

EXAMPLE: floating point literals
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;
float_literal :

| decimal_float_literal

| hex_float_literal

decimal_float_literal :

| /((([0-9]*\.[0-9]+|[0-9]+\.[0-9]*)([eE](\+|-)?[0-9]+)?)|([0-9]+[eE](\+|-)?[0-9]+))[fh]?|0[fh]|[1-9][0-9]*[fh]/

hex_float_literal :

| /0[xX]((([0-9a-fA-F]*\.[0-9a-fA-F]+|[0-9a-fA-F]+\.[0-9a-fA-F]*)([pP](\+|-)?[0-9]+[fh]?)?)|([0-9a-fA-F]+[pP](\+|-)?[0-9]+[fh]?))/

const_literal :

| int_literal

| float_literal

| bool_literal

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.

Mapping literals to types
Literal Suffix Type Examples
integer literal i i32 42i
integer literal u u32 42u
integer literal AbstractInt 124
floating point literal f f32 42f 1e5f 1.2f 0x1.0p10f
floating point literal h f16 42h 1e5h 1.2h 0x1.0p10h
floating point literal AbstractFloat 1e5 1.2 0x1.0p10

A shader-creation error results if:

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:

ident :

| /([_\p{XID_Start}][\p{XID_Continue}]+)|([\p{XID_Start}])/uy

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.

In particular, two identifiers may be distinct in WGSL, but considered the same under conventional normalization, mapping, and matching algorithms such as:

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.

attribute :

| attr ident paren_left ( literal_or_ident comma ) * literal_or_ident comma ? paren_right

| attr ident

literal_or_ident :

| float_literal

| int_literal

| ident

Attributes defined in WGSL
Attribute Valid Values Description
align positive i32 literal Must only be applied to a member of a structure type.

Must be a power of 2, and must satisfy the required-alignment for the member type:

If align(n) is applied to a member of S with type T, and S is the store type or contained in the store type for a variable in address space C, then n must satisfy: n = k × RequiredAlignOf(T,C) for some positive integer k.

See § 4.4.7 Memory Layout

binding non-negative i32 literal Must only be applied to a resource variable.

Specifies the binding number of the resource in a bind group. See § 9.3.2 Resource Interface.

builtin identifier name for a built-in value Must only be applied to an entry point function parameter, entry point return type, or member of a structure.

Declares a built-in value. See § 15 Built-in Values.

const None Must only be applied to function declarations.

Specifies that the function can be used as a creation-time function. It is a shader-creation error if this attribute is a applied to a user-defined function.

Note: This attribute is used as a notational convention to describe which built-in functions can be used in creation-time expressions.

group non-negative i32 literal Must only be applied to a resource variable.

Specifies the binding group of the resource. See § 9.3.2 Resource Interface.

id non-negative i32 literal Must only be applied to an override declaration of scalar type.

Specifies a numeric identifier as an alternate name for a pipeline-overridable constant.

interpolate One or two parameters.

The first parameter must be an interpolation type. The second parameter, if present, must specify the interpolation sampling.

Must only be applied to a declaration that is decorated with a location attribute.

Specifies how the user-defined IO must be interpolated. The attribute is only significant on user-defined vertex outputs and fragment inputs. See § 9.3.1.3 Interpolation.

invariant None Must only be applied to the position built-in value.

When applied to the position built-in output value of a vertex shader, the computation of the result is invariant across different programs and different invocations of the same entry point. That is, if the data and control flow match for two position outputs in different entry points, then the result values are guaranteed to be the same. There is no affect on a position built-in input value.

Note: this attribute maps to the precise qualifier in HLSL, and the invariant qualifier in GLSL.

location non-negative i32 literal Must only be applied to an entry point function parameter, entry point return type, or member of a structure type. Must only be applied to declarations of numeric scalar or numeric vector type. Must not be used with the compute shader stage.

Specifies a part of the user-defined IO of an entry point. See § 9.3.1.4 Input-output Locations.

size positive i32 literal Must only be applied to a member of a structure type.

The number of bytes reserved in the struct for this member.

This number must be at least the byte-size of the type of the member:

If size(n) is applied to a member with type T, then SizeOf(T) ≤ n.

See § 4.4.7 Memory Layout

workgroup_size One, two or three parameters.

Each parameter is either a literal or module-scope constant. All parameters must be of the same type, either i32 or u32.

Must be applied to a compute shader entry point function. Must not be applied to any other object.

Specifies the x, y, and z dimensions of the workgroup grid for the compute shader.

The first parameter specifies the x dimension. The second parameter, if provided, specifies the y dimension, otherwise is assumed to be 1. The third parameter, if provided, specifies the z dimension, otherwise is assumed to be 1. Each dimension must be at least 1 and at most an upper bound specified by the WebGPU API.

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

Pipeline Stage Attributes
Attribute Description
vertex
Declares the function to be an entry point for the vertex shader stage of a render pipeline.
fragment
Declares the function to be an entry point for the fragment shader stage of a render pipeline.
compute
Declares the function to be an entry point for the compute shader stage of a compute pipeline.

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.

global_directive :

| enable_directive

3.10. Declaration and Scope

A declaration associates an identifier with one of the following kinds of objects:

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:

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, there must be no cycles among the declarations:

Consider the directed graph where:

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.

EXAMPLE: Valid and invalid 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 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:

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:

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:

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:

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:

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:

A type rule applies to a syntactic phrase when:

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:

When analyzing a syntactic phrase, three cases may occur:

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:

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:

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.

ConversionRank from one type to another
Src Dest ConversionRank(Src,Dest) Notes
T T 0 Identity. No conversion performed.
ref<S,T,A>
where A is read or read_write
T 0 Apply the Load Rule to load a value from a memory reference.
AbstractFloat f32 1
AbstractFloat f16 2
AbstractInt i32 3
AbstractInt u32 4
AbstractInt AbstractFloat 5
AbstractInt f32 6 Behaves as AbstractInt to AbstractFloat, and then AbstractFloat to f32
AbstractInt f16 7 Behaves as AbstractInt to AbstractFloat, and then AbstractFloat to f16
S T
where above cases don’t apply
infinity There are no automatic conversions between other types.

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:

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:

An evaluation of an expression in one of these types must not overflow or produce undefined results. Otherwise, the result is a shader-creation error.

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:

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

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

Example: let x = 1 + 2.5;

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

EXAMPLE: Type inference for literals
// 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

// An integer literal without a suffix tends to be inferred as i32:
//   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.

Boolean literal type rules
Precondition Conclusion Description
true: bool The true value.
OpConstantTrue %bool
false: bool The false value.
OpConstantFalse %bool

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.

Type Description
vecN<T> Vector of N components of type T. N must be in {2, 3, 4} and T must be one of the scalar or abstract numeric types. We say T is the component type of the vector.

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

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

Key use cases of a vector include:

Many operations on vectors act component-wise, i.e. the result vector is formed by operating on each component independently.

EXAMPLE: Vector
vec2<f32>  // is a vector of two f32s.
EXAMPLE: Component-wise addition
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.

Type Description
matCxR<T> Matrix of C columns and R rows of type T, where C and R are both in {2, 3, 4}, and T must be f32, f16, or AbstractFloat. Equivalently, it can be viewed as C column vectors of type vecR<T>.

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:

See § 6.8 Arithmetic Expressions.

EXAMPLE: Matrix
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:

Type Description
atomic<T> Atomic of type T. T must be either u32 or i32.

An expression must not evaluate to an atomic type.

Atomic types may only be instantiated by variables in the workgroup address space or by storage buffer variables with a read_write access mode. The memory scope of operations on the type is determined by the address space it is instantiated in. Atomic types in the workgroup address space have a memory scope of Workgroup, while those in the storage address space have a memory scope of QueueFamily.

An atomic modification is any operation on an atomic object which sets the content of the object. The operation counts as a modification even if the new value is the same as the object’s existing value.

In WGSL, atomic modifications are mutually ordered, for each object. That is, during execution of a shader stage, for each atomic object A, all agents observe the same order of modification operations applied to A. The ordering for distinct atomic objects may not be related in any way; no causality is implied. Note that variables in workgroup space are shared within a workgroup, but are not shared between different workgroups.

4.3.8. Array Types

An array is an indexable grouping of element values.

Type Description
array<E,N> A fixed-size array with N elements of type E.
N is called the element count of the array.
array<E> A runtime-sized array of elements of type E. These may only appear in specific contexts.

The first element in an array is at index 0, and each successive element is at the next integer index. See § 6.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 must:

Note: The element count value is fully determined at pipeline creation time.

An array element type must be one of:

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:

EXAMPLE: Example fixed-size array types, non-overridable element count
// 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.

EXAMPLE: Workgroup variables sized by overridable constants
override blockSize = 16;

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

// An invalid example, because the overridable element count may only occur
// at the outer level.
// var<workgroup> both: array<array<i32,blockSize>,2>;

// An invalid example, because the overridable element count is only
// valid for workgroup variables.
// var<private> bad_address_space: array<i32,blockSize>;
array_type_decl :

| array less_than type_decl ( comma element_count_expression ) ? greater_than

element_count_expression :

| additive_expression

| bitwise_expression

4.3.9. Structure Types

A structure is a grouping of named member values.

Type Description
struct<T1,...,TN> An ordered tuple of N members of types T1 through TN, with N being an integer greater than 0. A structure type declaration specifies an identifier name for each member. Two members of the same structure type must not have the same name.

A structure member type must be one of:

Note: Each member type must be a plain type.

Some consequences of the restrictions on structure member and array element types are:

EXAMPLE: Structure
// A structure with four members.
struct Data {
  a: i32,
  b: vec2<f32>,
  c: array<i32,10>,
  d: array<f32>, // last comma is optional
}
struct_decl :

| struct ident struct_body_decl

struct_body_decl :

| brace_left ( struct_member comma ) * struct_member comma ? brace_right

struct_member :

| attribute * variable_ident_decl

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

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

EXAMPLE: Structure declaration
struct my_struct {
  a: f32,
  b: vec4<f32>
}
EXAMPLE: Structure used to declare a buffer
// 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:

For a composite type T, the nesting depth of T, written NestDepth(T) is:

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:

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:

Note: A constructible type has creation-fixed footprint.

The plain types with fixed footprint are any of:

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 may access padding between elements, but must not access padding at the end of the structure or array.

4.4.2. Memory Access Mode

A memory access is an operation that acts on 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.

access_mode :

| 'read'

| 'write'

| 'read_write'

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:

Note: That is, the storable types are the plain types, texture types, and sampler types.

4.4.4. IO-shareable Types

Pipeline input and output values must be of IO-shareable type.

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

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

Note: Only built-in pipeline inputs may have a boolean type. A user input or output data attribute must not be of bool type or contain a bool type. See § 9.3.1 Pipeline Input and Output Interface.

4.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 must 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:

WGSL defines the following attributes that affect memory layouts:

Note: An IO-shareable type T is host-shareable if T is not bool and does not contain bool. 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.

Address Spaces
Address space Sharing among invocations Supported access modes Variable scope Restrictions on stored values Notes
function Same invocation only read_write Function scope Constructible type
private Same invocation only read_write Module scope Constructible type
workgroup Invocations in the same compute shader workgroup read_write Module scope Plain type with fixed footprint The element count of an outermost array may be a pipeline-overridable constant.
uniform Invocations in the same shader stage read Module scope Constructible host-shareable types For uniform buffer variables
storage Invocations in the same shader stage read_write, read (default) Module scope Host-shareable For storage buffer variables
handle Invocations in the same shader stage read Module scope Sampler types or texture types For sampler and texture variables.

Note: The token handle is reserved: it is never used in a WGSL program.

Note: A texture variable holds an opaque handle which is used to access the underlying grid of texels. The handle itself is always read-only. In most cases the underlying texels are read-only. For a write-only storage texture, the underlying texels are write-only.

address_space :

| function

| private

| workgroup

| uniform

| storage

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, buffer producers and consumers must agree on the memory layout, which is the description of how the bytes in a buffer are organized into typed WGSL values.

The store type of a buffer variable must be 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:

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:

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:

Alignment and size for host-shareable types
Host-shareable type T AlignOf(T) SizeOf(T)
i32, u32, or f32 4 4
f16 2 2
atomic<|T|> 4 4
vec2<T>, T is i32, u32, or f32 8 8
vec2<f16> 4 4
vec3<T>, T is i32, u32, or f32 16 12
vec3<f16> 8 6
vec4<T>, T is i32, u32, or f32 16 16
vec4<f16> 8 8
matCxR (col-major)

(General form)

AlignOf(vecR) SizeOf(array<vecR, C>)
mat2x2<f32> 8 16
mat2x2<f16> 4 8
mat3x2<f32> 8 24
mat3x2<f16> 4 12
mat4x2<f32> 8 32
mat4x2<f16> 4 16
mat2x3<f32> 16 32
mat2x3<f16> 8 16
mat3x3<f32> 16 48
mat3x3<f16> 8 24
mat4x3<f32> 16 64
mat4x3<f16> 8 32
mat2x4<f32> 16 32
mat2x4<f16> 8 16
mat3x4<f32> 16 48
mat3x4<f16> 8 24
mat4x4<f32> 16 64
mat4x4<f16> 8 32
struct S with members M1...MN max(AlignOfMember(S,1), ... , AlignOfMember(S,N))
roundUp(AlignOf(S), justPastLastMember)

where justPastLastMember = OffsetOfMember(S,N) + SizeOfMember(S,N)
array<E, N>
AlignOf(E) N × roundUp(AlignOf(E), SizeOf(E))
array<E>
AlignOf(E) Nruntime × roundUp(AlignOf(E),SizeOf(E))

where Nruntime is the runtime-determined number of elements of T
4.4.7.2. Structure Member Layout

The i’th member of structure S has a size and alignment, denoted by SizeOfMember(S, i) and AlignOfMember(S, i), respectively. The member sizes and alignments are used to calculate each member’s byte offset from the start of the structure, as described in § 4.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 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))

EXAMPLE: Layout of structures using implicit member sizes and alignments
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;
EXAMPLE: Layout of structures with explicit member sizes and alignments
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
EXAMPLE: Fixed-size 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>;
EXAMPLE: Runtime-sized array layout examples
// 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:

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:

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:

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:

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

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

When a value of structure type S is placed at byte offset k of a host-shared memory buffer, then:

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.

Alignment requirements of a host-shareable type for storage and uniform address spaces
Host-shareable type S RequiredAlignOf(S, storage) RequiredAlignOf(S, uniform)
i32, u32, f32, or f16 AlignOf(S) AlignOf(S)
atomic<T> AlignOf(S) AlignOf(S)
vecN<T> AlignOf(S) AlignOf(S)
matCxR<T> AlignOf(S) AlignOf(S)
array<T, N> AlignOf(S) roundUp(16, AlignOf(S))
array<T> AlignOf(S) roundUp(16, AlignOf(S))
struct S AlignOf(S) roundUp(16, AlignOf(S))

Structure members of type T must have a byte offset from the start of the structure that is a multiple of the RequiredAlignOf(T, C) for the address space C:

OffsetOfMember(S, M) = k × RequiredAlignOf(T, C)
Where k is a positive integer and M is a member of structure S with type T

Arrays of element type T must have an element stride that is a multiple of the RequiredAlignOf(T, C) for the address space C:

StrideOf(array<T, N>) = k × RequiredAlignOf(T, C)
StrideOf(array<T>) = k × RequiredAlignOf(T, C)
Where k is a positive integer

Note: RequiredAlignOf(T, C) does not impose any additional restrictions on the values permitted for an align decoration, nor does it affect the rules of AlignOf(T). Data is laid out with the rules defined in previous sections and then the resulting layout is validated against the RequiredAlignOf(T, C) rules.

The uniform address space also requires that:

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.

EXAMPLE: Satisfying offset requirements for uniform address space
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;
EXAMPLE: Satisfying stride requirements for uniform address space
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:

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.

Constraint Type Description
S is a address space,
T is a storable type,
A is an access mode
ref<S,T,A> The reference type identified with the set of memory views for memory locations in S holding values of type T, supporting memory accesses described by mode A.
In this context T is known as the store type.
Reference types are not written in WGSL program source; instead they are used to analyze a WGSL program.
S is a address space,
T is a storable type,
A is an access mode
ptr<S,T,A> The pointer type identified with the set of memory views for memory locations in S holding values of type T, supporting memory accesses described by mode A.
In this context T is known as the pointee type.
Pointer types may appear in WGSL program source.

When analyzing a WGSL program, reference and pointer types are fully parameterized by an address space, a storable type, and an access mode. In code examples in this specification, the comments show this fully parameterized form.

However, in WGSL source text:

EXAMPLE: Pointer type
fn my_function(
  /* 'ptr<function,i32,read_write>' is the type of a pointer value that references
     memory for keeping an 'i32' value, using memory locations in the 'function'
     address space.  Here 'i32' is the pointee type.
     The implied access mode is 'read_write'. See below for access mode defaults. */
  ptr_int: ptr<function,i32>,

  // 'ptr<private,array<f32,50>,read_write>' is the type of a pointer value that
  // refers to memory for keeping an array of 50 elements of type 'f32', using
  // memory locations in the 'private' address space.
  // Here the pointee type is 'array<f32,50>'.
  // The implied access mode is 'read_write'. See below for access mode defaults.
  ptr_array: ptr<private, array<f32, 50>>
) { }

Reference types and pointer types are both sets of memory views: a particular memory view is associated with a unique reference value and also a unique pointer value:

Each pointer value p of type ptr<S,T,A> corresponds to a unique reference value r of type ref<S,T,A>, and vice versa, where p and r describe the same memory view.

4.5.1. Access Mode Defaults

The access mode for a memory view is often determined by context:

When writing a variable declaration or a pointer type in WGSL source:

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 originating variable for the reference value.

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

Note: The originating variable is a dynamic concept. The originating variable for a formal parameter of a function depends on the call sites for the function. Different call sites may supply pointers into different originating variables.

If a reference or pointer access is out of bounds, an invalid memory reference is produced. Loads from an invalid reference return one of:

Stores to an invalid reference may either: Read-modify-write atomics that operate on an invalid memory reference must load and store from the same memory locations if they access memory.

4.5.3. Use Cases for References and Pointers

References and pointers are distinguished by how they are used:

Defining references in this way enables simple idiomatic use of variables:

EXAMPLE: Reference types enable simple 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;
}
EXAMPLE: Returning a reference returns the value loaded via the reference
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:

Note: The following examples use WGSL features explained later in this specification.

EXAMPLE: Using a pointer as a short name for part of a variable
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;
}
EXAMPLE: Using a pointer as a formal parameter
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:

In all cases, the access mode of the result is the same as the access mode of the original reference.

EXAMPLE: Component reference from a composite 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:

In all cases, the access mode of the result is the same as the access mode of the original pointer.

EXAMPLE: Pointer from a variable
// 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 original variable, so
    // the fully specified type is ptr<private,f32,read_write>.  But read_write
    // is the default access mode for function address space, so read_write does not
    // have to be spelled in this case
    let x_ptr: ptr<private,f32> = &x;

    // The name 'y' resolves to the 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:

Note: From the above rules, it is not possible to form a "dangling" pointer, i.e. a pointer that does not reference the memory for a valid (or "live") originating variable.

4.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 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 origin toward the cube centered on the origin.

size

The extent of grid coordinates along each dimension.

mip level count

The mip level count is at least 1 for sampled textures, and equal to 1 for storage textures.
Mip level 0 contains a full size version of the texture. Each successive mip level contains a filtered version of the previous mip level at half the size (within rounding) of the previous mip level.
When sampling a texture, an explicit or 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:

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.

Channel Formats
Channel format Number of stored bits Interpretation of stored bits Shader type Shader value (Channel Transfer Function)
8unorm 8 unsigned integer v ∈ {0,...,255} f32 v ÷ 255
8snorm 8 signed integer v ∈ {-128,...,127} f32 max(-1, v ÷ 127)
8uint 8 unsigned integer v ∈ {0,...,255} u32 v
8sint 8 signed integer v ∈ {-128,...,127} i32 v
16uint 16 unsigned integer v ∈ {0,...,65535} u32 v
16sint 16 signed integer v ∈ {-32768,...,32767} i32 v
16float 16 IEEE-754 binary16 16-bit floating point value v, with 1 sign bit, 5 exponent bits, 10 mantissa bits f32 v
32uint 32 32-bit unsigned integer value v u32 v
32sint 32 32-bit signed integer value v i32 v
32float 32 IEEE-754 binary32 32-bit floating point value v f32 v

The texel formats listed in the Texel Formats for Storage Textures table correspond to the WebGPU plain color formats which support the WebGPU STORAGE usage. These texel formats are used to parameterize the storage texture types defined in § 4.6.5 Storage Texture Types.

When the texel format does not have all four channels, then:

The last column in the table below uses the format-specific channel transfer function from the channel formats table.

Texel Formats for Storage Textures
Texel format Channel format Channels in memory order Corresponding shader value
rgba8unorm 8unorm r, g, b, a vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a))
rgba8snorm 8snorm r, g, b, a vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a))
rgba8uint 8uint r, g, b, a vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a))
rgba8sint 8sint r, g, b, a vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a))
rgba16uint 16uint r, g, b, a vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a))
rgba16sint 16sint r, g, b, a vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a))
rgba16float 16float r, g, b, a vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a))
r32uint 32uint r vec4<u32>(CTF(r), 0u, 0u, 1u)
r32sint 32sint r vec4<i32>(CTF(r), 0, 0, 1)
r32float 32float r vec4<f32>(CTF(r), 0.0, 0.0, 1.0)
rg32uint 32uint r, g vec4<u32>(CTF(r), CTF(g), 0.0, 1.0)
rg32sint 32sint r, g vec4<i32>(CTF(r), CTF(g), 0.0, 1.0)
rg32float 32float r, g vec4<f32>(CTF(r), CTF(g), 0.0, 1.0)
rgba32uint 32uint r, g, b, a vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a))
rgba32sint 32sint r, g, b, a vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a))
rgba32float 32float r, g, b, a vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a))

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

4.6.3. Multisampled Texture Types

texture_multisampled_2d<type>

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

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:

Type Description
sampler Sampler. Mediates access to a sampled texture.
sampler_comparison Comparison sampler. Mediates access to a depth texture.

Samplers are parameterized when created in the WebGPU API. They cannot be modified by a WGSL program.

Samplers can only be used by the texture builtin functions.

sampler
sampler_comparison

4.6.8. Texture Types Grammar

texture_sampler_types :

| sampler_type

| depth_texture_type

| sampled_texture_type less_than type_decl greater_than

| multisampled_texture_type less_than type_decl greater_than

| storage_texture_type less_than texel_format comma access_mode greater_than

sampler_type :

| sampler

| sampler_comparison

sampled_texture_type :

| texture_1d

| texture_2d

| texture_2d_array

| texture_3d

| texture_cube

| texture_cube_array

multisampled_texture_type :

| texture_multisampled_2d

storage_texture_type :

| texture_storage_1d

| texture_storage_2d

| texture_storage_2d_array

| texture_storage_3d

depth_texture_type :

| texture_depth_2d

| texture_depth_2d_array

| texture_depth_cube

| texture_depth_cube_array

| texture_depth_multisampled_2d

texel_format :

| 'rgba8unorm'

| 'rgba8snorm'

| 'rgba8uint'

| 'rgba8sint'

| 'rgba16uint'

| 'rgba16sint'

| 'rgba16float'

| 'r32uint'

| 'r32sint'

| 'r32float'

| 'rg32uint'

| 'rg32sint'

| 'rg32float'

| 'rgba32uint'

| 'rgba32sint'

| 'rgba32float'

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_alias_decl :

| type ident equal type_decl

EXAMPLE: Type Alias
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

type_decl :

| ident

| type_decl_without_ident

type_decl_without_ident :

| bool

| float32

| float16

| int32

| uint32

| vec_prefix less_than type_decl greater_than

| mat_prefix less_than type_decl greater_than

| pointer less_than address_space comma type_decl ( comma access_mode ) ? greater_than

| array_type_decl

| atomic less_than type_decl greater_than

| texture_sampler_types

vec_prefix :

| vec2

| vec3

| vec4

mat_prefix :

| mat2x2

| mat2x3

| mat2x4

| mat3x2

| mat3x3

| mat3x4

| mat4x2

| mat4x3

| mat4x4

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.

EXAMPLE: Type Declarations
identifier
  Allows to specify types created by the type command

bool
f32
i32
u32
vec2<f32>
array<f32, 4>
array<f32>
mat2x3<f32>
EXAMPLE: Access modes for buffers
// 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:

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. The type of a let declaration is always concrete. 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.

EXAMPLE: let-declared constants at function scope
// '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:

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

EXAMPLE: Module constants, pipeline-overrideable
@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:

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

EXAMPLE: Creation-time constants
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 S, and access mode A, then its reference type is ref<S,T,A>. The store type of a variable is always concrete.

A variable declaration:

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:

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:

Consider the following snippet of WGSL:

EXAMPLE: Variable initial values
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; }
}
The loop body will execute five times. Variable i will take on values 0, 1, 2, 3, 4, 5, and variable twice will take on values 0, 2, 4, 6, 8.

Consider the following snippet of WGSL:

EXAMPLE: Reading a variable multiple times
var x: f32 = 1.0;
let y = x * x + x + 1;
Because x is a variable, all accesses to it turn into load and store operations. However, it is expected that either the browser or the driver optimizes this intermediate representation such that the redundant loads are eliminated.

5.3. Module Scope Variables

A variable declared outside all functions is at module scope. The variable name is in scope for the entire program.

Variables at module scope are restricted as follows:

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:

EXAMPLE: Module scope variable declarations
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.

EXAMPLE: Module constants
// 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 let-declared constant must be of constructible type, or of pointer type.

For a variable declared in function scope:

EXAMPLE: Function scope variables and constants
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

variable_statement :

| variable_decl

| variable_decl equal expression

| let ( ident | variable_ident_decl ) equal expression

| const ( ident | variable_ident_decl ) equal expression

variable_decl :

| var variable_qualifier ? ( ident | variable_ident_decl )

variable_ident_decl :

| ident colon type_decl

variable_qualifier :

| less_than address_space ( comma access_mode ) ? greater_than

global_variable_decl :

| attribute * variable_decl ( equal expression ) ?

global_constant_decl :

| const ( ident | variable_ident_decl ) equal expression

| attribute * override ( ident | variable_ident_decl ) ( equal expression ) ?

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:

Example: -5 is analyzed as follows:

Example: -2147483648 is analyzed as follows:

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

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

6.2. Literal Value Expressions

Scalar literal type rules
Precondition Conclusion Notes
true: bool true boolean value.
false: bool false boolean value.
e is an integer literal with no suffix e: AbstractInt Abstract integer literal value.
e is a floating point literal with no suffix e: AbstractFloat Abstract float literal value.
e is an integer literal with i suffix e: i32 32-bit signed integer literal value.
e is an integer literal with u suffix e: u32 32-bit unsigned integer literal value.
e is an floating point literal with f suffix e: f32 32-bit floating point literal value.
e is an floating point literal with h suffix e: f16 16-bit floating point literal value.

6.3. Parenthesized Expressions

Parenthesized expression type rules
Precondition Conclusion Description
e : T ( e ) : T Evaluates to e.
Use parentheses to isolate an expression from the surrounding text.

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

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

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.

Scalar constructor type rules
Precondition Conclusion Notes
e: bool bool(e): bool Identity.
e: i32 i32(e): i32 Identity.
e: u32 u32(e): u32 Identity.
e: f32 f32(e): f32 Identity.
e: f16 f16(e): f16 Identity.
Vector constructor type rules, where T is a scalar type
Precondition Conclusion Notes
e: T vecN<T>(e): vecN<T> Evaluates e once. Results in the N-component vector where each component has the value of e.
vecN(e): vecN<T>
e1: T
e2: T
vec2<T>(e1,e2): vec2<T>
vec2(e1,e2): vec2<T>
e: vec2<T> vec2<T>(e): vec2<T> Identity. The result is e.
vec2(e): vec2<T>
e1: T
e2: T
e3: T
vec3<T>(e1,e2,e3): vec3<T>
vec3(e1,e2,e3): vec3<T>
e1: T
e2: vec2<T>
vec3<T>(e1,e2): vec3<T>
vec3<T>(e2,e1): vec3<T>
vec3(e1,e2): vec3<T>
vec3(e2,e1): vec3<T>
e: vec3<T> vec3<T>(e): vec3<T> Identity. The result is e.
vec3(e): vec3<T>
e1: T
e2: T
e3: T
e4: T
vec4<T>(e1,e2,e3,e4): vec4<T>
vec4(e1,e2,e3,e4): vec4<T>
e1: T
e2: T
e3: vec2<T>
vec4<T>(e1,e2,e3): vec4<T>
vec4<T>(e1,e3,e2): vec4<T>
vec4<T>(e3,e1,e2): vec4<T>
vec4(e1,e2,e3): vec4<T>
vec4(e1,e3,e2): vec4<T>
vec4(e3,e1,e2): vec4<T>
e1: vec2<T>
e2: vec2<T>
vec4<T>(e1,e2): vec4<T>
vec4(e1,e2): vec4<T>
e1: T
e2: vec3<T>
vec4<T>(e1,e2): vec4<T>
vec4<T>(e2,e1): vec4<T>
vec4(e1,e2): vec4<T>
vec4(e2,e1): vec4<T>
e: vec4<T> vec4<T>(e): vec4<T> Identity. The result is e.
vec4(e): vec4<T>
Matrix constructor type rules
Precondition Conclusion Notes
e: mat2x2<T> mat2x2<T>(e): mat2x2<T>
mat2x2(e): mat2x2<T>
Identity type conversion. The result is e.
e: mat2x3<T> mat2x3<T>(e): mat2x3<T>
mat2x3(e): mat2x3<T>
e: mat2x4<T> mat2x4<T>(e): mat2x4<T>
mat2x4(e): mat2x4<T>
e: mat3x2<T> mat3x2<T>(e): mat3x2<T>
mat3x2(e): mat3x2<T>
e: mat3x3<T> mat3x3<T>(e): mat3x3<T>
mat3x3(e): mat3x3<T>
e: mat3x4<T> mat3x4<T>(e): mat3x4<T>
mat3x4(e): mat3x4<T>
e: mat4x2<T> mat4x2<T>(e): mat4x2<T>
mat4x2(e): mat4x2<T>
e: mat4x3<T> mat4x3<T>(e): mat4x3<T>
mat4x3(e): mat4x3<T>
e: mat4x4<T> mat4x4<T>(e): mat4x4<T>
mat4x4(e): mat4x4<T>
e1: T
...
eN: T
mat2x2<T>(e1,e2,e3,e4): mat2x2<T>
mat3x2<T>(e1,...,e6): mat3x2<T>
mat2x3<T>(e1,...,e6): mat2x3<T>
mat4x2<T>(e1,...,e8): mat4x2<T>
mat2x4<T>(e1,...,e8): mat2x4<T>
mat3x3<T>(e1,...,e9): mat3x3<T>
mat4x3<T>(e1,...,e12): mat4x3<T>
mat3x4<T>(e1,...,e12): mat3x4<T>
mat4x4<T>(e1,...,e16): mat4x4<T>
Column-major construction by elements.
mat2x2(e1,e2,e3,e4): mat2x2<T>
mat3x2(e1,...,e6): mat3x2<T>
mat2x3(e1,...,e6): mat2x3<T>
mat4x2(e1,...,e8): mat4x2<T>
mat2x4(e1,...,e8): mat2x4<T>
mat3x3(e1,...,e9): mat3x3<T>
mat4x3(e1,...,e12): mat4x3<T>
mat3x4(e1,...,e12): mat3x4<T>
mat4x4(e1,...,e16): mat4x4<T>
e1: vec2<T>
e2: vec2<T>
e3: vec2<T>
e4: vec2<T>
mat2x2<T>(e1,e2): mat2x2<T>
mat3x2<T>(e1,e2,e3): mat3x2<T>
mat4x2<T>(e1,e2,e3,e4): mat4x2<T>
Column by column construction.
mat2x2(e1,e2): mat2x2<T>
mat3x2(e1,e2,e3): mat3x2<T>
mat4x2(e1,e2,e3,e4): mat4x2<T>
e1: vec3<T>
e2: vec3<T>
e3: vec3<T>
e4: vec3<T>
mat2x3<T>(e1,e2): mat2x3<T>
mat3x3<T>(e1,e2,e3): mat3x3<T>
mat4x3<T>(e1,e2,e3,e4): mat4x3<T>
Column by column construction.
mat2x3(e1,e2): mat2x3<T>
mat3x3(e1,e2,e3): mat3x3<T>
mat4x3(e1,e2,e3,e4): mat4x3<T>
e1: vec4<T>
e2: vec4<T>
e3: vec4<T>
e4: vec4<T>
mat2x4<T>(e1,e2): mat2x4<T>
mat3x4<T>(e1,e2,e3): mat3x4<T>
mat4x4<T>(e1,e2,e3,e4): mat4x4<T>
Column by column construction.
mat2x4(e1,e2): mat2x4<T>
mat3x4(e1,e2,e3): mat3x4<T>
mat4x4(e1,e2,e3,e4): mat4x4<T>
Array constructor type rules
Precondition Conclusion Notes
e1: T
...
eN: T,
T is constructible
array<T,N>(e1,...,eN) : array<T,N> Construction of an array from elements.

Note: array<T,N> is constructible because its element count is equal to the number of arguments to the constructor, and hence fully determined at shader-creation time.

Structure constructor type rules
Precondition Conclusion Notes
e1: T1
...
eN: TN,
S is a constructible structure type with members having types T1 ... TN.
The expression is in the scope of declaration of S.
S(e1,...,eN): S Construction of a structure from members.

6.4.2. Zero Value Expressions

Each constructible T has a unique zero value written in WGSL as the type followed by an empty pair of parentheses: T ().

The zero values are as follows:

Note: WGSL does not have zero expression for atomic types, runtime-sized arrays, or other types that are not constructible.

Scalar zero value type rules
Precondition Conclusion Notes
bool(): bool false
Zero value
i32(): i32 0
Zero value
u32(): u32 0u
Zero value
f32(): f32 0.0
Zero value
f16(): f16 0.0
Zero value
Vector zero type rules, where T is a scalar type
Precondition Conclusion Notes
vec2<T>(): vec2<T> Zero value
vec3<T>(): vec3<T> Zero value
vec4<T>(): vec4<T> Zero value
EXAMPLE: Zero-valued vectors
vec2<f32>()                 // The zero-valued vector of two f32 components.
vec2<f32>(0.0, 0.0)         // The same value, written explicitly.

vec3<i32>()                 // The zero-valued vector of three i32 components.
vec3<i32>(0, 0, 0)          // The same value, written explicitly.
Matrix zero type rules
Precondition Conclusion Notes
T is f32 or f16 mat2x2<T>(): mat2x2<T>
mat3x2<T>(): mat3x2<T>
mat4x2<T>(): mat4x2<T>
Zero value
mat2x3<T>(): mat2x3<T>
mat3x3<T>(): mat3x3<T>
mat4x3<T>(): mat4x3<T>
Zero value
mat2x4<T>(): mat2x4<T>
mat3x4<T>(): mat3x4<T>
mat4x4<T>(): mat4x4<T>
Zero value
Array zero type rules
Precondition Conclusion Notes
T is a constructible array<T,N>(): array<T,N> Zero-valued array
EXAMPLE: Zero-valued arrays
array<bool, 2>()               // The zero-valued array of two booleans.
array<bool, 2>(false, false)   // The same value, written explicitly.
Structure zero type rules
Precondition Conclusion Notes
S is a constructible structure type.
The expression is in the scope of declaration of S.
S(): S Zero-valued structure: a structure of type S where each member is the zero value for its member type.
EXAMPLE: Zero-valued structures
struct Student {
  grade: i32,
  GPA: f32,
  attendance: array<bool,4>
}

fn func() {
  var s: Student;

  // The zero value for Student
  s = Student();

  // The same value, written explicitly.
  s = Student(0, 0.0, array<bool,4>(false, false, false, false));

  // The same value, written with zero-valued members.
  s = Student(i32(), f32(), array<bool,4>());
}

6.4.3. Conversion Expressions

WGSL does not implicitly convert or promote a numeric or boolean value to another type. Instead use a conversion expression as defined in the tables below.

For details on conversion to and from floating point types, see § 12.5.2 Floating Point Conversion.

Scalar conversion type rules
Precondition Conclusion Notes
e: u32 bool(e): bool Coercion to boolean.
The result is false if e is 0, and true otherwise.
e: i32 bool(e): bool Coercion to boolean.
The result is false if e is 0, and true otherwise.
e: f32 bool(e): bool Coercion to boolean.
The result is false if e is 0.0 or -0.0, and true otherwise. In particular NaN and infinity values map to true.
e: f16 bool(e): bool Coercion to boolean.
The result is false if e is 0.0 or -0.0, and true otherwise. In particular NaN and infinity values map to true.
e: bool i32(e): i32 Conversion of a boolean value to a signed integer
The result is 1 if e is true and 0 otherwise.
e: u32 i32(e): i32 Reinterpretation of bits.
The result is the unique value in i32 that has the same bit pattern as e.
e: f32 i32(e): i32 Value conversion, rounding toward zero.
e: f16 i32(e): i32 Value conversion, rounding toward zero.
e: bool u32(e): u32 Conversion of a boolean value to an unsigned integer.
The result is 1u if e is true and 0u otherwise.
e: AbstractInt or i32 u32(e): u32 Reinterpretation of bits.
The result is the unique value in u32 that has the same bit pattern as e.
e: f32 u32(e): u32 Value conversion, rounding toward zero.
e: f16 u32(e): u32 Value conversion, rounding toward zero.
e: bool f32(e): f32 Conversion of a boolean value to floating point.
The result is 1.0 if e is true and 0.0 otherwise.
e: i32 f32(e): f32 Value conversion, including invalid cases.
e: u32 f32(e): f32 Value conversion, including invalid cases.
e: f16 f32(e): f32 Exact value conversion.
e: bool f16(e): f16 Conversion of a boolean value to floating point
The result is 1.0 if e is true and 0.0 otherwise.
e: i32 f16(e): f16 Value conversion, including invalid cases.
e: u32 f16(e): f16 Value conversion, including invalid cases.
e: f32 f16(e): f16 Lossy value conversion.

Details of conversion to and from floating point are explained in § 12.5.2 Floating Point Conversion.

Vector conversion type rules
Precondition Conclusion Notes
e: vecN<u32> vecN<bool>(e): vecN<bool> Component-wise coercion of a unsigned integer vector to a boolean vector.
e: vecN<i32> vecN<bool>(e): vecN<bool> Component-wise coercion of a signed integer vector to a boolean vector.
e: vecN<f32> vecN<bool>(e): vecN<bool> Component-wise coercion of a binary32 floating point vector to a boolean vector.
e: vecN<f16> vecN<bool>(e): vecN<bool> Component-wise coercion of a binary16 floating point vector to a boolean vector.
e: vecN<bool> vecN<i32>(e): vecN<i32> Component-wise conversion of a boolean vector to signed.
Component i of the result is i32(e[i])
e: vecN<u32> vecN<i32>(e): vecN<i32> Component-wise reinterpretation of bits.
Component i of the result is i32(e[i])
e: vecN<f32> vecN<i32>(e): vecN<i32> Component-wise value conversion to signed integer, including invalid cases.
e: vecN<f16> vecN<i32>(e): vecN<i32> Component-wise value conversion to signed integer, including invalid cases.
e: vecN<bool> vecN<u32>(e): vecN<u32> Component-wise conversion of a boolean vector to unsigned.
Component i of the result is u32(e[i])
e: vecN<AbstractInt> or vecN<i32> vecN<u32>(e): vecN<u32> Component-wise reinterpretation of bits.
e: vecN<f32> vecN<u32>(e): vecN<u32> Component-wise value conversion to unsigned integer, including invalid cases.
e: vecN<f16> vecN<u32>(e): vecN<u32> Component-wise value conversion to unsigned integer, including invalid cases.
e: vecN<bool> vecN<f32>(e): vecN<f32> Component-wise conversion of a boolean vector to floating point.
Component i of the result is f32(e[i])
e: vecN<i32> vecN<f32>(e): vecN<f32> Component-wise value conversion to binary32 floating point, including invalid cases.
e: vecN<f32> vecN<f32>(e): vecN<f32> Component-wise value conversion to binary32 floating point, including invalid cases.
e: vecN<f16> vecN<f32>(e): vecN<f32> Component-wise exact value conversion to binary32 floating point.
e: vecN<bool> vecN<f16>(e): vecN<f16> Component-wise conversion of a boolean vector to binary16 floating point.
Component i of the result is f16(e[i])
e: vecN<i32> vecN<f16>(e): vecN<f16> Component-wise value conversion to binary16 floating point, including invalid cases.
e: vecN<u32> vecN<f16>(e): vecN<f> Component-wise value conversion to binary16 floating point, including invalid cases.
e: vecN<f32> vecN<f16>(e): vecN<f16> Component-wise lossy value conversion to binary16 floating point.
Matrix conversion type rules
Precondition Conclusion Notes
e: matCxR<f16> matCxR<f32>(e): matCxR<f32> Component-wise exact value conversion to binary32 floating point.
e: matCxR<f32> matCxR<f16>(e): matCxR<f16> Component-wise lossy value conversion to binary16 floating point.

6.5. Reinterpretation of Representation Expressions

A bitcast expression is used to reinterpet the bit representation of a value in one type as a value in another type.

Bitcast type rules
Precondition Conclusion Notes
e: T
T is a numeric scalar or numeric vector type
bitcast<T>(e): T Identity transform. Component-wise when T is a vector.
The result is e.
e: T1
T1 is i32, u32, or f32
T2 is not T1 and is i32, u32, or f32
bitcast<T2>(e): T2 Reinterpretation of bits as T2.
The result is the reinterpretation of the bits in e as a T2 value.
e: vecN<T1>
T1 is i32, u32, or f32
T2 is not T1 and is i32, u32, or f32
bitcast<vecN<T2>>(e): vecN<T2> Component-wise reinterpretation of bits as T2.
The result is the reinterpretation of the bits in e as a vecN<T2> value.
e: vec2<f16>
T is i32, u32, or f32
bitcast<T>(e): T Reinterpretation of bits as T.
The result is the reinterpretation of the 32 bits in e as a T value, following the internal layout rules.
e: T
T is i32, u32, or f32
bitcast<vec2<f16>>(e): vec2<f16> Reinterpretation of bits as vec2<f16>.
The result is the reinterpretation of the 32 bits in e as a vec2<f16> value, following the internal layout rules.
e: vec4<f16>
T is i32, u32, or f32
bitcast<vec2<T>>(e): vec2<T> Reinterpretation of bits as vec2<T>.
The result is the reinterpretation of the 64 bits in e as a vec2<T> value, following the internal layout rules.
e: vec2<T>
T is i32, u32, or f32
bitcast<vec4<f16>>(e): vec4<f16> Reinterpretation of bits as vec4<f16>.
The result is the reinterpretation of the 64 bits in e as a vec4<f16> value, following the internal layout rules.

The internal layout rules are described in § 4.4.7.4 Internal Layout of Values.

6.6. Composite Value Decomposition Expressions

6.6.1. Vector Access Expression

Accessing components of a vector can be done either using array subscripting (e.g. a[2]) or using a sequence of convenience names, each mapping to a component of the source vector.

The convenience names are accessed using the . notation. (e.g. color.bgra).

NOTE: the convenience letterings can not be mixed. (i.e. you can not use rybw).

Using a convenience letter, or array subscript, which accesses a component past the end of the vector is an error.

The convenience letterings can be applied in any order, including duplicating letters as needed. You can provide 1 to 4 letters when extracting components from a vector. Providing more then 4 letters is an error.

The result type depends on the number of letters provided. Assuming a vec4<f32>

Accessor Result type
r f32
rg vec2<f32>
rgb vec3<f32>
rgba vec4<f32>
var a: vec3<f32> = vec3<f32>(1., 2., 3.);
var b: f32 = a.y;          // b = 2.0
var c: vec2<f32> = a.bb;   // c = (3.0, 3.0)
var d: vec3<f32> = a.zyx;  // d = (3.0, 2.0, 1.0)
var e: f32 = a[1];         // e = 2.0
6.6.1.1. Vector Single Component Selection
Vector decomposition: single component selection
Precondition Conclusion Description
e: vecN<T>
e.x: T
e.r: T
Select the first component of e
e: vecN<T>
e.y: T
e.g: T
Select the second component of e
e: vecN<T>
N is 3 or 4
e.z: T
e.b: T
Select the third component of e
e: vec4<T> e.w: T
e.a: T
Select the fourth component of e
e: vecN<T>
i: i32 or u32
e[i]: T Select the ith component of vector
The first component is at index i=0.
If i is outside the range [0,N-1], then any valid value for T may be returned.
6.6.1.2. Vector Multiple Component Selection
Vector decomposition: multiple component selection
Precondition Conclusion Description
e: vecN<T>
I is the letter x, y, z, or w
J is the letter x, y, z, or w
e.IJ: vec2<T>
Computes the two-component vector with first component e.I, and second component e.J.
Letter z is valid only when N is 3 or 4.
Letter w is valid only when N is 4.
e: vecN<T>
I is the letter r, g, b, or a
J is the letter r, g, b, or a
e.IJ: vec2<T>
Computes the two-component vector with first component e.I, and second component e.J.
Letter b is valid only when N is 3 or 4.
Letter a is valid only when N is 4.
e: vecN<T>
I is the letter x, y, z, or w
J is the letter x, y, z, or w
K is the letter x, y, z, or w
e.IJK: vec3<T>
Computes the three-component vector with first component e.I, second component e.J, and third component e.K.
Letter z is valid only when N is 3 or 4.
Letter w is valid only when N is 4.
e: vecN<T>
I is the letter r, g, b, or a
J is the letter r, g, b, or a
K is the letter r, g, b, or a
e.IJK: vec3<T>
Computes the three-component vector with first component e.I, second component e.J, and third component e.K.
Letter b is only valid when N is 3 or 4.
Letter a is only valid when N is 4.
e: vecN<T>
I is the letter x, y, z, or w
J is the letter x, y, z, or w
K is the letter x, y, z, or w
L is the letter x, y, z, or w
e.IJKL: vec4<T>
Computes the four-component vector with first component e.I, second component e.J, third component e.K, and fourth component e.L.
Letter z is valid only when N is 3 or 4.
Letter w is valid only when N is 4.
e: vecN<T>
I is the letter r, g, b, or a
J is the letter r, g, b, or a
K is the letter r, g, b, or a
L is the letter r, g, b, or a
e.IJKL: vec4<T>
Computes the four-component vector with first component e.I, second component e.J, third component e.K, and fourth component e.L.
Letter b is only valid when N is 3 or 4.
Letter a is only valid when N is 4.
6.6.1.3. Component Reference from Vector Reference

A write access to component of a vector may access all of the memory locations associated with that vector.

Note: This means accesses to different components of a vector by different invocations must be synchronized if at least one access is a write access. See § 16.12 Synchronization Built-in Functions.

Getting a reference to a component from a reference to a vector
Precondition Conclusion Description
r: ref<S,vecN<T>>
r.x: ref<S,T>
r.r: ref<S,T>
Compute a reference to the first component of the vector referenced by the reference r.
The originating variable of the resulting reference is the same as the originating variable of r.
r: ref<S,vecN<T>>
r.y: ref<S,T>
r.g: ref<S,T>
Compute a reference to the second component of the vector referenced by the reference r.
The originating variable of the resulting reference is the same as the originating variable of r.
r: ref<S,vecN<T>>
N is 3 or 4
r.z: ref<S,T>
r.b: ref<S,T>
Compute a reference to the third component of the vector referenced by the reference r.
The originating variable of the resulting reference is the same as the originating variable of r.
r: ref<S,vec4<T>>
r.w: ref<S,T>
r.a: ref<S,T>
Compute a reference to the fourth component of the vector referenced by the reference r.
The originating variable of the resulting reference is the same as the originating variable of r.
r: ref<S,vecN<T>>
i: i32 or u32
r[i] : ref<S,T> Compute a reference to the ith component of the vector referenced by the reference r.

If i is outside the range [0,N-1], then the expression evaluates to invalid memory reference.

The originating variable of the resulting reference is the same as the originating variable of r.

6.6.2. Matrix Access Expression

Column vector extraction
Precondition Conclusion Description
e: matCxR<T>
i: i32 or u32
e[i]: vecR<T> The result is the ith column vector of e.

If i is outside the range [0,C-1], then any valid value for vecR<T> may be returned.

Getting a reference to a column vector from a reference to a matrix
Precondition Conclusion Description
r: ref<S,matCxR<T>>
i: i32 or u32
r[i] : ref<vecR<S,T>> Compute a reference to the ith column vector of the matrix referenced by the reference r.

If i is outside the range [0,C-1], then the expression evaluates to invalid memory reference.

The originating variable of the resulting reference is the same as the originating variable of r.

6.6.3. Array Access Expression

Array element extraction
Precondition Conclusion Description
e: array<T,N>
i: i32 or u32
e[i] : T The result is the value of the ith element of the array value e.

If i is outside the range [0,N-1], then any valid value for T may be returned.

Getting a reference to an array element from a reference to an array
Precondition Conclusion Description
r: ref<S,array<T,N>>
i: i32 or u32
r[i] : ref<S,T> Compute a reference to the ith element of the array referenced by the reference r.

If i is outside the range [0,N-1], then the expression evaluates to an invalid memory reference.

The originating variable of the resulting reference is the same as the originating variable of r.

r: ref<S,array<T>>
i: i32 or u32
r[i] : ref<S,T> Compute a reference to the ith element of the runtime-sized array referenced by the reference r.

If at runtime the array has N elements, and i is outside the range [0,N-1], then the expression evaluates to an invalid memory reference.

The originating variable of the resulting reference is the same as the originating variable of r.

6.6.4. Structure Access Expression

Structure member extraction
Precondition Conclusion Description
S is a structure type
M is the identifier name of a member of S, having type T
e: S
e.M: T The result is the value of the member with name M from the structure value e.
Getting a reference to a structure member from a reference to a structure
Precondition Conclusion Description
S is a structure type
M is the name of a member of S, having type T
r: ref<S,S>
r.M: ref<S,T> Given a reference to a structure, the result is a reference to the structure member with identifier name M.
The originating variable of the resulting reference is the same as the originating variable of r.

6.7. Logical Expressions

Unary logical operations
Precondition Conclusion Notes
e: bool
T is bool or vecN<bool>
!e: T Logical negation. The result is true when e is false and false when e is true. Component-wise when T is a vector.
Binary logical expressions
Precondition Conclusion Notes
e1: bool
e2: bool
e1 || e2: bool Short-circuiting "or". Yields true if either e1 or e2 are true; evaluates e2 only if e1 is false.
e1: bool
e2: bool
e1 && e2: bool Short-circuiting "and". Yields true if both e1 and e2 are true; evaluates e2 only if e1 is true.
e1: T
e2: T
T is bool or vecN<bool>
e1 | e2: T Logical "or". Component-wise when T is a vector. Evaluates both e1 and e2.
e1: T
e2: T
T is bool or vecN<bool>
e1 & e2: T Logical "and". Component-wise when T is a vector. Evaluates both e1 and e2.

6.8. Arithmetic Expressions

Unary arithmetic expressions
Precondition Conclusion Notes
e: T
T is AbstractInt, AbstractFloat, i32, f32, f16, vecN<AbstractInt>, vecN<AbstractFloat>, vecN<i32>, vecN<f32>, or vecN<f16>
-e: T Negation. Component-wise when T is a vector. If T is an integral type and e evaluates to the largest negative value, then the result is e.
Binary arithmetic expressions
Precondition Conclusion Notes
e1 : T
e2 : T
S is AbstractInt, AbstractFloat, i32, u32, f32, or f16
T is S, or vecN<S>
e1 + e2 : T Addition. Component-wise when T is a vector. If T is a concrete integral type, then the result is modulo 232.
e1 : T
e2 : T
S is AbstractInt, AbstractFloat, i32, u32, f32, or f16
T is S, or vecN<S>
e1 - e2 : T Subtraction Component-wise when T is a vector. If T is a concrete integral type, then the result is modulo 232.
e1 : T
e2 : T
S is AbstractInt, AbstractFloat, i32, u32, f32, or f16
T is S, or vecN<S>
e1 * e2 : T Multiplication. Component-wise when T is a vector. If T is a concrete integral type, then the result is modulo 232.
e1 : T
e2 : T
S is AbstractInt, AbstractFloat, i32, u32, f32, or f16
T is S, or vecN<S>
e1 / e2 : T Division. Component-wise when T is a vector.

If T is a signed integral type, the scalar case, evaluates to:

  • e1, when e2 is zero.

  • e1, when e1 is the most negative value in T, and e2 is -1.

  • truncate(x) otherwise, where x is the real-valued quotient e1 ÷ e2.

Note: The need to ensure truncation behaviour may require an implementation to perform more operations than when computing an unsigned division. Use unsigned division when both operands are known to have the same sign.

If T is an unsigned integral type, the scalar case, evaluates to:

  • e1, when e2 is zero.

  • Otherwise, the integer q such that e1 = q × e2 + r, where 0 ≤ r < e2.

e1 : T
e2 : T
S is AbstractInt, AbstractFloat, i32, u32, f32, or f16
T is S, or vecN<S>
e1 % e2 : T Remainder. Component-wise when T is a vector.

If T is a signed integral scalar type, evaluates e1 and e2 once, and evaluates to:

  • 0, when e2 is zero.

  • 0, when e1 is the most negative value in T, and e2 is -1.

  • Otherwise, e1 - truncate(e1 ÷ e2) × e2 where the quotient is computed as a real value.

Note: When non-zero, the result has the same sign as e1.

Note: The need to ensure consistent behaviour may require an implementation to perform more operations than when computing an unsigned remainder.

If T is an unsigned integral scalar type, evaluates to:

  • 0, when e2 is zero.

  • Otherwise, the integer r such that e1 = q × e2 + r, where q is an integer and 0 ≤ r < e2.

If T is a floating point type, the result is equal to:
e1 - e2 * trunc(e1 / e2)

Binary arithmetic expressions with mixed scalar and vector operands
Preconditions Conclusions Semantics
S is one of AbstractInt, AbstractFloat, f32, f16, i32, u32
V is vecN<S>
es: S
ev: V
ev + es: V ev + V(es)
es + ev: V V(es) + ev
ev - es: V ev - V(es)
es - ev: V V(es) - ev
ev * es: V ev * V(es)
es * ev: V V(es) * ev
ev / es: V ev / V(es)
es / ev: V V(es) / ev
ev % es: V ev % V(es)
es % ev: V V(es) % ev
Matrix arithmetic
Preconditions Conclusions Semantics
e1, e2: matCxR<T>
T is AbstractFloat, f32, or f16
e1 + e2: matCxR<T>
Matrix addition: column i of the result is e1[i] + e2[i]
e1 - e2: matCxR<T> Matrix subtraction: column i of the result is e1[i] - e2[i]
m: matCxR<T>
s: T
T is AbstractFloat, f32, or f16
m * s: matCxR<T>
Component-wise scaling: (m * s)[i][j] is m[i][j] * s
s * m: matCxR<T>
Component-wise scaling: (s * m)[i][j] is m[i][j] * s
m: matCxR<T>
v: vecC<T>
T is AbstractFloat, f32, or f16
m * v: vecR<T>
Linear algebra matrix-column-vector product: Component i of the result is dot(m[i],v)
m: matCxR<T>
v: vecR<T>
T is AbstractFloat, f32, or f16
v * m: vecC<T>
Linear algebra row-vector-matrix product:
transpose(transpose(m) * transpose(v))
e1: matKxR<T>
e2: matCxK<T>
T is AbstractFloat, f32, or f16
e1 * e2: matCxR<T>
Linear algebra matrix product.

6.9. Comparison Expressions

Comparisons
Precondtion Conclusion Notes
e1: T
e2: T
S is AbstractInt, AbstractFloat, bool, i32, u32, f32, or f16
T is S or vecN<S>
TB is bool if T is scalar or
vecN<bool> if T is a vector
e1 == e2: TB Equality. Component-wise when T is a vector.
e1: T
e2: T
S is AbstractInt, AbstractFloat, bool, i32, u32, or f32
T is S or vecN<S>
TB is bool if T is scalar or
vecN<bool> if T is a vector
e1 != e2: TB Inequality. Component-wise when T is a vector.
e1: T
e2: T
S is AbstractInt, AbstractFloat, i32, u32, f32, or f16
T is S, or vecN<S>
TB is bool if T is scalar, or
vecN<bool> if T is a vector
e1 < e2: TB Less than. Component-wise when T is a vector.
e1: T
e2: T
S is AbstractInt, AbstractFloat, i32, u32, f32, or f16
T is S, or vecN<S>
TB is bool if T is scalar, or
vecN<bool> if T is a vector
e1 <= e2: TB Less than or equal. Component-wise when T is a vector.
e1: T
e2: T
S is AbstractInt, AbstractFloat, i32, u32, f32, or f16
T is S, or vecN<S>
TB is bool if T is scalar, or
vecN<bool> if T is a vector
e1 > e2: TB Greater than. Component-wise when T is a vector.
e1: T
e2: T
S is AbstractInt, AbstractFloat, i32, u32, f32, or f16
T is S, or vecN<S>
TB is bool if T is scalar, or
vecN<bool> if T is a vector
e1 >= e2: TB Greater than or equal. Component-wise when T is a vector.

6.10. Bit Expressions

Unary bitwise operations
Precondition Conclusion Notes
e: T
S is AbstractInt, i32, or u32
T is S or vecN<S>
~e : T Bitwise complement on e. Each bit in the result is the opposite of the corresponding bit in e. Component-wise when T is a vector.
Binary bitwise operations
Precondition Conclusion Notes
e1: T
e2: T
S is AbstractInt, i32, or u32
T is S or vecN<S>
e1 | e2: T Bitwise-or. Component-wise when T is a vector.
e1: T
e2: T
S is AbstractInt, i32, or u32
T is S or vecN<S>
e1 & e2: T Bitwise-and. Component-wise when T is a vector.
e1: T
e2: T
S is AbstractInt, i32, or u32
T is S or vecN<S>
e1 ^ e2: T Bitwise-exclusive-or. Component-wise when T is a vector.

6.11. Function Call Expression

A function call expression executes a function call where the called function has a return type. If the called function does not return a value, a function call statement should be used instead. See § 7.5 Function Call Statement.

6.12. Variable Identifier Expression

Getting a reference from a variable name
Precondition Conclusion Description
v is an identifier resolving to an in-scope variable declared in address space S with store type T v: ref<S,T> Result is a reference to the memory for the named variable v.

6.13. Formal Parameter Expression

Getting the value of an identifier declared as a formal parameter to a function
Precondition Conclusion Description
a is an identifier resolving to an in-scope formal parameter declaration with type T a: T Result is the value supplied for the corresponding function call operand at the call site invoking this instance of the function.

6.14. Address-Of Expression

The address-of operator converts a reference to its corresponding pointer.

Getting a pointer from a reference
Precondition Conclusion Description
r: ref<S,T,A> &r: ptr<S,T,A> Result is the pointer value corresponding to the same memory view as the reference value r.

If r is an invalid memory reference, then the resulting pointer is also an invalid memory reference.

It is a shader-creation error if S is the handle address space.

It is a shader-creation error if r is a reference to a vector component.

6.15. Indirection Expression

The indirection operator converts a pointer to its corresponding reference.

Getting a reference from a pointer
Precondition Conclusion Description
p: ptr<S,T,A> *p: ref<S,T,A> Result is the reference value corresponding to the same memory view as the pointer value p.

If p is an invalid memory reference, then the resulting reference is also an invalid memory reference.

6.16. Constant Identifier Expression

Getting the value of a let-declared identifier
Precondition Conclusion Description
c is an identifier resolving to an in-scope override declaration with type T c: T If pipeline creation specified a value for the constant ID, then the result is that value. This value may be different for different pipeline instances.
Otherwise, the result is the value computed for the initializer expression. Pipeline-overridable constants appear at module-scope, so evaluation occurs before the shader begins execution.
Note: Pipeline creation fails if no initial value was specified in the API call and the let-declaration has no initializer expression.
c is an identifier resolving to an in-scope let declaration with type T c: T Result is the value computed for the initializer expression.
For a let declaration at module scope, evaluation occurs before the shader begins execution.
For a let declaration inside a function, evaluation occurs each time control reaches the declaration.

6.17. Expression Grammar Summary

When an identifier is used as a callable item, it is one of:

Declaration and scope rules ensure those names are always distinct.

primary_expression :

| ident

| callable argument_expression_list

| const_literal

| paren_expression

| bitcast less_than type_decl greater_than paren_expression

callable :

| ident

| type_decl_without_ident

| vec_prefix

| mat_prefix

paren_expression :

| paren_left expression paren_right

argument_expression_list :

| paren_left ( ( expression comma ) * expression comma ? ) ? paren_right

postfix_expression :

| bracket_left expression bracket_right postfix_expression ?

| period ident postfix_expression ?

unary_expression :

| singular_expression

| minus unary_expression

| bang unary_expression

| tilde unary_expression

| star unary_expression

| and unary_expression

singular_expression :

| primary_expression postfix_expression ?

lhs_expression :

| ( star | and ) * core_lhs_expression postfix_expression ?

core_lhs_expression :

| ident

| paren_left lhs_expression paren_right

multiplicative_expression :

| unary_expression

| multiplicative_expression star unary_expression

| multiplicative_expression forward_slash unary_expression

| multiplicative_expression modulo unary_expression

additive_expression :

| multiplicative_expression

| additive_expression plus multiplicative_expression

| additive_expression minus multiplicative_expression

relational_expression :

| additive_expression

| additive_expression less_than additive_expression

| additive_expression greater_than additive_expression

| additive_expression less_than_equal additive_expression

| additive_expression greater_than_equal additive_expression

| additive_expression equal_equal additive_expression

| additive_expression not_equal additive_expression

short_circuit_and_expression :

| relational_expression

| short_circuit_and_expression and_and relational_expression

short_circuit_or_expression :

| relational_expression

| short_circuit_or_expression or_or relational_expression

binary_or_expression :

| unary_expression

| binary_or_expression or unary_expression

binary_and_expression :

| unary_expression

| binary_and_expression and unary_expression

binary_xor_expression :

| unary_expression

| binary_xor_expression xor unary_expression

bitwise_expression :

| binary_and_expression and unary_expression

| binary_or_expression or unary_expression

| binary_xor_expression xor unary_expression

expression :

| relational_expression

| short_circuit_or_expression or_or relational_expression

| short_circuit_and_expression and_and relational_expression

| bitwise_expression

7. Statements

Statements are program fragments that control its execution. Statements are generally executed in sequential order; however, control flow statements may cause a program to execute in non-sequential order.

7.1. Compound Statement

A compound statement is a brace-enclosed sequence of zero or more statements. When a declaration is one of those statements, its identifier is in scope from the start of the next statement until the end of the compound statement.

compound_statement :

| brace_left statement * brace_right

There are two special forms of compound statement:

7.2. Assignment Statement

An assignment evaluates an expression, and optionally stores it in memory (thus updating the contents of a variable).

assignment_statement :

| lhs_expression ( equal | compound_assignment_operator ) expression

| underscore equal expression

The text to the left of the operator token is the left-hand side, and the expression to the right of the operator token is the right-hand side.

7.2.1. Simple Assignment

An assignment is a simple assignment when the left-hand side is an expression, and the operator is the equal token. In this case the value of the right-hand side is written to the memory referenced by the left-hand side.

Precondition Statement Description
r: ref<S,T,A>,
A is write or read_write
e: T,
T is a constructible type,
S is a writable address space
r = e Evaluates e, evaluates r, then writes the value computed for e into the memory locations referenced by r. Note: if the reference is an invalid memory reference, the write may not execute, or may write to a different memory location than expected.

In the simplest case, the left hand side is the name of a variable. See § 4.5.4 Forming Reference and Pointer Values for other cases.

EXAMPLE: Assignments
struct S {
    age: i32,
    weight: f32
}
var<private> person: S;

fn f() {
    var a: i32 = 20;
    a = 30;           // Replace the contents of 'a' with 30.

    person.age = 31;  // Write 31 into the age field of the person variable.

    var uv: vec2<f32>;
    uv.y = 1.25;      // Place 1.25 into the second component of uv.

    let uv_x_ptr: ptr<function,f32> = &uv.x;
    *uv_x_ptr = 2.5;   // Place 2.5 into the first component of uv.

    var friend: S;
    // Copy the contents of the 'person' variable into the 'friend' variable.
    friend = person;
}

7.2.2. Phony Assignment

An assignment is a phony assignment when the left-hand side is an underscore token. In this case the right-hand side is evaluated, and then ignored.

Precondition Statement Description
e: T,
T is constructible, a pointer type, a texture type, or a sampler type
_ = e Evaluates e.

Note: The resulting value is not stored. The _ token is not an identifier, and therefore cannot be used in an expression.

A phony-assignment is useful for:

EXAMPLE: Using phony-assignment to throw away an un-needed function result
var<private> counter: i32;

fn increment_and_yield_previous() -> i32 {
  let previous = counter;
  counter = counter + 1;
  return previous;
}

fn user() {
  // Increment the counter, but don’t use the result.
  _ = increment_and_yield_previous();
}
EXAMPLE: Using phony-assignment to occupy bindings without using them
struct BufferContents {
    counter: atomic<u32>,
    data: array<vec4<f32>>
}
@group(0) @binding(0) var<storage> buf: BufferContents;
@group(0) @binding(1) var t: texture_2d<f32>;
@group(0) @binding(2) var s: sampler;

@fragment
fn shade_it() -> @location(0) vec4<f32> {
  // Declare that buf, t, and s are part of the shader interface, without
  // using them for anything.
  _ = &buf;
  _ = t;
  _ = s;
  return vec4<f32>();
}

7.2.3. Compound assignment

An assignment is a compound assignment when the left-hand side is an expression, and the operator is one of the compound_assignment_operators.

compound_assignment_operator :

| plus_equal

| minus_equal

| times_equal

| division_equal

| modulo_equal

| and_equal

| or_equal

| xor_equal

The type requirements, semantics, and behavior of each statement is defined as if the compound assignment expands as in the following table, except that the reference expression e1 is evaluated only once.

Statement Expansion
e1 += e2 e1 = e1 + (e2)
e1 -= e2 e1 = e1 - (e2)
e1 *= e2 e1 = e1 * (e2)
e1 /= e2 e1 = e1 / (e2)
e1 %= e2 e1 = e1 % (e2)
e1 &= e2 e1 = e1 & (e2)
e1 |= e2 e1 = e1 | (e2)
e1 ^= e2 e1 = e1 ^ (e2)

Note: The syntax does not allow a compound assignment to also be a phony assignment.

Note: Even though the reference e1 is evaluated once, its underlying memory is accessed twice: first a read access gets the old value, and then a write access stores the updated value.

EXAMPLE: Compound assignment
var<private> next_item: i32 = 0;

fn advance_item() -> i32 {
   next_item += 1;   // Adds 1 to next_item.
   return next_item - 1;
}

fn bump_item() {
  var data: array<f32,10>;
  next_item = 0;
  // Adds 5.0 to data[0], calling advance_item() only once.
  data[advance_item()] += 5.0;
  // next_item will be 1 here.
}

fn precedence_example() {
  var value = 1;
  // The right-hand side of a compound assignment is its own expression.
  value *= 2 + 3; // Same as value = value * (2 + 3);
  // 'value' now holds 5.
}

Note: A compound assignment can rewritten as different WGSL code that uses a simple assignment instead. The idea is to use a pointer to hold the result of evaluating the reference once.

For example, when e1 is not a reference to a component inside a vector, then e1+=e2 can be rewritten as {let p = &(e1); *p = *p + (e2);}, where the identifier p is chosen to be different from all other identifiers in the program.

When e1 is a reference to a component inside a vector, the above technique needs to be modified because WGSL does not allow taking the address in that case. For example, if ev is a reference to a vector, the statement ev[c] += e2 can be rewritten as {let p = &(ev); let c0 = c; (*p)[c0] = (*p)[c0] + (e2);}, where identifiers c0 and p are chosen to be different from all other identifiers in the program.

7.3. Increment and Decrement Statements

An increment statement adds 1 to the contents of a variable. A decrement statement subtracts 1 from the contents of a variable.

increment_statement :

| lhs_expression plus_plus

decrement_statement :

| lhs_expression minus_minus

The expression must evaluate to a reference with an integer scalar store type and read_write access mode.

Precondition Statement Description
r : ref<SC,T,read_write>,
T is integer scalar
r++ Adds 1 to the contents of memory referenced by r.
Same as r += T(1)
r : ref<SC,T,read_write>,
T is integer scalar
r-- Subtracts 1 from the contents of memory referenced by r.
Same as r -= T(1)
EXAMPLE: Increment and decrement
fn f() {
    var a: i32 = 20;
    a++;
    // Now a contains 21
    a--;
    // Now a contains 20
}

7.4. Control Flow

Control flow statements may cause the program to execute in non-sequential order.

7.4.1. If Statement

if_statement :

| if expression compound_statement ( else else_statement ) ?

else_statement :

| compound_statement

| if_statement

An if statement conditionally executes at most one compound statement based on the evaluation of the condition expressions.

The if statements in WGSL use an if/else if/else structure, that contains a single required if clause, zero or more else if clauses and a single optional else clause.

Type rule precondition: Each of the expressions for the if and else if clause conditions must be scalar boolean expressions.

An if statement is executed as follows:

7.4.2. Switch Statement

switch_statement :

| switch expression brace_left switch_body + brace_right

switch_body :

| case case_selectors colon ? case_compound_statement

| default colon ? case_compound_statement

case_selectors :

| const_literal ( comma const_literal ) * comma ?

case_compound_statement :

| brace_left statement * fallthrough_statement ? brace_right

fallthrough_statement :

| fallthrough semicolon

A switch statement transfers control to one of a set of case clauses, or to the default clause, depending on the evaluation of a selector expression.

If the selector value equals a value in a case selector list, then control is transferred to the body of that case clause. If the selector value does not equal any of the case selector values, then control is transferred to the default clause.

Each switch statement must have exactly one default clause.

Type rule precondition: For a single switch statement, the selector expression and all case selector expressions must be of the same integer scalar type.

A literal value must not appear more than once in the case selectors for a switch statement.

Note: The value of the literal is what matters, not the spelling. For example 0 and 0x0000 both denote the zero value.

When control reaches the end of a case body, control normally transfers to the first statement after the switch statement. Alternately, executing a fallthrough statement transfers control to the body of the next case clause or default clause, whichever appears next in the switch body. A fallthrough statement must not appear as the last statement in the last clause of a switch. When a declaration appears in a case body, its identifier is in scope from the start of the next statement until the end of the case body.

Note: Identifiers declared in a case body are not in scope of case bodies which are reachable via a fallthrough statement.

EXAMPLE: WGSL Switch
var a : i32;
let x : i32 = generateValue();
switch x {
  case 0: {      // the colon is optional
    a = 1;
  }
  default {      // the default needn’t appear last
    a = 2;
  }
  case 1, 2 {    // multiple selector values can be used
    a = 3;       // a will be overridden in the next case
    fallthrough;
  }
  case 3 {
    a = 4;
  }
}

7.4.3. Loop Statement

loop_statement :

| loop brace_left statement * continuing_statement ? brace_right

A loop statement repeatedly executes a loop body; the loop body is specified as a compound statement. Each execution of the loop body is called an iteration.

The identifier of a declaration in a loop is in scope from the start of the next statement until the end of the loop body. The declaration is executed each time it is reached, so each new iteration creates a new instance of the variable or constant, and re-initializes it.

This repetition can be interrupted by a break, return, or discard statement.

Optionally, the last statement in the loop body may be a continuing statement.

Note: The loop statement is one of the biggest differences from other shader languages.

This design directly expresses loop idioms commonly found in compiled code. In particular, placing the loop update statements at the end of the loop body allows them to naturally use values defined in the loop body.

EXAMPLE: GLSL Loop
int a = 2;
for (int i = 0; i < 4; i++) {
  a *= 2;
}
EXAMPLE: WGSL Loop
let a: i32 = 2;
var i: i32 = 0;      // <1>
loop {
  if i >= 4 { break; }

  a = a * 2;

  i++;
}
EXAMPLE: GLSL Loop with continue
int a = 2;
let int step = 1;
for (int i = 0; i < 4; i += step) {
  if i % 2 == 0 continue;
  a *= 2;
}
EXAMPLE: WGSL Loop with continue
var a: i32 = 2;
var i: i32 = 0;
loop {
  if i >= 4 { break; }

  let step: i32 = 1;

  i = i + step;
  if i % 2 == 0 { continue; }

  a = a * 2;
}
EXAMPLE: WGSL Loop with continue and continuing
var a: i32 = 2;
var i: i32 = 0;
loop {
  if i >= 4 { break; }

  let step: i32 = 1;

  if i % 2 == 0 { continue; }

  a = a * 2;

  continuing {   // <2>
    i = i + step;
  }
}

7.4.4. For Statement

for_statement :

| for paren_left for_header paren_right compound_statement

for_header :

| for_init ? semicolon expression ? semicolon for_update ?

for_init :

| variable_statement

| increment_statement

| decrement_statement

| assignment_statement

| func_call_statement

for_update :

| increment_statement

| decrement_statement

| assignment_statement

| func_call_statement

The for statement takes the form for (initializer; condition; update_part) { body } and is syntactic sugar on top of a loop statement with the same body. Additionally:

Type rule precondition: The condition must be of bool type.

The initializer of a for loop is executed once prior to executing the loop. When a declaration appears in the initializer, its identifier is in scope until the end of the body. Unlike declarations in the body, the declaration is not re-initialized each iteration.

The condition, body and update_part execute in that order to form a loop iteration. The body is a special form of compound statement. The identifier of a declaration in the body is in scope from the start of the next statement until the end of the body. The declaration is executed each time it is reached, so each new iteration creates a new instance of the variable or constant, and re-initializes it.

EXAMPLE: For to Loop transformation
for(var i: i32 = 0; i < 4; i++) {
  if a == 0 {
    continue;
  }
  a = a + 2;
}

Converts to:

EXAMPLE: For to Loop transformation
{ // Introduce new scope for loop variable i
  var i: i32 = 0;
  var a: i32 = 0;
  loop {
    if !(i < 4) {
      break;
    }

    if a == 0 {
      continue;
    }
    a = a + 2;

    continuing {
      i++;
    }
  }
}

7.4.5. While Statement

while_statement :

| while expression compound_statement

The while statement is a kind of loop parameterized by a condition. At the start of each loop iteration, a boolean condition is evaluated. If the condition is false, the while loop ends execution. Otherwise, the rest of the iteration is executed.

Type rule precondition: The condition must be of bool type.

A while loop can be viewed as syntactic sugar over either a loop or for statement. The following statement forms are equivalent:

7.4.6. Break Statement

break_statement :

| break

A break statement transfers control to immediately after the body of the nearest-enclosing loop or switch statement, thus ending execution of the loop or switch statement.

A break statement must only be used within loop, for, while, and switch statements.

A break statement must not be placed such that it would exit from a loop’s continuing statement. Use a break-if statement instead.

EXAMPLE: WGSL Invalid loop break from a continuing clause
var a: i32 = 2;
var i: i32 = 0;
loop {
  let step: i32 = 1;

  if i % 2 == 0 { continue; }

  a = a * 2;

  continuing {
    i = i + step;
    if i >= 4 { break; } // Invalid.  Use break-if instead.
  }
}

7.4.7. Break-If Statement

break_if_statement :

| break if expression semicolon

A break-if statement evaluates a boolean condition; If the condition is true, control is transferred to immediately after the body of the nearest-enclosing loop statement, ending execution of that loop.

Type rule precondition: The condition must be of bool type.

Note: A break-if statement may only appear as the last statement in the body of a continuing statement.

EXAMPLE: WGSL Valid loop break-if from a continuing clause
var a: i32 = 2;
var i: i32 = 0;
loop {
  let step: i32 = 1;

  if i % 2 == 0 { continue; }

  a = a * 2;

  continuing {
    i = i + step;
    break if i >= 4;
  }
}

7.4.8. Continue Statement

continue_statement :

| continue

A continue statement transfers control in the nearest-enclosing loop:

A continue statement must only be used in a loop, for or while statement. A continue statement must not be placed such that it would transfer control to an enclosing continuing statement. (It is a forward branch when branching to a continuing statement.)

A continue statement must not be placed such that it would transfer control past a declaration used in the targeted continuing statement.

Note: A continue can only be used in a continuing statement if it is used for transferring control flow within another loop nested in the continuing statement. That is, a continue cannot be used to transfer control to the start of the currently executing continuing statement.

EXAMPLE: Invalid continue bypasses declaration
var i: i32 = 0;
loop {
  if i >= 4 { break; }
  if i % 2 == 0 { continue; } // <3>

  let step: i32 = 2;

  continuing {
    i = i + step;
  }
}

7.4.9. Continuing Statement

continuing_statement :

| continuing continuing_compound_statement

continuing_compound_statement :

| brace_left statement * break_if_statement ? brace_right

A continuing statement specifies a compound statement to be executed at the end of a loop iteration. The construct is optional.

The compound statement must not contain a return at any compound statement nesting level.

The compound statement must not contain a discard at any compound statement nesting level nor through function calls. See § 7.7 Statements Behavior Analysis for a more formal description of this rule.

7.4.10. Return Statement

return_statement :

| return expression ?

A return statement ends execution of the current function. If the function is an entry point, then the current shader invocation is terminated. Otherwise, evaluation continues with the next expression or statement after the evaluation of the call site of the current function invocation.

If the function does not have a return type, then the return statement is optional. If the return statement is provided for such a function, it must not supply a value. Otherwise the expression must be present, and is called the return value. In this case the call site of this function invocation evaluates to the return value. The type of the return value must match the return type of the function.

7.4.11. Discard Statement

A discard statement immediately ends execution of a fragment shader invocation and throws away the fragment. The discard statement must only be used in a fragment shader stage.

More precisely, executing a discard statement will:

Only statements executed prior to the discard statement will have observable effects.

Note: A discard statement may be executed by any function in a fragment stage and the effect is the same: immediate termination of the invocation.

After a discard statement is executed, control flow is [non-uniform] for the duration of the entry point.

EXAMPLE: Using the discard statement to throw away a fragment
var<private> will_emit_color: bool = false;

fn discard_if_shallow(pos: vec4<f32>) {
  if pos.z < 0.001 {
    // If this is executed, then the will_emit_color flag will
    // never be set to true.
    discard;
  }
  will_emit_color = true;
}

@fragment
fn main(@builtin(position) coord_in: vec4<f32>)
  -> @location(0) vec4<f32>
{
  discard_if_shallow(coord_in);

  // Set the flag and emit red, but only if the helper function
  // did not execute the discard statement.
  will_emit_color = true;
  return vec4<f32>(1.0, 0.0, 0.0, 1.0);
}

7.5. Function Call Statement

func_call_statement :

| ident argument_expression_list

A function call statement executes a function call.

Note: If the function returns a value, that value is ignored.

7.6. Statements Grammar Summary

The statement rule matches statements that can be used in most places inside a function body.

statement :

| semicolon

| return_statement semicolon

| if_statement

| switch_statement

| loop_statement

| for_statement

| while_statement

| func_call_statement semicolon

| variable_statement semicolon

| break_statement semicolon

| continue_statement semicolon

| discard semicolon

| assignment_statement semicolon

| compound_statement

| increment_statement semicolon

| decrement_statement semicolon

Additionally, certain statements may only be used in very specific contexts:

7.7. Statements Behavior Analysis

7.7.1. Rules

Some statements affecting control-flow are only valid in some contexts. For example, fallthrough is invalid outside of a switch, and continue is invalid outside of a loop, for, or while. Additionally, the uniformity analysis (see § 12.2 Uniformity) needs to know when control flow can exit a statement in multiple different ways.

Both goals are achieved by a system for summarizing execution behaviors of statements and expressions. Behavior analysis maps each statement and expression to the set of possible ways execution proceeds after evaluation of the statement or expression completes. As with type analysis for values and expressions, behavior analysis proceeds bottom up: first determine behaviors for certain basic statements, and then determine behavior for higher level constructs by applying combining rules.

A behavior is a set, whose elements may be:

Each of those correspond to a way to exit a compound statement: either through a keyword, or by falling to the next statement ("Next").

We note "s: B" to say that s respects the rules regarding behaviors, and has behavior B.

For each function:

We assign a behavior to each function: it is its body’s behavior (treating the body as a regular statement), with any "Return" replaced by "Next". As a consequence of the rules above, a function behavior is always one of {}, {Next}, {Discard}, or {Next, Discard}.

Similarly, we assign a behavior to each expression, since expressions can include function calls, which can discard. Like functions, expression behaviors are always one of {}, {Next}, {Discard}, or {Next, Discard}.

Note: There is currently no valid program with an expression that does not have Next in its behavior. The reason is that only functions without a return type can have such a behavior, and there is no compound expression in which such a function can be called.

Rules for analyzing and validating the behaviors of statements
Statement Preconditions Resulting behavior
empty statement {Next}
{s} s: B B
s1 s2

Note: s1 often ends in a semicolon.

s1: B1
Next in B1
s2: B2
(B1∖{Next}) ∪ B2
s1: B1
Next not in B1
s2: B2
B1
var x:T; {Next}
let x = e; e: B B
var x = e; e: B B
x = e; x: B1
e: B2
x is not _
B1B2
_ = e; e: B B
f(e1, ..., en); e1: B1
...
en: Bn
f has behavior B
B ∪ ((B1 ∪ ... ∪ Bn)∖{Next})
return; {Return}
return e; e: B
(B∖{Next}) ∪ {Return}
discard; {Discard}
break; {Break}
break if e; e: B B ∪ {Break}
continue; {Continue}
fallthrough; {Fallthrough}
if e s1 else s2 e: B
s1: B1
s2: B2
(B∖{Next}) ∪ B1B2
loop {s1 continuing {s2}} s1: B1
s2: B2
None of {Continue, Return, Discard} are in B2
Break is not in (B1B2)
(B1B2)∖{Continue, Next}
s1: B1
s2: B2
None of {Continue, Return, Discard} are in B2
Break is in (B1B2)
(B1B2 ∪ {Next})∖{Break, Continue}
switch e {case c1: s1 ... case cn: sn} e: B
s1: B1
...
sn: Bn
Fallthrough is not in Bn
Break is not in (B1 ∪ ... ∪ Bn)
((B∖{Next}) ∪ B1 ∪ ... ∪ Bn)∖{Fallthrough}
e: B
s1: B1
...
sn: Bn
Fallthrough is not in Bn
Break is in (B1 ∪ ... ∪ Bn)
(BB1 ∪ ... ∪ Bn ∪ {Next})∖{Break, Fallthrough}

Note: The empty statement case occurs when a loop has an empty body, or when a for loop lacks an initialization or update statement.

For the purpose of this analysis:

Rules for analyzing and validating the behaviors of expressions
Expression Preconditions Resulting behavior
f(e1, ..., en) e1: B1
...
en: Bn
f has behavior B
B ∪ ((B1 ∪ ... ∪ Bn)∖{Next})
Any literal {Next}
Any variable reference {Next}
e1[e2] e1: B1
e2: B2
B1B2
e.field e: B B
e1 || e2 e1: B1
e2: B2
B1B2
e1 && e2 e1: B1
e2: B2
B1B2

Each built-in function has a behavior of {Next}. And each operator application not listed in the table above has the same behavior as if it were a function call with the same operands and with a function’s behavior of {Next}.

A shader-creation error results if behavior analysis fails:

7.7.2. Notes

This section is informative, non-normative.

Here is the full list of ways that these rules can cause a program to be rejected (this is just restating information already listed above):

This analysis can be run in linear time, by analyzing the call-graph bottom-up (since the behavior of a function call can depend on the function’s code).

7.7.3. Examples

Here are some examples showing this analysis in action:

EXAMPLE: Trivially dead code is allowed
fn simple() -> i32 {
  var a: i32;
  return 0;  // Behavior: {Return}
  a = 1;     // Valid, statically unreachable code.
             //   Statement behavior: {Next}
             //   Overall behavior (due to sequential statements): {Return}
  return 2;  // Valid, statically unreachable code. Behavior: {Return}
} // Function behaviour: {Return}
EXAMPLE: Compound statements are supported
fn nested() -> i32 {
  var a: i32;
  {             // The start of a compound statement.
    a = 2;      // Behavior: {Next}
    return 1;   // Behavior: {Return}
  }             // The compound statement as a whole has behavior {Return}
  a = 1;        // Valid, statically unreachable code.
                //   Statement behavior: {Next}
                //   Overall behavior (due to sequential statements): {Return}
  return 2;     // Valid, statically unreachable code. Behavior: {Return}
}
EXAMPLE: if/then behaves as if there is an empty else
fn if_example() {
  var a: i32 = 0;
  loop {
    if a == 5 {
      break;      // Behavior: {Break}
    }             // Behavior of the whole if compound statement: {Break, Next},
                  //   as the if has an implicit empty else
    a = a + 1;    // Valid, as the previous statement had "Next" in its behavior
  }
}
EXAMPLE: if/then/else has the behavior of both sides
fn if_example() {
  var a: i32 = 0;
  loop {
    if a == 5 {
      break;      // Behavior: {Break}
    } else {
      continue;   // Behavior: {Continue}
    }             // Behavior of the whole if compound statement: {Break, Continue}
    a = a + 1;    // Valid, statically unreachable code.
                  //   Statement behavior: {Next}
                  //   Overall behavior: {Break, Continue}
  }
}
EXAMPLE: if/else if/else behaves like a nested if/else
fn if_example() {
  var a: i32 = 0;
  loop {
    // if e1 s1 else if e2 s2 else s3
    // is identical to
    // if e1 else { if e2 s2 else s3 }
    if a == 5 {
      break;      // Behavior: {Break}
    } else if a == 42 {
      continue;   // Behavior: {Continue}
    } else {
      return;     // Behavior {Return}
    }             // Behavior of the whole if compound statement:
                  //   {Break, Continue, Return}
  }               // Behavior of the whole loop compound statement {Next, Return}
}                 // Behavior of the whole function {Next}
EXAMPLE: Break in switch becomes Next
fn switch_example() {
  var a: i32 = 0;
  switch a {
    default: {
      break;   // Behavior: {Break}
    }
  }            // Behavior: {Next}, as switch replaces Break by Next
  a = 5;       // Valid, as the previous statement had Next in its behavior
}
EXAMPLE: Obviously infinite loops
fn invalid_infinite_loop() {
  loop { }     // Behavior: { }.  Invalid because it’s empty.
}
EXAMPLE: A conditional continue with continuing statement
fn conditional_continue() {
  var a: i32;
  loop {
    if a == 5 { break; } // Behavior: {Break, Next}
    if a % 2 == 1 {      // Valid, as the previous statement has Next in its behavior
      continue;          // Behavior: {Continue}
    }                    // Behavior: {Continue, Next}
    a = a * 2;           // Valid, as the previous statement has Next in its behavior
    continuing {         // Valid as the continuing statement has behavior {Next}
                         //  which does not include any of:
                         //  {Break, Continue, Discard, Return}
      a = a + 1;
    }
  }                      // The loop as a whole has behavior {Next},
                         //  as it absorbs "Continue" and "Next",
                         //  then replaces "Break" with "Next"
}
EXAMPLE: A redundant continue with continuing statement
fn redundant_continue_with_continuing() {
  var a: i32;
  loop {
    if a == 5 { break; }
    continue;   // Valid. This is redundant, branching to the next statement.
    continuing {
      a = a + 1;
    }
  }
}
EXAMPLE: A continue at the end of a loop body
fn continue_end_of_loop_body() {
  for (var i: i32 = 0; i < 5; i++ ) {
    continue;   // Valid. This is redundant,
                //   branching to the end of the loop body.
  }             // Behavior: {Next},
                //   as loops absorb "Continue",
                //   and "for" loops always add "Next"
}
for loops desugar to loop with a conditional break. As shown in a previous example, the conditional break has behavior {Break, Next}, which leads to adding "Next" to the loop’s behavior.
EXAMPLE: Effect of a function that discards unconditionally
fn always_discard() {
  discard;
}                   // The whole function has behavior {Discard}
fn code_after_discard() {
  var a: i32;
  always_discard(); // Behavior: {Discard}
  a = a + 1;        // Valid, statically unreachable code.
                    //   Statement behavior: {Next}
                    //   Overall behavior: {Discard}
}
EXAMPLE: Effect of a function that discards conditionally
fn sometimes_discard(a: i32) {
  if a {
    discard;        // Behavior: {Discard}
  }                 // Behavior: {Next, Discard}
}                   // The whole function has behavior {Next, Discard}
fn code_after_discard() {
  var a: i32;
  a = 42;
  sometimes_discard(a);  // Behavior: {Next, Discard}
  a = a + 1;             // Valid
}                        // The whole function has behavior {Next, Discard}
EXAMPLE: return required in functions that have a return type
fn missing_return () -> i32 {
  var a: i32 = 0;
  if a == 42 {
    return a;       // Behavior: {Return}
  }                 // Behavior: {Next, Return}
}                   // Error: Next is invalid in the body of a
                    //   function with a return type
EXAMPLE: continue must be in a loop
fn continue_out_of_loop () {
  var a: i32 = 0;
  if a > 0  {
    continue;       // Behavior: {Continue}
  }                 // Behavior: {Next, Continue}
}                   // Error: Continue is invalid in the body of a function
The same example would also be invalid for the same reason if continue was replaced by break or fallthrough.

8. Functions

A function performs computational work when invoked.

A function is invoked in one of the following ways:

There are two kinds of functions:

8.1. Declaring a User-defined Function

A function declaration creates a user-defined function, by specifying:

A function declaration must only occur at module scope. A function name is in scope for the entire program.

A formal parameter declaration specifies an identifier name and a type for a value that must be provided when invoking the function. A formal parameter may have attributes. See § 8.2 Function Calls. The identifier is in scope until the end of the function. Two formal parameters for a given function must not have the same name.

Note: Some built-in functions may allow parameters to be abstract numeric types; however, this functionality is not currently supported for user-declared functions.

The return type, if specified, must be constructible.

function_decl :

| attribute * function_header compound_statement

function_header :

| fn ident paren_left param_list ? paren_right ( arrow attribute * type_decl ) ?

param_list :

| ( param comma ) * param comma ?

param :

| attribute * variable_ident_decl

WGSL defines the following attributes that can be applied to function declarations:

WGSL defines the following attributes that can be applied to function parameters and return types:

EXAMPLE: Simple functions
// Declare the add_two function.
// It has two formal paramters, i and b.
// It has a return type of i32.
// It has a body with a return statement.
fn add_two(i: i32, b: f32) -> i32 {
  return i + 2;  // A formal parameter is available for use in the body.
}

// A compute shader entry point function, 'main'.
// It has no specified return type.
// It invokes the ordinary_two function, and captures
// the resulting value in the named value 'six'.
@compute fn main() {
   let six: i32 = add_two(4, 5.0);
}

8.2. Function Calls

A function call is a statement or expression which invokes a function.

The function containing the function call is the calling function, or caller. The function being invoked is the called function, or callee.

The function call:

The function call must supply the same number of argument values as there are formal parameters in the called function. Each argument value must evaluate to the same type as the corresponding formal parameter, by position.

In summary, when calling a function:

  1. Execution of the calling function is suspended.

  2. The called function executes until it returns.

  3. Execution of the calling function resumes.

A called function returns as follows:

In detail, when a function call is executed the following steps occur:

  1. Function call argument values are evaluated. The relative order of evaluation is left-to-right.

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

  3. If the called function is user-defined, memory is allocated for each function scope variable in the called function.

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

  5. If the called function is user-defined, control is transferred to the first statement in its body.

  6. The called function is executed, until it returns.

  7. Control is transferred back to the calling function, and the called function’s execution is unsuspended. If the called function returns a value, that value is supplied for the value of the function call expression.

Note: The current function will not resume execution if the called function or any descendent called function executes a discard statement.

The location of a function call is referred to as a call site. Call sites are a dynamic context. As such, the same textual location may represent multiple call sites.

8.3. Creation-time Functions

A function declared with a const attribute can be evaluated at shader-creation time. These functions are called creation-time functions. Calls to these functions can part of creation-time expressions.

It is a shader-creation error if the function contains any expressions that are not creation-time expressions, or any declarations that are not creation-time constants.

Note: The const attribute cannot be applied to user-declared functions.

EXAMPLE: Creation-time functions
const first_one = firstLeadingBit(1234 + 4567); // Evaluates to 12
                                                // first_one has the type i32, because
                                                // firstLeadingBit cannot operate on
                                                // AbstractInt

@id(1) override x : i32;
override y = firstLeadingBit(x); // Creation-time expressions can be
                                 // used in override expressions.
                                 // firstLeadingBit(x) is not a
                                 // creation-time expression in this context.

fn foo() {
  var a : array<i32, firstLeadingBit(257)>; // Creation-time functions can be used in
                                            // creation-time expressions if all their
                                            // parameters are creation-time expressions.
}

8.4. Restrictions on Functions

Note: Recursion is disallowed because cycles are not permitted among any kinds of declarations.

8.4.1. Aliasing Memory Views

Memory locations can be accessed during the execution of a function using memory views. Within a function, each memory view has a particular root identifier. The root identifier can be an originating variable or a formal parameter of pointer type.

Locally derived expressions of reference or [=pointer type|pointer] type may introduce new names for a particular root identifier, but each expression has a statically determinable root identifier. While the originating variable of a root identifier is a dynamic concept that depends on the call sites for the function, WGSL programs can be statically analyzed to determine the set of all possible originating variables for each root identifier.

Two root identifiers alias when they have the same originating variable. It is a dynamic error if:

Note: This aliasing restriction applies to memory locations written by function calls made in the function.

Note: Originating variables cannot have aliased memory locations. See § 5.2 var Declarations and § 9.3.2 Resource Interface.

EXAMPLE: Aliased memory views
var x : i32 = 0;

fn foo() {
  bar(&x, &x); // Both p and q parameters are aliases of x.
}

// This function produces a dynamic error because of the aliased
// memory accesses.
fn bar(p : ptr<private, i32>, q : ptr<private, i32>) {
  if (x == 0) {
    *p = 1;
  } else {
    *q = 2;
  }
}
EXAMPLE: Aliasing in a helper function
var x : i32 = 0;

fn baz(p : ptr<private, i32>) {
  *p = 2;
}

// This function produces a dynamic error if x == 0, because x is read and
// written through different root identifiers even though the write occurs
// in the scope of baz.
fn bar(p : ptr<private, i32>) {
  if (x == 0) {
    baz(p);
  }
}

fn foo() {
  bar(&x); // p in bar is aliased to x.
}

9. Entry Points

An entry point is a user-defined function that performs the work for a particular shader stage.

9.1. Shader Stages

WebGPU issues work to the GPU in the form of draw or dispatch commands. These commands execute a pipeline in the context of a set of inputs, outputs, and attached resources.

A pipeline describes the work to be performed on the GPU, as a sequence of stages, some of which are programmable. In WebGPU, a pipeline is created before scheduling a draw or dispatch command for execution. There are two kinds of pipelines: GPUComputePipeline, and GPURenderPipeline.

A dispatch command uses a GPUComputePipeline to run a compute shader stage over a logical grid of points with a controllable amount of parallelism, while reading and possibly updating buffer and image resources.

A draw command uses a GPURenderPipeline to run a multi-stage process with two programmable stages among other fixed-function stages:

The WebGPU specification describes pipelines in greater detail.

WGSL defines three shader stages, corresponding to the programmable parts of pipelines:

Each shader stage has its own set of features and constraints, described elsewhere.

9.2. Entry Point Declaration

To create an entry point, declare a user-defined function with a pipeline stage attribute.

When configuring a pipeline in the WebGPU API, the entry point’s function name maps to the entryPoint attribute of the WebGPU § GPUProgrammableStage object.

The entry point’s formal parameters form the stage’s pipeline inputs. The entry point’s return type, if specified, forms the stage’s pipeline output. Each input and output must be an entry point IO type.

Note: Compute entry points never have a return type.

EXAMPLE: Entry Point
@vertex
fn vert_main() -> @builtin(position) vec4<f32> {
  return vec4<f32>(0.0, 0.0, 0.0, 1.0);
}

@fragment
fn frag_main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
  return vec4<f32>(coord_in.x, coord_in.y, 0.0, 1.0);
}

@compute
fn comp_main() { }

The set of functions in a shader stage is the union of:

The union is applied repeatedly until it stabilizes. It will stabilize in a finite number of steps.

9.2.1. Function Attributes for Entry Points

WGSL defines the following attributes that can be applied to entry point declarations:

Can we query upper bounds on workgroup size dimensions? Is it independent of the shader, or a property to be queried after creating the shader module?

EXAMPLE: workgroup_size Attribute
@compute @workgroup_size(8,4,1)
fn sorter() { }

@compute @workgroup_size(8u)
fn reverser() { }

// Using an pipeline-overridable constant.
@id(42) override block_width = 12u;
@compute @workgroup_size(block_width)
fn shuffler() { }

// Error: workgroup_size must be specified on compute shader
@compute
fn bad_shader() { }

9.3. Shader Interface

The shader interface is the set of objects through which the shader accesses data external to the shader stage, either for reading or writing. The interface includes:

These objects are represented by module-scope variables in certain address spaces.

When an identifier used in a function declaration resolves to a module-scope variable, then we say the variable is statically accessed by the function. Static access of a let-declared constant is defined similarly. Note that being statically accessed is independent of whether an execution of the shader will actually evaluate the expression referring to the variable, or even execute the statement that may enclose the expression.

More precisely, the interface of a shader stage consists of:

9.3.1. Pipeline Input and Output Interface

The entry point IO types include the following:

A pipeline input is data provided to the shader stage from upstream in the pipeline. A pipeline input is denoted by the arguments of the entry point.

A pipeline output is data the shader provides for further processing downstream in the pipeline. A pipeline output is denoted by the return type of the entry point.

Each pipeline input or output is one of:

9.3.1.1. Built-in Inputs and Outputs

A built-in input value provides access to system-generated control information. The set of built-in inputs are listed in § 15 Built-in Values.

A built-in input for stage S with name X and type TX is accessed via a formal parameter to an entry point for shader stage S, in one of two ways:

  1. The parameter has attribute builtin(X) and is of type TX.

  2. The parameter has structure type, where one of the structure members has attribute builtin(X) and is of type TX.

Conversely, when a parameter or member of a parameter for an entry point has a builtin attribute, the corresponding builtin must be an input for the entry point’s shader stage.

A built-in output value is used by the shader to convey control information to later processing steps in the pipeline. The set of built-in outputs are listed in § 15 Built-in Values.

A built-in output for stage S with name Y and type TY is set via the return value for an entry point for shader stage S, in one of two ways:

  1. The entry point return type has attribute builtin(Y) and is of type TY.

  2. The entry point return type has structure type, where one of the structure members has attribute builtin(Y) and is of type TY.

Conversely, when the return type or member of a return type for an entry point has a builtin attribute, the corresponding builtin must be an output for the entry point’s shader stage.

Note: The position built-in is both an output of a vertex shader, and an input to the fragement shader.

9.3.1.2. User-defined Inputs and Outputs

User-defined data can be passed as input to the start of a pipeline, passed between stages of a pipeline or output from the end of a pipeline. User-defined IO must not be passed to compute shader entry points. User-defined IO must be of numeric scalar or numeric vector type, or of a structure type whose members are numeric scalars or vectors. All user-defined IO must be assigned locations (See § 9.3.1.4 Input-output Locations).

9.3.1.3. Interpolation

Authors can control how user-defined IO data is interpolated through the use of the interpolate attribute. WGSL offers two aspects of interpolation to control: the type of interpolation, and the sampling of the interpolation.

The interpolation type must be one of:

The interpolation sampling must be one of:

For user-defined IO of scalar or vector floating-point type:

User-defined IO of scalar or vector integer type must always be specified as @interpolate(flat).

Interpolation attributes must match between vertex outputs and fragment inputs with the same location assignment within the same pipeline.

9.3.1.4. Input-output Locations

Each location can store a value up to 16 bytes in size. The byte size of a type is defined using the SizeOf column in § 4.4.7.1 Alignment and Size. For example, a four-component vector of floating-point values occupies a single location.

Locations are specified via the location attribute.

Every user-defined input and output must have a fully specified set of locations. Each structure member in the entry point IO must be one of either a built-in value (see § 9.3.1.1 Built-in Inputs and Outputs), or assigned a location.

Locations must not overlap within each of the following sets:

Note: Location numbering is distinct between inputs and outputs: Location numbers for an entry point’s pipeline inputs do not conflict with location numbers for the entry point’s pipeline outputs.

Note: No additional rule is required to prevent location overlap within an entry point’s outputs. When the output is a structure, the first rule above prevents overlap. Otherwise, the output is a scalar or a vector, and can have only a single location assigned to it.

Note: The number of available locations for an entry point is defined by the WebGPU API.

EXAMPLE: Applying location attributes
struct A {
  @location(0) x: f32,
  // Despite locations being 16-bytes, x and y cannot share a location
  @location(1) y: f32
}

// in1 occupies locations 0 and 1.
// in2 occupies location 2.
// The return value occupies location 0.
@fragment
fn fragShader(in1: A, @location(2) in2: f32) -> @location(0) vec4<f32> {
 // ...
}

User-defined IO can be mixed with built-in values in the same structure. For example,

EXAMPLE: Mixing builtins and user-defined IO
// Mixed builtins and user-defined inputs.
struct MyInputs {
  @location(0) x: vec4<f32>,
  @builtin(front_facing) y: bool,
  @location(1) @interpolate(flat) z: u32
}

struct MyOutputs {
  @builtin(frag_depth) x: f32,
  @location(0) y: vec4<f32>
}

@fragment
fn fragShader(in1: MyInputs) -> MyOutputs {
  // ...
}
EXAMPLE: Invalid location assignments
struct A {
  @location(0) x: f32,
  // Invalid, x and y cannot share a location.
  @location(0) y: f32
}

struct B {
  @location(0) x: f32
}

struct C {
  // Invalid, structures with user-defined IO cannot be nested.
  b: B
}

struct D {
  x: vec4<f32>
}

@fragment
// Invalid, location cannot be applied to a structure type.
fn fragShader1(@location(0) in1: D) {
  // ...
}

@fragment
// Invalid, in1 and in2 cannot share a location.
fn fragShader2(@location(0) in1: f32, @location(0) in2: f32) {
  // ...
}

@fragment
// Invalid, location cannot be applied to a structure.
fn fragShader3(@location(0) in1: vec4<f32>) -> @location(0) D {
  // ...
}

9.3.2. Resource Interface

A resource is an object, other than a pipeline input or output, which provides access to data external to a shader stage. Resources are shared by all invocations of the shader.

There are four kinds of resources:

The resource interface of a shader is the set of module-scope resource variables statically accessed by functions in the shader stage.

Each resource variable must be declared with both group and binding attributes. Together with the shader’s stage, these identify the binding address of the resource on the shader’s pipeline. See WebGPU § GPUPipelineLayout.

Bindings must not alias within a shader stage: two different variables in the resource interface of a given shader must not have the same group and binding values, when considered as a pair of values.

9.3.3. Resource Layout Compatibility

WebGPU requires that a shader’s resource interface match the layout of the pipeline using the shader.

Each WGSL variable in a resource interface must be bound to a WebGPU resource with a compatible resource type and binding type, where compatibility is defined by the following table.

WebGPU binding type compatibility
WGSL resource WebGPU
Resource type
WebGPU Binding type
uniform buffer GPUBufferBinding GPUBufferBindingType uniform
storage buffer with read_write access storage
storage buffer with read access read-only-storage
sampler GPUSampler GPUSamplerBindingType filtering
non-filtering
sampler_comparison comparison
sampled texture GPUTextureView GPUTextureSampleType float
unfilterable-float
sint
uint
depth
write-only storage texture GPUTextureView GPUStorageTextureAccess write-only

See the WebGPU API specification for interface validation requirements.

10. Language Extensions

The WGSL language is expected to evolve over time.

An extension is a named grouping for a coherent set of modifications to a particular version of the WGSL specification, consisting of any combination of:

Hypothetically, extensions could be used to:

10.1. Enable Directive

An enable directive indicates that the functionality described by a particular named extension may be used. The grammar rules imply that all enable directives must appear before any declarations.

The directive uses an identifier, keyword, or reserved word to name the extension. The valid extension name are listed in § 10.2 Extensions list.

If the name is an identifier, the directive does not create a scope for the identifier. Use of the identifier by the directive does not conflict with the use of that identifier as the name in any declaration.

enable_directive :

| enable ident semicolon

Note: The grammar rule includes the terminating semicolon token, ensuring the additional functionality is usable only after that semicolon. Therefore any WGSL implementation can parse the entire enable directive. When an implementation encounters an enable directive for an unsupported extension, the implementation can issue a clear diagnostic.

EXAMPLE: Using hypothetical extensions
// Enable a hypothetical extension for arbitrary precision floating point types.
enable aribtrary_precision_float;
enable arbitrary_precision_float; // A redundant enable directive is ok.

// Enable a hypothetical extension to control the rounding mode.
enable rounding_mode;

// Assuming arbitrary_precision_float enables use of:
//    - a type f<E,M>
//    - as a type in function return, formal parameters and let-declarations
//    - as a type constructor from AbstractFloat
//    - operands to division operator: /
// Assuming @rounding_mode attribute is enabled by the rounding_mode enable directive.
@rounding_mode(round_to_even)
fn halve_it(x : f<8, 7>) -> f<8, 7> {
  let two = f<8, 7>(2);
  return x / 2; // uses round to even rounding mode.
}

10.2. Extensions list

Extension identifier
Identifier WebGPU extension name Description
f16 "shader-f16" Keyword f16 and any floating point literal with a h suffix is valid if and only if this extension is enabled. Otherwise, using f16 keyword or any floating point literal with a h suffix will result in a shader-creation error.

11. WGSL Program

A WGSL program is a sequence of optional directives followed by module scope declarations.

translation_unit :

| global_directive * global_decl *

global_decl :

| semicolon

| global_variable_decl semicolon

| global_constant_decl semicolon

| type_alias_decl semicolon

| struct_decl

| function_decl

11.1. Limits

A program must satisfy the following limits:

Quantifiable shader complexity limits
Limit Maximum value
Number of members in a structure type 16383
Nesting depth of a composite type 255
Number of parameters for a function 255
Number of case selector values in a switch statement 16383

12. Execution

§ 1.1 Technical Overview describes how a shader is invoked and partitioned into invocations. This section describes further constraints on how invocations execute, individually and collectively.

12.1. Program Order Within an Invocation

Each statement in a WGSL program may be executed zero or more times during execution. For a given invocation, each execution of a given statement represents a unique dynamic statement instance.

When a statement includes an expression, the statement’s semantics determines:

Expression nesting defines data dependencies which must be satisfied to complete evaluation. That is, a nested expression must be evaluated before the enclosing expression can be evaluated. The order of evaluation for operands of an expression is left-to-right in WGSL. For example, foo() + bar() must evaluate foo() before bar(). See § 6 Expressions.

Statements in a WGSL program are executed in control flow order. See § 7 Statements and § 8.2 Function Calls.

12.2. Uniformity

12.2.1. Terminology and concepts

The following definitions are merely informative, trying to give an intuition for what the analysis in the next subsection is computing. The analysis is what actually defines these concepts, and when a program is valid or breaks the uniformity rules.

For a given group of invocations:

12.2.2. Uniformity analysis overview

Some functions (e.g. barriers and derivatives) are only safe to call in uniform control flow. In this section we specify an analysis that verifies that these functions are only called in such a context.

Note: This analysis has the following desirable properties:

The analysis analyzes each function, verifying that there is a context where it is safe to call this function. It rejects the program as invalid if there is no such context.

At the same time, it computes metadata about the function to help analyze its callers in turn. This means that the call graph must first be built, and functions must be analyzed from the leaves upwards, i.e. from functions that call no function outside the standard library toward the entry point. This way, whenever a function is analyzed, the metadata for all of its callees has already been computed. There is no risk of being trapped in a cycle, as recurrence is forbidden in the language.

Note: another way of saying the same thing is that we do a topological sort of functions ordered by the "is a (possibly indirect) callee of" partial order, and analyze them in that order.

12.2.3. Analyzing the uniformity requirements of a function

Each function is analyzed in two phases.

The first phase walks over the syntax of the function, building a directed graph along the way based on the rules in the following subsections. The second phase explores that graph, resulting in either rejecting the program, or computing the constraints on calling this function.

Note: apart from two special nodes RequiredToBeUniform and MayBeNonUniform, all nodes can be understood as having one of the following meanings:

An edge can be understood as an implication from the statement corresponding to its source node to the statement corresponding to its target node.

To express that something must always be uniform (e.g. the control flow at the call site of a derivative), we add an edge from RequiredToBeUniform to the corresponding node. One way to understand this, is that RequiredToBeUniform corresponds to the proposition True, so that RequiredToBeUniform -> X is the same as saying that X is true.

Reciprocally, to express that we cannot ensure the uniformity of something (e.g. a variable which holds the thread id), we add an edge from the corresponding node to MayBeNonUniform. One way to understand this, is that MayBeNonUniform corresponds to the proposition False, so that X -> MayBeNonUniform is the same as saying that X is false.

A consequence of this interpretation is that every node reachable from RequiredToBeUniform corresponds to something which must be uniform for the program to be valid, and every node from which MayBeNonUniform is reachable corresponds to something whose uniformity we cannot guarantee. It follows that we have a uniformity violation (and thus reject the program) if there is any path from RequiredToBeUniform to MayBeNonUniform.

For each function, two tags are computed:

Additionally, for each formal parameter of a function, a parameter tag is computed that describes the uniformity requirement of the parameter value.

Call site tag values
Call Site Tag Description
CallSiteRequiredToBeUniform The function must only be called from uniform control flow.
CallSiteNoRestriction The function may be called from non-uniform control flow.
Function tag values
Function Tag Description
SubsequentControlFlowMayBeNonUniform Calling this function may cause control flow to be non-uniform immediately after the call site.
ReturnValueMayBeNonUniform The return value of the function may be non-uniform.
NoRestriction The function does not introduce non-uniformity.
Parameter tag values
Parameter Tag Description
ParameterRequiredToBeUniform The parameter must be a uniform value.
ParameterRequiredToBeUniformForSubsequentControlFlow The parameter must be a uniform value for control flow after the function call to be uniform.
ParameterRequiredToBeUniformForReturnValue The parameter must be a uniform value in order for the return value to be a uniform value.
ParameterNoRestriction The parameter value has no uniformity requirement.

The following algorithm describes how to compute these tags for a given function:

Note: The entire graph can be destroyed at this point. The tags listed above are all that we need to remember to analyze callers of this function.

12.2.4. Uniformity rules for statements

The rules for analyzing statements take as argument both the statement itself and the node corresponding to control flow at the beginning of it (which we’ll note "CF" below) and return both of the following:

In the table below, (CF1, S) => CF2 means "run the analysis on S starting with control flow CF1, apply the required changes to the graph, and name the resulting control flow CF2". Similarly, (CF1, E) => (CF2, V) means "run the analysis on expression E, starting with control flow CF1, apply the required changes to the graph, and name the resulting control flow node CF2 and the resulting value node V" (see next section for the analysis of expressions).

We have a similar set of rules for expressions in left-value positions, that we denote by LValue: (CF, E) => (CF, L). Instead of computing the node which corresponds to the uniformity of the value, it computes the node which corresponds to the uniformity of the variable we are addressing.

When several edges have to be created we use X -> {Y, Z} as a short-hand for X -> Y, X -> Z.

Uniformity rules for statements
Statement New nodes Recursive analyses Resulting control flow node New edges
{s} (CF, s) => CF' CF'
s1 s2,
with Next in behavior of s1

Note: s1 often ends in a semicolon.

(CF, s1) => CF1
(CF1, s2) => CF2
CF2
s1 s2,
without Next in behavior of s1

Note: s1 often ends in a semicolon.

(CF, s1) => CF1

Note: s2 is statically unreachable and not recursively analyzed. s2 does not contribute to the uniformity analysis.

CF1
if e s1 else s2
with behavior {Next}
(CF, e) => (CF', V)
(V, s1) => CF1
(V, s2) => CF2
CF
if e s1 else s2
with another behavior
CFend CFend CFend -> {CF1, CF2}
loop {s1 continuing {s2}}
with behavior {Next}
CF' (CF', s1) => CF1
(CF1, s2) => CF2
CF CF' -> {CF2, CF}
loop {s1 continuing {s2}}
with another behavior
CF'
loop {s1}
with behavior {Next}
CF' (CF', s1) => CF1 CF CF' -> {CF1, CF}
loop {s1}
with another behavior
CF'
switch e case _: s_1 .. case _: s_n
with behavior {Next}
(CF, e) => (CF', V)
(V, s_1) => CF_1
...
if s_(n-1) may fallthrough, (CF_(n-1), s_n) => CF_n
else (V, s_n) => CF_n
CF
switch e case _: s_1 .. case _: s_n
with another behavior
CFend CFend CFend -> {CF_1, ..., CF_n}
var x: T; CF
break;
break if e; (CF, e) => (CF', V) CF'
continue; CF
fallthrough;
discard;
return; CF CF_return -> CF
return e; (CF, e) => (CF', V) CF' CF_return -> CF'
Value_return -> V
e2 = e1; (CF, e1) => (CF1, V1)
LValue: (CF1, e2) => (CF2, L2)
CF2 L2 -> V1
_ = e (CF, e) => (CF', V) CF'
let x = e; (CF, e) => (CF', V) CF'
var x = e;

Analysis of for and while loops follows from their respective desugaring translations to loop statements.

In switch, a default block is treated exactly like a case block with regards to uniformity.

Note: If the set of behaviors (see § 7.7 Statements Behavior Analysis) for an if, switch, or loop statement is {Next}, this means that we either did not diverge within the statement, or we reconverged, so we pick the node corresponding to control flow at the start of the statement as the node corresponding to control flow at the exit of the statement.

12.2.5. Uniformity rules for function calls

The most complex rule is for function calls:

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

Most built-in functions have tags of:

Here is the list of exceptions:

12.2.6. Uniformity rules for expressions

The rules for analyzing expressions take as argument both the expression itself and the node corresponding to control flow at the beginning of it (which we’ll note "CF" below) and return the following:

Uniformity rules for expressions (in normal rvalue position)
Expression New nodes Recursive analyses Resulting control flow node, value node New edges
e1 || e2 (CF, e1) => (CF1, V1)
(V1, e2) => (CF2, V2)
CF2, V2
e1 && e2
Literal CF, CF
reference to function-scope variable, creation-time constant, let-declaration, or non-built-in parameter "x" Result X is the node corresponding to "x" CF, Result Result -> {CF, X}
reference to uniform built-in value "x" CF, CF
reference to non-uniform built-in value "x" CF, MayBeNonUniform
reference to read-only module-scope variable "x" CF, CF
reference to non-read-only module-scope variable "x" CF, MayBeNonUniform
op e,
where op is a unary operator
(CF, e) => (CF', V) CF', V
e.field
e1 op e2,
where op is a non-short-circuiting binary operator
Result (CF, e1) => (CF1, V1)
(CF1, e2) => (CF2, V2)
CF2, Result Result -> {V1, V2}
e1[e2]

The following built-in input variables are considered uniform:

All other ones (see § 15 Built-in Values) are considered non-uniform.

Uniformity rules for expressions in lvalue positions
Expression New nodes Recursive analyses Resulting control flow node, variable node New edges
reference to function-scope variable, creation-time constant, let-declaration, or parameter "x" X is the node corresponding to "x" CF, X
reference to module-scope variable "x" CF, MayBeNonUniform
e.field LValue: (CF, e) => (CF1, L1) CF1, L1
e1[e2] LValue: (CF, e1) => (CF1, L1)
(CF1, e2) => (CF2, V2)
CF2, L1 L1 -> V2

12.2.7. Annotating the uniformity of every point in the control-flow

This entire subsection is non-normative.

If implementers want to provide developers with a diagnostic mode that shows for each point in the control-flow of the entire shader whether it is uniform or not (and thus whether it would be valid to call a function that requires uniformity there), we suggest the following:

Any node which is not visited by these reachability analyses can be proven to be uniform by the analysis (and so it would be safe to call a derivative or similar function there).

Note: The bottom-up analysis is still required, as it lets us know what edges to add to the graphs when encountering calls.

12.3. Compute Shaders and Workgroups

A workgroup is a set of invocations which concurrently execute a compute shader stage entry point, and share access to shader variables in the workgroup address space.

The workgroup grid for a compute shader is the set of points with integer coordinates (i,j,k) with:

where (workgroup_size_x, workgroup_size_y, workgroup_size_z) is the value specified for the workgroup_size attribute of the entry point.

There is exactly one invocation in a workgroup for each point in the workgroup grid.

An invocation’s local invocation ID is the coordinate triple for the invocation’s corresponding workgroup grid point.

When an invocation has local invocation ID (i,j,k), then its local invocation index is

i + (j * workgroup_size_x) + (k * workgroup_size_x * workgroup_size_y)

Note that if a workgroup has W invocations, then each invocation I the workgroup has a unique local invocation index L(I) such that 0 ≤ L(I) < W, and that entire range is covered.

A compute shader begins execution when a WebGPU implementation removes a dispatch command from a queue and begins the specified work on the GPU. The dispatch command specifies a dispatch size, which is an integer triple (group_count_x, group_count_y, group_count_z) indicating the number of workgroups to be executed, as described in the following.

The compute shader grid for a particular dispatch is the set of points with integer coordinates (CSi,CSj,CSk) with:

where workgroup_size_x, workgroup_size_y, and workgroup_size_z are as above for the compute shader entry point.

The work to be performed by a compute shader dispatch is to execute exactly one invocation of the entry point for each point in the compute shader grid.

An invocation’s global invocation ID is the coordinate triple for the invocation’s corresponding compute shader grid point.

The invocations are organized into workgroups, so that each invocation (CSi, CSj, CSk) is identified with the workgroup grid point

( CSi mod workgroup_size_x , CSj mod workgroup_size_y , CSk mod workgroup_size_z )

in workgroup ID

( ⌊ CSi ÷ workgroup_size_x ⌋, ⌊ CSj ÷ workgroup_size_y ⌋, ⌊ CSk ÷ workgroup_size_z ⌋).

WebGPU provides no guarantees about:

12.4. Collective Operations

12.4.1. Barriers

A barrier is a synchronization built-in function that orders memory operations in a program. A control barrier is executed by all invocations in the same workgroup as if it were executed concurrently. As such, control barriers must only be executed in uniform control flow in a compute shader.

12.4.2. Derivatives

A partial derivative is the rate of change of a value along an axis.

Fragment shader invocations operating on neighbouring fragments (in screen-space coordinates) collaborate to compute approximate partial derivatives. These neighbouring fragments are referred to as a quad.

Partial derivatives of the fragment coordinate are computed implicitly as part of operation of the following built-in functions:

For these, the derivatives help determine the mip levels of texels to be sampled, or in the case of textureSampleCompare, sampled and compared against a reference value.

Partial derivatives of invocation-specified values are computed by the built-in functions described in § 16.7 Derivative Built-in Functions:

Because neighbouring invocations must collaborate to compute derivatives, these functions must only be invoked in uniform control flow in a fragment shader.

12.5. Floating Point Evaluation

WGSL follows the IEEE-754 standard for floating point computation with the following exceptions:

12.5.1. Floating Point Accuracy

Let x be the exact real-valued or infinite result of an operation when computed with unbounded precision. The correctly rounded result of the operation for floating point type T is:

That is, the result may be rounded up or down: WGSL does not specify a rounding mode.

Note: Floating point types include positive and negative infinity, so the correctly rounded result may be finite or infinite.

The units in the last place, ULP, for a floating point number x is the minimum distance between two non-equal floating point numbers a and b such that axb (i.e. ulp(x) = mina,b|b - a|).

In the following tables, the accuracy of an operation is provided among five possibilities:

For any accuracy values specified over a range, the accuracy is undefined for results outside that range.

If an allowable return value for any operation is greater in magnitude than the largest representable finite floating-point value, then that operation may additionally return either the infinity with the same sign or the largest finite value with the same sign.

Accuracy of expressions
Expression Accuracy for f32 Accuracy for f16
x + y Correctly rounded
x - y Correctly rounded
x * y Correctly rounded
x / y 2.5 ULP for |y| in the range [2-126, 2126] 2.5 ULP for |y| in the range [2-14, 214]
x % y Derived from x - y * trunc(x/y)
-x Correctly rounded
x == y Correct result
x != y Correct result
x < y Correct result
x <= y Correct result
x > y Correct result
x >= y Correct result
Accuracy of built-in functions
Built-in Function Accuracy for f32 Accuracy for f16
abs(x) Correctly rounded
acos(x) Inherited from atan2(sqrt(1.0 - x * x), x)
acosh(x) Inherited from log(x + sqrt(x * x - 1.0))
asin(x) Inherited from atan2(x, sqrt(1.0 - x * x))
asinh(x) Inherited from log(x + sqrt(x * x + 1.0))
atan(x) 4096 ULP 5 ULP
atan2(y, x) 4096 ULP 5 ULP
atanh(x) Inherited from log( (1.0 + x) / (1.0 - x) ) * 0.5
ceil(x) Correctly rounded
clamp(x,low,high) Correctly rounded
cos(x) Absolute error ≤ 2-11 inside the range of [-π, π] Absolute error ≤ 2-7 inside the range of [-π, π]
cosh(x) Inherited from (exp(x) - exp(-x)) * 0.5
cross(x, y) Inherited from (x[i] * y[j] - x[j] * y[i])
degrees(x) Inherited from x * 57.295779513082322865
distance(x, y) Inherited from length(x - y)
exp(x) 3 + 2 * |x| ULP 1 + 2 * |x| ULP
exp2(x) 3 + 2 * |x| ULP 1 + 2 * |x| ULP
faceForward(x, y, z) Inherited from select(-x, x, dot(z, y) < 0.0)
floor(x) Correctly rounded
fma(x, y, z) Inherited from x * y + z
fract(x) Correctly rounded
frexp(x) Correctly rounded
inverseSqrt(x) 2 ULP
ldexp(x, y) Correctly rounded
length(x) Inherited from sqrt(dot(x, x))
log(x) 3 ULP outside the range [0.5, 2.0].
Absolute error < 2-21 inside the range [0.5, 2.0]
3 ULP outside the range [0.5, 2.0].
Absolute error < 2-7 inside the range [0.5, 2.0]
log2(x) 3 ULP outside the range [0.5, 2.0].
Absolute error < 2-21 inside the range [0.5, 2.0]
3 ULP outside the range [0.5, 2.0].
Absolute error < 2-7 inside the range [0.5, 2.0]
max(x, y) Correctly rounded
min(x, y) Correctly rounded
mix(x, y, z) Inherited from x * (1.0 - z) + y * z
modf(x) Correctly rounded
normalize(x) Inherited from x / length(x)
pow(x, y) Inherited from exp2(y * log2(x))
radians(x) Inherited from x * 0.017453292519943295474
reflect(x, y) Inherited from x - 2.0 * dot(x, y) * y
refract(x, y, z) Inherited from z * x - (z * dot(y, x) + sqrt(k)) * y,
where k = 1.0 - z * z * (1.0 - dot(y, x) * dot(y, x))
If k < 0.0 the result is precisely 0.0
round(x) Correctly rounded
sign(x) Correctly rounded
sin(x) Absolute error ≤ 2-11 inside the range [-π, π] Absolute error ≤ 2-7 inside the range [-π, π]
sinh(x) Inherited from (exp(x) - exp(-x)) * 0.5
smoothstep(low, high, x) Inherited from t * t * (3.0 - 2.0 * t),
where t = clamp((x - low) / (high - low), 0.0, 1.0)
sqrt(x) Inherited from 1.0 / inverseSqrt(x)
step(edge, x) Correctly rounded
tan(x) Inherited from sin(x) / cos(x)
tanh(x) Inherited from sinh(x) / cosh(x)
trunc(x) Correctly rounded

Reassociation is the reordering of operations in an expression such that the answer is the same if computed exactly. For example:

However, the result may not be the same when computed in floating point. The reassociated result may be inaccurate due to approximation, or may trigger an overflow or NaN when computing intermediate results.

An implementation may reassociate operations.

An implementation may fuse operations if the transformed expression is at least as accurate as the original formulation. For example, some fused multiply-add implementations can be more accurate than performing a multiply followed by an addition.

12.5.2. Floating Point Conversion

In this section, a floating point type may be any of:

Note: Recall that the f32 WGSL type corresponds to the IEEE-754 binary32 format, and the f16 WGSL type corresponds to the IEEE-754 binary16 format.

When converting a floating point scalar value to an integral type:

Note: In other words, floating point to integer conversion rounds toward zero, then saturates.

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

When converting a value to a floating point type:

NOTE: An integer value may lie between two adjacent representable floating point values. In particular, the f32 type uses 23 explicit fractional bits. Additionally, when the floating point value is in the normal range (the exponent is neither extreme value), then the mantissa is the set of fractional bits together with an extra 1-bit at the most significant position at bit position 23. Then, for example, integers 228 and 1+228 both map to the same floating point value: the difference in the least significant 1 bit is not representable by the floating point format. This kind of collision occurs for pairs of adjacent integers with a magnitude of at least 225.

Check behaviour of the f32 to f16 conversion for numbers just beyond the max normal f16 values. I’ve written what an NVIDIA GPU does. See https://github.com/google/amber/pull/918 for an executable test case.

13. Memory Model

In general, WGSL follows the Vulkan Memory Model. The remainder of this section describes how WGSL programs map to the Vulkan Memory Model.

Note: The Vulkan Memory Model is a textual version of a formal Alloy model.

13.1. Memory Operation

In WGSL, a read access is equivalent to a memory read operation in the Vulkan Memory Model. A WGSL, a write access is equivalent to a memory write operation in the Vulkan Memory Model.

A read access occurs when an invocation executes one of the following:

A write access occurs when an invocation executes one of the following:

Atomic read-modify-write built-in functions perform a single memory operation that is both a read access and a write access.

Read and write accesses do not occur under any other circumstances. Read and write accesses are collectively known as memory operations in the Vulkan Memory Model.

A memory operation accesses exactly the set of locations associated with the particular memory view used in the operation. For example, a memory read that accesses a u32 from a struct containing multiple members, only reads the memory locations associated with that u32 member.

EXAMPLE: Accessing memory locations
struct S {
  a : f32,
  b : u32,
  c : f32
}

@group(0) @binding(0)
var<storage> v : S;

fn foo() {
  let x = v.b; // Does not access memory locations for v.a or v.c.
}

13.2. Memory Model Reference

Each module-scope variable in WGSL forms a unique memory model reference for the lifetime of a given entry point. Each function-scope variable in WGSL forms a unique memory model reference for the lifetime of the variable.

13.3. Scoped Operations

When an invocation performs a scoped operation, it will affect one or two sets of invocations. These sets are the memory scope and the execution scope. The memory scope specifies the set of invocations that will see any updates to memory contents affected by the operation. For synchronization built-in functions, this also means that all affected memory operations program ordered before the function are visible to affected operations program ordered after the function. The execution scope specifies the set of invocations which may participate in an operation (see § 12.4 Collective Operations).

Atomic built-in functions map to atomic operations whose memory scope is:

Synchronization built-in functions map to control barriers whose execution and memory scopes are Workgroup.

Implicit and explicit derivatives have an implicit quad execution scope.

Note: If the Vulkan memory model is not enabled in generated shaders, Device scope should be used instead of QueueFamily.

13.4. Memory Semantics

All Atomic built-in functions use Relaxed memory semantics and, thus, no address space semantics.

workgroupBarrier uses AcquireRelease memory semantics and WorkgroupMemory semantics. storageBarrier uses AcquireRelease memory semantics and UniformMemory semantics.

Note: A combined workgroupBarrier and storageBarrier uses AcquireRelease ordering semantics and both WorkgroupMemory and UniformMemory memory semantics.

Note: No atomic or synchronization built-in functions use MakeAvailable or MakeVisible semantics.

13.5. Private vs Non-private

All non-atomic read accesses in the storage or workgroup address spaces are considered non-private and correspond to read operations with NonPrivatePointer | MakePointerVisible memory operands with the Workgroup scope.

All non-atomic write accesses in the storage or workgroup address spaces are considered non-private and correspond to write operations with NonPrivatePointer | MakePointerAvailable memory operands with the Workgroup scope.

https://github.com/gpuweb/gpuweb/issues/1621

14. Keyword and Token Summary

14.1. Keyword Summary

14.1.1. Type-defining Keywords

array :

| 'array'

atomic :

| 'atomic'

bool :

| 'bool'

float32 :

| 'f32'

float16 :

| 'f16'

int32 :

| 'i32'

mat2x2 :

| 'mat2x2'

mat2x3 :

| 'mat2x3'

mat2x4 :

| 'mat2x4'

mat3x2 :

| 'mat3x2'

mat3x3 :

| 'mat3x3'

mat3x4 :

| 'mat3x4'

mat4x2 :

| 'mat4x2'

mat4x3 :

| 'mat4x3'

mat4x4 :

| 'mat4x4'

override :

| 'override'

pointer :

| 'ptr'

sampler :

| 'sampler'

sampler_comparison :

| 'sampler_comparison'

struct :

| 'struct'

texture_1d :

| 'texture_1d'

texture_2d :

| 'texture_2d'

texture_2d_array :

| 'texture_2d_array'

texture_3d :

| 'texture_3d'

texture_cube :

| 'texture_cube'

texture_cube_array :

| 'texture_cube_array'

texture_multisampled_2d :

| 'texture_multisampled_2d'

texture_storage_1d :

| 'texture_storage_1d'

texture_storage_2d :

| 'texture_storage_2d'

texture_storage_2d_array :

| 'texture_storage_2d_array'

texture_storage_3d :

| 'texture_storage_3d'

texture_depth_2d :

| 'texture_depth_2d'

texture_depth_2d_array :

| 'texture_depth_2d_array'

texture_depth_cube :

| 'texture_depth_cube'

texture_depth_cube_array :

| 'texture_depth_cube_array'

texture_depth_multisampled_2d :

| 'texture_depth_multisampled_2d'

uint32 :

| 'u32'

vec2 :

| 'vec2'

vec3 :

| 'vec3'

vec4 :

| 'vec4'

14.1.2. Other Keywords

bitcast :

| 'bitcast'

break :

| 'break'

case :

| 'case'

const :

| 'const'

continue :

| 'continue'

continuing :

| 'continuing'

default :

| 'default'

discard :

| 'discard'

else :

| 'else'

enable :

| 'enable'

fallthrough :

| 'fallthrough'

false :

| 'false'

fn :

| 'fn'

for :

| 'for'

function :

| 'function'

if :

| 'if'

let :

| 'let'

loop :

| 'loop'

private :

| 'private'

return :

| 'return'

storage :

| 'storage'

switch :

| 'switch'

true :

| 'true'

type :

| 'type'

uniform :

| 'uniform'

var :

| 'var'

while :

| 'while'

workgroup :

| 'workgroup'

14.2. Reserved Words

A reserved word is a token which is reserved for future use. A WGSL program must not contain a reserved word.

The following are reserved words:

_reserved :

| 'AppendStructuredBuffer'

| 'BlendState'

| 'Buffer'

| 'ByteAddressBuffer'

| 'CompileShader'

| 'ComputeShader'

| 'ConsumeStructuredBuffer'

| 'DepthStencilState'

| 'DepthStencilView'

| 'DomainShader'

| 'GeometryShader'

| 'Hullshader'

| 'InputPatch'

| 'LineStream'

| 'NULL'

| 'OutputPatch'

| 'PixelShader'

| 'PointStream'

| 'RWBuffer'

| 'RWByteAddressBuffer'

| 'RWStructuredBuffer'

| 'RWTexture1D'

| 'RWTexture1DArray'

| 'RWTexture2D'

| 'RWTexture2DArray'

| 'RWTexture3D'

| 'RasterizerState'

| 'RenderTargetView'

| 'SamplerComparisonState'

| 'SamplerState'

| 'Self'

| 'StructuredBuffer'

| 'Texture1D'

| 'Texture1DArray'

| 'Texture2D'

| 'Texture2DArray'

| 'Texture2DMS'

| 'Texture2DMSArray'

| 'Texture3D'

| 'TextureCube'

| 'TextureCubeArray'

| 'TriangleStream'

| 'VertexShader'

| 'abstract'

| 'active'

| 'alignas'

| 'alignof'

| 'as'

| 'asm'

| 'asm_fragment'

| 'async'

| 'atomic_uint'

| 'attribute'

| 'auto'

| 'await'

| 'become'

| 'bf16'

| 'buffer'

| 'cast'

| 'catch'

| 'cbuffer'

| 'centroid'

| 'char'

| 'class'

| 'co_await'

| 'co_return'

| 'co_yield'

| 'coherent'

| 'column_major'

| 'common'

| 'compile'

| 'compile_fragment'

| 'concept'

| 'const_cast'

| 'consteval'

| 'constexpr'

| 'constinit'

| 'crate'

| 'debugger'

| 'decltype'

| 'delete'

| 'demote'

| 'demote_to_helper'

| 'do'

| 'dword'

| 'dynamic_cast'

| 'enum'

| 'explicit'

| 'export'

| 'extends'

| 'extern'

| 'external'

| 'f64'

| 'filter'

| 'final'

| 'finally'

| 'fixed'

| 'flat'

| 'friend'

| 'from'

| 'fvec2'

| 'fvec3'

| 'fvec4'

| 'fxgroup'

| 'get'

| 'goto'

| 'groupshared'

| 'handle'

| 'highp'

| 'hvec2'

| 'hvec3'

| 'hvec4'

| 'i16'

| 'i64'

| 'i8'

| 'iimage1D'

| 'iimage1DArray'

| 'iimage2D'

| 'iimage2DArray'

| 'iimage2DMS'

| 'iimage2DMSArray'

| 'iimage2DRect'

| 'iimage3D'

| 'iimageBuffer'

| 'iimageCube'

| 'iimageCubeArray'

| 'image1D'

| 'image1DArray'

| 'image2D'

| 'image2DArray'

| 'image2DMS'

| 'image2DMSArray'

| 'image2DRect'

| 'image3D'

| 'imageBuffer'

| 'imageCube'

| 'imageCubeArray'

| 'impl'

| 'implements'

| 'import'

| 'in'

| 'inline'

| 'inout'

| 'input'

| 'instanceof'

| 'interface'

| 'invariant'

| 'isampler1D'

| 'isampler1DArray'

| 'isampler2D'

| 'isampler2DArray'

| 'isampler2DMS'

| 'isampler2DMSArray'

| 'isampler2DRect'

| 'isampler3D'

| 'isamplerBuffer'

| 'isamplerCube'

| 'isamplerCubeArray'

| 'isubpassInput'

| 'isubpassInputMS'

| 'itexture1D'

| 'itexture1DArray'

| 'itexture2D'

| 'itexture2DArray'

| 'itexture2DMS'

| 'itexture2DMSArray'

| 'itexture2DRect'

| 'itexture3D'

| 'itextureBuffer'

| 'itextureCube'

| 'itextureCubeArray'

| 'layout'

| 'line'

| 'lineadj'

| 'linear'

| 'lowp'

| 'macro'

| 'macro_rules'

| 'mat'

| 'match'

| 'matrix'

| 'mediump'

| 'meta'

| 'mod'

| 'module'

| 'move'

| 'mut'

| 'mutable'

| 'namespace'

| 'new'

| 'nil'

| 'noexcept'

| 'noinline'

| 'nointerpolation'

| 'noperspective'

| 'null'

| 'nullptr'

| 'of'

| 'operator'

| 'out'

| 'output'

| 'package'

| 'packoffset'

| 'partition'

| 'pass'

| 'patch'

| 'pixelfragment'

| 'point'

| 'precise'

| 'precision'

| 'premerge'

| 'priv'

| 'protected'

| 'pub'

| 'public'

| 'readonly'

| 'ref'

| 'regardless'

| 'register'

| 'reinterpret_cast'

| 'requires'

| 'resource'

| 'restrict'

| 'row_major'

| 'samper'

| 'sample'

| 'sampler1D'

| 'sampler1DArray'

| 'sampler1DArrayShadow'

| 'sampler1DShadow'

| 'sampler2D'

| 'sampler2DArray'

| 'sampler2DArrayShadow'

| 'sampler2DMS'

| 'sampler2DMSArray'

| 'sampler2DRect'

| 'sampler2DRectShadow'

| 'sampler2DShadow'

| 'sampler3D'

| 'sampler3DRect'

| 'samplerBuffer'

| 'samplerCube'

| 'samplerCubeArray'

| 'samplerCubeArrayShadow'

| 'samplerCubeShadow'

| 'samplerShadow'

| 'self'

| 'set'

| 'shared'

| 'signed'

| 'sizeof'

| 'smooth'

| 'snorm'

| 'stateblock'

| 'stateblock_state'

| 'static'

| 'static_assert'

| 'static_cast'

| 'std'

| 'string'

| 'subpassInput'

| 'subpassInputMS'

| 'subroutine'

| 'super'

| 'superp'

| 'target'

| 'tbuffer'

| 'technique'

| 'technique10'

| 'technique11'

| 'template'

| 'texture'

| 'texture1D'

| 'texture1DArray'

| 'texture2D'

| 'texture2DArray'

| 'texture2DMS'

| 'texture2DMSArray'

| 'texture2DRect'

| 'texture3D'

| 'textureBuffer'

| 'textureCube'

| 'textureCubeArray'

| 'this'

| 'thread_local'

| 'throw'

| 'trait'

| 'triangle'

| 'triangleadj'

| 'try'

| 'typedef'

| 'typeid'

| 'typename'

| 'typeof'

| 'u16'

| 'u64'

| 'u8'

| 'uimage1D'

| 'uimage1DArray'

| 'uimage2D'

| 'uimage2DArray'

| 'uimage2DMS'

| 'uimage2DMSArray'

| 'uimage2DRect'

| 'uimage3D'

| 'uimageBuffer'

| 'uimageCube'

| 'uimageCubeArray'

| 'union'

| 'unless'

| 'unorm'

| 'unsafe'

| 'unsigned'

| 'unsized'

| 'usampler1D'

| 'usampler1DArray'

| 'usampler2D'

| 'usampler2DArray'

| 'usampler2DMS'

| 'usampler2DMSArray'

| 'usampler2DRect'

| 'usampler3D'

| 'usamplerBuffer'

| 'usamplerCube'

| 'usamplerCubeArray'

| 'use'

| 'using'

| 'usubpassInput'

| 'usubpassInputMS'

| 'utexture1D'

| 'utexture1DArray'

| 'utexture2D'

| 'utexture2DArray'

| 'utexture2DMS'

| 'utexture2DMSArray'

| 'utexture2DRect'

| 'utexture3D'

| 'utextureBuffer'

| 'utextureCube'

| 'utextureCubeArray'

| 'varying'

| 'vec'

| 'vector'

| 'vertexfragment'

| 'virtual'

| 'void'

| 'volatile'

| 'wchar_t'

| 'wgsl'

| 'where'

| 'with'

| 'writeonly'

| 'yield'

14.3. Syntactic Tokens

A syntactic token is a sequence of special code points, used:

and :

| '&' (Code point: U+0026)

and_and :

| '&&' (Code points: U+0026 U+0026)

arrow :

| '->' (Code points: U+002D U+003E)

attr :

| '@' (Code point: U+0040)

forward_slash :

| '/' (Code point: U+002F)

bang :

| '!' (Code point: U+0021)

bracket_left :

| '[' (Code point: U+005B)

bracket_right :

| ']' (Code point: U+005D)

brace_left :

| '{' (Code point: U+007B)

brace_right :

| '}' (Code point: U+007D)

colon :

| ':' (Code point: U+003A)

comma :

| ',' (Code point: U+002C)

equal :

| '=' (Code point: U+003D)

equal_equal :

| '==' (Code points: U+003D U+003D)

not_equal :

| '!=' (Code points: U+0021 U+003D)

greater_than :

| '>' (Code point: U+003E)

greater_than_equal :

| '>=' (Code points: U+003E U+003D)

less_than :

| '<' (Code point: U+003C)

less_than_equal :

| '<=' (Code points: U+003C U+003D)

modulo :

| '%' (Code point: U+0025)

minus :

| '-' (Code point: U+002D)

minus_minus :

| '--' (Code points: U+002D U+002D)

period :

| '.' (Code point: U+002E)

plus :

| '+' (Code point: U+002B)

plus_plus :

| '++' (Code points: U+002B U+002B)

or :

| '|' (Code point: U+007C)

or_or :

| '||' (Code points: U+007C U+007C)

paren_left :

| '(' (Code point: U+0028)

paren_right :

| ')' (Code point: U+0029)

semicolon :

| ';' (Code point: U+003B)

star :

| '*' (Code point: U+002A)

tilde :

| '~' (Code point: U+007E)

underscore :

| '_' (Code point: U+005F)

xor :

| '^' (Code point: U+005E)

plus_equal :

| '+=' (Code points: U+002B U+003D)

minus_equal :

| '-=' (Code points: U+002D U+003D)

times_equal :

| '*=' (Code points: U+002A U+003D)

division_equal :

| '/=' (Code points: U+002F U+003D)

modulo_equal :

| '%=' (Code points: U+0025 U+003D)

and_equal :

| '&=' (Code points: U+0026 U+003D)

or_equal :

| '|=' (Code points: U+007C U+003D)

xor_equal :

| '^=' (Code points: U+005E U+003D)

15. Built-in Values

The following table lists the available built-in input values and built-in output values.

See § 9.3.1.1 Built-in Inputs and Outputs for how to declare a built-in value.

Built-in input and output values
Name Stage Input or Output Type Description
vertex_index vertex input u32 Index of the current vertex within the current API-level draw command, independent of draw instancing.

For a non-indexed draw, the first vertex has an index equal to the firstVertex argument of the draw, whether provided directly or indirectly. The index is incremented by one for each additional vertex in the draw instance.

For an indexed draw, the index is equal to the index buffer entry for the vertex, plus the baseVertex argument of the draw, whether provided directly or indirectly.

instance_index vertex input u32 Instance index of the current vertex within the current API-level draw command.

The first instance has an index equal to the firstInstance argument of the draw, whether provided directly or indirectly. The index is incremented by one for each additional instance in the draw.

position vertex output vec4<f32> Output position of the current vertex, using homogeneous coordinates. After homogeneous normalization (where each of the x, y, and z components are divided by the w component), the position is in the WebGPU normalized device coordinate space. See WebGPU § Coordinate Systems.
position fragment input vec4<f32> Framebuffer position of the current fragment in framebuffer space. (The x, y, and z components have already been scaled such that w is now 1.) See WebGPU § Coordinate Systems.
front_facing fragment input bool True when the current fragment is on a front-facing primitive. False otherwise. See WebGPU § Front-facing.
frag_depth fragment output f32 Updated depth of the fragment, in the viewport depth range. See WebGPU § Coordinate Systems.
local_invocation_id compute input vec3<u32> The current invocation’s local invocation ID, i.e. its position in the workgroup grid.
local_invocation_index compute input u32 The current invocation’s local invocation index, a linearized index of the invocation’s position within the workgroup grid.
global_invocation_id compute input vec3<u32> The current invocation’s global invocation ID, i.e. its position in the compute shader grid.
workgroup_id compute input vec3<u32> The current invocation’s workgroup ID, i.e. the position of the workgroup in the workgroup grid.
num_workgroups compute input vec3<u32> The dispatch size, vec<u32>(group_count_x, group_count_y, group_count_z), of the compute shader dispatched by the API.
sample_index fragment input u32 Sample index for the current fragment. The value is least 0 and at most sampleCount-1, where sampleCount is the number of MSAA samples specified for the GPU render pipeline.
See WebGPU § GPURenderPipeline.
sample_mask fragment input u32 Sample coverage mask for the current fragment. It contains a bitmask indicating which samples in this fragment are covered by the primitive being rendered.
See WebGPU § Sample Masking.
sample_mask fragment output u32 Sample coverage mask control for the current fragment. The last value written to this variable becomes the shader-output mask. Zero bits in the written value will cause corresponding samples in the color attachments to be discarded.
See WebGPU § Sample Masking.
EXAMPLE: Declaring built-in values
 struct VertexOutput {
   @builtin(position) my_pos: vec4<f32>
 }

 @vertex
 fn vs_main(
   @builtin(vertex_index) my_index: u32,
   @builtin(instance_index) my_inst_index: u32,
 ) -> VertexOutput {}

 struct FragmentOutput {
   @builtin(frag_depth) depth: f32,
   @builtin(sample_mask) mask_out: u32
 }

 @fragment
 fn fs_main(
   @builtin(front_facing) is_front: bool,
   @builtin(position) coord: vec4<f32>,
   @builtin(sample_index) my_sample_index: u32,
   @builtin(sample_mask) mask_in: u32,
 ) -> FragmentOutput {}

 @compute
 fn cs_main(
   @builtin(local_invocation_id) local_id: vec3<u32>,
   @builtin(local_invocation_index) local_index: u32,
   @builtin(global_invocation_id) global_id: vec3<u32>,
) {}

16. Built-in Functions

Certain functions are predeclared, provided by the implementation, and therefore always available for use in a WGSL program. These are called built-in functions.

A built-in function is a family of functions, all with the same name, but distinguished by the number, order, and types of their formal parameters. Each of these distinct function variations is an overload.

Note: Each user-defined function only has one overload.

Each overload is described below via:

Since a built-in function is always in scope, it is an error to attempt to redefine one or to use the name of a built-in function as an identifier for any other module-scope declaration.

When calling a built-in function, all arguments to the function are evaluated before function evaluation begins. See § 8.2 Function Calls.

16.1. Logical Built-in Functions

Parameterization Overload Description
@const fn
all(e: vecN<bool>) -> bool
Returns true if each component of e is true.
e: bool @const fn
all(e) -> bool
Returns e.
@const fn
any(e: vecN<bool>) -> bool
Returns true if any component of e is true.
e: bool @const fn
any(e) -> bool
Returns e.
T is scalar, abstract numeric type, or vector @const fn
select(f: T, t: T, cond: bool) -> T
Returns t when cond is true, and f otherwise.
T is scalar or abstract numeric type @const fn
select(f: vecN<T>, t: vecN<T>, cond: vecN<bool>) -> vecN<T>
Component-wise selection. Result component i is evaluated as select(f[i],t[i],cond[i]).

16.2. Array Built-in Functions

Parameterization Overload Description
fn arrayLength(e: ptr<storage,array<T>> ) -> u32 Returns the number of elements in the runtime-sized array.

16.3. Float Built-in Functions

Parameterization Overload Description
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
abs(e: T ) -> T
Returns the absolute value of e (e.g. e with a positive sign bit). Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
acos(e: T ) -> T
Returns the arc cosine of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
acosh(e: T ) -> T
Returns the hyperbolic arc cosine of e. The result is 0 when e < 1.
Computes the non-negative functional inverse of cosh.
Component-wise when T is a vector.

Note: The result is not mathematically meaningful when e < 1.

T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
asin(e: T ) -> T
Returns the arc sine of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
asinh(e: T ) -> T
Returns the hyperbolic arc sine of e.
Computes the functional inverse of sinh.
Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
atan(e: T ) -> T
Returns the arc tangent of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
atanh(e: T ) -> T
Returns the hyperbolic arc tangent of e. The result is 0 when abs(e) ≥ 1.
Computes the functional inverse of tanh.
Component-wise when T is a vector.

Note: The result is not mathematically meaningful when abs(e) ≥ 1.

T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
atan2(e1: T ,e2: T ) -> T
Returns the arc tangent of e1 over e2. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
ceil(e: T ) -> T
Returns the ceiling of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> clamp(e: T , low: T , high: T) -> T Returns either min(max(e,low),high), or the median of the three values e, low, high. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
cos(e: T ) -> T
Returns the cosine of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
cosh(e: T ) -> T
Returns the hyperbolic cosine of e. Component-wise when T is a vector
T is AbstractFloat, f32, or f16 @const fn
cross(e1: vec3<T> ,e2: vec3<T>) -> vec3<T>
Returns the cross product of e1 and e2.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
degrees(e1: T ) -> T
Converts radians to degrees, approximating e1 × 180 ÷ π. Component-wise when T is a vector
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
distance(e1: T ,e2: T ) -> f32
Returns the distance between e1 and e2 (e.g. length(e1-e2)).
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
exp(e1: T ) -> T
Returns the natural exponentiation of e1 (e.g. ee1). Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
exp2(e: T ) -> T
Returns 2 raised to the power e (e.g. 2e). Component-wise when T is a vector.
T is vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
faceForward(e1: T ,e2: T ,e3: T ) -> T
Returns e1 if dot(e2,e3) is negative, and -e1 otherwise.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
floor(e: T ) -> T
Returns the floor of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
fma(e1: T ,e2: T ,e3: T ) -> T
Returns e1 * e2 + e3. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
fract(e: T ) -> T
Returns the fractional part of e, computed as e - floor(e).
Component-wise when T is a vector.
T is f32 @const fn
frexp(e:T) -> __frexp_result
Splits e into a significand and exponent of the form significand * 2exponent. Returns the __frexp_result built-in structure, defined as follows:
struct __frexp_result {
  sig : f32, // significand part
  exp : i32  // exponent part
}
The magnitude of the significand is in the range of [0.5, 1.0) or 0.

Note: A value cannot be explicitly declared with the type __frexp_result, but a value may infer the type.

EXAMPLE: frexp usage
// Infers result type
let sig_and_exp = frexp(1.5);
// Sets fraction_direct to 0.75
let fraction_direct = frexp(1.5).sig;
T is f16 @const fn frexp(e:T) -> __frexp_result_f16
Splits e into a significand and exponent of the form significand * 2exponent. Returns the __frexp_result_f16 built-in structure, defined as if as follows:
struct __frexp_result_f16 {
  sig : f16, // significand part
  exp : i32  // exponent part
}
The magnitude of the significand is in the range of [0.5, 1.0) or 0.

Note: A value cannot be explicitly declared with the type __frexp_result_f16, but a value may infer the type.

T is vecN<f32> @const fn
frexp(e:T) -> __frexp_result_vecN
Splits the components of e into a significand and exponent of the form significand * 2exponent. Returns the __frexp_result_vecN built-in structure, defined as follows:
struct __frexp_result_vecN {
  sig : vecN<f32>, // significand part
  exp : vecN<i32>  // exponent part
}
The magnitude of each component of the significand is in the range of [0.5, 1.0) or 0.

Note: A value cannot be explicitly declared with the type __frexp_result_vecN, but a value may infer the type.

T is vecN<f16> @const fn frexp(e:T) -> __frexp_result_vecN_f16
Splits the components of e into a significand and exponent of the form significand * 2exponent. Returns the __frexp_result_vecN built-in structure, defined as if as follows:
struct __frexp_result_vecN_f16 {
  sig : vecN<f16>, // significand part
  exp : vecN<i32>  // exponent part
}
The magnitude of each component of the significand is in the range of [0.5, 1.0) or 0.

Note: A value cannot be explicitly declared with the type __frexp_result_vecN_f16, but a value may infer the type.

T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
inverseSqrt(e: T ) -> T
Returns the reciprocal of sqrt(e). Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16>
I is AbstractInt, i32, vecN<AbstractInt>, or vecN<i32>, where
I is a scalar if T is a scalar, or
a vector when T is a vector
@const fn
ldexp(e1: T ,e2: I ) -> T
Returns e1 * 2e2. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
length(e: T ) -> f32
Returns the length of e (e.g. abs(e) if T is a scalar, or sqrt(e[0]2 + e[1]2 + ...) if T is a vector).
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
log(e: T ) -> T
Returns the natural logarithm of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
log2(e: T ) -> T
Returns the base-2 logarithm of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
max(e1: T ,e2: T ) -> T
Returns e2 if e1 is less than e2, and e1 otherwise. If one operand is a NaN, the other is returned. If both operands are NaNs, a NaN is returned. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
min(e1: T ,e2: T ) -> T
Returns e2 if e2 is less than e1, and e1 otherwise. If one operand is a NaN, the other is returned. If both operands are NaNs, a NaN is returned. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
mix(e1: T ,e2: T ,e3: T) -> T
Returns the linear blend of e1 and e2 (e.g. e1*(1-e3)+e2*e3). Component-wise when T is a vector.
T is AbstractFloat, f32, or f16
T2 is vecN<T>
@const fn
mix(e1: T2 ,e2: T2 ,e3: f32 ) -> T
Returns the component-wise linear blend of e1 and e2, using scalar blending factor e3 for each component.
Same as mix(e1,e2,T2(e3)).
T is f32 @const fn
modf(e:T) -> __modf_result
Splits e into fractional and whole number parts. Returns the __modf_result built-in structure, defined as follows:
struct __modf_result {
  fract : f32, // fractional part
  whole : f32  // whole part
}

Note: A value cannot be explicitly declared with the type __modf_result, but a value may infer the type.

EXAMPLE: modf usage
// Infers result type
let fract_and_whole = modf(1.5);
// Sets fract_direct to 0.5
let fract_direct = modf(1.5).fract;
T is f16 @const fn modf(e:T) -> __modf_result_f16
Splits e into fractional and whole number parts. Returns the __modf_result_f16 built-in structure, defined as if as follows:
struct __modf_result_f16 {
  fract : f16, // fractional part
  whole : f16  // whole part
}

Note: A value cannot be explicitly declared with the type __modf_result_f16, but a value may infer the type.

T is vecN<f32> @const fn
modf(e:T) -> __modf_result_vecN
Splits the components of e into fractional and whole number parts. Returns the __modf_result_vecN built-in structure, defined as follows:
struct __modf_result_vecN {
  fract : vecN<f32>, // fractional part
  whole : vecN<f32>  // whole part
}

Note: A value cannot be explicitly declared with the type __modf_result_vecN, but a value may infer the type.

T is vecN<f16> @const fn modf(e:T) -> __modf_result_vecN_f16
Splits the components of e into fractional and whole number parts. Returns the __modf_result_vecN_f16 built-in structure, defined as if as follows:
struct __modf_result_vecN_f16 {
  fract : vecN<f16>, // fractional part
  whole : vecN<f16>  // whole part
}

Note: A value cannot be explicitly declared with the type __modf_result_vecN_f16, but a value may infer the type.

T is AbstractFloat, f32, or f16 @const fn
normalize(e: vecN<T> ) -> vecN<T>
Returns a unit vector in the same direction as e.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
pow(e1: T ,e2: T ) -> T
Returns e1 raised to the power e2. Component-wise when T is a vector.
T is f32 or vecN<f32> @const fn quantizeToF16(e: T ) -> T Quantizes a 32-bit floating point value e as if e were converted to a IEEE 754 binary16 value, and then converted back to a IEEE 754 binary32 value.
See § 12.5.2 Floating Point Conversion.
Component-wise when T is a vector.

Note: The vec2<f32> case is the same as unpack2x16float(pack2x16float(e)).

T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
radians(e1: T ) -> T
Converts degrees to radians, approximating e1 × π ÷ 180. Component-wise when T is a vector
T is vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
reflect(e1: T ,e2: T ) -> T
For the incident vector e1 and surface orientation e2, returns the reflection direction e1-2*dot(e2,e1)*e2.
T is vecN<I>
I is AbstractFloat, f32, or f16
@const fn
refract(e1: T ,e2: T ,e3: I ) -> T
For the incident vector e1 and surface normal e2, and the ratio of indices of refraction e3, let k = 1.0 -e3*e3* (1.0 - dot(e2,e1) * dot(e2,e1)). If k < 0.0, returns the refraction vector 0.0, otherwise return the refraction vector e3*e1- (e3* dot(e2,e1) + sqrt(k)) *e2.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
round(e: T ) -> T
Result is the integer k nearest to e, as a floating point value.
When e lies halfway between integers k and k+1, the result is k when k is even, and k+1 when k is odd.
Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
sign(e: T ) -> T
Returns the sign of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
sin(e: T ) -> T
Returns the sine of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
sinh(e: T ) -> T
Returns the hyperbolic sine of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
smoothstep(low: T , high: T , x: T ) -> T
Returns the smooth Hermite interpolation between 0 and 1. Component-wise when T is a vector.

For scalar T, the result is t * t * (3.0 - 2.0 * t), where t = clamp((x - low) / (high - low), 0.0, 1.0).

T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
sqrt(e: T ) -> T
Returns the square root of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
step(edge: T ,x: T ) -> T
Returns 1.0 if edgex, and 0.0 otherwise. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
tan(e: T ) -> T
Returns the tangent of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
tanh(e: T ) -> T
Returns the hyperbolic tangent of e. Component-wise when T is a vector.
T is AbstractFloat, f32, f16, vecN<AbstractFloat>, vecN<f32>, or vecN<f16> @const fn
trunc(e: T ) -> T
Returns the nearest whole number whose absolute value is less than or equal to e. Component-wise when T is a vector.

16.4. Integer Built-in Functions

Parameterization Overload Description
S is AbstractInt, i32, or u32
T is S or vecN<S>
@const fn
abs(e: T ) -> T
The absolute value of e. Component-wise when T is a vector. If e is a signed integral scalar type and evaluates to the largest negative value, then the result is e. If e is an unsigned integral type, then the result is e.
S is AbstractInt, i32, or u32
T is S or vecN<S>
@const fn
clamp(e: T , low: T, high: T) -> T
Returns min(max(e,low),high). Component-wise when T is a vector.
T is i32, u32, vecN<i32>, or vecN<u32> @const fn
countLeadingZeros(e: T ) -> T
The number of consecutive 0 bits starting from the most significant bit of e, when T is a scalar type.
Component-wise when T is a vector.
Also known as "clz" in some languages.
T is i32, u32, vecN<i32>, or vecN<u32> @const fn
countOneBits(e: T ) -> T
The number of 1 bits in the representation of e.
Also known as "population count".
Component-wise when T is a vector.
T is i32, u32, vecN<i32>, or vecN<u32> @const fn
countTrailingZeros(e: T ) -> T
The number of consecutive 0 bits starting from the least significant bit of e, when T is a scalar type.
Component-wise when T is a vector.
Also known as "ctz" in some languages.
T is i32 or vecN<i32> @const fn
firstLeadingBit(e: T ) -> T
For scalar T, the result is:
  • -1 if e is 0 or -1.
  • Otherwise the position of the most significant bit in e that is different from e’s sign bit.

Note: Since signed integers use twos-complement representation, the sign bit appears in the most significant bit position.

Component-wise when T is a vector.

T is u32 or vecN<u32> @const fn
firstLeadingBit(e: T ) -> T
For scalar T, the result is:
  • T(-1) if e is zero.
  • Otherwise the position of the most significant 1 bit in e.
Component-wise when T is a vector.
T is i32, u32, vecN<i32>, or vecN<u32> @const fn
firstTrailingBit(e: T ) -> T
For scalar T, the result is:
  • T(-1) if e is zero.
  • Otherwise the position of the least significant 1 bit in e.
Component-wise when T is a vector.
T is i32 or vecN<i32> @const fn
extractBits(
 e: T,
 offset: u32,
 count: u32) -> T
Reads bits from an integer, with sign extension.

When T is a scalar type, then:

  • w is the bit width of T
  • o = min(offset,w)
  • c = min(count, w - o)
  • The result is 0 if c is 0.
  • Otherwise, bits 0..c-1 of the result are copied from bits o..o+c-1 of e. Other bits of the result are the same as bit c-1 of the result.
Component-wise when T is a vector.
T is u32 or vecN<u32> @const fn
extractBits(
 e: T,
 offset: u32,
 count: u32) -> T
Reads bits from an integer, without sign extension.

When T is a scalar type, then:

  • w is the bit width of T
  • o = min(offset,w)
  • c = min(count, w - o)
  • The result is 0 if c is 0.
  • Otherwise, bits 0..c-1 of the result are copied from bits o..o+c-1 of e. Other bits of the result are 0.
Component-wise when T is a vector.
T is i32, u32, vecN<i32>, or vecN<u32> @const fn
insertBits(
 e: T,
 newbits:T,
 offset: u32,
 count: u32) -> T
Sets bits in an integer.

When T is a scalar type, then:

  • w is the bit width of T
  • o = min(offset,w)
  • c = min(count, w - o)
  • The result is e if c is 0.
  • Otherwise, bits o..o+c-1 of the result are copied from bits 0..c-1 of newbits. Other bits of the result are copied from e.
Component-wise when T is a vector.
S is AbstractInt, i32, or u32
T is S or vecN<S>
@const fn
max(e1: T ,e2: T) -> T
Returns e2 if e1 is less than e2, and e1 otherwise. Component-wise when T is a vector.
S is AbstractInt, i32, or u32
T is S or vecN<S>
@const fn
min(e1: T ,e2: T) -> T
Returns e1 if e1 is less than e2, and e2 otherwise. Component-wise when T is a vector.
T is i32, u32, vecN<i32>, or vecN<u32> @const fn
reverseBits(e: T ) -> T
Reverses the bits in e: The bit at position k of the result equals the bit at position 31-k of e.
Component-wise when T is a vector.
S is AbstractInt, i32, or u32
T is S or vecN<S>
TS is AbstractInt or u32 if T is scalar, or
vecN<AbstractInt>, or vecN<u32> otherwise
@const fn shiftLeft(e1: T, e2: TS) -> T Logical shift left.
Shift e1 left, inserting zero bits at the least significant positions, and discarding the most significant bits.

The number of bits to shift is the value of e2. If e1 has a concrete type, the shift value is modulo the bit width of e1.
Component-wise when T is a vector.

If e2 is an AbstractInt or vecN<AbstractInt>, it is a shader-creation error if any of the values are less than 0.

S is AbstractInt, i32, or u32
T is S or vecN<S>>br> TS is AbstractInt or u32 if T is scalar, or
vecN<AbstractInt>, or vecN<u32> otherwise
@const fn shiftRight(e1: T, e2: TS) -> T Logical shift right.
If e1 is signed, shift e1 right, inserting zero bits at the most significant positions, and discarding the least significant bits. If e1 is unsigned, shift e1 right, copying the sign bit of e1 into the most significant positions, and discarding the least significant bits.

The number of bits to shift is the value of e2. If e1 has a concrete type, the shift value is modulo the bit width of e1.
Component-wise when T is a vector.

If e2 is an AbstractInt or vecN<AbstractInt>, it is a shader-creation error if any of the values are less than 0.

16.5. Matrix Built-in Functions

Parameterization Overload Description
T is AbstractFloat, f32, or f16 determinant(e: matCxC<T> ) -> T Returns the determinant of e.
T is AbstractFloat, f32, or f16 transpose(e: matRxC<T> ) -> matCxR<T> Returns the transpose of e.

16.6. Vector Built-in Functions

Parameterization Overload Description
T is AbstractInt, AbstractFloat, i32, u32, f32, or f16 @const fn
dot(e1: vecN<T>,e2: vecN<T>) -> T
Returns the dot product of e1 and e2.

16.7. Derivative Built-in Functions

See § 12.4.2 Derivatives.

These functions:

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

16.8. Texture Built-in Functions

In this section, texture types are shown with the following parameters:

Parameter values must be valid for the respective texture types.

16.8.1. textureDimensions

Returns the dimensions of a texture, or texture’s mip level in texels.

fn textureDimensions(t: texture_1d<T>) -> i32
fn textureDimensions(t: texture_1d<T>, level: i32) -> i32
fn textureDimensions(t: texture_2d<T>) -> vec2<i32>
fn textureDimensions(t: texture_2d<T>, level: i32) -> vec2<i32>
fn textureDimensions(t: texture_2d_array<T>) -> vec2<i32>
fn textureDimensions(t: texture_2d_array<T>, level: i32) -> vec2<i32>
fn textureDimensions(t: texture_3d<T>) -> vec3<i32>
fn textureDimensions(t: texture_3d<T>, level: i32) -> vec3<i32>
fn textureDimensions(t: texture_cube<T>) -> vec2<i32>
fn textureDimensions(t: texture_cube<T>, level: i32) -> vec2<i32>
fn textureDimensions(t: texture_cube_array<T>) -> vec2<i32>
fn textureDimensions(t: texture_cube_array<T>, level: i32) -> vec2<i32>
fn textureDimensions(t: texture_multisampled_2d<T>)-> vec2<i32>
fn textureDimensions(t: texture_depth_2d) -> vec2<i32>
fn textureDimensions(t: texture_depth_2d, level: i32) -> vec2<i32>
fn textureDimensions(t: texture_depth_2d_array) -> vec2<i32>
fn textureDimensions(t: texture_depth_2d_array, level: i32) -> vec2<i32>
fn textureDimensions(t: texture_depth_cube) -> vec2<i32>
fn textureDimensions(t: texture_depth_cube, level: i32) -> vec2<i32>
fn textureDimensions(t: texture_depth_cube_array) -> vec2<i32>
fn textureDimensions(t: texture_depth_cube_array, level: i32) -> vec2<i32>
fn textureDimensions(t: texture_depth_multisampled_2d)-> vec2<i32>
fn textureDimensions(t: texture_storage_1d<F,A>) -> i32
fn textureDimensions(t: texture_storage_2d<F,A>) -> vec2<i32>
fn textureDimensions(t: texture_storage_2d_array<F,A>) -> vec2<i32>
fn textureDimensions(t: texture_storage_3d<F,A>) -> vec3<i32>
fn textureDimensions(t: texture_external) -> vec2<i32>

Parameters:

t The sampled, multisampled, depth, storage, or external texture.
level The mip level, with level 0 containing a full size version of the texture.
If omitted, the dimensions of level 0 are returned.

Returns:

The dimensions of the texture in texels.

For textures based on cubes, the results are the dimensions of each face of the cube. Cube faces are square, so the x and y components of the result are equal.

If level is outside the range [0, textureNumLevels(t)) then any valid value for the return type may be returned.

16.8.2. textureGather

A texture gather operation reads from a 2D, 2D array, cube, or cube array texture, computing a four-component vector as follows:

TODO: The four texels are the "sample footprint" that should be described by the WebGPU spec. https://github.com/gpuweb/gpuweb/issues/2343

fn textureGather(component: i32, t: texture_2d<T>, s: sampler, coords: vec2<f32>) -> vec4<T>
fn textureGather(component: i32, t: texture_2d<T>, s: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<T>
fn textureGather(component: i32, t: texture_2d_array<T>, s: sampler, coords: vec2<f32>, array_index: i32) -> vec4<T>
fn textureGather(component: i32, t: texture_2d_array<T>, s: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<T>
fn textureGather(component: i32, t: texture_cube<T>, s: sampler, coords: vec3<f32>) -> vec4<T>
fn textureGather(component: i32, t: texture_cube_array<T>, s: sampler, coords: vec3<f32>, array_index: i32) -> vec4<T>
fn textureGather(t: texture_depth_2d, s: sampler, coords: vec2<f32>) -> vec4<f32>
fn textureGather(t: texture_depth_2d, s: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<f32>
fn textureGather(t: texture_depth_2d_array, s: sampler, coords: vec2<f32>, array_index: i32) -> vec4<f32>
fn textureGather(t: texture_depth_2d_array, s: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<f32>
fn textureGather(t: texture_depth_cube, s: sampler, coords: vec3<f32>) -> vec4<f32>
fn textureGather(t: texture_depth_cube_array, s: sampler, coords: vec3<f32>, array_index: i32) -> vec4<f32>

Parameters:

component Only applies to non-depth textures.
The index of the channel to read from the selected texels.
When provided, the component expression must a creation-time expression (e.g. 1).
Its value must be at least 0 and at most 3. Values outside of this range will result in a shader-creation error.
t The sampled or depth texture to read from.
s The sampler type.
coords The texture coordinates.
array_index The 0-based texture array index.
offset The optional texel offset applied to the unnormalized texture coordinate before sampling the texture. This offset is applied before applying any texture wrapping modes.
The offset expression must be a creation-time expression (e.g. vec2<i32>(1, 2)).
Each offset component must be at least -8 and at most 7. Values outside of this range will result in a shader-creation error.

Returns:

A four component vector with components extracted from the specified channel from the selected texels, as described above.

EXAMPLE: Gather components from texels in 2D texture
@group(0) @binding(0) var t: texture_2d<f32>;
@group(0) @binding(1) var dt: texture_depth_2d;
@group(0) @binding(2) var s: sampler;

fn gather_x_components(c: vec2<f32>) -> vec4<f32> {
  return textureGather(0,t,s,c);
}
fn gather_y_components(c: vec2<f32>) -> vec4<f32> {
  return textureGather(1,t,s,c);
}
fn gather_z_components(c: vec2<f32>) -> vec4<f32> {
  return textureGather(2,t,s,c);
}
fn gather_depth_components(c: vec2<f32>) -> vec4<f32> {
  return textureGather(dt,s,c);
}

16.8.3. textureGatherCompare

A texture gather compare operation performs a depth comparison on four texels in a depth texture and collects the results into a single vector, as follows:

fn textureGatherCompare(t: texture_depth_2d, s: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> vec4<f32>
fn textureGatherCompare(t: texture_depth_2d, s: sampler_comparison, coords: vec2<f32>, depth_ref: f32, offset: vec2<i32>) -> vec4<f32>
fn textureGatherCompare(t: texture_depth_2d_array, s: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32) -> vec4<f32>
fn textureGatherCompare(t: texture_depth_2d_array, s: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, offset: vec2<i32>) -> vec4<f32>
fn textureGatherCompare(t: texture_depth_cube, s: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> vec4<f32>
fn textureGatherCompare(t: texture_depth_cube_array, s: sampler_comparison, coords: vec3<f32>, array_index: i32, depth_ref: f32) -> vec4<f32>
Parameters:
t The depth texture to read from.
s The sampler comparison.
coords The texture coordinates.
array_index The 0-based texture array index.
depth_ref The reference value to compare the sampled depth value against.
offset The optional texel offset applied to the unnormalized texture coordinate before sampling the texture. This offset is applied before applying any texture wrapping modes.
The offset expression must be a creation-time expression (e.g. vec2<i32>(1, 2)).
Each offset component must be at least -8 and at most 7. Values outside of this range will result in a shader-creation error.

Returns:

A four component vector with comparison result for the selected texels, as described above.

EXAMPLE: Gather depth comparison
@group(0) @binding(0) var dt: texture_depth_2d;
@group(0) @binding(1) var s: sampler;

fn gather_depth_compare(c: vec2<f32>, depth_ref: f32) -> vec4<f32> {
  return textureGatherCompare(dt,s,c,depth_ref);
}

16.8.4. textureLoad

Reads a single texel from a texture without sampling or filtering.

fn textureLoad(t: texture_1d<T>, coords: i32, level: i32) -> vec4<T>
fn textureLoad(t: texture_2d<T>, coords: vec2<i32>, level: i32) -> vec4<T>
fn textureLoad(t: texture_2d_array<T>, coords: vec2<i32>, array_index: i32, level: i32) -> vec4<T>
fn textureLoad(t: texture_3d<T>, coords: vec3<i32>, level: i32) -> vec4<T>
fn textureLoad(t: texture_multisampled_2d<T>, coords: vec2<i32>, sample_index: i32)-> vec4<T>
fn textureLoad(t: texture_depth_2d, coords: vec2<i32>, level: i32) -> f32
fn textureLoad(t: texture_depth_2d_array, coords: vec2<i32>, array_index: i32, level: i32) -> f32
fn textureLoad(t: texture_depth_multisampled_2d, coords: vec2<i32>, sample_index: i32)-> f32
fn textureLoad(t: texture_external, coords: vec2<i32>) -> vec4<f32>

Parameters:

t The sampled, multisampled, depth, or external texture.
coords The 0-based texel coordinate.
array_index The 0-based texture array index.
level The mip level, with level 0 containing a full size version of the texture.
sample_index The 0-based sample index of the multisampled texture.

Returns:

The unfiltered texel data.

An out of bounds access occurs if:

If an out of bounds access occurs, the built-in function returns one of:

16.8.5. textureNumLayers

Returns the number of layers (elements) of an array texture.

fn textureNumLayers(t: texture_2d_array<T>) -> i32
fn textureNumLayers(t: texture_cube_array<T>) -> i32
fn textureNumLayers(t: texture_depth_2d_array) -> i32
fn textureNumLayers(t: texture_depth_cube_array) -> i32
fn textureNumLayers(t: texture_storage_2d_array<F,A>) -> i32

Parameters:

t The sampled, multisampled, depth or storage array texture.

Returns:

The number of layers (elements) of the array texture.

16.8.6. textureNumLevels

Returns the number of mip levels of a texture.

fn textureNumLevels(t: texture_1d<T>) -> i32
fn textureNumLevels(t: texture_2d<T>) -> i32
fn textureNumLevels(t: texture_2d_array<T>) -> i32
fn textureNumLevels(t: texture_3d<T>) -> i32
fn textureNumLevels(t: texture_cube<T>) -> i32
fn textureNumLevels(t: texture_cube_array<T>) -> i32
fn textureNumLevels(t: texture_depth_2d) -> i32
fn textureNumLevels(t: texture_depth_2d_array) -> i32
fn textureNumLevels(t: texture_depth_cube) -> i32
fn textureNumLevels(t: texture_depth_cube_array) -> i32

Parameters:

t The sampled or depth texture.

Returns:

The number of mip levels for the texture.

16.8.7. textureNumSamples

Returns the number samples per texel in a multisampled texture.

textureNumSamples(t: texture_multisampled_2d<T>) -> i32
textureNumSamples(t: texture_depth_multisampled_2d) -> i32

Parameters:

t The multisampled texture.

Returns:

The number of samples per texel in the multisampled texture.

16.8.8. textureSample

Samples a texture.

Must only be used in a fragment shader stage. Must only be invoked in uniform control flow.

fn textureSample(t: texture_1d<f32>, s: sampler, coords: f32) -> vec4<f32>
fn textureSample(t: texture_2d<f32>, s: sampler, coords: vec2<f32>) -> vec4<f32>
fn textureSample(t: texture_2d<f32>, s: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<f32>
fn textureSample(t: texture_2d_array<f32>, s: sampler, coords: vec2<f32>, array_index: i32) -> vec4<f32>
fn textureSample(t: texture_2d_array<f32>, s: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<f32>
fn textureSample(t: texture_3d<f32>, s: sampler, coords: vec3<f32>) -> vec4<f32>
fn textureSample(t: texture_3d<f32>, s: sampler, coords: vec3<f32>, offset: vec3<i32>) -> vec4<f32>
fn textureSample(t: texture_cube<f32>, s: sampler, coords: vec3<f32>) -> vec4<f32>
fn textureSample(t: texture_cube_array<f32>, s: sampler, coords: vec3<f32>, array_index: i32) -> vec4<f32>
fn textureSample(t: texture_depth_2d, s: sampler, coords: vec2<f32>) -> f32
fn textureSample(t: texture_depth_2d, s: sampler, coords: vec2<f32>, offset: vec2<i32>) -> f32
fn textureSample(t: texture_depth_2d_array, s: sampler, coords: vec2<f32>, array_index: i32) -> f32
fn textureSample(t: texture_depth_2d_array, s: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> f32
fn textureSample(t: texture_depth_cube, s: sampler, coords: vec3<f32>) -> f32
fn textureSample(t: texture_depth_cube_array, s: sampler, coords: vec3<f32>, array_index: i32) -> f32

Parameters:

t The sampled, depth, or external texture to sample.
s The sampler type.
coords The texture coordinates used for sampling.
array_index The 0-based texture array index to sample.
offset The optional texel offset applied to the unnormalized texture coordinate before sampling the texture. This offset is applied before applying any texture wrapping modes.
The offset expression must be a creation-time expression (e.g. vec2<i32>(1, 2)).
Each offset component must be at least -8 and at most 7. Values outside of this range will result in a shader-creation error.

Returns:

The sampled value.

16.8.9. textureSampleBias

Samples a texture with a bias to the mip level.

Must only be used in a fragment shader stage. Must only be invoked in uniform control flow.

fn textureSampleBias(t: texture_2d<f32>, s: sampler, coords: vec2<f32>, bias: f32) -> vec4<f32>
fn textureSampleBias(t: texture_2d<f32>, s: sampler, coords: vec2<f32>, bias: f32, offset: vec2<i32>) -> vec4<f32>
fn textureSampleBias(t: texture_2d_array<f32>, s: sampler, coords: vec2<f32>, array_index: i32, bias: f32) -> vec4<f32>
fn textureSampleBias(t: texture_2d_array<f32>, s: sampler, coords: vec2<f32>, array_index: i32, bias: f32, offset: vec2<i32>) -> vec4<f32>
fn textureSampleBias(t: texture_3d<f32>, s: sampler, coords: vec3<f32>, bias: f32) -> vec4<f32>
fn textureSampleBias(t: texture_3d<f32>, s: sampler, coords: vec3<f32>, bias: f32, offset: vec3<i32>) -> vec4<f32>
fn textureSampleBias(t: texture_cube<f32>, s: sampler, coords: vec3<f32>, bias: f32) -> vec4<f32>
fn textureSampleBias(t: texture_cube_array<f32>, s: sampler, coords: vec3<f32>, array_index: i32, bias: f32) -> vec4<f32>

Parameters:

t The texture to sample.
s The sampler type.
coords The texture coordinates used for sampling.
array_index The 0-based texture array index to sample.
bias The bias to apply to the mip level before sampling. bias must be between -16.0 and 15.99.
offset The optional texel offset applied to the unnormalized texture coordinate before sampling the texture. This offset is applied before applying any texture wrapping modes.
The offset expression must be a creation-time expression (e.g. vec2<i32>(1, 2)).
Each offset component must be at least -8 and at most 7. Values outside of this range will result in a shader-creation error.

Returns:

The sampled value.

16.8.10. textureSampleCompare

Samples a depth texture and compares the sampled depth values against a reference value.

Must only be used in a fragment shader stage. Must only be invoked in uniform control flow.

fn textureSampleCompare(t: texture_depth_2d, s: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> f32
fn textureSampleCompare(t: texture_depth_2d, s: sampler_comparison, coords: vec2<f32>, depth_ref: f32, offset: vec2<i32>) -> f32
fn textureSampleCompare(t: texture_depth_2d_array, s: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32) -> f32
fn textureSampleCompare(t: texture_depth_2d_array, s: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, offset: vec2<i32>) -> f32
fn textureSampleCompare(t: texture_depth_cube, s: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> f32
fn textureSampleCompare(t: texture_depth_cube_array, s: sampler_comparison, coords: vec3<f32>, array_index: i32, depth_ref: f32) -> f32

Parameters:

t The depth texture to sample.
s The sampler comparision type.
coords The texture coordinates used for sampling.
array_index The 0-based texture array index to sample.
depth_ref The reference value to compare the sampled depth value against.
offset The optional texel offset applied to the unnormalized texture coordinate before sampling the texture. This offset is applied before applying any texture wrapping modes.
The offset expression must be a creation-time expression (e.g. vec2<i32>(1, 2)).
Each offset component must be at least -8 and at most 7. Values outside of this range will result in a shader-creation error.

Returns:

A value in the range [0.0..1.0].

Each sampled texel is compared against the reference value using the comparision operator defined by the sampler_comparison, resulting in either a 0 or 1 value for each texel.

If the sampler uses bilinear filtering then the returned value is the filtered average of these values, otherwise the comparision result of a single texel is returned.

16.8.11. textureSampleCompareLevel

Samples a depth texture and compares the sampled depth values against a reference value.

fn textureSampleCompareLevel(t: texture_depth_2d, s: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> f32
fn textureSampleCompareLevel(t: texture_depth_2d, s: sampler_comparison, coords: vec2<f32>, depth_ref: f32, offset: vec2<i32>) -> f32
fn textureSampleCompareLevel(t: texture_depth_2d_array, s: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32) -> f32
fn textureSampleCompareLevel(t: texture_depth_2d_array, s: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, offset: vec2<i32>) -> f32
fn textureSampleCompareLevel(t: texture_depth_cube, s: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> f32
fn textureSampleCompareLevel(t: texture_depth_cube_array, s: sampler_comparison, coords: vec3<f32>, array_index: i32, depth_ref: f32) -> f32

The textureSampleCompareLevel function is the same as textureSampleCompare, except that:

16.8.12. textureSampleGrad

Samples a texture using explicit gradients.

fn textureSampleGrad(t: texture_2d<f32>, s: sampler, coords: vec2<f32>, ddx: vec2<f32>, ddy: vec2<f32>) -> vec4<f32>
fn textureSampleGrad(t: texture_2d<f32>, s: sampler, coords: vec2<f32>, ddx: vec2<f32>, ddy: vec2<f32>, offset: vec2<i32>) -> vec4<f32>
fn textureSampleGrad(t: texture_2d_array<f32>, s: sampler, coords: vec2<f32>, array_index: i32, ddx: vec2<f32>, ddy: vec2<f32>) -> vec4<f32>
fn textureSampleGrad(t: texture_2d_array<f32>, s: sampler, coords: vec2<f32>, array_index: i32, ddx: vec2<f32>, ddy: vec2<f32>, offset: vec2<i32>) -> vec4<f32>
fn textureSampleGrad(t: texture_3d<f32>, s: sampler, coords: vec3<f32>, ddx: vec3<f32>, ddy: vec3<f32>) -> vec4<f32>
fn textureSampleGrad(t: texture_3d<f32>, s: sampler, coords: vec3<f32>, ddx: vec3<f32>, ddy: vec3<f32>, offset: vec3<i32>) -> vec4<f32>
fn textureSampleGrad(t: texture_cube<f32>, s: sampler, coords: vec3<f32>, ddx: vec3<f32>, ddy: vec3<f32>) -> vec4<f32>
fn textureSampleGrad(t: texture_cube_array<f32>, s: sampler, coords: vec3<f32>, array_index: i32, ddx: vec3<f32>, ddy: vec3<f32>) -> vec4<f32>

Parameters:

t The texture to sample.
s The sampler type.
coords The texture coordinates used for sampling.
array_index The 0-based texture array index to sample.
ddx The x direction derivative vector used to compute the sampling locations.
ddy The y direction derivative vector used to compute the sampling locations.
offset The optional texel offset applied to the unnormalized texture coordinate before sampling the texture. This offset is applied before applying any texture wrapping modes.
The offset expression must be a creation-time expression (e.g. vec2<i32>(1, 2)).
Each offset component must be at least -8 and at most 7. Values outside of this range will result in a shader-creation error.

Returns:

The sampled value.

16.8.13. textureSampleLevel

Samples a texture using an explicit mip level, or at mip level 0.

fn textureSampleLevel(t: texture_2d<f32>, s: sampler, coords: vec2<f32>, level: f32) -> vec4<f32>
fn textureSampleLevel(t: texture_2d<f32>, s: sampler, coords: vec2<f32>, level: f32, offset: vec2<i32>) -> vec4<f32>
fn textureSampleLevel(t: texture_2d_array<f32>, s: sampler, coords: vec2<f32>, array_index: i32, level: f32) -> vec4<f32>
fn textureSampleLevel(t: texture_2d_array<f32>, s: sampler, coords: vec2<f32>, array_index: i32, level: f32, offset: vec2<i32>) -> vec4<f32>
fn textureSampleLevel(t: texture_3d<f32>, s: sampler, coords: vec3<f32>, level: f32) -> vec4<f32>
fn textureSampleLevel(t: texture_3d<f32>, s: sampler, coords: vec3<f32>, level: f32, offset: vec3<i32>) -> vec4<f32>
fn textureSampleLevel(t: texture_cube<f32>, s: sampler, coords: vec3<f32>, level: f32) -> vec4<f32>
fn textureSampleLevel(t: texture_cube_array<f32>, s: sampler, coords: vec3<f32>, array_index: i32, level: f32) -> vec4<f32>
fn textureSampleLevel(t: texture_depth_2d, s: sampler, coords: vec2<f32>, level: i32) -> f32
fn textureSampleLevel(t: texture_depth_2d, s: sampler, coords: vec2<f32>, level: i32, offset: vec2<i32>) -> f32
fn textureSampleLevel(t: texture_depth_2d_array, s: sampler, coords: vec2<f32>, array_index: i32, level: i32) -> f32
fn textureSampleLevel(t: texture_depth_2d_array, s: sampler, coords: vec2<f32>, array_index: i32, level: i32, offset: vec2<i32>) -> f32
fn textureSampleLevel(t: texture_depth_cube, s: sampler, coords: vec3<f32>, level: i32) -> f32
fn textureSampleLevel(t: texture_depth_cube_array, s: sampler, coords: vec3<f32>, array_index: i32, level: i32) -> f32
fn textureSampleLevel(t: texture_external, s: sampler, coords: vec2<f32>) -> vec4<f32>

Parameters:

t The sampled or depth texture to sample.
s The sampler type.
coords The texture coordinates used for sampling.
array_index The 0-based texture array index to sample.
level The mip level, with level 0 containing a full size version of the texture. For the functions where level is a f32, fractional values may interpolate between two levels if the format is filterable according to the Texture Format Capabilities.
When not specified, mip level 0 is sampled.
offset The optional texel offset applied to the unnormalized texture coordinate before sampling the texture. This offset is applied before applying any texture wrapping modes.
The offset expression must be a creation-time expression (e.g. vec2<i32>(1, 2)).
Each offset component must be at least -8 and at most 7. Values outside of this range will result in a shader-creation error.

Returns:

The sampled value.

16.8.14. textureStore

Writes a single texel to a texture.

fn textureStore(t: texture_storage_1d<F,write>, coords: i32, value: vec4<T>)
fn textureStore(t: texture_storage_2d<F,write>, coords: vec2<i32>, value: vec4<T>)
fn textureStore(t: texture_storage_2d_array<F,write>, coords: vec2<i32>, array_index: i32, value: vec4<T>)
fn textureStore(t: texture_storage_3d<F,write>, coords: vec3<i32>, value: vec4<T>)

The channel format T depends on the storage texel format F. See the texel format table for the mapping of texel format to channel format.

Parameters:

t The write-only storage texture.
coords The 0-based texel coordinate.
array_index The 0-based texture array index.
value The new texel value.

Note:

An out-of-bounds access occurs if:

If an out-of-bounds access occurs, the built-in function may do any of the following:

16.9. Atomic Built-in Functions

Atomic built-in functions can be used to read/write/read-modify-write atomic objects. They are the only operations allowed on § 4.3.7 Atomic Types.

All atomic built-in functions use a relaxed memory ordering. This means synchronization and ordering guarantees only apply among atomic operations acting on the same memory locations. No synchronization or ordering guarantees apply between atomic and non-atomic memory accesses, or between atomic accesses acting on different memory locations.

Atomic built-in functions must not be used in a vertex shader stage.

The address space SC of the atomic_ptr parameter in all atomic built-in functions must be either storage or workgroup.

The access mode A in all atomic built-in functions must be read_write.

16.9.1. Atomic Load

fn atomicLoad(atomic_ptr: ptr<SC, atomic<T>, A>) -> T

Returns the atomically loaded the value pointed to by atomic_ptr. It does not modify the object.

16.9.2. Atomic Store

fn atomicStore(atomic_ptr: ptr<SC, atomic<T>, A>, v: T)

Atomically stores the value v in the atomic object pointed to by atomic_ptr.

16.9.3. Atomic Read-modify-write

fn atomicAdd(atomic_ptr: ptr<SC, atomic<T>, A>, v: T) -> T
fn atomicSub(atomic_ptr: ptr<SC, atomic<T>, A>, v: T) -> T
fn atomicMax(atomic_ptr: ptr<SC, atomic<T>, A>, v: T) -> T
fn atomicMin(atomic_ptr: ptr<SC, atomic<T>, A>, v: T) -> T
fn atomicAnd(atomic_ptr: ptr<SC, atomic<T>, A>, v: T) -> T
fn atomicOr(atomic_ptr: ptr<SC, atomic<T>, A>, v: T) -> T
fn atomicXor(atomic_ptr: ptr<SC, atomic<T>, A>, v: T) -> T
Each function performs the following steps atomically:
  1. Load the original value pointed to by atomic_ptr.

  2. Obtains a new value by performing the operation (e.g. max) from the function name with the value v.

  3. Store the new value using atomic_ptr.

Each function returns the original value stored in the atomic object.

fn atomicExchange(atomic_ptr: ptr<SC, atomic<T>, A>, v: T) -> T

Atomically stores the value v in the atomic object pointed to atomic_ptr and returns the original value stored in the atomic object.

fn atomicCompareExchangeWeak(atomic_ptr: ptr<SC, atomic<T>, A>, cmp: T, v: T) -> __atomic_compare_exchange_result<T>

struct __atomic_compare_exchange_result<T> {
  old_value : T; // old value stored in the atomic
  exchanged : bool; // true if the exchange was done
}

Note: A value cannot be explicitly declared with the type __atomic_compare_exchange_result, but a value may infer the type.

Performs the following steps atomically:

  1. Load the original value pointed to by atomic_ptr.

  2. Compare the original value to the value v using an equality operation.

  3. Store the value v only if the result of the equality comparison was true.

Returns a two member structure, where the first member, old_value, is the original value of the atomic object and the second member, exchanged, is whether or not the comparison succeeded.

Note: the equality comparison may spuriously fail on some implementations. That is, the second component of the result vector may be false even if the first component of the result vector equals cmp.

16.10. Data Packing Built-in Functions

Data packing builtin functions can be used to encode values using data formats that do not correspond directly to types in WGSL. This enables a program to write many densely packed values to memory, which can reduce a shader’s memory bandwidth demand.

Overload Description
@const fn pack4x8snorm(e: vec4<f32>) -> u32 Converts four normalized floating point values to 8-bit signed integers, and then combines them into one u32 value.
Component e[i] of the input is converted to an 8-bit twos complement integer value ⌊ 0.5 + 127 × min(1, max(-1, e[i])) ⌋ which is then placed in bits 8 × i through 8 × i + 7 of the result.
@const fn pack4x8unorm(e: vec4<f32>) -> u32 Converts four normalized floating point values to 8-bit unsigned integers, and then combines them into one u32 value.
Component e[i] of the input is converted to an 8-bit unsigned integer value ⌊ 0.5 + 255 × min(1, max(0, e[i])) ⌋ which is then placed in bits 8 × i through 8 × i + 7 of the result.
@const fn pack2x16snorm(e: vec2<f32>) -> u32 Converts two normalized floating point values to 16-bit signed integers, and then combines them into one u32 value.
Component e[i] of the input is converted to a 16-bit twos complement integer value ⌊ 0.5 + 32767 × min(1, max(-1, e[i])) ⌋ which is then placed in bits 16 × i through 16 × i + 15 of the result.
@const fn pack2x16unorm(e: vec2<f32>) -> u32 Converts two normalized floating point values to 16-bit unsigned integers, and then combines them into one u32 value.
Component e[i] of the input is converted to a 16-bit unsigned integer value ⌊ 0.5 + 65535 × min(1, max(0, e[i])) ⌋ which is then placed in bits 16 × i through 16 × i + 15 of the result.
@const fn pack2x16float(e: vec2<f32>) -> u32 Converts two floating point values to half-precision floating point numbers, and then combines them into one u32 value.
Component e[i] of the input is converted to a IEEE-754 binary16 value, which is then placed in bits 16 × i through 16 × i + 15 of the result. See § 12.5.2 Floating Point Conversion.

16.11. Data Unpacking Built-in Functions

Data unpacking builtin functions can be used to decode values in data formats that do not correspond directly to types in WGSL. This enables a program to read many densely packed values from memory, which can reduce a shader’s memory bandwidth demand.

Overload Description
@const fn unpack4x8snorm(e: u32) -> vec4<f32> Decomposes a 32-bit value into four 8-bit chunks, then reinterprets each chunk as a signed normalized floating point value.
Component i of the result is max(v ÷ 127, -1), where v is the interpretation of bits 8×i through 8×i+7 of e as a twos-complement signed integer.
@const fn unpack4x8unorm(e: u32) -> vec4<f32> Decomposes a 32-bit value into four 8-bit chunks, then reinterprets each chunk as an unsigned normalized floating point value.
Component i of the result is v ÷ 255, where v is the interpretation of bits 8×i through 8×i+7 of e as an unsigned integer.
@const fn unpack2x16snorm(e: u32) -> vec2<f32> Decomposes a 32-bit value into two 16-bit chunks, then reinterprets each chunk as a signed normalized floating point value.
Component i of the result is max(v ÷ 32767, -1), where v is the interpretation of bits 16×i through 16×i+15 of e as a twos-complement signed integer.
@const fn unpack2x16unorm(e: u32) -> vec2<f32> Decomposes a 32-bit value into two 16-bit chunks, then reinterprets each chunk as an unsigned normalized floating point value.
Component i of the result is v ÷ 65535, where v is the interpretation of bits 16×i through 16×i+15 of e as an unsigned integer.
@const fn unpack2x16float(e: u32) -> vec2<f32> Decomposes a 32-bit value into two 16-bit chunks, and reinterpets each chunk as a floating point value.
Component i of the result is the f32 representation of v, where v is the interpretation of bits 16×i through 16×i+15 of e as an IEEE-754 binary16 value. See § 12.5.2 Floating Point Conversion.

16.12. Synchronization Built-in Functions

WGSL provides the following synchronization functions:

fn storageBarrier()
fn workgroupBarrier()

All synchronization functions execute a control barrier with Acquire/Release memory ordering. That is, all synchronization functions, and affected memory and atomic operations are ordered in program order relative to the synchronization function. Additionally, the affected memory and atomic operations program-ordered before the synchronization function must be visible to all other threads in the workgroup before any affected memory or atomic operation program-ordered after the synchronization function is executed by a member of the workgroup. All synchronization functions use the Workgroup memory scope. All synchronization functions have a Workgroup execution scope. All synchronization functions must only be used in the compute shader stage.

storageBarrier affects memory and atomic operations in the storage address space.

workgroupBarrier affects memory and atomic operations in the workgroup address space.

Conformance

Document conventions

Conformance requirements are expressed with a combination of descriptive assertions and RFC 2119 terminology. The key words “MUST”, “MUST NOT”, “REQUIRED”, “SHALL”, “SHALL NOT”, “SHOULD”, “SHOULD NOT”, “RECOMMENDED”, “MAY”, and “OPTIONAL” in the normative parts of this document are to be interpreted as described in RFC 2119. However, for readability, these words do not appear in all uppercase letters in this specification.

All of the text of this specification is normative except sections explicitly marked as non-normative, examples, and notes. [RFC2119]

Examples in this specification are introduced with the words “for example” or are set apart from the normative text with class="example", like this:

This is an example of an informative example.

Informative notes begin with the word “Note” and are set apart from the normative text with class="note", like this:

Note, this is an informative note.

Conformant Algorithms

Requirements phrased in the imperative as part of algorithms (such as "strip any leading space characters" or "return false and abort these steps") are to be interpreted with the meaning of the key word ("must", "should", "may", etc) used in introducing the algorithm.

Conformance requirements phrased as algorithms or specific steps can be implemented in any manner, so long as the end result is equivalent. In particular, the algorithms defined in this specification are intended to be easy to understand and are not intended to be performant. Implementers are encouraged to optimize.

Index

Terms defined by this specification

Terms defined by reference

References

Normative References

[CHARMOD-NORM]
Addison Phillips; et al. Character Model for the World Wide Web: String Matching. 11 August 2021. NOTE. URL: https://www.w3.org/TR/charmod-norm/
[IEEE-754]
IEEE Standard for Floating-Point Arithmetic. 29 August 2008. URL: http://ieeexplore.ieee.org/servlet/opac?punumber=4610933
[RFC2119]
S. Bradner. Key words for use in RFCs to Indicate Requirement Levels. March 1997. Best Current Practice. URL: https://datatracker.ietf.org/doc/html/rfc2119
[UAX14]
Christopher Chapman. Unicode Line Breaking Algorithm. 23 August 2021. Unicode Standard Annex #14. URL: https://www.unicode.org/reports/tr14/tr14-47.html
[UAX15]
Ken Whistler. Unicode Normalization Forms. 27 August 2021. Unicode Standard Annex #15. URL: https://www.unicode.org/reports/tr15/tr15-51.html
[UAX31]
Mark Davis. Unicode Identifier and Pattern Syntax. 26 August 2021. Unicode Standard Annex #31. URL: https://www.unicode.org/reports/tr31/tr31-35.html
[UnicodeVersion14]
The Unicode Standard, Version 14.0.0. URL: http://www.unicode.org/versions/Unicode14.0.0/
[UTS46]
Mark Davis; Michel Suignard. Unicode IDNA Compatibility Processing. 24 August 2021. Unicode Technical Standard #46. URL: https://www.unicode.org/reports/tr46/tr46-27.html
[VulkanMemoryModel]
Jeff Bolz; et al. Vulkan Memory Model. URL: https://www.khronos.org/registry/vulkan/specs/1.2-extensions/html/vkspec.html#memory-model
[WebGPU]
Dzmitry Malyshau; Kai Ninomiya; Brandon Jones. WebGPU. Editor's Draft. URL: https://gpuweb.github.io/gpuweb/

Issues Index

Can we query upper bounds on workgroup size dimensions? Is it independent of the shader, or a property to be queried after creating the shader module?
Check behaviour of the f32 to f16 conversion for numbers just beyond the max normal f16 values. I’ve written what an NVIDIA GPU does. See https://github.com/google/amber/pull/918 for an executable test case.
https://github.com/gpuweb/gpuweb/issues/1621