WebGPU Shading Language

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.

// A fragment shader which lights textured geometry with point lights.

// Lights from a storage buffer binding.
struct PointLight {
  position : vec3f,
  color : vec3f,
}

struct LightStorage {
  pointCount : u32,
  point : array<PointLight>,
}
@group(0) @binding(0) var<storage> lights : LightStorage;

// Texture and sampler.
@group(1) @binding(0) var baseColorSampler : sampler;
@group(1) @binding(1) var baseColorTexture : texture_2d<f32>;

// Function arguments are values from the vertex shader.
@fragment
fn fragmentMain(@location(0) worldPos : vec3f,
                @location(1) normal : vec3f,
                @location(2) uv : vec2f) -> @location(0) vec4f {
  // Sample the base color of the surface from a texture.
  let baseColor = textureSample(baseColorTexture, baseColorSampler, uv);

  let N = normalize(normal);
  var surfaceColor = vec3f(0);

  // Loop over the scene point lights.
  for (var i = 0u; i < lights.pointCount; i++) {
    let worldToLight = lights.point[i].position - worldPos;
    let dist = length(worldToLight);
    let dir = normalize(worldToLight);

    // Determine the contribution of this light to the surface color.
    let radiance = lights.point[i].color * (1 / pow(dist, 2));
    let nDotL = max(dot(N, dir), 0);

    // Accumulate light contribution to the surface color.
    surfaceColor += baseColor.rgb * radiance * nDotL;
  }

  // Return the accumulated surface color.
  return vec4(surfaceColor, baseColor.a);
}

1.1. Overview

WebGPU issues a unit of work to the GPU in the form of a GPU command. WGSL is concerned with two kinds of GPU commands:

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

• a dispatch command executes a compute pipeline in the context of inputs and attached resources.

Both kinds of pipelines use shaders written in WGSL.

A shader is the portion of a WGSL program that executes a shader stage in a pipeline. A shader comprises:

• An entry point function.

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

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

• The set of types used to define or analyze all those functions, variables, and constants.

Note: A WGSL program does not require an entry point; however, such a program cannot be executed by the API because an entry point is required to create a GPUProgrammableStage.

When executing a shader stage, the implementation:

• Computes the values of constants declared at module-scope.

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

• Allocates memory for other module-scope variables, and populates that memory with the specified initial values.

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

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

• Then it invokes the entry point.

A WGSL program is organized into:

• Directives, which specify module-level behavior controls.

• Functions, which specify execution behavior.

• Statements, which are declarations or units of executable behavior.

• Literals, which are text representations for pure mathematical values.

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

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

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

• Types, each of which describes:

  • A set of values.

  • Constraints on supported expressions.

  • The semantics of those expressions.

• Attributes, which modify an object to specify extra information such as:

  • Specifying the interfaces to entry points.

  • Specifying diagnostic filters.

Note: A WGSL program is currently composed of a single WGSL module.

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

• Declare constants or variables.

• Modify the contents of variables.

• Modify execution order using structured programming constructs:

  • Selective execution: if (with optional else if and else clauses), switch.

  • Repetition: loop, while, for.

  • Escaping a nested execution construct: continue, break, break if.

  • Refactoring: function call and return.

• Evaluate expressions to compute values as part of the above behaviors.

• Check assumptions at shader creation time on constant expressions.

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 describing booleans and numbers (integers and floating point). These types can be aggregated into composites (vectors, matrices, arrays, and structures). WGSL has special types (e.g. atomics) that provide unique operations. WGSL describes the types that can be stored in memory as memory views. WGSL provides commonly used rendering types in the form of textures and samplers. These types have associated built-in functions to expose commonly provided GPU hardware for graphics rendering.

WGSL does not have implicit conversions or promotions from concrete types, but does provide implicit conversions and promotions from abstract types. Converting a value from one concrete numeric or boolean type to another requires an explicit conversion, value constructor, or reinterpretation of bits; however, WGSL does provide some limited facility to promote scalar types to vector types. This also applies to composite types.

The work of a shader stage is partitioned into one or more invocations, each of which executes the entry point, but under slightly different conditions. Invocations in a shader stage share access to certain variables:

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

• In a compute shader, invocations in the same workgroup share variables in the workgroup address space. Invocations in different workgroups do not share those variables.

However, the invocations act on different sets of shader stage 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 behavior of the invocations in a shader stage:

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

• Coordinate potentially conflicting accesses to shared variables, to avoid data races.

WGSL sometimes permits several possible behaviors for a given feature. This is a portability hazard, as different implementations may exhibit the different behaviors. 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.

Behavioral requirements are actions the implementation will perform when processing or executing a WGSL program. They describe the implementation’s obligations in the contract with the programmer. The specification explicitly states these obligations when they might not be otherwise obvious.

1.2. Syntax Notation

Following syntax notation describes the conventions of the syntactic grammar of WGSL:

• Italic text on both sides of a rule indicates a syntax rule.

• Bold monospace text starting and ending with single quotes (') on the right-hand side of a rule indicates keywords and tokens.

• A colon (:) in regular text registers a syntax rule.

• A vertical bar (|) in regular text indicates alternatives.

• An question mark (?) in regular text indicates that previous keyword, token, rule, or group occurs zero or one times (is optional).

• An asterisk (*) in regular text indicates that previous keyword, token, rule, or group occurs zero or more times.

• A plus (+) in regular text indicates that previous keyword, token, rule, or group occurs one or more times.

• A matching pair of opening parenthesis (() and closing parenthesis ()) in regular text indicates a group of elements.

1.3. Mathematical Terms and Notation

Angles:

• By convention, angles are measured in radians.

• The reference ray for measuring angles is the ray from the origin (0,0) toward (+∞,0).

• Let θ be the angle subtended by a comparison ray and the reference ray. Then θ increases as the comparison ray moves counterclockwise.

• There are 2 π radians in a complete circle.

• Examples:

  • The angle 0 points from the origin to the right, toward (1,0)

  • The angle 2π points from the origin to the right, toward (1,0)

  • The angle π/4 points from the origin to the point (1,1)

  • The angle π/2 points from the origin to the point (0,1)

  • The angle π points from the origin to the point (-1,0)

  • The angle (3/2)π points from the origin to the point (0,-1)

A hyperbolic angle is a unitless area, not an angle in the traditional sense. Specifically:

• Consider the hyperbola x² - y² = 1 for x > 0.

• Let R be a ray from the origin to some point (x, y) on the hyperbola.

• Let a be twice the area enclosed by R, the x axis, and the curve of the hyperbola itself.

• Consider a to be positive when R is above the x axis, and negative when below.

Then the area a is a hyperbolic angle such that x is the hyperbolic cosine of a, and y is the hyperbolic sine of a.

Positive infinity, denoted by +∞, is a unique value strictly greater than all real numbers.

Negative infinity, denoted by −∞, is a unique value strictly lower than all real numbers.

The extended real numbers (also known as the affinely extended real numbers) is the set of real numbers together with +∞ and −∞. Computers use floating point types to approximately represent the extended reals, including values for both infinities. See § 15.7 Floating Point Evaluation.

An interval is a contiguous set of numbers with a lower and upper bound. Depending on context, they are sets of integers, floating point numbers, real numbers, or extended real numbers.

• The closed interval [a,b] is the set of numbers x such that a ≤ x ≤ b.

• The half-open interval [a,b) is the set of numbers x such that a ≤ x < b.

• The half-open interval (a,b] is the set of numbers x such that a < x ≤ b.

The floor expression is defined for extended real numbers x:

• ⌊ +∞ ⌋ = +∞

• ⌊ −∞ ⌋ = −∞

• for real number x, ⌊x⌋ = k, where k is the unique integer such that k ≤ x < k+1

The ceiling expression is defined for extended real numbers x:

• ⌈ +∞ ⌉ = +∞

• ⌈ −∞ ⌉ = −∞

• for real number x, ⌈x⌉ = k, where k is the unique integer such that k-1 < x ≤ k

The truncate function is defined for extended real numbers x:

• truncate(+∞) = +∞

• truncate(−∞) = −∞

• for real number x, computes the nearest whole number whose absolute value is less than or equal to the absolute value of x:

  • truncate(x) = ⌊x⌋ if x ≥ 0, and ⌈x⌉ if x < 0.

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

• roundUp(k, n) = ⌈n ÷ k⌉ × k

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

• transpose(A) = A^T

• transpose(A)_i,j = A_j,i

The transpose of a column vector is defined by interpreting the column vector as a 1-row matrix. Similarly, the transpose of a row vector is defined by interpreting the row vector as a 1-column matrix.

2. WGSL Module

A WGSL program is composed of a single WGSL module.

A module is a sequence of optional directives followed by module scope declarations and assertions. A module is organized into:

• Directives, which specify module-level behavior controls.

• Functions, which specify execution behavior.

• Statements, which are declarations or units of executable behavior.

• Literals, which are text representations for pure mathematical values.

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

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

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

• Types, each of which describes:

  • A set of values.

  • Constraints on supported expressions.

  • The semantics of those expressions.

• Attributes, which modify an object to specify extra information such as:

  • Specifying the interfaces to entry points.

  • Specifying diagnostic filters.

2.1. Shader Lifecycle

There are four key events in the lifecycle of a WGSL program and the shaders it may contain. The first two correspond to the WebGPU API methods used to prepare a WGSL program for execution. The last two are the start and end of execution of a shader.

The events are:

1. Shader module creation

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

2. Pipeline creation

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

   • Only the code forming the shader of the specified entry point of the GPUProgrammableStage is considered during pipeline creation. That is, code unrelated to the entry point is effectively stripped before compilation.

   • Note: Each shader stage is considered to be compiled separately and, thus, might include different portions of the module.

3. Shader execution start

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

4. Shader execution end

   • This occurs when all work in the shader completes:

     • all its invocations terminate, and

     • all accesses to resources complete, and

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

The events are ordered due to:

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

• causality: the shader must start executing before it can finish executing.

2.2. Errors

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

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

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

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

  • A defect in the WebGPU implementation.

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

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

• A pipeline-creation error is an error detectable at pipeline creation time. Detection relies on the WGSL module source text and other information available to the particular pipeline creation API method. These errors are only triggered for code present in the shader of the entry point being compiled for the GPUProgrammableStage.

• A dynamic error is an error occurring during shader execution. These errors may or may not be detectable.

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

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

• A shader-creation error results when failing to meet a requirement detectable at shader-creation time.

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

When unclear from context, this specification indicates whether failure to meet a particular requirement results in a shader-creation, pipeline-creation, or dynamic error.

The consequences of an error are as follows:

• A WGSL module with a shader-creation error or pipeline-creation error error will not be incorporated into a pipeline and hence will not be executed.

• Detectable errors will trigger a diagnostic.

• If a dynamic error occurs:

  • Memory accesses will be restricted to:

    • shader stage inputs,

    • shader stage outputs,

    • any part of a resource bound to a variable in the WGSL module, and

    • other variables declared in the WGSL module.

  • Otherwise, the program may not behave as described in the rest of this specification. Note: These effects may be non-local.

2.3. Diagnostics

An implementation can generate diagnostics during shader module creation or pipeline creation. A diagnostic is a message produced by the implementation for the benefit of the application author.

A diagnostic is created, or triggered, when a particular condition is met, known as the triggering rule. The place in the source text where the condition is met, expressed as a point or range in the source text, is known as the triggering location.

A diagnostic has the following properties:

• A severity.

• A triggering rule.

• A triggering location.

The severity of a diagnostic is one of the following, ordered from greatest to least:

error

The diagnostic is an error. This corresponds to a shader-creation error or to a pipeline-creation error.

warning

The diagnostic describes an anomaly that merits the attention of the application developer, but is not an error.

info

The diagnostic describes a notable condition that merits attention of the application developer, but is not an error or warning.

off

The diagnostic is disabled. It will not be conveyed to the application.

The name of a triggering rule is either:

• a diagnostic_name_token, or

• two diagnostic_name_token tokens, separated by a period '.' (U+002E).

2.3.1. Diagnostic Processing

Triggered diagnostics will be processed as follows:

1. For each diagnostic D, find the diagnostic filter with the smallest affected range that contains D’s triggering location, and which has the same triggering rule.

   • If such a filter exists, apply it to D, updating D's severity.

   • Otherwise D remains unchanged.

2. Discard diagnostics that have severity off.

3. If at least one remaining diagnostic DI has severity info, then:

   • Other info diagnostics with same triggering rule may be discarded, leaving only the original diagnostic DI.

4. If at least one remaining diagnostic DW has severity warning, then:

   • Other info or warning diagnostics with same triggering rule may be discarded, leaving only the original diagnostic DW.

5. If at least one remaining diagnostic has error severity, then:

   • Other diagnostics may be discarded, including other diagnostics with error severity.

   • A program error is generated.

     • The error is a shader-creation error if the diagnostic was triggered at shader module creation time.

     • The error is a pipeline-creation error if the diagnostic was triggered at pipeline creation time.

6. If processing during shader module creation time, the remaining diagnostics populate the messages member of the WebGPU GPUCompilationInfo object.

7. If processing during pipeline creation, error diagnostics result in WebGPU validation failure when validating GPUProgrammableStage.

Note: The rules allow an implementation to stop processing a WGSL module as soon as an error is detected. Additionally, an analysis for a particular kind of warning can stop on the first warning, and an analysis for a particular kind of info diagnostic can stop on the first occurrence. WGSL does not specify the order to perform different kinds of analyses, or an ordering within a single analysis. Therefore, for the same WGSL module, different implementations may report different instances of diagnostics with the same severity.

2.3.2. Filterable Triggering Rules

Most diagnostics are unconditionally reported to the WebGPU application. Some kinds of diagnostics can be filtered, in part by naming their triggering rule. The following table lists the standard set of triggering rules that can be filtered.

Using an unrecognized triggering rule consisting of a single diagnostic name-token should trigger a warning from the user agent.

An implementation may support triggering rules not specified here, provided they are spelled using the multiple-token form of diagnostic_rule_name. Using an unrecognized triggering rule spelled in the multiple-token form may itself trigger a diagnostic.

Future versions of this specification may remove a particular rule or weaken its default severity (i.e. replace its current default with a less severe default) and still be deemed as satisfying backward compatibility. For example, a future version of WGSL may change the default severity for derivative_uniformity from error to either warning or info. After such a change to the specification, previously valid programs would remain valid.

2.3.3. Diagnostic Filtering

Once a diagnostic with a filterable triggering rule is triggered, WGSL provides mechanisms to discard the diagnostic, or to modify its severity.

A diagnostic filter DF has three parameters:

• AR: a range of source text known as the affected range

• NS: a new severity

• TR: a triggering rule

Applying a diagnostic filter DF(AR,NS,TR) to a diagnostic D has the following effect:

• If D's triggering location is in AR and D's triggering rule is TR, then set D's severity property to NS.

• Otherwise, D is unchanged.

A range diagnostic filter is a diagnostic filter whose affected range is a specified range of source text. A range diagnostic filter is specified as a @diagnostic attribute at the start of the affected source range, as specified in the following table. A @diagnostic attribute must not appear anywhere else.

Note: The following are also compound statements: a function body, a case clause, a default-alone clause, the bodies of while and for loops, and the bodies of if_clause, else_if_clause, and else_clause.

var<private> d: f32;
fn helper() -> vec4<f32> {
  // Disable the derivative_uniformity diagnostic in the
  // body of the "if".
  if (d < 0.5) @diagnostic(off,derivative_uniformity) {
    return textureSample(t,s,vec2(0,0));
  }
  return vec4(0.0);
}

A global diagnostic filter can be used to apply a diagnostic filter to the entire WGSL module.

diagnostic(off,derivative_uniformity);
var<private> d: f32;
fn helper() -> vec4<f32> {
  if (d < 0.5) {
    // The derivative_uniformity diagnostic is disabled here
    // by the global diagnostic filter.
    return textureSample(t,s,vec2(0,0));
  } else {
    // The derivative_uniformity diagnostic is set to 'warning' severity.
    @diagnostic(warning,derivative_uniformity) {
      return textureSample(t,s,vec2(0,0));
    }
  }
  return vec4(0.0);
}

Two diagnostic filters DF(AR1,NS1,TR1) and DF(AR2,NS2,TR2) conflict when:

• (AR1 = AR2), and

• (TR1 = TR2), and

• (NS1 ≠ NS2).

Diagnostic filters must not conflict.

Note: Multiple global diagnostic filters are permitted when they do not conflict.

WGSL’s diagnostic filters are designed so their affected ranges nest perfectly. If the affected range of DF1 overlaps with the affected range of DF2, then either DF1’s affected range is fully contained in DF2’s affected range, or the other way around.

The nearest enclosing diagnostic filter for source location L and triggering rule TR, if one exists, is the diagnostic filter DF(AR,NS,TR) where:

• L falls in the affected range AR, and

• If there is another filter DF'(AR',NS',TR) where L falls in AR', then AR is contained in AR'.

Because affected ranges nest, the nearest enclosing diagnostic:

• is a unique range diagnostic filter,

• otherwise is one of a set of duplicate (non-conflicting) global diagnostic filters,

• otherwise does not exist.

2.4. Limits

A WGSL implementation will support shaders that satisfy the following limits. A WGSL implementation may support shaders that go beyond the specified limits.

Note: A WGSL implementation should issue an error if it does not support a shader that goes beyond the specified limits.

3. Textual Structure

The text/wgsl media type is used to identify content as a WGSL module. See Appendix A: The text/wgsl Media Type.

A WGSL module is Unicode text using the UTF-8 encoding, with no byte order mark (BOM).

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

• comments

• tokens

• blankspaces

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

3.1. Parsing

To parse a WGSL module:

1. Remove comments:

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

   • Repeat until no comments remain.

2. Find template lists, using the algorithm in § 3.9 Template Lists.

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

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

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

     • If the next code point is the start of a template list, consume it and return _template_args_start.

     • If the next code point is the end of a template list, consume it and return _template_args_end.

     • Otherwise:

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

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

A shader-creation error results if:

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

• the translation_unit grammar rule does not match the entire token sequence.

3.2. Blankspace and Line Breaks

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

• space (U+0020)

• horizontal tab (U+0009)

• line feed (U+000A)

• vertical tab (U+000B)

• form feed (U+000C)

• carriage return (U+000D)

• next line (U+0085)

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

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

• line separator (U+2028)

• paragraph separator (U+2029)

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

• line feed (U+000A)

• vertical tab (U+000B)

• form feed (U+000C)

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

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

• next line (U+0085)

• line separator (U+2028)

• paragraph separator (U+2029)

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

3.3. Comments

A comment is a span of text that does not influence the validity or meaning of a WGSL program, except that a comment can separate tokens. Shader authors can use comments to document their programs.

A line-ending comment is a kind of comment consisting of the two code points // (U+002F followed by U+002F) and the code points that follow, up until but not including:

• the next line break, or

• the end of the program.

A block comment is a kind of comment consisting of:

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

• Then any sequence of:

  • A block comment, or

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

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

Note: Block comments can be nested. Since a block comment requires matching start and end text sequences, and allows arbitrary nesting, a block comment cannot be recognized with a regular expression. This is a consequence of the Pumping Lemma for Regular Languages.

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

3.4. Tokens

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

• a literal.

• a keyword.

• a reserved word.

• a syntactic token.

• an identifier.

• a context-dependent name.

3.5. Literals

A literal is one of:

• A boolean literal: either true or false.

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

3.5.1. Boolean Literals

const a = true;
const b = false;

3.5.2. Numeric Literals

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

An integer literal is:

• An integer specified as any of:

  • 0

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

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

• Then an optional i or u suffix.

Note: A leading zero on a non-zero integer literal (e.g. 012) is forbidden, so as to avoid confusion with other languages' leading-zero-means-octal notation.

const a = 1u;
const b = 123;
const c = 0;
const d = 0i;

const a = 0x123;
const b = 0X123u;
const c = 0x3f;

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

A floating point literal has two logical parts: a significand to represent a fraction, and an optional exponent. Roughly, the value of the literal is the significand multiplied by a base value raised to the given exponent. A significand digit is significant if it is non-zero, or if there are significand digits to its left and to its right that are both non-zero. Significant digits are counted from left-to-right: the N'th significant digit has N-1 significant digits to its left.

A decimal floating point literal is:

• A significand, specified as a sequence of digits, with an optional decimal point (.) somewhere among them. The significand represents a fraction in base 10 notation.

• Then an optional exponent suffix consisting of:

  • e or E.

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

  • Then an optional f or h suffix.

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

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;

Note: The decimal significand is truncated after 20 decimal digits, preserving approximately log(10)/log(2)×20 ≈ 66.4 significant bits in the fraction.

A hexadecimal floating point literal is:

• A 0x or 0X prefix

• Then a significand, specified as a sequence of hexadecimal digits, with an optional hexadecimal point (.) somewhere among them. The significand represents a fraction in base 16 notation.

• Then an optional exponent suffix consisting of:

  • p or P

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

  • Then an optional f or h suffix.

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

const a = 0xa.fp+2;
const b = 0x1P+4f;
const c = 0X.3;
const d = 0x3p+2h;
const e = 0X1.fp-4;
const f = 0x3.2p+2h;

Note: The hexadecimal significand is truncated after 16 hexadecimal digits, preserving approximately 4 ×16 = 64 significant bits in the fraction.

When a numeric literal has a suffix, the literal denotes a value in a specific concrete scalar type. Otherwise, the literal denotes a value one of the abstract numeric types defined below. In either case, the value denoted by the literal is its mathematical value after conversion to the target type, following the rules in § 15.7.6 Floating Point Conversion.

A shader-creation error results if:

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

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

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

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

Note: The hexadecimal float value 0x1.00000001p0 requires 33 significand bits to be represented exactly, but f32 only has 23 explicit significand 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 refers to a predefined language concept. See § 16.1 Keyword Summary for the list of WGSL keywords.

3.7. Identifiers

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

WGSL uses two grammar nonterminals to separate use cases:

• An ident is used to name a declared object.

• A member_ident is used to name a member of a structure type.

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

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

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

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

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

With the following exceptions:

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

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

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

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

Note: The return type for some built-in functions are structure types whose name cannot be used WGSL source. Those structure types are described as if they were predeclared with a name starting with two underscores. The result value can be saved into newly declared let or var using type inferencing, or immediately have one of its members immediately extracted by name. See example usages in the description of frexp and modf.

3.7.1. Identifier Comparison

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

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

Note: A user agent should issue developer-visible warnings when the meaning of a WGSL module 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. Context-Dependent Names

A context-dependent name is a token used to name a concept, but only in specific grammatical contexts. The spelling of the token may be the same as an identifier, but the token does not resolve to a declared object. This section lists the tokens used as context-dependent names. The token must not be a keyword or reserved word.

3.8.1. Attribute Names

See § 12 Attributes.

The attribute names are:

• 'align'

• 'binding'

• 'builtin'

• 'compute'

• 'const'

• 'diagnostic'

• 'fragment'

• 'group'

• 'id'

• 'interpolate'

• 'invariant'

• 'location'

• 'blend_src'

• 'must_use'

• 'size'

• 'vertex'

• 'workgroup_size'

3.8.2. Built-in Value Names

A built-in value name-token is a token used in the name of a built-in value.

See § 13.3.1.1 Built-in Inputs and Outputs.

The built-in value names are:

• 'vertex_index'

• 'instance_index'

• 'position'

• 'front_facing'

• 'frag_depth'

• 'sample_index'

• 'sample_mask'

• 'local_invocation_id'

• 'local_invocation_index'

• 'global_invocation_id'

• 'workgroup_id'

• 'num_workgroups'

• 'subgroup_invocation_id'

• 'subgroup_size'

3.8.3. Diagnostic Rule Names

A diagnostic name-token is a token used in the name of a diagnostic triggering rule.

See § 2.3 Diagnostics.

The pre-defined diagnostic rule names are:

• 'derivative_uniformity'

• 'subgroup_uniformity'

3.8.4. Diagnostic Severity Control Names

The valid diagnostic filter severity control names are listed in § 2.3 Diagnostics but have the same form as an identifier:

The diagnostic filter severity control names are:

• 'error'

• 'warning'

• 'info'

• 'off'

3.8.5. Extension Names

The valid enable-extension names are listed in § 4.1.1 Enable Extensions but in general have the same form as an identifier:

The enable-extension names are:

• 'f16'

• 'clip_distances'

• 'dual_source_blending'

• 'subgroups'

• 'primitive_index'

The valid language extension names are listed in § 4.1.2 Language Extensions but in general have the same form as an identifier:

The language extension names are:

• 'readonly_and_readwrite_storage_textures'

• 'packed_4x8_integer_dot_product'

• 'unrestricted_pointer_parameters'

• 'pointer_composite_access'

3.8.6. Interpolation Type Names

An interpolation type name-token is a token used in the name of an interpolation type for interpolate_type_name.

See § 13.3.1.4 Interpolation.

The interpolation type names are:

• 'perspective'

• 'linear'

• 'flat'

3.8.7. Interpolation Sampling Names

An interpolation sampling name-token is a token used in the name of an interpolation sampling.

See § 13.3.1.4 Interpolation.

The interpolation sampling names are:

• 'center'

• 'centroid'

• 'sample'

• 'first'

• 'either'

3.8.8. Swizzle Names

The swizzle names are used in vector access expressions:

3.9. Template Lists

Template parameterization is a way to specify parameters that modify a general concept. To write a template parameterization, write the general concept, followed by a template list.

Ignoring comments and blankspace, a template list is:

• An initial '<' (U+003C) code point, then

• A comma-separated list of one or more template parameters, then

• An optional trailing comma, then

• A terminating '>' (U+003E) code point.

The form of a template parameter is implicitly defined by the template list discovery algorithm below. Generally, they are names, expressions, or types.

Note: For example, the phrase vec3<f32> is a template parameterization where vec3 is the general concept being modified, and <f32> is a template list containing one parameter, the f32 type. Together, vec3<f32> denotes a specific vector type.

Note: For example, the phrase var<storage,read_write> modifies the general var concept with template parameters storage and read_write.

The '<' (U+003C) and '>' (U+003E) code points that delimit a template list are also used when spelling:

• A comparison operator in a relational_expression.

• A shift operator in a shift_expression.

• A compound_assignment_operator for performing a shift operation followed by an assignment.

The syntactic ambiguity is resolved in favour of template lists:

• Template lists are discovered in an early phase of parsing, before declarations, expressions, statements are parsed.

• During tokenization in a later phase, the intial '<' (U+003C) of a template list is mapped to a _template_args_start token, and the terminating '>' (U+003E) of a template list is mapped to a _template_args_end token.

The template list discovery algorithm is given below. It uses the following assumptions and properties:

1. A template parameter is an expression, and therefore does not start with either a '<' (U+003C) or a '=' (U+003D) code point.

2. An expression does not contain code points ';' (U+003B), '{' (U+007B), or ':' (U+003A).

3. An expression does not contain an assignment.

4. The only time a '=' (U+003D) code point appears is as part of a comparison operation, i.e. in one of '<=', '>=', '==', or '!='. Otherwise, a '=' (U+003D) code point appears as part of an assignment.

5. Template list delimiters respect nested expressions formed by parentheses '(...)', and array indexing '[...]'. The start and end of a template list must appear at the same nesting level.

Algorithm: Template list discovery

Input: The program source text.

Record types:

Let UnclosedCandidate be a record type containing:

• position, a location in the source text

• depth, an integer, the expression nesting depth at position

Let TemplateList be a record type containing:

• start_position, the source location of the '<' (U+003C) code point that starts this template list.

• end_position, the source location of the '>' (U+003E) code point that ends this template list.

Output: DiscoveredTemplateLists, a list of TemplateList records.

Procedure:

• Initialize DiscoveredTemplateLists to an empty list.

• Initialize a Pending variable to be an empty stack of UnclosedCandidate records.

• Initialize a CurrentPosition integer variable to 0. It encodes the position of the code point currently being examined, as a count of the number of code points after the start of the source text.

  • This variable will advance forward in the text while executing the algorithm. When the end of text is reached, terminate the algorithm immediately and have it return DiscoveredTemplateLists.

• Initialize a NestingDepth integer variable to 0.

• Repeat the following steps:

  • Advance CurrentPosition past blankspace, comments, and literals.

  • If ident_pattern_token matches the text at CurrentPosition, then:

    • Advance CurrentPosition past the ident_pattern_token.

    • Advance CurrentPosition past blankspace and comments, if present.

    • If '<' (U+003C) appears at CurrentPosition, then:

      • Note: This code point is a candidate for being the start of a template list. Save enough state so it can be matched against a terminating '>' (U+003E) appearing later in the input.

      • Push UnclosedCandidate(position=CurrentPosition,depth=NestingDepth) onto the Pending stack.

      • Advance CurrentPosition to the next code point.

      • If '<' (U+003C) appears at CurrentPosition, then:

        • Note: From assumption 1, no template parameter starts with '<' (U+003C), so the previous code point cannot be the start of a template list. Therefore the current and previous code point must be '<<' operator.

        • Pop the top entry from the Pending stack.

        • Advance CurrentPosition past this code point, and start the next iteration of the loop.

      • If '=' (U+003D) appears at CurrentPosition, then:

        • Note: From assumption 1, no template parameter starts with '=' (U+003C), so the previous code point cannot be the start of a template list. Assume the current and previous code point form a '<=' comparison operator. Skip over the '=' (U+003D) code point so a later step does not mistake it for an assignment.

        • Pop the top entry from the Pending stack.

        • Advance CurrentPosition past this code point, and start the next iteration of the loop.

    • Start the next iteration of the loop.

  • If '>' (U+003E) appears at CurrentPosition then:

    • Note: This code point is a candidate for being the end of a template list.

    • If Pending is not empty, then let T be its top entry, and if T.depth equals NestingDepth then:

      • Note: This code point ends the current template list whose start is recorded in T.

      • Add TemplateList(start_position=T.position, end_position=CurrentPosition) to DiscoveredTemplateLists.

      • Pop T off the Pending stack.

      • Advance CurrentPosition past this code point, and start the next iteration of the loop.

    • Otherwise, this code point does not end a template list:

      • Advance CurrentPosition past this code point.

      • If '=' (U+003D) appears at CurrentPosition then:

        • Note: Assume the current and previous code points form a '>=' comparison operator. Skip over the '=' (U+003D) code point so a later step does not mistake it for an assignment.

        • Advance CurrentPosition past this code point.

      • Start the next iteration of the loop.

  • If '(' (U+0028) or '[' (U+005B) appears at CurrentPosition then:

    • Note: Enter a nested expression.

    • Add 1 to NestingDepth.

    • Advance CurrentPosition past this code point, and start the next iteration of the loop.

  • If ')' (U+0029) or ']' (U+005D) appears at CurrentPosition then:

    • Note: Exit a nested expression.

    • Pop entries from the Pending stack until it is empty, or until the its top entry has depth < NestingDepth.

    • Set NestingDepth to 0 or NestingDepth − 1, whichever is larger.

    • Advance CurrentPosition past this code point, and start the next iteration of the loop.

  • If '!' (U+0021) appears at CurrentPosition then:

    • Advance CurrentPosition past this code point.

    • If '=' (U+003D) appears at CurrentPosition then:

      • Note: Assume the current and previous code points form a '!=' comparison operator. Skip over the '=' (U+003D) code point so a later step does not mistake it for an assignment.

      • Advance CurrentPosition past this code point.

    • Start the next iteration of the loop.

  • If '=' (U+003D) appears at CurrentPosition then:

    • Advance CurrentPosition past this code point.

    • If '=' (U+003D) appears at CurrentPosition then:

      • Note: Assume the current and previous code points form a '==' comparison operator. Skip over the '=' (U+003D) code point so a later step does not mistake it for an assignment.

      • Advance CurrentPosition past this code point, and start the next iteration of the loop.

    • Note: Assume this code point is part of an assignment, which cannot appear as part of an expression, and therefore cannot appear in a template list. Clear pending unclosed candidates.

    • Set NestingDepth to 0.

    • Remove all entries from the Pending stack.

    • Advance CurrentPosition past this code point, and start the next iteration of the loop.

  • If ';' (U+003B) or '{' (U+007B) or ':' (U+003A) appears at CurrentPosition then:

    • Note: These cannot appear in the middle of an expression, and therefore cannot appear in a template list. Clear pending unclosed candidates.

    • Set NestingDepth to 0.

    • Remove all entries from the Pending stack.

    • Advance CurrentPosition past this code point, and start the next iteration of the loop.

  • If '&&' or '||' matches the text at CurrentPosition then:

    • Note: These are operators that have lower precedence than comparisons. Reject any pending unclosed candidates at the current expression level.

    • Note: With this rule, no template list will be found in the program fragment a<b || c>d. Instead it will be recognized as the short-circuiting disjunction of two comparisons.

    • Pop entries from the Pending stack until it is empty, or until the its top entry has depth < NestingDepth.

    • Advance CurrentPosition past the two code points, and start the next iteration of the loop.

  • Advance CurrentPosition past the current code point.

Note: The algorithm explicitly skips past literals because some numeric literals end in a letter, for example 1.0f. The terminating f should not be mistaken as the start of an ident_pattern_token.

Note: In the phrase A ( B < C, D > ( E ) ), the segment < C, D > is a template list.

Note: The algorithm respects expression nesting: The start and end of a particular template list cannot appear at different expression nesting levels. For example, in array<i32,select(2,3,a>b)>, the template list has three parameters, where the last one is select(2,3,a>b). The '>' in a>b does not terminate the template list because it is enclosed in a parenthesized part of the expression calling the select function.

Note: Both ends of a template list must appear within the same indexing expression. For example a[b<d]>() does not contain a valid template list.

Note: In the phrase A<B<<C>, the phrase B<<C is parsed as B followed by the left-shift operator '<<' followed by C. The template discovery algorithm starts examining B then '<' (U+003C) but then sees that the next '<' (U+003C) code point cannot start a template argument, and so the '<' immediately after the B is not the start of a template list. The initial '<' and final '>' are the only template list delimiters, and it has template parameter B<<C.

Note: The phrase A<B<=C> is analyzed similarly to the previous note, so the phrase B<=C is parsed as B followed by the less-than-or-equal operator '<=' followed by C. The template discovery algorithm starts examining B then '<' (U+003C) but then sees that the next '=' (U+003D) code point cannot start a template argument, and so the '<' immediately after the B is not the start of a template list. The initial '<' and final '>' are the only template list delimiters, and it has template parameter B<=C.

Note: When examining the phrase A<(B>=C)>, there is one template list, starting at the first '<' (U+003C) code point and ending at the last '>' (U+003E) code point, and having template argument B>=C. After examining the first '>' (U+003C) code point (after B), the '=' (U+003D) code point needs to be recognized specially so it isn’t assumed to be part of an assignment.

Note: When examining the phrase A<(B!=C)>, there is one template list, starting at the first '<' (U+003C) code point and ending at the last '>' (U+003E) code point, and having template argument B!=C. After examining the '!' (U+0021) code point (after 'B'), the '=' (U+003D) code point needs to be recognized specially so it isn’t assumed to be part of an assignment.

Note: When examining the phrase A<(B==C)>, there is one template list, starting at the first '<' (U+003C) code point and ending at the last '>' (U+003E) code point, and having template argument B==C. After examining the first '=' (U+003D) code point (after 'B'), the second '=' (U+003D) code point needs to be recognized specially so neither are assumed to be part of an assignment.

After template list discovery completes, parsing will attempt to match each template list to the template_list grammar rule.

4. 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 or const assertions.

4.1. Extensions

WGSL is expected to evolve over time.

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

• Addition of new concepts and behaviors via new syntax, including:

  • declarations, statements, attributes, and built-in functions.

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

• Syntax for reducing the set of permissible behaviors.

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

• A description of how the extension interacts with the existing specification, and optionally with other extensions.

Hypothetically, extensions could:

• Add numeric scalar types, such as different bit width integers.

• Add syntax to constrain floating point rounding mode.

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

• Add new kinds of statements.

• Add new built-in functions.

• Add syntax to constrain how shader invocations execute.

• Add new shader stages.

There are two kinds of extensions: enable-extensions and language extensions.

4.1.1. Enable Extensions

An enable-extension is an extension whose functionality is available only if:

• The implementation supports it, and

• The shader explicitly requests it via an enable directive, and

• The corresponding WebGPU GPUFeatureName was one of the required features requested when creating the GPUDevice.

Enable-extensions are intended to expose hardware functionality that is not universally available.

An enable directive is a directive that turns on support for one or more enable-extensions. A shader-creation error results if the implementation does not support all the listed enable-extensions.

Like other directives, if an enable directive is present, it must appear before all declarations and const assertions. Extension names are not identifiers: they do not resolve to declarations.

The valid enable-extensions are listed in the following table.

// Enable a hypothetical extension for arbitrary precision floating point types.
enable arbitrary_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 value 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.
}

4.1.2. Language Extensions

A language extension is an extension which is automatically available if the implementation supports it. The program does not have to explicitly request it.

Language extensions embody functionality which could reasonably be supported on any WebGPU implementation. If the feature is not universally available, that it is because some WebGPU implementation has not yet implemented it.

Note: For example, do-while loops could be a language extension.

The wgslLanguageFeatures member of the WebGPU GPU object lists the set of language extensions supported by the implementation.

A requires-directive is a directive that documents the program’s use of one or more language extensions. It does not change the functionality exposed by the implementation. A shader-creation error results if the implementation does not support one of the required extensions.

A WGSL module can use a requires-directive to signal the potential for non-portability, and to signal the intended minimum bar for portability.

Note: Tooling outside of a WebGPU implementation could check whether all the language extensions used by a program are covered by requires-directives in the program.

Like other directives, if a requires-directive is present, it must appear before all declarations and const assertions. Extension names are not identifiers: they do not resolve to declarations.

Note: The intent is that, over time, WGSL will define language extensions embodying all functionality in language extensions commonly supported at that time. In a requires-directive, these serve as a shorthand for listing all those common features. They represent progressively increasing sets of functionality, and can be thought of as language versions, of a sort.

4.2. Global Diagnostic Filter

A global diagnostic filter is a diagnostic filter whose affected range is the whole WGSL module. It is a directive, thus appearing before any module-scope declarations. It is spelled like the attribute form, but without the leading @ (U+0040) code point, and with a terminating semicolon.

5. Declaration and Scope

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

• a type

• a type-generator

• a value

• a variable

• a function

• a formal parameter

• an enumerant

In other words, a declaration introduces a name for an object.

A declaration is at module scope if the declaration appears in the program source, but outside the text of any other declaration.

A function declaration appears at module-scope. A function declaration contains declarations for formal parameters, if it has any, and it may contain variable and value declarations inside its body. Those contained declarations are therefore not at module-scope.

Note: The only kind of declaration that contain another declaration is a function declaration.

Certain objects are provided by the WebGPU implementation, and are treated as if they have been declared before the start of the WGSL module source. We say such objects are predeclared. For example, WGSL predeclares:

• built-in functions,

• built-in types such as i32 and f32,

• built-in type-generators such as array, ptr, and texture_2d, and

• enumerants such as read_write, workgroup, and rgba8unorm.

The scope of a declaration is the set of program source locations where a declared identifier potentially denotes its associated object. We say the identifier is in scope (of the declaration) at those source locations.

Where a declaration appears determines its scope:

• Predeclared objects, and objects declared at module-scope, are in scope across the entire program source.

• Each formal parameter of a user-declared function is in scope across the corresponding function body. See § 11.1 Declaring a User-defined Function.

• Otherwise, the scope is a span of text beginning immediately after the end of the declaration. For details, see § 7 Variable and Value Declarations.

Two declarations in the same WGSL source program must not simultaneously:

• introduce the same identifier name, and

• have the same end-of-scope.

Note: A predeclared object does not have a declaration in the WGSL source. So a user-specified declaration at module-scope or inside a function can have the same name as a predeclared object.

Identifiers are used as follows, distinguished by grammatical context:

• A token matching the ident grammar element is:

  • Used in a declaration, as the name of the object being declared, or

  • Used as a name, denoting an object declared elsewhere. This is the common case.

• A token matching the member_ident grammar element is:

  • Used in a structure type declaration, as the name of a member, or

  • Used as a name, denoting a member of a structure value, or denoting a reference to a member of a structure. See § 8.5.4 Structure Access Expression.

When an ident token appears as a name denoting an object declared elsewhere, it must be in scope for some declaration. The object denoted by the identifier token is determined as follows:

• If the token is in scope for at least one non-module-scope declaration, then the token denotes the object associated with the nearest of those declarations.

  Note: The nearest such declaration appears before the identifier token.

• Otherwise, if there is a module-scope declaration with that name, then the token denotes that declared object.

  Note: The module-scope declaration may appear before or after the identifier token.

• Otherwise, if there is a predeclared object with that name, then the token denotes that object.

When the above algorithm is used to map an identifier to a declaration, we say the identifier resolves to that declaration. Similarly, we also say the identifier resolves to the declared object.

It is a shader-creation error if any module scope declaration is recursive. That is, no cycles can exist among the declarations:

Consider the directed graph where:

• Each node corresponds to a declaration D.

• There is an edge from declaration D to declaration T when the definition for D mentions an identifier which resolves to T.

This graph must not have a cycle.

Note: The function body is part of the function declaration, thus functions must not be recursive, either directly or indirectly.

Note: Non-module scope identifier declarations must precede their uses in the text.

// Valid, user-defined variables can have the same name as a built-in function.
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, modf resolves to the module-scope variable.
  let res = modf(foo);

  // 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, foo_4 is in scope until the end of the compound statement.
    var foo : f32; // foo_4

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

// This declaration hides the predeclared 'min' built-in function.
// Since this declaration is at module-scope, it is in scope over the entire
// source.  The built-in function is no longer accessible.
fn min() -> u32 { return 0; }

const rgba8unorm = 12; // This shadows the predeclared 'rgba8unorm' enumerant.

6. Types

Programs calculate values.

In WGSL, a type is a set of values, and each value belongs to exactly one type. A value’s type determines the syntax and semantics of operations that can be performed on that value.

For example, the mathematical number 1 corresponds to these distinct values in WGSL:

• the 32-bit signed integer value 1i,

• the 32-bit unsigned integer value 1u,

• the 32-bit floating point value 1.0f,

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

• the AbstractInt value 1, and

• the AbstractFloat value 1.0

WGSL treats these as different because their machine representation and operations differ.

A type is either predeclared, or created in WGSL source via a declaration.

Some types are expressed as template parameterizations. A type-generator is a predeclared object which, when parameterized with a template list, denotes a type. For example, the type atomic<u32> combines the type-generator atomic with template list <u32>.

We distinguish between the concept of a type and the syntax in WGSL to denote that type. In many cases the spelling of a type in this specification is the same as its WGSL syntax. For example:

• the set of 32-bit unsigned integer values is spelled u32 in this specification, and also in a WGSL module.

• the spelling is different for structure types, or types containing structures.

Some WGSL types are only used for analyzing a source program and for determining the program’s runtime behavior. This specification will describe such types, but they do not appear in WGSL source text.

Note: Reference types are not written in WGSL modules. See § 6.4.3 Reference and Pointer Types.

6.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 § 8.18 Expression Grammar Summary.

The particular value produced by an expression evaluation depends on:

• static context: the source text surrounding the expression, and

• dynamic context: the state of the invocation evaluating the expression, and the execution context in which the invocation is running.

The values that may result from evaluating a particular expression will always belong to a specific WGSL type, known as the static type of the expression. The rules of WGSL are designed so that the static type of an expression depends only on the expression’s static context.

A type assertion is a mapping from some WGSL source expression to a WGSL type. The notation

e : T

is a type assertion meaning T is the static type of WGSL expression e.

Note: A type assertion is a statement of fact about the text of a program. It is not a runtime check.

Statements often use expressions, and may place requirements on the static types of those expressions. For example:

• The condition expression of an if statement must be of type bool.

• In a let declaration with an explicit type specified, the initializer expression must evaluate to that type.

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

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

• A conclusion.

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

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

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

• Preconditions, consisting of:

  • For expressions:

    • Type assertions for subexpressions, when it has subexpressions. Each may be satisfied directly, or via a feasible automatic conversion (as defined in § 6.1.2 Conversion Rank). WGSL will evaluate all const-expressions required to determine the static types of a program.

    • How the expression is used in a statement.

  • For statements:

    • The syntactic form of the statement, and

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

  • Conditions on the other schematic parameters, if any.

  • Optionally, other static context.

Type rules may have type parameters in their preconditions and conclusions. When a type rule’s conclusion or preconditions contain type parameters, we say it is parameterized. When they do not, we say the rule is fully elaborated. We can make a fully elaborated type rule from a parameterized one by substituting a type for each of its type parameters, using the same type for all occurrences of a given parameter in the rule. An assignment of types to a rule’s type parameters is called a substitution.

For example, here is the type rule for logical negation (an expression of the form !e):

This is a parameterized rule, because it contains the type parameter T, which can represent any one of four types bool, vec2<bool>, vec3<bool>, or vec4<bool>. Applying the substitution that maps T to vec3<bool> produces the fully elaborated type rule:

Each fully elaborated rule we can produce from a parameterized rule by applying some substitution that meets the rule’s other conditions is called an overload of the parameterized rule. For example, the boolean negation rule has four overloads, because there are four possible ways to assign a type to its type parameter T.

Note: In other words, a parameterized type rule provides the pattern for a collection of fully elaborated type rules, each one produced by applying a different substitution to the parameterized rule.

A type rule applies to a syntactic phrase when:

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

• The rule’s preconditions are satisfied.

A parameterized type rule applies to an expression if there exists a substitution producing a fully elaborated type rule that applies to the expression.

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 rules in § 8.7 Arithmetic Expressions, the type rule for addition applies to the expression, because:

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

• e1 is of type u32, and

• e2 is of type u32, and

• we can substitute u32 for the type parameter T in the type rule, resulting in a fully elaborated rule that applies to the entire expression.

When analyzing a syntactic phrase, three cases may occur:

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

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

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

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

  • If overload resolution fails, a type error results.

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

A WGSL source program is well-typed when:

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

• The type requirements for each statement are satisfied.

Otherwise there is a type error and the source program is not a valid WGSL module.

WGSL is a statically typed language because type checking a WGSL module will either succeed or discover a type error, while only having to inspect the program source text.

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

fn foo(p : ptr<function, i32>) -> i32 {
  let x = *p;
  *p += 1;
  return x;
}

fn bar() {
  var a: i32;
  let x = foo(&a); // the call to foo returns a value
                   // and updates the value of a
}

6.1.2. Conversion Rank

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

• e is already of type T, or

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

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

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

Note: Automatic conversions only occur in two kinds of situations. First, when converting a const-expression 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.

The type T is the concretization of type S if:

• T is concrete, and

• T is not a reference type, and

• ConversionRank(S, T) is finite, and

• For any other non-reference type T2, ConversionRank(S, T2) > ConversionRank(S, T).

The concretization of a value e of type T is the value resulting from applying, to e, the feasible conversion that maps T to the concretization of T.

Note: Conversion to f32 is always preferred over f16, therefore automatic conversion will only ever produce an f16 if extension is enabled in the module.

6.1.3. Overload Resolution

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

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

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

• Its conclusion has:

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

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

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

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

   • Its static type has been computed.

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

2. Eliminate any candidate where one of its subexpressions resolves to an abstract type after feasible automatic conversions, but another of the candidate’s subexpressions is not a const-expression.

   Note: As a consequence, if any subexpression in the phrase is not a const-expression, then all subexpressions in the phrase must have a concrete type.

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

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

6.2. 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 and abstract numeric types.

6.2.1. Abstract Numeric Types

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

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

WGSL defines two abstract numeric types for these evaluations:

• The AbstractInt type is the set of integers representable in the 64-bit two’s complement format, with the sign bit in the most significant bit position.

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

An evaluation of an expression in one of these types must not overflow or produce an infinite or NaN value.

A type is abstract if it is an abstract numeric type or contains an abstract numeric type. A type is concrete if it is not abstract.

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

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

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

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

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

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

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

  • AbstractInt to AbstractFloat. (Conversion rank 4)

  • AbstractInt to f32. (Conversion rank 5)

  • AbstractInt to f16. (Conversion rank 6)

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

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

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

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

• However, using feasible automatic conversions, there are three potential overloads:

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

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

  • 1 is converted to f16 value 1.0f (rank 6) and 2.5 is converted to f16 value 2.5h (rank 2).

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

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

Example: let x = 1 + 2.5;

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

• Therefore, there are two viable overload candidates: addition using f32 or f16.

• The preferable candidate uses f32.

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

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

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

• The 2.5 term is an expression of type AbstractFloat.

• There are no valid overload candidates:

  • There is no feasible automatic conversion from a GPU-materialized integer scalar type to a floating point type.

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

Example: -1 * i32(-2147483648) does not result in a shader-creation error:

• The -1 term is an expression of type AbstractInt.

• The i32(-2147483648) term is an expression of type i32.

• There is no overload of the multiply operator for these two types and the i32 term cannot be up-converted to AbstractInt.

• The only feasible automatic conversion is converting the AbstractInt to i32 so:

  • The resulting computation is a multiplication operation between i32.

  • i32 operations use two’s complement arithmetic and have defined overflow behavior.

  • Wrapping occurs.

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

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

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

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

// When a concrete type is required, but no part of the statement or
// expression forces a particular concrete type, an integer literal is
// interpreted as an i32 value:
//   Initializer for a let-declaration must be constructible (or pointer).
//   The most preferred automatic conversion from AbstractInt to a constructible type
//   is AbstractInt to i32, with conversion rank 2.  So '1' is inferred as i32.
let some_i32 = 1; // like let some_i32: i32 = 1i;

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

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

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

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

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

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

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

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

// Inference based on built-in function parameters.

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

// The following examples all promote to f32 with an initial value of 10f.
let f32_promotion1 = 1.0 + 2 + 3 + 4;
let f32_promotion2 = 2 + 1.0 + 3 + 4;
let f32_promotion3 = 1f + ((2 + 3) + 4);
let f32_promotion4 = ((2 + (3 + 1f)) + 4);

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

6.2.2. Boolean Type

The bool type contains the values true and false.

6.2.3. 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 the two’s complement representation, with the sign bit in the most significant bit position.

Expressions on concrete integer types that overflow produce a result that is modulo 2^bitwidth

Note: AbstractInt is also an integer type.

6.2.4. Floating Point Types

The f32 type is the set of 32-bit floating point values of the IEEE-754 binary32 (single precision) format. See § 15.7 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 § 15.7 Floating Point Evaluation for details.

The following table lists certain extreme values for floating point types. Each has a corresponding negative value.

Note: AbstractFloat is also a floating point type.

6.2.5. Scalar Types

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

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

The integer scalar types are AbstractInt, i32, and u32.

A scalar conversion maps a value in one scalar type to a value in a different scalar type. Generally the result value is close to the original value, within the limitations of the destination type. Scalar conversions occur either:

• By explicitly invoking a value constructor, or

• When converting a const-expression in an abstract numeric type to another type, via a feasible automatic conversion.

6.2.6. Vector Types

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

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

Key use cases of a vector include:

• to express both a direction and a magnitude.

• to express a position in space.

• to express a color in some color space. For example, the components could be intensities of red, green, and blue, while the fourth component could be an alpha (opacity) value.

Many operations on vectors (and matrices) act component-wise, i.e. the result is formed by operating on each scalar component independently.

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

let x : vec3<f32> = a + b; // a and b are vec3<f32>
// x[0] = a[0] + b[0]
// x[1] = a[1] + b[1]
// x[2] = a[2] + b[2]

WGSL also predeclares the following type aliases:

6.2.7. Matrix Types

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

The key use case for a matrix is to embody a linear transformation. In this interpretation, the vectors of a matrix are treated as column vectors.

The product operator (*) is used to either:

• scale the transformation by a scalar magnitude.

• apply the transformation to a vector.

• combine the transformation with another matrix.

See § 8.7 Arithmetic Expressions.

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

WGSL also predeclares the following type aliases:

6.2.8. Atomic Types

An atomic type encapsulates a concrete integer scalar type such that:

• atomic objects provide certain guarantees to concurrent observers, and

• the only valid operations on atomic objects are the atomic builtin functions.

An expression must not evaluate to an atomic type.

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

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

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

6.2.9. Array Types

An array is an indexable sequence of element values.

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

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

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

• It must be an override-expression.

• It must evaluate to a concrete integer scalar.

• If N is not greater than zero:

  • It is a shader-creation error if N is a const-expression.

  • Otherwise, it is a pipeline-creation error.

Note: The element count value is fully determined at pipeline creation time if N depends on any override-declarations, and shader module creation otherwise.

Note: To qualify for type-equivalency, any override expression that is not a const expression must be an identifier. See Workgroup variables sized by overridable constants

The number of elements in a runtime-sized array is determined by the size of buffer binding associated with the corresponding storage buffer variable. See § 13.3.4 Buffer Binding Determines Runtime-Sized Array Element Count.

An array element type must be one of:

• a scalar type

• a vector type

• a matrix type

• an atomic type

• an array type having a creation-fixed footprint

• a structure type having a creation-fixed footprint.

Note: The element type must be a plain type.

Two array types are the same if and only if all of the following are true:

• They have the same element type.

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

  • They are both runtime-sized.

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

  • They are both fixed-sized with element counts specified as identifiers resolving to the same declaration of a pipeline-overridable constant.

// array<f32,8> and array<i32,8> are different types:
// different element types
var<private> a: array<f32,8>;
var<private> b: array<i32,8>;
var<private> c: array<i32,8u>;  // array<i32,8> and array<i32,8u> are the same type

const width = 8;
const height = 8;

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

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

Note: The only valid use of an array type sized by an overridable constant is as a memory view in the workgroup address space. This includes the store type of a workgroup variable. See § 7 Variable and Value Declarations.

override blockSize = 16;

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

// None of the following have the same type as 'odds' and 'evens'.

// Different type: Not the identifier 'blockSize'
var<workgroup> evens_0: array<i32,16>;
// Different type: Uses arithmetic to express the element count.
var<workgroup> evens_1: array<i32,(blockSize * 2 / 2)>;
// Different type: Uses parentheses, not just an identifier.
var<workgroup> evens_2: 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>;

6.2.10. Structure Types

A structure is a named grouping of named member values.

Structure types are declared at module scope. Elsewhere in the program source, a structure type is denoted by its identifier name. See § 5 Declaration and Scope.

Two structure types are the same if and only if they have the same name.

A structure member type must be one of:

• a scalar type

• a vector type

• a matrix type

• an atomic type

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

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

• a structure type that has a creation-fixed footprint

Note: All user-declared structure types are concrete.

Note: Each member type must be a plain type.

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

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

• When a runtime-sized array is part of a larger type, it may only appear as the last element of a structure, which itself cannot be part of an enclosing array or structure.

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

// Declare a variable storing a value of type Data.
var<private> some_data: Data;

The following attributes can be applied to structure members:

• align

• builtin

• location

• blend_src

• interpolate

• invariant

• size

Attributes builtin, location, blend_src, interpolate, and invariant are IO attributes. An IO attribute on a member of a structure S has effect only when S is used as the type of a formal parameter or return type of an entry point. See § 13.3.1 Inter-stage Input and Output Interface.

Attributes align and size are layout attributes, and may be required if the structure type is used to define a uniform buffer or a storage buffer. See § 14.4 Memory Layout.

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

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

6.2.11. 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. A composite value may be decomposed into its components. See § 8.5 Composite Value Decomposition Expressions.

The composite types are:

• vector type

• matrix type

• array type

• structure type

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

• 1 for a vector type

• 2 for a matrix type

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

• 1 + max(NestDepth(M₁),..., NestDepth(M_N)) if T is a structure type with member types M₁,...,M_N

6.2.12. Constructible Types

Many kinds of values can be created, loaded, stored, passed into functions, and returned from functions. We call these constructible.

A type is constructible if it is one of:

• a scalar type

• a vector type

• a matrix type

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

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

Note: All constructible types have a 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.

6.2.13. 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 type has a creation-fixed footprint if its concretization has a size that is fully determined at shader creation time.

A type has a fixed footprint if its size is fully determined at pipeline creation time.

Note: All concrete creation-fixed footprint and fixed footprint types are storable.

Note: Pipeline creation depends on shader creation, so a type with creation-fixed footprint also has fixed footprint.

The types with creation-fixed footprint are:

• a scalar type

• a vector type

• a matrix type

• an atomic type

• a fixed-size array type, when:

  • its element count is a const-expression.

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

Note: A constructible type has a creation-fixed footprint.

The plain types with fixed footprint are any of:

• a type with creation-fixed footprint

• a fixed-size array type (without further constraining its element count)

Note: The only valid use of a fixed-size array with an element count that is an override-expression that is not a const-expression is as a memory view in the workgroup address space. This includes the store type of a workgroup variable.

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

Note: Fixed-footprint types exclude runtime-sized arrays, and any structure that contains a runtime-sized array.

6.3. Enumeration Types

An enumeration type is a limited set of named values. An enumeration is used to distinguish among the set of possibilities for a specific concept, such as the set of valid texel formats.

An enumerant is one of the named values in an enumeration. Each enumerant is distinct from all other enumerants, and distinct from all other kinds of values.

There is no mechanism for declaring new enumerants or new enumeration types in WGSL source.

Note: Enumerants are used as template parameters.

6.3.1. Predeclared enumerants

The following table lists the enumeration types in WGSL, and their predeclared enumerants. The enumeration types exist, but cannot be spelled in WGSL source.

6.4. Memory Views

In addition to calculating with plain values, a WGSL program will also often read values from memory or write values to memory, via memory access operations. Each memory access is performed via a memory view.

A memory view comprises:

• a set of memory locations in a particular address space,

• a memory model reference,

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

• an access mode.

The access mode of a memory view must be supported by the address space. See § 7 Variable and Value Declarations.

6.4.1. 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 § 14.4.4 Internal Layout of Values, or it may be opaque, such as for textures and samplers.

A type is storable if it is both concrete and one of:

• a scalar type

• a vector type

• a matrix type

• an atomic type

• an array type

• a structure type

• a texture type

• a sampler type

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

6.4.2. Host-shareable Types

Host-shareable types are used to describe the contents of buffers which are shared between the host and the GPU, or copied between host and GPU without format translation. When used for this purpose, the type may additionally have layout attributes applied as described in § 14.4 Memory Layout. As described in § 7.3 var Declarations, the store type of uniform buffer and storage buffer variables must be host-shareable.

A type is host-shareable if it is both concrete and one of:

• a numeric scalar type

• a numeric vector type

• a matrix type

• an atomic type

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

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

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

Note: Restrictions on the types of inter-stage inputs and outputs are described in § 13.3.1 Inter-stage Input and Output Interface and subsequent sections. Those types are also sized, but the counting differs.

Note: Textures and samplers can also be shared between the host and the GPU, but their contents are opaque. The host-shareable types in this section are specifically for use in storage and uniform buffers.

6.4.3. Reference and Pointer Types

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

Two pointer types are the same if and only if they have the same address space, store type, and access mode.

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

However, in WGSL source text:

• Reference types must not appear.

• Pointer types may appear.

  • A pointer type is spelled with parameterization by:

    • address space,

    • store type, and

    • sometimes by access mode, as specified in § 14.3 Address Spaces.

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

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

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

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

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

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

6.4.4. Valid and Invalid Memory References

A reference value is either valid or invalid.

References are formed as described in detail in § 6.4.8 Forming Reference and Pointer Values. Generally, a valid reference is formed by:

• naming a variable, or

• applying the indirection (unary *) operation to a valid pointer, or

• a named component expression where the base is a valid memory view, or

• an indexing expression where the base is a valid memory view, and using an in-bounds index.

Generally, an invalid memory reference is formed by:

• applying the indirection operator to an invalid pointer, or

• a named component expression where the base is an invalid memory reference, or

• an indexing expression where the base is a memory view, and either:

  • the base is an invalid memory reference, or

  • the index is out-of-bounds.

A valid pointer is a pointer that corresponds to a valid reference. An invalid pointer is a pointer that corresponds to an invalid memory reference.

6.4.5. Originating Variable

The originating variable of a pointer value is defined as the originating variable of the corresponding reference value.

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.

A valid reference always corresponds to a non-empty memory view for some or all of the memory locations for some variable.

struct Particle {
   position: vec2f,
   velocity: vec2f,
   color_index: i32,
}

@group(0) @binding(0)
var<storage,read_write> the_particle: Particle;

fn particle_velocity_component(p: Particle, i: i32) -> f32 {
  return the_particle.velocity[i]; // A valid reference when i is 0 or 1.
}

6.4.6. Out-of-Bounds Access

An operation that accesses an invalid memory reference is an out-of-bounds access.

An out-of-bounds access is a program defect, because if it were performed as written, it would typically:

• read or write memory locations outside of a variable, or

• interpret the contents of those locations as the wrong store type, or

• cause an unintended data race.

For this reason, an implementation will not perform the access as written. Executing an out-of-bounds access generates a dynamic error.

Note: An example of interpreting the store type incorrectly occurs in the example from the previous section. When i is 2, the expression the_particle.velocity[i] has type ref<storage,f32,read_write>, meaning it is a memory view with f32 as its store type. However, the memory locations are allocated to for the color_index member, so the stored value is actually of type i32.

6.4.7. Use Cases for References and Pointers

References and pointers are distinguished by how they are used:

• The type of a variable is a reference type.

• The address-of operation (unary &) converts a reference value to its corresponding pointer value.

• The indirection operation (unary *) converts a pointer value to its corresponding reference value.

• A let-declaration can be of pointer type, but not of reference type.

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

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

  • The left-hand side of the assignment statement must be of reference type, with access mode write or read_write.

  • The right-hand side of the assignment statement must evaluate to the store type of the left-hand side.

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

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

  • r has an access mode of read or read_write, and

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

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

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

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

@compute @workgroup_size(1)
fn main() {
  // 'i' has reference type ref<function,i32,read_write>
  // The memory locations for 'i' store the i32 value 0.
  var i: i32 = 0;

  // 'i + 1' can only match a type rule where the 'i' subexpression is of type i32.
  // So the expression 'i + 1' has type i32, and at evaluation, the 'i' subexpression
  // evaluates to the i32 value stored in the memory locations for 'i' at the time
  // of evaluation.
  let one: i32 = i + 1;

  // Update the value in the locations referenced by 'i' so they hold the value 2.
  i = one + 1;

  // Update the value in the locations referenced by 'i' so they hold the value 5.
  // The evaluation of the right-hand-side occurs before the assignment takes effect.
  i = i + 3;
}

var<private> age: i32;
fn get_age() -> i32 {
  // The type of the expression in the return statement must be 'i32' since it
  // must match the declared return type of the function.
  // The 'age' expression is of type ref<private,i32,read_write>.
  // Apply the Load Rule, since the store type of the reference matches the
  // required type of the expression, and no other type rule applies.
  // The evaluation of 'age' in this context is the i32 value loaded from the
  // memory locations referenced by 'age' at the time the return statement is
  // executed.
  return age;
}

fn caller() {
  age = 21;
  // The copy_age constant will get the i32 value 21.
  let copy_age: i32 = get_age();
}

Defining pointers in this way enables two key use cases:

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

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

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

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 @workgroup_size(1)
fn main() {
  // Form a pointer to a specific Particle in storage memory.
  let active_particle = &system.particles[system.active_index];

  let delta_position: vec3<f32> = (*active_particle).velocity * system.timestep;
  let current_position: vec3<f32>  = (*active_particle).position;
  (*active_particle).position = delta_position + current_position;
}

fn add_one(x: ptr<function,i32>) {
  /* Update the locations for 'x' to contain the next higher integer value,
     (or to wrap around to the largest negative i32 value).
     On the left-hand side, unary '*' converts the pointer to a reference that
     can then be assigned to. It has a read_write access mode, by default.
     /* On the right-hand side:
        - Unary '*' converts the pointer to a reference, with a read_write
          access mode.
        - The only matching type rule is for addition (+) and requires '*x' to
          have type i32, which is the store type for '*x'.  So the Load Rule
          applies and '*x' evaluates to the value stored in the memory for '*x'
          at the time of evaluation, which is the i32 value for 0.
        - Add 1 to 0, to produce a final value of 1 for the right-hand side. */
     Store 1 into the memory for '*x'. */
  *x = *x + 1;
}

@compute @workgroup_size(1)
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.
}

6.4.8. Forming Reference and Pointer Values

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

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

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

• Use a named component expression on a memory view to a composite:

  • Given a memory view with a vector store type, appending a single-letter vector access phrase results in a reference to the named component of the vector. See § 8.5.1.3 Component Reference from Vector Memory View.

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

• Use an indexing expression on a memory view to a composite:

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

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

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

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

struct S {
    age: i32,
    weight: f32
}
var<private> person: S;
// Elsewhere, 'person' denotes the reference to the memory underlying the variable,
// and will have type ref<private,S,read_write>.

fn f() {
    var uv: vec2<f32>;
    // For the remainder of this function body, 'uv' denotes 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;

    // Alternatively, references can also be formed from pointers using
    // the same syntax.

    let uv_ptr = &uv;
    // For the remainder of this function body, 'uv_ptr' denotes a pointer
    // to the memory underlying 'uv', and will have the type
    // ptr<function,vec2<f32>,read_write>.

    // Evaluate the left-hand side of the assignment:
    //   Evaluate '*uv_ptr' to yield a reference:
    //   1. First evaluate 'uv_ptr', yielding a pointer to the memory for
    //      the 'uv' variable. The result has type ptr<function,vec2<f32>,read_write>.
    //   2. Then apply the indirection expression operator, yielding a
    //      reference to memory for 'uv'.
    // Evaluating the right-hand side of the assignment yields the vec2<f32> value (1.0, 2.0).
    // Store the value (1.0, 2.0) into the storage memory locations referenced by uv.
    *uv_ptr = vec2f(1.0, 2.0);

    // Evaluate the left-hand side of the assignment:
    //   Evaluate 'uv_ptr.x' to yield a reference:
    //   1. First evaluate 'uv_ptr', yielding a pointer to the memory for
    //      the 'uv' variable. The result has type ptr<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_ptr.x = 1.0;

    // Evaluate the left-hand side of the assignment:
    //   Evaluate 'uv_ptr[1]' to yield a reference:
    //   1. First evaluate 'uv_ptr', yielding a pointer to the memory for
    //      the 'uv' variable. The result has type ptr<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_ptr[1] = 2.0;

    let m_ptr = &m;
    // When evaluating 'm_ptr[2]':
    // 1. First evaluate 'm_ptr', yielding a pointer to the memory for
    //    the 'm' variable. The result has type ptr<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_ptr[2];

    let A_ptr = &A;
    // When evaluating 'A[4]'
    // 1. First evaluate 'A', yielding a pointer to the memory for
    //    the 'A' variable. The result has type ptr<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_ptr[4];

    let person_ptr = &person;
    // When evaluating 'person.weight'
    // 1. First evaluate 'person_ptr', yielding a pointer to the memory for
    //    the 'person' variable declared at module scope.
    //    The result has type ptr<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_ptr.weight;
}

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

• Use the address-of (unary &) operator on a reference.

  • The result is a valid pointer if and only if the original reference is valid.

  • The originating variable of a valid result is defined as the originating variable of the reference.

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

  • The value denoted by the formal parameter (at runtime) is a valid pointer if and only if the pointer value at the call site is valid.

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

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

// Declare a variable in the private address space, for storing an f32 value.
var<private> x: f32;

fn f() {
    // Declare a variable in the function address space, for storing an i32 value.
    var y: i32;

    // The name 'x' resolves to the 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;
}