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:
-
Modify the contents of variables.
-
Modify execution order using structured programming constructs:
-
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)
-
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, or 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 over real numbers x extended with +∞ and −∞:
-
⌊ + ∞ ⌋ = +∞
-
⌊ − ∞ ⌋ = −∞
-
for real number x, ⌊x⌋ = k, where k is the unique integer such that k ≤ x < k+1
The ceiling expression is defined over real numbers x extended with +∞ and −∞:
-
⌈ +∞ ⌉ = +∞
-
⌈ −∞ ⌉ = −∞
-
for real number x, ⌈x⌉ = k, where k is the unique integer such that k-1 < x ≤ k
The truncate function is defined over real numbers x extended with +∞ and −∞:
-
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 AT formed by copying the rows of A as the columns of AT:
-
transpose(A) = AT
-
transpose(A)i,j = Aj,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 const_assert statements. 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:
-
Shader module creation
-
This occurs when the WebGPU
createShaderModule()
method is called. The source text for a WGSL program is provided at this time.
-
-
Pipeline creation
-
This occurs when the WebGPU
createComputePipeline()
method or the WebGPUcreateRenderPipeline()
method is invoked. These methods use one or more previously created shader modules, together with other configuration information.
-
-
Shader execution start
-
This occurs when a draw or dispatch command is issued to the GPU, begins executing the pipeline, and invokes the shader stage entry point function.
-
-
-
This occurs when all work in the shader completes:
-
all its invocations terminate, and
-
all accesses to resources complete, 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.
-
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:
-
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.
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:
-
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.
-
-
Discard diagnostics that have severity off.
-
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.
-
-
If at least one remaining diagnostic DW has severity warning, then:
-
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.
-
-
-
If processing during shader module creation time, the remaining diagnostics populate the
messages
member of the WebGPUGPUCompilationInfo
object. -
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.
Filterable Triggering Rule | Default Severity | Triggering Location | Description |
---|---|---|---|
derivative_uniformity | error | The location of the call site for any builtin function that computes a derivative. That is, the location of a call to any of: |
A call to a builtin function computes derivatives, but uniformity analysis cannot prove that the call occurs in uniform control flow.
See § 14.2 Uniformity. |
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.
Placement | Affected Range |
---|---|
Beginning of a compound statement. | The compound statement. |
Beginning of a function declaration. | The function declaration. |
Beginning of an if statement. | The if statement: the if_clause and all associated else_if_clause and else_clause clauses, including all controlling condition expressions. |
Beginning of a switch statement. | The switch statement: the selector expression and the switch_body. |
Beginning of a switch_body. | The switch_body. |
Beginning of a loop statement. | The loop statement. |
Beginning of a while statement. | The while statement: both the condition expression and the loop body. |
Beginning of a for statement. | The for statement: the for_header and the loop body. |
Immediately before the opening brace ('{' ) of the loop body of a loop, while, or for loop.
| The loop body. |
Beginning of a continuing_compound_statement. | The continuing_compound_statement. |
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.
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 unique, or 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.
Limit | Minimum supported value |
---|---|
Maximum number of members in a structure type | 16383 |
Maximum nesting depth of a composite type | 255 |
Maximum nesting depth of brace-enclosed statements in a function | 127 |
Maximum number of parameters for a function | 255 |
Maximum number of case selector values in a switch statement | 16383 |
Maximum byte-size of an array type instantiated in the function or private address spaces
For the purposes of this limit, bool has a size of 1 byte. | 65535 |
Maximum byte-size of an array type instantiated in the workgroup address space.
For the purposes of this limit, bool has a size of 1 byte and a fixed-footprint array is treated as a creation-fixed footprint array when substituting the override value. This maps the WebGPU maxComputeWorkgroupStorageSize limit into a standalone WGSL limit. Note: Several workgroup variables that individually satisfy this limit can still combine to exceed the API limit. | 16384 |
Maximum number of elements in const-expression of array type | 65535 |
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:
The program text must not include a null code point (U+0000
).
3.1. Parsing
To parse a WGSL module:
Remove comments:
Replace the first comment with a space code point (
U+0020
).Repeat until no comments remain.
Find template lists, using the algorithm in § 3.10 Template Lists.
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 byU+002A
) -
Then any sequence of:
-
A block comment, or
-
Text that does not contain either
*/
(U+002A
followed byU+002F
) or/*
(U+002F
followed byU+002A
)
-
-
Then the two code points
*/
(U+002A
followed byU+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.
-
an identifier.
3.5. Literals
A literal is one of:
-
A numeric literal: either an integer literal or a floating point literal, and is used to represent a number.
3.5.1. Boolean Literals
`'true'`
| `'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
or0X
followed by a sequence of hexadecimal digits.
-
-
Then an optional
i
oru
suffix.
`/0[iu]?/`
| `/[1-9][0-9]*[iu]?/`
`/0[xX][0-9a-fA-F]+[iu]?/`
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 mantissa to representing a fraction, and an optional exponent. Roughly, the value of the literal is the mantissa multiplied by a base value raised to the given exponent. A mantissa digit is significant if it is non-zero, or if there are mantissa 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 mantissa, specified as a sequence of digits, with an optional decimal point (
.
) somewhere among them. The mantissa represents a fraction in base 10 notation. -
Then an optional exponent suffix consisting of:
-
e
orE
. -
Then an exponent specified as an decimal number with an optional leading sign (
+
or-
). -
Then an optional
f
orh
suffix.
-
-
At least one of the decimal point, or the exponent, or the
f
orh
suffix must be present. If none are, then the token is instead an integer literal.
`/0[fh]/`
| `/[1-9][0-9]*[fh]/`
| `/[0-9]*\.[0-9]+([eE][+-]?[0-9]+)?[fh]?/`
| `/[0-9]+\.[0-9]*([eE][+-]?[0-9]+)?[fh]?/`
| `/[0-9]+[eE][+-]?[0-9]+[fh]?/`
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 ;
-
Compute effective_mantissa from mantissa:
-
If mantissa has 20 or fewer significant digits, then effective_mantissa is mantissa.
-
Otherwise:
-
Let truncated_mantissa be the same as mantissa except each digit to the right of the 20th significant digit is replaced with 0.
-
Let truncated_mantissa_next be the same as mantissa except:
-
the 20th significant digit is incremented by 1, and carries are propagated to the left as needed needed to ensure each digit remains in the range 0 through 9, and
-
each digit to the right of the 20th significant digit is replaced with 0.
-
-
Set effective_mantissa to either truncated_mantissa or truncated_mantissa_next. This is an implementation choice.
-
-
-
The mathematical value of the literal is the mathematical value of effective_mantissa as a decimal fraction, multiplied by 10 to the power of the exponent. When no exponent is specified, an exponent of 0 is assumed.
Note: The decimal mantissa 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
or0X
prefix -
Then a mantissa, specified as a sequence of hexadecimal digits, with an optional hexadecimal point (
.
) somewhere among them. The mantissa represents a fraction in base 16 notation. -
Then an optional exponent suffix consisting of:
-
p
orP
-
Then an exponent specified as an decimal number with an optional leading sign (
+
or-
). -
Then an optional
f
orh
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.
`/0[xX][0-9a-fA-F]*\.[0-9a-fA-F]+([pP][+-]?[0-9]+[fh]?)?/`
| `/0[xX][0-9a-fA-F]+\.[0-9a-fA-F]*([pP][+-]?[0-9]+[fh]?)?/`
| `/0[xX][0-9a-fA-F]+[pP][+-]?[0-9]+[fh]?/`
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 ;
-
Compute effective_mantissa from mantissa:
-
If mantissa has 16 or fewer significant digits, then effective_mantissa is mantissa.
-
Otherwise:
-
Let truncated_mantissa be the same as mantissa except each digit to the right of the 16th significant digit is replaced with 0.
-
Let truncated_mantissa_next be the same as mantissa except:
-
the 16th significant digit is incremented by 1, and carries are propagated to the left as needed needed to ensure each digit remains in the range 0 through
f
, and -
each digit to the right of the 16th significant digit is replaced with 0.
-
-
Set effective_mantissa to either truncated_mantissa or truncated_mantissa_next. This is an implementation choice.
-
-
-
The mathematical value of the literal is the mathematical value of effective_mantissa as a hexadecimal fraction, multiplied by 2 to the power of the exponent. When no exponent is specified, an exponent of 0 is assumed.
Note: The hexadecimal mantissa 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 § 14.6.4 Floating Point Conversion.
Numeric Literal | Suffix | Type | Examples |
---|---|---|---|
integer literal | i
| i32 | 42i |
integer literal | u
| u32 | 42u |
integer literal | AbstractInt | 124 | |
floating point literal | f
| f32 | 42f 1e5f 1.2f 0x1.0p10f |
floating point literal | h
| f16 | 42h 1e5h 1.2h 0x1.0p10h |
floating point literal | AbstractFloat | 1e5 1.2 0x1.0p10 |
A shader-creation error results if:
-
An integer literal with a
i
oru
suffix cannot be represented by the target type. -
A hexadecimal floating point literal with a
f
orh
suffix overflows or cannot be exactly represented by the target type. -
A decimal floating point literal with a
f
orh
suffix overflows the target type. -
A floating point literal with a
h
suffix is used while the f16 extension is not enabled.
Note: The hexadecimal float value 0x1.00000001p0 requires 33 mantissa bits to be represented exactly, but f32 only has 23 explicit mantissa bits.
Note: If you want to use an f
suffix to force a hexadecimal float literal to be of type, the literal must also
use a binary exponent. For example, write 0x1p0f
. In comparison, 0x1f
is a hexadecimal integer literal.
3.6. Keywords
A keyword is a token which refers to a predefined language concept. See § 15.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 byU+005F
).
`/([_\p{XID_Start}][\p{XID_Continue}]+)|([\p{XID_Start}])/u`
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.
Section § 15.4 Context-Dependent Name Tokens lists all such tokens.
3.9. Diagnostic Rule Names
A diagnostic name-token is a token used in the name of a diagnostic triggering rule. The spelling of the token may be the same as an identifier but does not resolve to a declared object. The token must not be a keyword or reserved word.
See § 2.3 Diagnostics.
3.10. 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
.
array<vec4<f32>>
has two template parameterizations:
-
vec4<f32>
modifies the generalvec4
concept with template parameterf32
. -
array<vec4<f32>>
modifies the generalarray
concept with template parametervec4<f32>
.
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:
-
A template parameter is an expression, and therefore does not start with either a
'<'
(U+003C) or a'='
(U+003D) code point. -
An expression does not contain code points
';'
(U+003B),'{'
(U+007B), or':'
(U+003A). -
An expression does not contain an assignment.
-
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. -
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 discoveryInput: 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.
Algorithm:
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.
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.
-
Modify UnclosedCandidate to add the following fields:
-
parameters, a list of source ranges of template parameters.
-
parameter_start_position, a source location.
-
-
Modify TemplateList to add a field:
-
parameters, a list of source ranges of template parameters.
-
-
When pushing a new UnclosedCandidate onto the Pending stack:
-
Set its parameters field to an empty list.
-
Set parameter_start_position to one code point past CurrentPosition.
-
-
When adding a TemplateList, TL, to DiscoveredTemplateLists:
-
Let T be the top of the Pending stack, as in the original algorithm.
-
Push the source range starting at T.parameter_start_position and ending at CurrentPosition−1 onto T.parameters.
-
Prepare TL as in the original algorithm.
-
Set TL.parameters to T.parameters.
-
-
Insert a check at the end the loop, just before advancing past the current code point:
-
If '
,
' (U+002C) appears at CurrentPosition and Pending is not empty, then:-
Let T be the top of the Pending stack.
-
Push the source range starting at T.parameter_start_position and ending at CurrentPosition−1 onto T.parameters.
-
Set T.parameter_start_position to CurrentPosition+1
-
-
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.
_template_args_start template_arg_comma_list _template_args_end
template_arg_expression ( `','` template_arg_expression ) * `','` ?
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 theGPUDevice
.
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.
`'enable'` enable_extension_list `';'`
enable_extension_name ( `','` enable_extension_name ) * `','` ?
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.
WGSL enable-extension | WebGPU GPUFeatureName
| Description |
---|---|---|
f16
| "shader-f16"
| The f16 type is valid to use in the WGSL module. Otherwise, using f16 (directly or indirectly) will result in a shader-creation error. |
// 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.
`'requires'` software_extension_list `';'`
software_extension_name ( `','` software_extension_name ) * `','` ?
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.
WGSL language extension | Description |
---|---|
readonly_and_readwrite_storage_textures | Allows the use of read and read_write access modes with storage textures. Additionally, adds the textureBarrier built-in function. |
packed_4x8_integer_dot_product | Supports using 32-bit integer scalars packing 4-component vectors of 8-bit integers as inputs to the dot product instructions, packing and unpacking instructions with packed 4-component vectors of 8-bit integers |
unrestricted_pointer_parameters |
Removes the following restrictions from user-defined functions:
For user-defined functions, a parameter of pointer type must be in one of the following address spaces: Each argument of pointer type to a user-defined function must have the same memory view as its root identifier.
|
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.
`'diagnostic'` diagnostic_control `';'`
5. Declaration and Scope
A declaration associates an identifier with one of the following kinds of objects:
In other words, a declaration introduces a name for an object.
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 type-generators such as
array
,ptr
, andtexture_2d
, and -
enumerants such as read_write, perspective, 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 § 10.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).
-
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):
Precondition | Conclusion |
---|---|
e: T T is bool or vecN<bool> | ! e: T
|
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:
Precondition | Conclusion |
---|---|
e: vec3<bool> | ! e: vec3<bool>
|
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 for1u
and e2 standing for2u
, 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 § 14.6 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.
Src | Dest | ConversionRank(Src,Dest) | Description |
---|---|---|---|
T | T | 0 | Identity. No conversion performed. |
ref<AS,T,AM> for address space AS, and where access mode AM is read or read_write. | T | 0 | Apply the Load Rule to load a value from a memory reference. |
AbstractFloat | f32 | 1 | See § 14.6.4 Floating Point Conversion |
AbstractFloat | f16 | 2 | See § 14.6.4 Floating Point Conversion |
AbstractInt | i32 | 3 | Identity if the value is in i32. Produces a shader-creation error otherwise. |
AbstractInt | u32 | 4 | Identity if the value is in u32. Produces a shader-creation error otherwise. |
AbstractInt | AbstractFloat | 5 | See § 14.6.4 Floating Point Conversion |
AbstractInt | f32 | 6 | Behaves as AbstractInt to AbstractFloat, and then AbstractFloat to f32 |
AbstractInt | f16 | 7 | Behaves as AbstractInt to AbstractFloat, and then AbstractFloat to f16 |
vecN<S> | vecN<T> | ConversionRank(S,T) | Inherit conversion rank from component type. |
matCxR<S> | matCxR<T> | ConversionRank(S,T) | Inherit conversion rank from component type. |
array<S,N> | array<T,N> | ConversionRank(S,T) | Inherit conversion rank from component type. Note: Only fixed-size arrays may have an abstract component type. |
__frexp_result_abstract | __frexp_result_f32 | 1 | |
__frexp_result_abstract | __frexp_result_f16 | 2 | |
__frexp_result_vecN_abstract | __frexp_result_vecN_f32 | 1 | |
__frexp_result_vecN_abstract | __frexp_result_vecN_f16 | 2 | |
__modf_result_abstract | __modf_result_f32 | 1 | |
__modf_result_abstract | __modf_result_f16 | 2 | |
__modf_result_vecN_abstract | __modf_result_vecN_f32 | 1 | |
__modf_result_vecN_abstract | __modf_result_vecN_f16 | 2 | |
S | T where above cases don’t apply | infinity | There are no automatic conversions between other types. |
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:
-
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.
-
-
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.
-
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.
-
-
-
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 i, with -263 ≤ i < 263.
-
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
oru
suffix denotes an AbstractInt value. -
A floating point literal without an
f
orh
suffix denotes a AbstractFloat value.
Example: The expression log2(32)
is analyzed as follows:
-
log2(32)
is parsed as a function call to thelog2
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 value1.0
(rank 4) and2.5
remains an AbstractFloat (rank 0). -
1
is converted to f32 value1.0f
(rank 5) and2.5
is converted to f32 value2.5f
(rank 1). -
1
is converted to f16 value1.0f
(rank 6) and2.5
is converted to f16 value2.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.
-
// 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
.
Precondition | Conclusion | Description |
---|---|---|
true : bool
| The true value. | |
false : bool
| The false value. |
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 a two’s complementation representation, with the sign bit in the most significant bit position.
Type | Lowest value | Highest value |
---|---|---|
i32 | i32(-2147483648) | 2147483647i |
i32(-0x80000000) | 0x7fffffffi | |
u32 | 0u | 4294967295u |
0x0u | 0xffffffffu |
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 § 14.6 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 § 14.6 Floating Point Evaluation for details.
The following table lists certain extreme values for floating point types. Each has a corresponding negative value.
Type | Smallest positive denormal | Smallest positive normal | Largest positive finite | Largest finite power of 2 |
---|---|---|---|---|
f32 | 1.40129846432481707092e-45f | 1.17549435082228750797e-38f | 3.40282346638528859812e+38f | 0x1p+127f |
0x1p-149f | 0x1p-126f | 0x1.fffffep+127f | ||
f16 | 5.9604644775390625e-8h | 0.00006103515625h | 65504.0h | 0x1p+15h |
0x1p-24h | 0x1p-14h | 0x1.ffcp+15h |
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.
6.2.6. Vector Types
A vector is a grouped sequence of 2, 3, or 4 scalar components.
Type | Description |
---|---|
vecN<T> | Vector of N components of type T. N must be in {2, 3, 4} and T must be one of the scalar types. We say T is the component type of the vector. |
A vector is a numeric vector if its component type is a numeric scalar.
Key use cases of a vector include:
-
to express both a direction and a magnitude.
-
to express a position in space.
-
to express a color in some color space. For example, the components could be intensities of red, green, and blue, while the fourth component could be an alpha (opacity) value.
Many operations on vectors act component-wise, i.e. the result vector is formed by operating on each component independently.
let x : vec3< f32> = a + b ; // a and b are vec3<f32> // x[0] = a[0] + b[0] // x[1] = a[1] + b[1] // x[2] = a[2] + b[2]
WGSL also predeclares the following type aliases:
Predeclared alias | Original type | Restrictions |
---|---|---|
vec2i | vec2<i32> | |
vec3i | vec3<i32> | |
vec4i | vec4<i32> | |
vec2u | vec2<u32> | |
vec3u | vec3<u32> | |
vec4u | vec4<u32> | |
vec2f | vec2<f32> | |
vec3f | vec3<f32> | |
vec4f | vec4<f32> | |
vec2h | vec2<f16> | Requires the f16 extension. |
vec3h | vec3<f16> | |
vec4h | vec4<f16> |
6.2.7. Matrix Types
A matrix is a grouped sequence of 2, 3, or 4 floating point vectors.
Type | Description |
---|---|
matCxR<T> | Matrix of C columns and R rows of type T, where C and R are both in {2, 3, 4}, and T must be f32, f16, or AbstractFloat. Equivalently, it can be viewed as C column vectors of type vecR<T>. |
The key use case for a matrix is to embody a linear transformation. In this interpretation, the vectors of a matrix are treated as column vectors.
The product operator (*
) is used to either:
-
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:
Predeclared alias | Original type | Restrictions |
---|---|---|
mat2x2f | mat2x2<f32> | |
mat2x3f | mat2x3<f32> | |
mat2x4f | mat2x4<f32> | |
mat3x2f | mat3x2<f32> | |
mat3x3f | mat3x3<f32> | |
mat3x4f | mat3x4<f32> | |
mat4x2f | mat4x2<f32> | |
mat4x3f | mat4x3<f32> | |
mat4x4f | mat4x4<f32> | |
mat2x2h | mat2x2<f16> | Requires the f16 extension. |
mat2x3h | mat2x3<f16> | |
mat2x4h | mat2x4<f16> | |
mat3x2h | mat3x2<f16> | |
mat3x3h | mat3x3<f16> | |
mat3x4h | mat3x4<f16> | |
mat4x2h | mat4x2<f16> | |
mat4x3h | mat4x3<f16> | |
mat4x4h | mat4x4<f16> |
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.
Type | Description |
---|---|
atomic<T> | Atomic of type T. T must be either u32 or i32. |
An expression must not evaluate to an atomic type.
Atomic types may only be instantiated by variables in the workgroup address space or by storage buffer variables with a read_write access mode.
The memory scope of operations on the type is determined by the address space it is instantiated in.
Atomic types in the workgroup address space have a memory scope of Workgroup
, while those in the storage address space have a memory scope of QueueFamily
.
An atomic modification is any operation on an atomic object which sets the content of the object. The operation counts as a modification even if the new value is the same as the object’s existing value.
In WGSL, atomic modifications are mutually ordered, for each object. That is, during execution of a shader stage, for each atomic object A, all agents observe the same order of modification operations applied to A. The ordering for distinct atomic objects may not be related in any way; no causality is implied. Note that variables in workgroup space are shared within a workgroup, but are not shared between different workgroups.
6.2.9. Array Types
An array is an indexable sequence of element values.
Type | Description |
---|---|
array<E,N> | A fixed-size array with N elements of type E. N is called the element count of the array. |
array<E> | A runtime-sized array of elements of type E.
These may only appear in specific contexts. |
The first element in an array is at index 0, and each successive element is at the next integer index. See § 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.
-
It is a pipeline-creation error if N is not greater than zero.
Note: The element count value is fully determined at pipeline creation time.
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 § 12.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.
Type | Description |
---|---|
struct AStructName {M1 : T1, ... MN : TN, } |
A declaration of a structure type named by the identifier AStructName and having N members,
where member i is named by the identifier Mi and is of the type Ti.
N must be at least 1. Two members of the same structure type must not have the same name. |
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 ;
`'struct'` ident struct_body_decl
`'{'` struct_member ( `','` struct_member ) * `','` ? `'}'`
The following attributes can be applied to structure members:
Attributes builtin, location, 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 § 12.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 § 13.4 Memory Layout.
// 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:
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(M1),..., NestDepth(MN)) if T is a structure type with member types M1,...,MN
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.
-
A variable or value declaration cannot have an enumeration as its store type or its effective-value-type.
-
A function formal parameter cannot be an enumeration type, in part because enumerations are not constructible.
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.
Enumeration (Cannot be spelled in WGSL) | Predeclared enumerant |
---|---|
access mode | read |
write | |
read_write | |
address space
Note: The | function |
private | |
workgroup | |
uniform | |
storage | |
interpolation type | perspective |
linear | |
flat | |
interpolation sampling | center |
centroid | |
sample | |
built-in value | vertex_index |
instance_index | |
position | |
front_facing | |
frag_depth | |
local_invocation_id | |
local_invocation_index | |
global_invocation_id | |
workgroup_id | |
num_workgroups | |
sample_index | |
sample_mask | |
texel format | rgba8unorm |
rgba8snorm | |
rgba8uint | |
rgba8sint | |
rgba16uint | |
rgba16sint | |
rgba16float | |
r32uint | |
r32sint | |
r32float | |
rg32uint | |
rg32sint | |
rg32float | |
rgba32uint | |
rgba32sint | |
rgba32float | |
bgra8unorm |
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,
-
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 § 13.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 § 13.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 § 12.3.1 Inter-stage Input and Output Interface and subsequent sections. Those types are also sized, but the counting is 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.
Constraint | Type | Description |
---|---|---|
AS is an address space, T is a storable type, AM is an access mode | ref<AS,T,AM> |
The reference type identified with the set of memory views for memory locations in AS holding values of type T,
supporting memory accesses described by mode AM.
Here, T is the store type. Reference types are not written in WGSL source; instead they are used to analyze a WGSL module. |
AS is an address space, T is a storable type, AM is an access mode | ptr<AS,T,AM> |
The pointer type identified with the set of memory views for memory locations in AS holding values of type T,
supporting memory accesses described by mode AM.
Here, T is the store type. Pointer types may appear in WGSL source. |
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:
-
store type, and
-
sometimes by access mode, as specified in § 13.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 reference, or
-
an indexing expression where the base is a valid reference, 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 reference, 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
-
It is the variable, when R is a variable.
-
It is the originating variable of the pointer value P, when R is the application of the indirection operator (unary *) on P.
-
It is the originating variable of the base, when R is a named component expression or an indexing expression.
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.
In the following example, the reference the_particle.position[i]
is valid if and only if i
is 0 or 1.
When i
is 2, the reference will be an invalid memory reference, but would otherwise correspond
the memory locations for the_particle.color_index
.
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.
Those outcomes include, but are not limited to, the following:
- Trap
-
The shader invocation immediately terminates, and shader stage outputs are set to zero values.
- Invalid Load
-
Loads from an invalid reference may return one of:
-
when the originating variable is a uniform buffer or a storage buffer, the value from any memory location(s) of the WebGPU
GPUBuffer
bound to the originating variable -
when the originating variable is not a uniform buffer or storage buffer, a value from any memory location(s) in the originating variable
-
the zero value for store type of the reference
-
if the loaded value is a vector, the value (0, 0, 0, x), where x is:
-
0, 1, or the maximum positive value for integer components
-
0.0 or 1.0 for floating-point components
-
-
- Invalid Store
-
Stores to an invalid reference may do one of:
-
when the originating variable is a storage buffer, store the value to any memory location(s) of the WebGPU
GPUBuffer
bound to the originating variable -
when the originating variable is not a storage buffer, store the value to any memory locations(s) in the originating variable
-
not be executed.
-
A data race may occur if an invalid load or store is redirected to access different locations inside a variable in a shared address space. For example, the accesses of several concurrently executing invocations may be redirected to the first element in an array. If at least one access is a write, and they are not otherwise synchronized, then the result is a data race, and hence a dynamic error.
An out-of-bounds access invalidates the assumptions of uniformity analysis. For example, if an invocation terminates early due to an out-of-bounds access, then it can no longer particpate in collective operations. In particular, a call to workgroupBarrier may hang the shader, and derivatives may yield invalid results.
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 : ptr< storage, Particle > = & system . particles [ system . active_index ]; let delta_position : vec3< f32> = ( * active_particle ). velocity * system . timestep ; let current_position : vec3< f32> = ( * active_particle ). position; ( * active_particle ). position= delta_position + current_position ; }
fn add_one ( x : ptr< function, i32> ) { /* Update the locations for 'x' to contain the next higher integer value, (or to wrap around to the largest negative i32 value). On the left-hand side, unary '*' converts the pointer to a reference that can then be assigned to. It has a read_write access mode, by default. /* On the right-hand side: - Unary '*' converts the pointer to a reference, with a read_write access mode. - The only matching type rule is for addition (+) and requires '*x' to have type i32, which is the store type for '*x'. So the Load Rule applies and '*x' evaluates to the value stored in the memory for '*x' at the time of evaluation, which is the i32 value for 0. - Add 1 to 0, to produce a final value of 1 for the right-hand side. */ Store 1 into the memory for '*x'. */ * x = * x + 1 ; } @compute @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 reference to a composite:
-
Given a reference with a vector store type, appending a single-letter vector access phrase results in a reference to the named component of the vector. See § 8.5.1.3 Component Reference from Vector Reference.
-
Given a reference with a structure store type, appending a member access phrase results in a reference to the named member of the structure. See § 8.5.4 Structure Access Expression.
-
-
Use an indexing expression on a reference to a composite:
-
Given a reference with a vector store type, appending an array index access phrase results in a reference to the indexed component of the vector. See § 8.5.1.3 Component Reference from Vector Reference.
-
Given a reference with a matrix store type, appending an array index access phrase results in a reference to the indexed column vector of the matrix. See § 8.5.2 Matrix Access Expression.
-
Given a reference with an array store type, appending an array index access phrase results in a reference to the indexed element of the array. See § 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 ; }
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 ; }
6.4.9. Comparison with References and Pointers in Other Languages
This section is informative, not normative.
References and pointers in WGSL are more restricted than in other languages. In particular:
-
In WGSL a reference can’t directly be declared as an alias to another reference or variable, either as a variable or as a formal parameter.
-
In WGSL pointers and references are not storable. That is, the content of a WGSL variable declaration may not contain a pointer or a reference.
-
In WGSL a function must not return a pointer or reference.
-
In WGSL there is no way to convert between integer values and pointer values.
-
In WGSL there is no way to forcibly change the type of a pointer value into another pointer type.
-
A composite component reference expression is different: it takes a reference to a composite value and yields a reference to one of the components or elements inside the composite value. These are considered different references in WGSL, even though they may have the same machine address at a lower level of implementation abstraction.
-
-
In WGSL there is no way to forcibly change the type of a reference value into another reference type.
-
In WGSL there is no way to change the access mode of a pointer or reference.
-
By comparison, C++ automatically converts a non-const pointer to a const pointer, and has a
const_cast
to convert a const value to a non-const value.
-
-
In WGSL there is no way to allocate new memory from a "heap".
-
In WGSL there is no way to explicitly destroy a variable. The memory for a WGSL variable becomes inaccessible only when the variable goes out of scope.
Note: From the above rules, it is not possible to form a "dangling" pointer, i.e. a pointer that does not reference the memory for a "live" originating variable. A memory view may be an invalid memory reference, but it will never access memory locations not associated with the originating variable or buffer.
6.5. Texture and Sampler Types
A texel is a scalar or vector used as the smallest independently accessible element of a texture. The word texel is short for texture element.
A texture is a collection of texels supporting special operations useful for rendering. In WGSL, those operations are invoked via texture builtin functions. See § 16.7 Texture Built-in Functions for a complete list.
A WGSL texture corresponds to a WebGPU GPUTexture
.
A texture has the following features:
- texel format
-
The data representation of each texel. See § 6.5.1 Texel Formats.
- dimensionality
-
The number of dimensions in the grid coordinates, and how the coordinates are interpreted. The number of dimensions is 1, 2, or 3. Most textures use cartesian coordinates. Cube textures have six square faces, and are sampled with a three dimensional coordinate interpreted as a direction vector from the origin toward the cube centered on the origin.
- size
-
The extent of grid coordinates along each dimension. This is a function of mip level.
- mip level count
-
The mip level count is at least 1 for sampled textures and depth textures, and equal to 1 for storage textures.
Mip level 0 contains a full size version of the texture. Each successive mip level contains a filtered version of the previous mip level at half the size (within rounding) of the previous mip level.
When sampling a texture, an explicit or implicitly-computed level-of-detail is used to select the mip levels from which to read texel data. These are then combined via filtering to produce the sampled value. - arrayed
-
Whether the texture is arrayed.
-
A non-arrayed texture is a grid of texels.
-
An arrayed texture is a homogeneous array of grids of texels.
-
- array size
-
The number of homogeneous grids, if the texture is arrayed.
- sample count
-
The number of samples, if the texture is multisampled.
Each texel in a texture is associated with a unique logical texel address, which is an integer tuple having:
-
A mip level in [0, mip level count).
-
A number of components, controlled by the dimensionality, with each component value in [0, Si), where Si is the size in the i'th component.
-
An array index in [0, array size), if the texture is arrayed. Note that size is a function of mip level.
-
A sample index in [0, sample count), if the texture is multisampled.
A texture’s physical organization is typically optimized for rendering operations. To achieve this, many details are hidden from the programmer, including data layouts, data types, and internal operations that cannot be expressed directly in the shader language.
As a consequence, a shader does not have direct access to the texel memory within a texture variable. Instead, access is mediated through an opaque handle:
-
Within the shader:
-
Declare a module-scope variable where the store type is one of the texture types described in later sections. The variable stores an opaque handle to the underlying texture memory, and is automatically placed in the handle address space.
-
Inside a function, call one of the texture builtin functions, and provide the texture variable or function parameter as the builtin function’s first parameter.
-
-
When constructing the WebGPU pipeline, the texture variable’s store type and binding must be compatible with the corresponding bind group layout entry.
In this way, the set of supported operations for a texture type is determined by the availability of texture built-in functions having a formal parameter with that texture type.
Note: The handle stored by a texture variable cannot be changed by the shader. That is, the variable is read-only, even if the underlying texture to which it provides access may be mutable (e.g. a write-only storage texture).
The texture types are the set of types defined in:
A sampler is an opaque handle that controls how texels are accessed from a sampled texture or a depth texture.
A WGSL sampler maps to a WebGPU GPUSampler
.
Texel access is controlled via several properties of the sampler:
- addressing mode
-
Controls how texture boundaries and out-of-bounds coordinates are resolved. The addressing mode for each texture dimension can be set independently. See WebGPU
GPUAddressMode
. - filter mode
-
Controls which texels are accessed to produce the final result. Filtering can either use the nearest texel or interpolate between multiple texels. Multiple filter modes can be set independently. See WebGPU
GPUFilterMode
. - LOD clamp
-
Controls the min and max levels of details that are accessed.
- comparison
-
Controls the type of comparison done for comparison sampler. See WebGPU
GPUCompareFunction
. - max anisotropy
-
Controls the maximum anisotropy value used by the sampler.
Samplers cannot be created in WGSL modules and their state (e.g. the properties listed above) are immutable within a shader and can only be set by the WebGPU API.
It is a pipeline-creation error if a filtering sampler (i.e. any sampler using interpolative filtering) is used with texture that has a non-filterable format.
Note: The handle stored by a sampler variable cannot be changed by the shader.
6.5.1. Texel Formats
In WGSL, certain texture types are parameterized by texel format.
A texel format is characterized by:
- channels
-
Each channel contains a scalar. A texel format has up to four channels:
r
,g
,b
, anda
, normally corresponding to the concepts of red, green, blue, and alpha channels. - channel format
-
The number of bits in the channel, and how those bits are interpreted.
Each texel format in WGSL corresponds to a WebGPU GPUTextureFormat
with the same name.
Only certain texel formats are used in WGSL source code. The channel formats used to define those texel formats are listed in the Channel Formats table. The last column specifies the conversion from the stored channel bits to the value used in the shader. This is also known as the channel transfer function, or CTF.
Note: The channel transfer function for 8unorm maps {0,...,255} to the floating point interval [0.0, 1.0].
Note: The channel transfer function for 8snorm maps {-128,...,127} to the floating point interval [-1.0, 1.0].
Channel format | Number of stored bits | Interpretation of stored bits | Shader type | Shader value (Channel Transfer Function) |
---|---|---|---|---|
8unorm | 8 | unsigned integer v ∈ {0,...,255} | f32 | v ÷ 255 |
8snorm | 8 | signed integer v ∈ {-128,...,127} | f32 | max(-1, v ÷ 127) |
8uint | 8 | unsigned integer v ∈ {0,...,255} | u32 | v |
8sint | 8 | signed integer v ∈ {-128,...,127} | i32 | v |
16uint | 16 | unsigned integer v ∈ {0,...,65535} | u32 | v |
16sint | 16 | signed integer v ∈ {-32768,...,32767} | i32 | v |
16float | 16 | IEEE-754 binary16 16-bit floating point value v, with 1 sign bit, 5 exponent bits, 10 mantissa bits | f32 | v |
32uint | 32 | 32-bit unsigned integer value v | u32 | v |
32sint | 32 | 32-bit signed integer value v | i32 | v |
32float | 32 | IEEE-754 binary32 32-bit floating point value v | f32 | v |
The texel formats listed in the Texel Formats for Storage Textures table
correspond to the WebGPU plain color formats which support the WebGPU STORAGE_BINDING
usage.
These texel formats are used to parameterize the storage texture types defined
in § 6.5.5 Storage Texture Types.
When the texel format does not have all four channels, then:
-
When reading the texel:
-
If the texel format has no green channel, then the second component of the shader value is 0.
-
If the texel format has no blue channel, then the third component of the shader value is 0.
-
If the texel format has no alpha channel, then the fourth component of the shader value is 1.
-
-
When writing the texel, shader value components for missing channels are ignored.
The last column in the table below uses the format-specific channel transfer function from the channel formats table.
Texel format | Channel format | Channels in memory order | Corresponding shader value |
---|---|---|---|
rgba8unorm | 8unorm | r, g, b, a | vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) |
rgba8snorm | 8snorm | r, g, b, a | vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) |
rgba8uint | 8uint | r, g, b, a | vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a)) |
rgba8sint | 8sint | r, g, b, a | vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a)) |
rgba16uint | 16uint | r, g, b, a | vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a)) |
rgba16sint | 16sint | r, g, b, a | vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a)) |
rgba16float | 16float | r, g, b, a | vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) |
r32uint | 32uint | r | vec4<u32>(CTF(r), 0u, 0u, 1u) |
r32sint | 32sint | r | vec4<i32>(CTF(r), 0, 0, 1) |
r32float | 32float | r | vec4<f32>(CTF(r), 0.0, 0.0, 1.0) |
rg32uint | 32uint | r, g | vec4<u32>(CTF(r), CTF(g), 0.0, 1.0) |
rg32sint | 32sint | r, g | vec4<i32>(CTF(r), CTF(g), 0.0, 1.0) |
rg32float | 32float | r, g | vec4<f32>(CTF(r), CTF(g), 0.0, 1.0) |
rgba32uint | 32uint | r, g, b, a | vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a)) |
rgba32sint | 32sint | r, g, b, a | vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a)) |
rgba32float | 32float | r, g, b, a | vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) |
bgra8unorm | 8unorm | b, g, r, a | vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) |
WGSL predeclares an enumerant for each of the texel formats in the table.
6.5.2. Sampled Texture Types
A sampled texture is capable of being accessed in conjunction with a sampler. It can also be accessed without the use of a sampler. Sampled textures only allow read accesses.
The texel format is the format
attribute of the GPUTexture
bound to the texture variable.
WebGPU validates compatibility between the
texture, the sampleType
of the bind group layout,
and the sampled type of the texture variable.
The texture is parameterized by a sampled type and must be f32
, i32
, or u32
.
Type | Dimensionality | Arrayed |
---|---|---|
texture_1d<T> | 1D
| No |
texture_2d<T> | 2D
| No |
texture_2d_array<T> | 2D
| Yes |
texture_3d<T> | 3D
| No |
texture_cube<T> | Cube
| No |
texture_cube_array<T> | Cube
| Yes |
-
T is the sampled type.
-
The parameterized type for the images is the type after conversion from sampling. E.g. you can have an image with texels with 8bit unorm components, but when you sample them you get a 32-bit float result (or vec-of-f32).
6.5.3. Multisampled Texture Types
A multisampled texture has a sample count of 1 or more. Despite the name, it cannot be used with a sampler. It effectively stores multiple texels worth of data per logical texel address if the sample index is ignored.
The texel format is the format
attribute of the GPUTexture
bound to the texture variable.
WebGPU validates compatibility between the
texture, the sampleType
of the bind group layout,
and the sampled type of the texture variable.
The texture is parameterized by a sampled type and must be f32
, i32
, or u32
.
Type | Dimensionality | Arrayed |
---|---|---|
texture_multisampled_2d<T> | 2D
| No |
texture_depth_multisampled_2d | 2D
| No |
-
T is the sampled type.
6.5.4. External Sampled Texture Types
An External texture is an opaque
two-dimensional float-sampled texture type similar to texture_2d<f32>
but
potentially with a different representation.
It can be read using textureLoad or textureSampleBaseClampToEdge built-in
functions, which handle these different representations.
See WebGPU § 6.4 GPUExternalTexture.
Type | Dimensionality | Arrayed |
---|---|---|
texture_external | 2D
| No |
6.5.5. Storage Texture Types
A storage texture supports accessing individual texel values without the use of a sampler.
-
A write-only storage texture supports writing individual texels, with automatic conversion of the shader value to a stored texel value.
-
A read-only storage texture supports reading individual texels, with automatic conversion of the stored texel value to a shader texel value.
-
A read-write storage texture supports reading and writing individual texels, with automatic conversion between shader and stored texel values.
A storage texture type must be parameterized by one of the texel formats for storage textures. The texel format determines the conversion function as specified in § 6.5.1 Texel Formats.
When writing texels to a storage texture the inverse of the conversion function is used to convert the shader value to the stored texel.
Type | Dimensionality | Arrayed |
---|---|---|
texture_storage_1d<Format, Access> | 1D
| No |
texture_storage_2d<Format, Access> | 2D
| No |
texture_storage_2d_array<Format, Access> | 2D
| Yes |
texture_storage_3d<Format, Access> | 3D
| No |
-
Format must be an enumerant for one of the texel formats for storage textures
-
Access must be an enumerant for one of the access modes.
6.5.6. Depth Texture Types
A depth texture is capable of being accessed in conjunction with a sampler_comparison. It can also be accessed without the use of a sampler. Depth textures only allow read accesses.
The texel format of the texture is defined in the GPUTextureBindingLayout
.
Type | Dimensionality | Arrayed |
---|---|---|
texture_depth_2d | 2D
| No |
texture_depth_2d_array | 2D
| Yes |
texture_depth_cube | Cube
| No |
texture_depth_cube_array | Cube
| Yes |
6.5.7. Sampler Type
A sampler mediates access to a sampled texture or a depth texture, by performing a combination of:
-
coordinate transformation.
-
optionally modifying mip-level selection.
-
for a sampled texture, optionally filtering retrieved texel values.
-
for a depth texture, determining the comparison function applied to the retrieved texel.
A sampler types are:
Type | Description |
---|---|
sampler | Sampler. Mediates access to a sampled texture. |
sampler_comparison | Comparison sampler. Mediates access to a depth texture. |
Samplers are parameterized when created in the WebGPU API. They cannot be modified by a WGSL module.
Samplers can only be used by the texture built-in functions.
sampler sampler_comparison
6.6. AllTypes Type
The AllTypes type is the set of all WGSL types.
There is no way to write the AllTypes type in WGSL source.
See § 6.9 Predeclared Types and Type-Generators Summary for the list of all predeclared types and type-generators.
Instead, the AllTypes type exists so type checking rules will apply to any phrase that may contain an ordinary value. WGSL makes the rules consistent by defining a type to be a kind of value, and allowing an expression to denote a type.
The motivating case is a template parameter, which in various contexts may denote several kinds of things, including a type, an enumerant, or a plain value. In particular, the template_arg_expression grammar rule expands to the expression grammar nonterminal.
6.7. Type Aliases
A type alias declares a new name for an existing type. The declaration must appear at module scope, and its scope is the entire program.
When type T is defined as a type alias for a structure type S, all properties of the members of S, including attributes, carry over to the members of T.
`'alias'` ident `'='` type_specifier
alias Arr = array< i32, 5 > ; alias RTArr = array< vec4< f32>> ; alias single = f32; // Declare an alias for f32 const pi_approx : single = 3.1415 ; fn two_pi () -> single { return single ( 2 ) * pi_approx ; }
6.8. Type Specifier Grammar
Note: An expression can also denote a type, by expanding via the primary_expression grammar rule to template_elaborated_ident, and via parenthesization.
6.9. Predeclared Types and Type-Generators Summary
The predeclared types that can be spelled in WGSL source are:
WGSL also predeclares the return types for the frexp, modf, and atomicCompareExchangeWeak built-in functions. However, they cannot be spelled in WGSL source.
The predeclared type-generators are listed in the following table:
Predeclared type-generator | Cross-reference |
---|---|
array | See § 6.2.9 Array Types |
atomic | See § 6.2.8 Atomic Types |
mat2x2 |
See § 6.2.7 Matrix Types, which also lists
predeclared aliases for matrix types.
Note: These are also used in value constructor expressions to create matrices. |
mat2x3 | |
mat2x4 | |
mat3x2 | |
mat3x3 | |
mat3x4 | |
mat4x2 | |
mat4x3 | |
mat4x4 | |
ptr | See § 6.4.3 Reference and Pointer Types |
texture_1d | See § 6.5.2 Sampled Texture Types |
texture_2d | |
texture_2d_array | |
texture_3d | |
texture_cube | |
texture_cube_array | |
texture_multisampled_2d | See § 6.5.3 Multisampled Texture Types |
texture_storage_1d | See § 6.5.5 Storage Texture Types |
texture_storage_2d | |
texture_storage_2d_array | |
texture_storage_3d | |
vec2 |
See § 6.2.6 Vector Types, which also lists
predeclared aliases for vector types.
Note: These are also used in value constructor expressions to create vectors. |
vec3 | |
vec4 |
7. Variable and Value Declarations
Variable and value declarations provide names for data values.
A value declaration creates a name for a value, and that
value is immutable once it has been declared.
The four kinds of value declarations are const
, override
, let
, and formal parameter declarations,
further described below (see § 7.2 Value Declarations).
A variable declaration creates a name for memory locations for storing a value; the value stored there may be updated, if the variable has
a read_write access mode.
There is one kind of variable declaration, var
, but it has options for address space and access modes in various combinations, described
below (see § 7.3 var Declarations).
Note: A value declaration does not have associated memory locations. For example, no WGSL expression can form a pointer to the value.
A declaration appearing outside of any function definition is at module scope. Its name is in scope for the entire program.
A declaration appearing within a function definition is in function scope. The name is available for use in the statement immediately after its declaration until the end of the brace-delimited list of statements immediately enclosing the declaration. A function-scope declaration is a dynamic context.
Variable and value declarations have a similar overall syntax:
// Specific value declarations. const name [: type ] = initializer ; [ attribute ] * override name [: type ] [ = initializer ]; let name [: type ] = initializer ; // General variable form. [ attribute ] * var [ < address_space [, access_mode ] > ] name [: type ] [ = initializer ]; // Specific variable declarations. // Function scope. var [ < function> ] name [: type ] [ = initializer ]; // Module scope. var < private> name [: type ] [ = initializer ]; var < workgroup> name : type ; [ attribute ] + var < uniform> name : type ; [ attribute ] + var name : texture_type ; [ attribute ] + var name : sampler_type ; [ attribute ] + var < storage[, access_mode ] > name : type ;
Each such declaration must have an explicitly specified type or an initializer. Both a type and an initializer may be specified. Each such declaration determines the type for the associated data value, known as the effective-value-type for the declaration. The effective-value-type of the declaration is:
-
The declared type, if explicitly specified.
-
Otherwise, if the initializer expression has type
T
:-
For a
const
declaration, the effective-value-type isT
itself. -
For a
override
,let
, orvar
declaration, the effective-value-type is the concretization ofT
.
-
Each kind of value or variable declaration may place additional constraints on the form of the initializer expression, if present, and on the effective-value-type.
-
Only const-declarations can be abstract types, and only when the type is not explicitly specified.
-
The type of the expression must be feasibly converted to the effective-value-type.
-
If an initializer is not specified, a value must be provided at pipeline-creation time.
-
Override-declarations are part of the shader interface, but are not bound resources.
-
Storage buffers with an access mode other than read and storage textures cannot be statically accessed in a vertex shader stage. See WebGPU
createBindGroupLayout()
. -
Atomic types can only appear in mutable storage buffers or workgroup variables.
-
The data in storage textures with a write or read_write access mode is mutable, but can only be modified via textureStore built-in function. The variable itself cannot be modified.
-
Variables in the workgroup address space can only be statically accessed in a compute shader stage.
-
The element count of the outermost array may be an override-expression.
-
If there is no initializer, the variable is default initialized.
7.1. Variables vs Values
Variable declarations are the only mutable data in a WGSL module. Value declarations are always immutable. Variables can be the basis of reference and pointer values because variables have associated memory locations, whereas a value declaration cannot be the basis of a pointer or reference value.
Using variables is generally more expensive than using value declarations, because using a variable requires extra operations to read or write to the memory locations associated with the variable.
Generally speaking, an author should prefer using declarations in the following order, with the most preferred option listed first:
This will generally result in the best overall performance of a shader.
7.2. Value Declarations
When an identifier resolves to a value declaration, the identifier denotes that value.
WGSL provides multiple kinds of value declarations. The value for each kind of declaration is fixed at a different point in the shader lifecycle. The different kinds of value declarations and when their values are fixed are:
-
let-declarations, when they are executed
-
formal parameter declarations, when the associated function call argument is executed
Note: Formal parameters are described in § 10 Functions.
7.2.1. const
Declarations
A const-declaration specifies a name for a data value that is fixed at shader-creation time. Each const-declaration requires an initializer. A const-declaration can be declared in module or function scope. The initializer expression must be a const-expression. The type of a const-declaration must be a concrete or abstract constructible type. const-declarations are the only declarations where the effective-value-type may be abstract.
Note: Since abstract numeric types cannot be spelled in WGSL, they can only be used via type inference.
const a = 4 ; // AbstractInt with a value of 4. const b : i32= 4 ; // i32 with a value of 4. const c : u32= 4 ; // u32 with a value of 4. const d : f32= 4 ; // f32 with a value of 4. const e = vec3( a , a , a ); // vec3 of AbstractInt with a value of (4, 4, 4). const f = 2.0 ; // AbstractFloat with a value of 2. const g = mat2x2( a , f , a , f ); // mat2x2 of AbstractFloat with a value of: // ((4.0, 2.0), (4.0, 2.0)). // The AbstractInt a converts to AbstractFloat. // An AbstractFloat cannot convert to AbstractInt. const h = array( a , f , a , f ); // array of AbstractFloat with 4 components: // (4.0, 2.0, 4.0, 2.0).
7.2.2. override
Declarations
An override-declaration specifies a name for a pipeline-overridable constant value. The value of a pipeline-overridable constant is fixed at pipeline-creation time. The value is one provided by the WebGPU pipeline-creation method, if specified, and otherwise is the value of its concretized initializer expression. The effective-value-type of an override-declaration must be a concrete scalar type.
An initializer expression is optional. If present, it must be an override-expression and represents the pipeline-overridable constant default value. If no initializer is specified, it is a pipeline-creation error if a value is not provided at pipeline-creation time.
If the declaration has an id attribute applied, the literal operand is known as the pipeline constant ID, and must be a unique integer between 0 and 65535 inclusive. That is, two override-declarations must not use the same pipeline constant ID.
The application can specify its own value for an override-declaration at pipeline-creation time. The pipeline creation API accepts a mapping from overridable constants to a value of the constant’s type. The constant is identified by a pipeline-overridable constant identifier string, which is the base-10 representation of the pipeline constant ID if specified, and otherwise the declared name of the constant.
@id ( 0 ) override has_point_light : bool= true ; // Algorithmic control @id ( 1200 ) override specular_param : f32= 2.3 ; // Numeric control @id ( 1300 ) override gain : f32; // Must be overridden override width : f32= 0.0 ; // Specified at the API level using // the name "width". override depth : f32; // Specified at the API level using // the name "depth". // Must be overridden. override height = 2 * depth ; // The default value // (if not set at the API level), // depends on another // overridable constant.
7.2.3. let
Declarations
A let-declaration specifies a name for a value that is fixed each time the statement is executed at runtime. A let-declaration must only be declared in function scope, and as such, is a dynamic context. A let-declaration must have an initializer expression. The value is the concretized value of the initializer. The effective-value-type of a let-declaration must be either a concrete constructible type or a pointer type.
// 'blockSize' denotes the i32 value 1024. let blockSize : i32= 1024 ; // 'row_size' denotes the u32 value 16u. The type is inferred. let row_size = 16u ;
7.3. var
Declarations
A variable is a named reference to memory that can contain a value of a particular storable type.
Two types are associated with a variable: its store type (the type of value
that may be placed in the referenced memory) and its reference type (the type
of the variable itself).
If a variable has store type T
, address space AS
, and access mode AM
, then its reference type is ref<AS,T,AM>
.
The store type of a variable is always concrete.
A variable declaration:
-
Specifies the variable’s name.
-
Determines the variable’s address space, store type, and access mode. Together these comprise the variable’s reference type.
-
The store type is the effective-value-type of the variable’s declaration.
-
-
Ensures the execution environment allocates memory for a value of the store type, in the specified address space, supporting the given access mode, for the lifetime of the variable.
-
Optionally has an initializer expression if the variable is in the private or function address spaces. If present, the initializer must evaluate to the variable’s store type. If present, the initializer for a private variable must be a const-expression or an override-expression. Variables in address spaces other than function or private must not have an initializer.
When an identifier resolves to a variable declaration, the identifier is an expression denoting the reference memory view for the variable’s memory, and its type is the variable’s reference type. See § 8.11 Variable Identifier Expression.
If the address space or access mode for a variable declaration are specified
in program source, they are written as a template list after the var
keyword:
-
The address space is specified first, as one of the predeclared address space enumerants.
-
The access mode is specified second, if present, as one of the predeclared address mode enumerants.
-
The address space must be specified if the access mode is specified.
-
Variables in the private, storage, uniform, workgroup, and handle address spaces must only be declared in module scope, while variables in the function address space must only be declared in function scope. The address space must be specified for all address spaces except handle and function. The handle address space must not be specified. Specifying the function address space is optional.
The access mode always has a default value, and except for variables in the storage address space, must not be specified in the WGSL source. See § 13.3 Address Spaces.
A variable in the uniform address space is a uniform buffer variable. Its store type must be a host-shareable constructible type, and must satisfy the address space layout constraints.
A variable in the storage address space is a storage buffer variable. Its store type must be a host-shareable type and must satisfy the address space layout constraints. The variable may be declared with a read or read_write access mode; the default is read.
A texture resource is a variable whose effective-value-type is a texture type. It is declared at module scope. It holds an opaque handle which is used to access the underlying grid of texels in a texture. The handle itself is in the handle address space and is always read-only. In many cases the underlying texels are read-only, and we say the texture variable immutable. For write-only storage textures and read-write storage textures, the underlying texels are mutable, and by convention we say the texture variable is mutable.
A sampler resource is a variable whose effective-value-type is a sampler type. It is declared at module scope, exists in the handle address space, and is immutable.
As described in § 12.3.2 Resource Interface, uniform buffers, storage buffers, textures, and samplers form the resource interface of a shader.
The lifetime of a variable is the period during shader execution for which the memory locations are associated with the variable. The lifetime of a module scope variable is the entire execution of the shader stage. There is an independent version of a variable in the private and function address spaces for each invocation. Function-scope variables are a dynamic context. The lifetime of a function-scope variable is determined by its scope:
-
It starts when control enters the variable’s declaration.
-
It ends when the name is no longer in scope of any part of the dynamic context. That is, the lifetime includes any functions called while the name is in scope.
Two resource variables may have overlapping memory locations, but it is a dynamic error if either of those variables is mutable. Other variables with overlapping lifetimes will not have overlapping memory locations. When a variable’s lifetime ends, its memory may be used for another variable.
Note: WGSL ensures the contents of a variable are only observable during the variable’s lifetime.
When a variable in the private, function, or workgroup address spaces is created, it will have an initial value. If no initializer is specified the initial value is the default initial value. The initial values are computed as follows:
-
For variables in the function address space:
-
The zero value of the store type, if the variable declaration did not specify an initializer.
-
Otherwise it is the result of evaluating the concretized initializer expression at that point in program execution.
-
-
For variables in the private address space:
-
The zero value of the store type, if the variable declaration did not specify an initializer.
-
Otherwise it is the result of evaluating the concretized initializer expression. The initializer must be an override-expression, and so its value is fixed no later than pipeline-creation time.
-
-
For variables in the workgroup address space:
-
When the store type is constructible, the zero value for the store type.
-
If the store type is an atomic type, the zero value is that of the underlying type (concrete integer scalar).
-
Otherwise, if the store type is not constructible, the zero value is determined by recursively applying these rules to each component of the composite until a constructible type is encountered.
-
Note: This commonly occurs when using an array with a pipeline-overridable element count or a composite that contains an atomic type.
-
-
Variables in other address spaces are resources set by bindings in the draw command or dispatch command.
Consider the following snippet of WGSL:
var i : i32; // Initial value is 0. Not recommended style. loop { var twice : i32= 2 * i ; // Re-evaluated each iteration. i ++ ; if i == 5 { break ; } }
i
will take on values 0, 1, 2, 3, 4, 5, and variable twice
will take on values 0, 2, 4, 6, and 8.
Consider the following snippet of WGSL:
Becausex
is a variable, all accesses to it turn into load and store operations.
However, it is expected that either the browser or the driver optimizes this intermediate representation
such that the redundant loads are eliminated.
var < private> decibels : f32; var < workgroup> worklist : array< i32, 10 > ; struct Params { specular : f32, count : i32} // Uniform buffer. Always read-only, and has more restrictive layout rules. @group ( 0 ) @binding ( 2 ) var < uniform> param : Params ; // A uniform buffer // A storage buffer, for reading and writing @group ( 0 ) @binding ( 0 ) var < storage, read_write> pbuf : array< vec2< f32>> ; // Textures and samplers are always in "handle" space. @group ( 0 ) @binding ( 1 ) var filter_params : sampler;
// Storage buffers @group ( 0 ) @binding ( 0 ) var < storage, read> buf1 : Buffer ; // Can read, cannot write. @group ( 0 ) @binding ( 0 ) var < storage> buf2 : Buffer ; // Can read, cannot write. @group ( 0 ) @binding ( 1 ) var < storage, read_write> buf3 : Buffer ; // Can both read and write. struct ParamsTable { weight : f32} // Uniform buffer. Always read-only, and has more restrictive layout rules. @group ( 0 ) @binding ( 2 ) var < uniform> params : ParamsTable ; // Can read, cannot write.
fn f () { var < function> count : u32; // A variable in function address space. var delta : i32; // Another variable in the function address space. var sum : f32= 0.0 ; // A function address space variable with initializer. var pi = 3.14159 ; // Infer the f32 store type from the initializer. }
7.4. Variable and Value Declaration Grammar Summary
| variable_decl `'='` expression
| `'let'` optionally_typed_ident `'='` expression
| `'const'` optionally_typed_ident `'='` expression
`'var'` _disambiguate_template template_list ? optionally_typed_ident
ident ( `':'` type_specifier ) ?
attribute * variable_decl ( `'='` expression ) ?
`'const'` optionally_typed_ident `'='` expression
| attribute * `'override'` optionally_typed_ident ( `'='` expression ) ?
8. Expressions
Expressions specify how values are computed.
The different kinds of value expressions provide a tradeoff between when they are evaluated and how expressive they can be. The sooner the evaluation, the more constrained the operations, but also the more places the value can be used. This tradeoff leads to different flexibility with each kind of value declaration. const-expressions and override-expressions are evaluated prior to execution on the GPU, so only the result of the computation of the expression is necessary in the final GPU code. Additionally, because const-expressions are evaluated at shader-creation time they can be used in more situations than override-expressions, for example, to size arrays in function scope variables. A runtime expression is an expression that is neither a const-expression nor an override-expression. A runtime expression is computed on the GPU during shader execution. While runtime expressions can be used by fewer grammar elements, they can be computed from a larger class of expressions, for example, other runtime values.
8.1. Early Evaluation Expressions
WGSL defines two types of expressions that can be evaluated before runtime:
8.1.1. const
Expressions
Expressions that can be evaluated at shader-creation time are called const-expressions. An expression is a const-expression if all its identifiers resolve to:
-
const-functions, or
-
type aliases, or
-
structure names.
The type of a const
expression must resolve to a type with a creation-fixed footprint.
Note: Abstract types can be the inferred type of a const-expression.
A const-expression E will be evaluated if and only if:
-
E is top-level expression, or
-
E is a subexpression of an expression OuterE, and OuterE will be evaluated, and evaluation of OuterE requires E to be evaluated.
Note: The evaluation rule implies that short-circuiting operators &&
and ||
guard evaluation of their right-hand
side subexpressions.
A const-expression may be evaluated by the CPU implementing the WebGPU API methods. Therefore accuracy requirements for operations on AbstractFloat values are no more strict than required for common WebGPU runtime environments, such as WebAssembly [WASM-CORE-2] and ECMAScript [ECMASCRIPT]. Accuracy requirements for concrete floating point types (such as f32) are specified in § 14.6.2.1 Accuracy of Concrete Floating Point Expressions.
Example: (42)
is analyzed as follows:
-
The term
42
is the AbstractInt value 42. -
Surrounding that term with parentheses produces a new expression
(42)
that is of type AbstractInt with value 42.
Example: -5
is analyzed as follows:
-
The term
5
is the AbstractInt value 5. -
Preceding that term with '
-
' produces a new expression-5
that is of type AbstractInt with value -5.
Example: -2147483648
is analyzed as follows:
-
The term
2147483648
is the AbstractInt value 2147483648. Note that this value does not fit in a 32-bit signed integer. -
Preceding that term with '
-
' produces a new expression-2147483648
that is of type AbstractInt with value -2147483648.
Example: const minint = -2147483648;
is analyzed as follows:
-
As above,
-2147483648
evaluates to a AbstractInt value -2147483648. -
A const-declaration allows the initializer to be an abstract numeric type.
-
The result is that
minint
is declared to be the AbstractInt value -2147483648.
Example: let minint = -2147483648;
is analyzed as follows:
-
As above,
-2147483648
evaluates to a AbstractInt value -2147483648. -
A let-declaration requires the initializer to be a concrete constructible type or a pointer type.
-
The let-declaration does not have an explicit type, so overload resolution is used. The overload candidates that apply use feasible automatic conversions from AbstractInt to either i32, u32, or f32. The one of lowest rank is to i32, and so AbstractInt -2147483648 value is converted to the i32 value -2147483648.
-
The result is that
minint
is declared to be the i32 value -2147483648.
Example: false && (10i < i32(5 * 1000 * 1000 * 1000))
is analyzed as follows:
-
The entire expression is a const-expression.
-
However, the short-circuiting rules of the
&&
operator apply: the left-hand side evaluates tofalse
, and so the right-hand side is not evaluated. -
Evaluation of i32(5 * 1000 * 1000 * 1000) would have caused a shader-creation error because the AbstractInt value 5000000000 overflows the i32 type.
8.1.2. override
Expressions
Expressions that can be evaluated at pipeline creation time are called override-expressions. An expression is an override-expression if all its identifiers resolve to:
-
const-functions, or
-
type aliases, or
-
structure names.
Note: All const-expressions are also override-expressions.
An override-expression E will be evaluated if and only if:
-
E is top-level expression, or
-
E is a subexpression of an expression OuterE, and OuterE will be evaluated, and evaluation of OuterE requires E to be evaluated.
Note: Not all override-expressions may be usable as the initializer for an override-declaration, because such initializers must resolve to a concrete scalar type.
Example: override x = 42;
is analyzed as follows:
-
The term
42
is the AbstractInt value 42. -
An override-declaration requires a concrete scalar type.
-
42
is converted to i32 via a feasible automatic conversion.
Example: let y = x + 1;
is analyzed as follows:
-
From above,
x
has a type of i32. -
The expression
x + 1
is an override-expression because it is composed of an override-declaration and an integer literal. -
The expression has a type of i32 and is evaluated at pipeline creation time. Its value depends on whether or not
x
is overridden at pipeline creation time.
Example: vec3(x,x,x)
is analyzed as follows:
-
From above,
x
is an override-declaration with the type i32. -
vec3(x,x,x)
is an override-expression because the only identifiers resolve to override-declarations. -
The type of the expression is a vector of 3 components of i32 (
vec3<i32>
).
8.2. Indeterminate values
In limited cases, an evaluation of a runtime expression can occur using unsupported values for its subexpressions.
In such a case, the result of that evaluation is an indeterminate value of the expression’s static type, meaning some arbitrary implementation-chosen value of the static type.
A distinct value may be produced for each unique dynamic context in which the expression is evaluated. For example, if the evaluation occurs once per iteration of a loop, a distinct value may be computed for each loop iteration.
Note: If the type is a floating point type and the implementation supports NaN values, then the indeterminate value produced at runtime may be a NaN value.
fn fun () { var extracted_values : array< i32, 2 > ; const v = vec2< i32> ( 0 , 1 ); for ( var i : i32= 0 ; i < 2 ; i ++ ) { // A runtime-expression used to index a vector, but outside the // indexing bounds of the vector, produces an indeterminate value // of the vector component type. let extract = v [ i + 5 ]; // Now 'extract' is any value of type i32. // Save it for later. extracted_values [ i ] = extract ; if extract == extract { // This is always executed } if extract < 2 { // This might be executed, but might not be executed. // Even though the original vector components are 0 and 1, // the extracted value might not be either of those values. } } if extracted_values [ 0 ] == extracted_values [ 1 ] { // This might be executed, but might not be executed. } } fn float_fun ( runtime_index : u32) { const v = vec2< f32> ( 0 , 1 ); // A vector of floating point values // As in the previous example, 'float_extract' is an indeterminate value. // Since it is a floating point type, it may be a NaN. let float_extract : f32= v [ runtime_index + 5 ]; if float_extract == float_extract { // This *might not* be executed, because: // - 'float_extract' may be NaN, and // - a NaN is never equal to any other floating point number, // even another NaN. } }
8.3. Literal Value Expressions
Precondition | Conclusion | Description |
---|---|---|
true : bool
| true boolean value.
| |
false : bool
| false boolean value.
| |
e is an integer literal with no suffix | e: AbstractInt | Abstract integer literal value. |
e is a floating point literal with no suffix | e: AbstractFloat | Abstract float literal value. |
e is an integer literal with i suffix
| e: i32 | 32-bit signed integer literal value. |
e is an integer literal with u suffix
| e: u32 | 32-bit unsigned integer literal value. |
e is an floating point literal with f suffix
| e: f32 | 32-bit floating point literal value. |
e is an floating point literal with h suffix
| e: f16 | 16-bit floating point literal value. |
8.4. Parenthesized Expressions
Precondition | Conclusion | Description |
---|---|---|
e : T | ( e ) : T
| Evaluates to e. Use parentheses to isolate an expression from the surrounding text. |
8.5. Composite Value Decomposition Expressions
This section describes expressions for getting a component of a composite value, and for getting a reference to a component from a reference to the containing composite value. For this discussion, the composite value, or the reference to composite value, is known as the base.
There are two ways of doing so:
- named component expression
-
The expression for the base B is followed by a period
'.'
(U+002D), and then the name of the component. - indexing expression
-
The expression for the base is followed by
'['
(U+005B), then the expression for an index then']'
(U+005D).-
The base may be a vector, matrix, or fixed-size array type, or or a reference to a vector, matrix, fixed-size array, or runtime-sized array type.
-
The index expression must be of integer scalar type.
-
Syntactically, these two forms are embodied by uses of the component_or_swizzle_specifier grammar rule.
-
N is the number of components of the vector type, when the base is a vector or a reference to a vector.
-
N is the number of columns of the matrix type, when the base is a matrix or a reference to a matrix.
-
N is the element count of the fixed-size array type, when the base is a fixed-size array or a reference to a fixed-size array.
-
N is NRuntime for the base, when the base is a reference to a runtime-sized array.
The index value is an out-of-bounds index when it is not an in-bounds index. An out-of-bounds index is often a programming defect, and will often cause a error. See below for details.
Additionally, vector types support a swizzling syntax for creating a new vector value from the components of another vector.
8.5.1. Vector Access Expression
Accessing components of a vector can be done either:
-
Using array subscripting (e.g.
v[2]
), or -
Using a swizzle name, a context-dependent name written as a sequence of convenience names, each mapping to a component of the source vector.
-
The color set of convenience names:
r
,g
,b
,a
for vector components 0, 1, 2, and 3 respectively. -
The dimensional set of convenience names:
x
,y
,z
,w
for vector components 0, 1, 2, and 3, respectively.
-
The convenience names are accessed using the .
notation. (e.g. color.bgra
).
The convenience letterings must not be mixed. For example, you cannot use .rybw
.
A convenience letter must not access a component past the end of the vector.
The convenience letterings can be applied in any order, including duplicating letters as needed. The provided number of letters must be between 1 and 4. That is, using convenience letters can only produce a scalar type or a valid vector type.
The result type depends on the number of letters provided. Assuming a vec4<f32>
Accessor | Result type |
---|---|
r | f32
|
rg | vec2<f32>
|
rgb | vec3<f32>
|
rgba | vec4<f32>
|
var a : vec3< f32> = vec3< f32> ( 1. , 2. , 3. ); var b : f32= a . y ; // b = 2.0 var c : vec2< f32> = a . bb ; // c = (3.0, 3.0) var d : vec3< f32> = a . zyx ; // d = (3.0, 2.0, 1.0) var e : f32= a [ 1 ]; // e = 2.0
8.5.1.1. Vector Single Component Selection
Precondition | Conclusion | Description |
---|---|---|
e: vecN<T> | e.x : Te .r : T
| Select the first component of e |
e: vecN<T> | e.y : Te .g : T
| Select the second component of e |
e: vecN<T> N is 3 or 4 | e.z : Te .b : T
| Select the third component of e |
e: vec4<T> | e.w : Te .a : T
| Select the fourth component of e |
e: vecN<T> i: i32 or u32 T is concrete | e[i]: T |
Select the i’th component of vector The first component is at index i=0. If i is outside the range [0,N-1]:
|
e: vecN<T> i: i32 or u32 T is abstract i is a const-expression | e[i]: T |
Select the i’th component of vector The first component is at index i=0. It is a shader-creation error if i is outside the range [0,N-1]. Note: When an abstract vector value e is indexed by an expression that is not a const-expression, then the vector is concretized before the index is applied. |
8.5.1.2. Vector Multiple Component Selection
Precondition | Conclusion | Description |
---|---|---|
e: vecN<T> I is the letter x , y , z , or w J is the letter x , y , z , or w | e. IJ: vec2<T> | Computes the two-component vector with first component e.I, and second component e.J. Letter z is valid only when N is 3 or 4.Letter w is valid only when N is 4.
|
e: vecN<T> I is the letter r , g , b , or a J is the letter r , g , b , or a | e. IJ: vec2<T> | Computes the two-component vector with first component e.I, and second component e.J. Letter b is valid only when N is 3 or 4.Letter a is valid only when N is 4.
|
e: vecN<T> I is the letter x , y , z , or w J is the letter x , y , z , or w K is the letter x , y , z , or w | e. IJK: vec3<T> | Computes the three-component vector with first component e.I, second component e.J, and third component e.K. Letter z is valid only when N is 3 or 4.Letter w is valid only when N is 4.
|
e: vecN<T> I is the letter r , g , b , or a J is the letter r , g , b , or a K is the letter r , g , b , or a | e. IJK: vec3<T> | Computes the three-component vector with first component e.I, second component e.J, and third component e.K. Letter b is only valid when N is 3 or 4.Letter a is only valid when N is 4.
|
e: vecN<T> I is the letter x , y , z , or w J is the letter x , y , z , or w K is the letter x , y , z , or w L is the letter x , y , z , or w | e. IJKL: vec4<T> | Computes the four-component vector with first component e.I, second component e.J, third component e.K, and fourth component e.L. Letter z is valid only when N is 3 or 4.Letter w is valid only when N is 4.
|
e: vecN<T> I is the letter r , g , b , or a J is the letter r , g , b , or a K is the letter r , g , b , or a L is the letter r , g , b , or a | e. IJKL: vec4<T> | Computes the four-component vector with first component e.I, second component e.J, third component e.K, and fourth component e.L. Letter b is only valid when N is 3 or 4.Letter a is only valid when N is 4.
|
8.5.1.3. Component Reference from Vector Reference
A write access to component of a vector may access all of the memory locations associated with that vector.
Note: This means accesses to different components of a vector by different invocations must be synchronized if at least one access is a write access. See § 16.11 Synchronization Built-in Functions.
Precondition | Conclusion | Description |
---|---|---|
r: ref<AS,vecN<T>,AM> | r.x : ref<AS,T,AM>r .r : ref<AS,T,AM> | Compute a reference to the first component of the vector referenced by the reference r. The originating variable of the resulting reference is the same as the originating variable of r. |
r: ref<AS,vecN<T>,AM> | r.y : ref<AS,T,AM>r .g : ref<AS,T,AM> | Compute a reference to the second component of the vector referenced by the reference r. The originating variable of the resulting reference is the same as the originating variable of r. |
r: ref<AS,vecN<T>,AM> N is 3 or 4 | r.z : ref<AS,T,AM>r .b : ref<AS,T,AM> | Compute a reference to the third component of the vector referenced by the reference r. The originating variable of the resulting reference is the same as the originating variable of r. |
r: ref<AS,vec4<T>,AM> | r.w : ref<AS,T,AM>r .a : ref<AS,T,AM> | Compute a reference to the fourth component of the vector referenced by the reference r. The originating variable of the resulting reference is the same as the originating variable of r. |
r: ref<AS,vecN<T>,AM> i: i32 or u32 | r[i] : ref<AS,T,AM> |
Compute a reference to the i’th component of the vector
referenced by the reference r.
If i is outside the range [0,N-1]:
The originating variable of the resulting reference is the same as the originating variable of r. |
8.5.2. Matrix Access Expression
Precondition | Conclusion | Description |
---|---|---|
e: matCxR<T> i: i32 or u32 T is concrete | e[i]: vecR<T> |
The result is the i’th column vector of e.
If i is outside the range [0,C-1]:
|
e: matCxR<T> i: i32 or u32 T is abstract i is a const-expression | e[i]: vecR<T> |
The result is the i’th column vector of e.
It is a shader-creation error if i is outside the range [0,C-1]. Note: When an abstract matrix value e is indexed by an expression that is not a const-expression, then the matrix is concretized before the index is applied. |
Precondition | Conclusion | Description |
---|---|---|
r: ref<AS,matCxR<T>,AM> i: i32 or u32 | r[i] : ref<AS,vecR<T>,AM> |
Compute a reference to the i’th column vector of the
matrix referenced by the reference r.
If i is outside the range [0,C-1]:
The originating variable of the resulting reference is the same as the originating variable of r. |
8.5.3. Array Access Expression
Precondition | Conclusion | Description |
---|---|---|
e: array<T,N> i: i32 or u32 T is concrete | e[i] : T |
The result is the value of the i’th element of the array value e.
If i is outside the range [0,N-1]:
|
e: array<T,N> i: i32 or u32 T is abstract i is a const-expression | e[i] : T |
The result is the value of the i’th element of the array value e.
It is a shader-creation error if i is outside the range [0,N-1]. Note: When an abstract array value e is indexed by an expression that is not a const-expression, then the array is concretized before the index is applied. |
Precondition | Conclusion | Description |
---|---|---|
r: ref<AS,array<T,N>,AM> i: i32 or u32 | r[i] : ref<AS,T,AM> |
Compute a reference to the i’th element of the array
referenced by the reference r.
If i is outside the range [0,N-1]:
The originating variable of the resulting reference is the same as the originating variable of r. |
r: ref<AS,array<T>,AM> i: i32 or u32 | r[i] : ref<AS,T,AM> |
Compute a reference to the i’th element of the
runtime-sized array referenced by the reference r.
If at runtime the array has N elements, and i is outside the range [0,N-1], then the expression evaluates to an invalid memory reference. If i is a signed integer, and i is less than 0:
The originating variable of the resulting reference is the same as the originating variable of r. |
8.5.4. Structure Access Expression
Precondition | Conclusion | Description |
---|---|---|
S is a structure type M is the identifier name of a member of S, having type T e: S | e.M: T | The result is the value of the member with name M from the structure value e. |
Precondition | Conclusion | Description |
---|---|---|
S is a structure type M is the identifier name of a member of S, having type T r: ref<AS,S,AM> | r.M: ref<AS,T,AM> | Given a reference to a structure, the result is a reference to the structure member with identifier name M. The originating variable of the resulting reference is the same as the originating variable of r. |
8.6. Logical Expressions
Precondition | Conclusion | Description |
---|---|---|
e: T T is bool or vecN<bool> | ! e: T
| Logical negation.
The result is true when e is false and false when e is true . Component-wise when T is a vector.
|
Precondition | Conclusion | Description |
---|---|---|
e1: bool e2: bool | e1 || e2: bool
| Short-circuiting "or". Yields true if either e1 or e2 are true;
evaluates e2 only if e1 is false.
|
e1: bool e2: bool | e1 && e2: bool
| Short-circuiting "and". Yields true if both e1 and e2 are true;
evaluates e2 only if e1 is true.
|
e1: T e2: T T is bool or vecN<bool> | e1 | e2: T
| Logical "or". Component-wise when T is a vector. Evaluates both e1 and e2. |
e1: T e2: T T is bool or vecN<bool> | e1 & e2: T
| Logical "and". Component-wise when T is a vector. Evaluates both e1 and e2. |
8.7. Arithmetic Expressions
Precondition | Conclusion | Description |
---|---|---|
e: T T is AbstractInt, AbstractFloat, i32, f32, f16, vecN<AbstractInt>, vecN<AbstractFloat>, vecN<i32>, vecN<f32>, or vecN<f16> | - e: T
| Negation. Component-wise when T is a vector. If T is an integer scalar type and e evaluates to the largest negative value, then the result is e. |
Precondition | Conclusion | Description |
---|---|---|
e1 : T e2 : T S is AbstractInt, AbstractFloat, i32, u32, f32, or f16 T is S, or vecN<S> | e1 + e2 : T
| Addition. Component-wise when T is a vector. If T is a concrete integer scalar type, then the result is modulo 232. |
e1 : T e2 : T S is AbstractInt, AbstractFloat, i32, u32, f32, or f16 T is S, or vecN<S> | e1 - e2 : T
| Subtraction Component-wise when T is a vector. If T is a concrete integer scalar type, then the result is modulo 232. |
e1 : T e2 : T S is AbstractInt, AbstractFloat, i32, u32, f32, or f16 T is S, or vecN<S> | e1 * e2 : T
| Multiplication. Component-wise when T is a vector. If T is a concrete integer scalar type, then the result is modulo 232. |
e1 : T e2 : T S is AbstractInt, AbstractFloat, i32, u32, f32, or f16 T is S, or vecN<S> | e1 / e2 : T
|
Division. Component-wise when T is a vector.
If T is a signed integer scalar type, evaluates to:
Note: The need to ensure truncation behavior may require an implementation to perform more operations than when computing an unsigned division. Use unsigned division when both operands are known to have the same sign. If T is an unsigned integer scalar type, evaluates to:
|
e1 : T e2 : T S is AbstractInt, AbstractFloat, i32, u32, f32, or f16 T is S, or vecN<S> | e1 % e2 : T
|
Remainder. Component-wise when T is a vector.
If T is a signed integer scalar type, evaluates e1 and e2 once, and evaluates to:
Note: When non-zero, the result has the same sign as e1. Note: The need to ensure consistent behavior may require an implementation to perform more operations than when computing an unsigned remainder. If T is an unsigned integer scalar type, evaluates to:
If T is a floating point type, the result is equal to: |
Preconditions | Conclusions | Semantics |
---|---|---|
S is one of AbstractInt, AbstractFloat, f32, f16, i32, u32 V is vecN<S> es: S ev: V | ev + es: V
| ev + V(es)
|
es + ev: V
| V(es) + ev
| |
ev - es: V
| ev - V(es)
| |
es - ev: V
| V(es) - ev
| |
ev * es: V
| ev * V(es)
| |
es * ev: V
| V(es) * ev
| |
ev / es: V
| ev / V(es)
| |
es / ev: V
| V(es) / ev
| |
ev % es: V
| ev % V(es)
| |
es % ev: V
| V(es) % ev
|
Preconditions | Conclusions | Semantics |
---|---|---|
e1, e2: matCxR<T> T is AbstractFloat, f32, or f16 | e1 + e2: matCxR<T> | Matrix addition: column i of the result is e1[i] + e2[i] |
e1 - e2: matCxR<T>
| Matrix subtraction: column i of the result is e1[i] - e2[i] | |
m: matCxR<T> s: T T is AbstractFloat, f32, or f16 | m * s: matCxR<T> | Component-wise scaling: (m * s)[i][j] is m[i][j] * s
|
s * m: matCxR<T> | Component-wise scaling: (s * m)[i][j] is m[i][j] * s
| |
m: matCxR<T> v: vecC<T> T is AbstractFloat, f32, or f16 | m * v: vecR<T> | Linear algebra matrix-column-vector product:
Component i of the result is dot (transpose(m)[i],v)
|
m: matCxR<T> v: vecR<T> T is AbstractFloat, f32, or f16 | v * m: vecC<T> | Linear algebra row-vector-matrix product: transpose(transpose(m) * transpose(v))
|
e1: matKxR<T> e2: matCxK<T> T is AbstractFloat, f32, or f16 | e1 * e2: matCxR<T> | Linear algebra matrix product. |
8.8. Comparison Expressions
Precondition | Conclusion | Description |
---|---|---|
e1: T e2: T S is AbstractInt, AbstractFloat, bool, i32, u32, f32, or f16 T is S or vecN<S> TB is vecN<bool> if T is a vector, otherwise TB is bool | e1 == e2: TB
| Equality. Component-wise when T is a vector. |
e1: T e2: T S is AbstractInt, AbstractFloat, bool, i32, u32, f32, or f16 T is S or vecN<S> TB is vecN<bool> if T is a vector, otherwise TB is bool | e1 != e2: TB
| Inequality. Component-wise when T is a vector. |
e1: T e2: T S is AbstractInt, AbstractFloat, i32, u32, f32, or f16 T is S, or vecN<S> TB is vecN<bool> if T is a vector, otherwise TB is bool | e1 < e2: TB
| Less than. Component-wise when T is a vector. |
e1: T e2: T S is AbstractInt, AbstractFloat, i32, u32, f32, or f16 T is S, or vecN<S> TB is vecN<bool> if T is a vector, otherwise TB is bool | e1 <= e2: TB
| Less than or equal. Component-wise when T is a vector. |
e1: T e2: T S is AbstractInt, AbstractFloat, i32, u32, f32, or f16 T is S, or vecN<S> TB is vecN<bool> if T is a vector, otherwise TB is bool | e1 > e2: TB
| Greater than. Component-wise when T is a vector. |
e1: T e2: T S is AbstractInt, AbstractFloat, i32, u32, f32, or f16 T is S, or vecN<S> TB is vecN<bool> if T is a vector, otherwise TB is bool | e1 >= e2: TB
| Greater than or equal. Component-wise when T is a vector. |
8.9. Bit Expressions
Precondition | Conclusion | Description |
---|---|---|
e: T S is AbstractInt, i32, or u32 T is S or vecN<S> | ~ e : T
| Bitwise complement on e. Each bit in the result is the opposite of the corresponding bit in e. Component-wise when T is a vector. |
Precondition | Conclusion | Description |
---|---|---|
e1: T e2: T S is AbstractInt, i32, or u32 T is S or vecN<S> | e1 | e2: T
| Bitwise-or. Component-wise when T is a vector. |
e1: T e2: T S is AbstractInt, i32, or u32 T is S or vecN<S> | e1 & e2: T
| Bitwise-and. Component-wise when T is a vector. |
e1: T e2: T S is AbstractInt, i32, or u32 T is S or vecN<S> | e1 ^ e2: T
| Bitwise-exclusive-or. Component-wise when T is a vector. |
Precondition | Conclusion | Description |
---|---|---|
e1: T e2: TS S is i32 or u32 T is S or vecN<S> TS is u32 when T is S, otherwise TS is vecN<u32> | e1 << e2: T
|
Shift left (shifted value is concrete):
Shift e1 left, inserting zero bits at the least significant positions, and discarding the most significant bits. The number of bits to shift is the value of e2, modulo the bit width of e1.
When both e1 and e2 are known before shader execution start, the result must not overflow:
Component-wise when T is a vector. |
e1: T e2: TS T is AbstractInt or vecN<AbstractInt> TS is u32 when T is AbstractInt, otherwise TS is vecN<u32> | e1 << e2: T
|
Shift left (shifted value abstract):
Shift e1 left, inserting zero bits at the least significant positions, and discarding the most significant bits. The number of bits to shift is the value of e2. The e2+1 most significant bits of e1 must have the same bit value. Otherwise overflow would occur. Note: This condition means all the discarded bits must be the same as the sign bit of the original value, and the same as the sign bit of the final value. Component-wise when T is a vector. |
e1: T e2: TS S is i32 or u32 T is S or vecN<S> TS is u32 when T is S, otherwise TS is vecN<u32> | e1 >> e2: T |
Shift right (shifted value is concrete).
Shift e1 right, discarding the least significant bits. If S is an unsigned type, insert zero bits at the most significant positions. If S is a signed type:
The number of bits to shift is the value of e2, modulo the bit width of e1. If e2 is greater than or equal to the bit width or e1, then:
Component-wise when T is a vector. |
e1: T e2: TS T is AbstractInt or vecN<AbstractInt> TS is u32 when T is AbstractInt, otherwise TS is vecN<u32> | e1 >> e2: T |
Shift right (abstract).
Shift e1 right, discarding the least significant bits. If e1 is negative, each inserted bit is 1, and so the result is also negative. Otherwise, each inserted bit is 0. The number of bits to shift is the value of e2. Component-wise when T is a vector. |
8.10. Function Call Expression
A function call expression executes a function call where the called function has a return type. If the called function does not return a value, a function call statement should be used instead. See § 9.5 Function Call Statement.
8.11. Variable Identifier Expression
Precondition | Conclusion | Description |
---|---|---|
v is an identifier resolving to an in-scope variable declared in address space AS with store type T and access mode AM | v: ref<AS,T,AM> | Result is a reference to the memory for the named variable v. |
8.12. Formal Parameter Expression
Precondition | Conclusion | Description |
---|---|---|
a is an identifier resolving to an in-scope formal parameter declaration with type T | a: T | Result is the value supplied for the corresponding function call operand at the call site invoking this instance of the function. |
8.13. Address-Of Expression
The address-of operator converts a reference to its corresponding pointer.
Precondition | Conclusion | Description |
---|---|---|
r: ref<AS,T,AM> | & r: ptr<AS,T,AM>
|
Result is the pointer value corresponding to the
same memory view as the reference value r.
If r is an invalid memory reference, then the resulting pointer is also an invalid memory reference. It is a shader-creation error if AS is the handle address space. It is a shader-creation error if r is a reference to a vector component. |
8.14. Indirection Expression
The indirection operator converts a pointer to its corresponding reference.
Precondition | Conclusion | Description |
---|---|---|
p: ptr<AS,T,AM> | * p: ref<AS,T,AM>
|
Result is the reference value corresponding to the
same memory view as the pointer value p.
If p is an invalid memory reference, then the resulting reference is also an invalid memory reference. |
8.15. Identifier Expressions for Value Declarations
Precondition | Conclusion | Description |
---|---|---|
c is an identifier resolving to an in-scope const-declaration with type T | c: T | Result is the value computed for the initializer expression. The expression is a const-expression, and is evaluated at shader-creation time. |
c is an identifier resolving to an in-scope override-declaration with type T | c: T |
If pipeline creation specified a value for the constant ID,
then the result is that value.
This value may be different for different pipeline instances.
Otherwise, the result is the value computed for the initializer expression. Pipeline-overridable constants appear at module-scope, so evaluation occurs before the shader begins execution. Note: Pipeline creation fails if no initial value was specified in the API call
and the |
c is an identifier resolving to an in-scope let-declaration with type T | c: T | Result is the value computed for the initializer expression.
A let-declaration appears inside a function body, and its initializer
is evaluated each time control flow reaches the declaration. |
8.16. Enumeration Expressions
Precondition | Conclusion | Description |
---|---|---|
e is an identifier resolving to a predeclared enumerant belonging to enumeration type E | e : E | See § 6.3.1 Predeclared enumerants |
8.17. Type Expressions
Precondition | Conclusion | Description |
---|---|---|
t is an identifier resolving to a predeclared type | t : AllTypes | See § 6.9 Predeclared Types and Type-Generators Summary |
a is an identifier resolving to a type alias. | a : AllTypes | Additionally, a denotes the type to which it is aliased. |
s is an identifier resolving to the declaration of a structure type. | s : AllTypes | Additionally, s denotes the structure type. |
tg is an identifier resolving to a type-generator
e1: T1 | tg _template_args_start e1, ..., eN _template_args_end : AllTypes |
Each type-generator has its own requirements on the template parameters it requires and accepts,
and defines how the template paramters help determine the resulting type.
The expressions e1 through eN are the template parameters for the type-generator. For example, the type expression See § 6.9 Predeclared Types and Type-Generators Summary for the list of predeclared type-generators. Note: The two variants here differ only in whether they have a trailing comma after eN. |
tg _template_args_start e1, ..., eN, _template_args_end : AllTypes |
8.18. Expression Grammar Summary
When an identifier is the first token in a call_phrase, it is one of:
-
The name of a user-defined function or built-in function, as part of a function call.
-
The name of a type, type alias, or type-generator, as part of a value constructor expression.
Declaration and scope rules ensure those names are always distinct.
Note: The call_expression rule exists to ensure type checking applies to the call expression.
expression ( `','` expression ) * `','` ?
`'['` expression `']'` component_or_swizzle_specifier ?
| multiplicative_expression multiplicative_operator unary_expression
| additive_expression additive_operator multiplicative_expression
| shift_expression _less_than shift_expression
| shift_expression _greater_than shift_expression
| shift_expression _less_than_equal shift_expression
| shift_expression _greater_than_equal shift_expression
binary_and_expression `'&'` unary_expression
| short_circuit_or_expression `'||'` relational_expression
8.19. Operator Precedence and Associativity
This entire subsection is non-normative.
Operator precedence and associativity in right-hand side WGSL expressions emerge from their grammar in summary. Right-hand expressions group operators to organize them, as illustrated by the following diagram:
To promote readability through verbosity, the following groups do not associate with other groups:
-
Short-circuit OR (can associate with self and relational weakly),
-
Short-circuit AND (can associate with self and relational weakly),
-
Binary AND (can associate with self and unary weakly),
-
Binary XOR (can associate with self and unary weakly).
And the following groups do not associate with themselves:
-
Relational (can associate with additive and shift weakly).
Associating both group sections above requires parentheses to set the relationship explicitly. The following exemplifies where these rules render expressions invalid in comments:
let a = x & ( y ^ ( z | w )); // Invalid: x & y ^ z | w let b = ( x + y ) << ( z >= w ); // Invalid: x + y << z >= w let c = x < ( y > z ); // Invalid: x < y > z let d = x && ( y || z ); // Invalid: x && y || z
Emergent precedence controls the implicit parentheses of an expression, where the stronger
binding operator will act as if it is surrounded by parentheses when together with operators
of weaker precedence. For example, stronger binding multiplicative operators than additive will
infer (a + (b * c))
from a + b * c
expression. Similarly, the emergent associativity controls
the direction of these implicit parentheses. For example, a left-to-right association will
infer ((a + b) + c)
from a + b + c
expression, whereas a right-to-left association
will infer (* (* a))
from * * a
expression.
The following table summarizes operator precedence, associativity, and binding, sorting by starting with strongest to weakest. The binding column contains the stronger expression of the given operator, meaning, for example, if "All above" is the value, then this operator can include any of the stronger expressions. But, for example, if "Unary" is the value, then anything weaker than unary but stronger than the operator at row would require parentheses to bind with this operator. This column is necessary for linearly listing operators.
Name | Operators | Associativity | Binding |
---|---|---|---|
Parenthesized | (...)
| ||
Primary | a() , a[] , a.b
| Left-to-right | |
Unary | -a , !a , ~a , *a , &a
| Right-to-left | All above |
Multiplicative | a*b , a/b , a%b
| Left-to-right | All above |
Additive | a+b , a-b
| Left-to-right | All above |
Shift | a<<b , a>>b
| Requires parentheses | Unary |
Relational | a<b , a>b , a<=b , a>=b , a==b , a!=b
| Requires parentheses | All above |
Binary AND | a&b
| Left-to-right | Unary |
Binary XOR | a^b
| Left-to-right | Unary |
Binary OR | a|b
| Left-to-right | Unary |
Short-circuit AND | a&&b
| Left-to-right | Relational |
Short-circuit OR | a||b
| Left-to-right | Relational |
9. Statements
A statement is a program fragment that controls execution. Statements are generally executed in sequential order; however, control flow statements may cause a program to execute in non-sequential order.
9.1. Compound Statement
A compound statement is a brace-enclosed sequence of zero or more statements. When a declaration is one of those statements, its identifier is in scope from the start of the next statement until the end of the compound statement.
The continuing_compound_statement is a special form of compound statement that forms the body of a continuing statement, and allows an optional break-if statement at the end.
9.2. Assignment Statement
An assignment evaluates an expression, and optionally stores it in memory (thus updating the contents of a variable).
lhs_expression ( `'='` | compound_assignment_operator ) expression
The text to the left of the operator token is the left-hand side, and the expression to the right of the operator token is the right-hand side.
9.2.1. Simple Assignment
An assignment is a simple assignment when the left-hand side is an expression, and the operator is the equal ('='
) token.
In this case the value of the right-hand side is written to the memory referenced by the left-hand side.
Precondition | Statement | Description |
---|---|---|
e: T, T is a concrete constructible type, r: ref<AS,T,AM>, AS is a writable address space, access mode AM is write or read_write | r = e |
Evaluates r, then evaluates e, then writes the value computed for e into
the memory locations referenced by r.
Note: If the reference is an invalid memory reference, the write may not execute, or may write to a different memory location than expected. |
In the simplest case, the left hand side is the name of a variable. See § 6.4.8 Forming Reference and Pointer Values for other cases.
struct S { age : i32, weight : f32} var < private> person : S ; fn f () { var a : i32= 20 ; a = 30 ; // Replace the contents of 'a' with 30. person . age = 31 ; // Write 31 into the age field of the person variable. var uv : vec2< f32> ; uv . y = 1.25 ; // Place 1.25 into the second component of uv. let uv_x_ptr : ptr< function, f32> = & uv . x ; * uv_x_ptr = 2.5 ; // Place 2.5 into the first component of uv. var sibling : S ; // Copy the contents of the 'person' variable into the 'sibling' variable. sibling = person ; }
9.2.2. Phony Assignment
An assignment is a phony assignment when the left-hand side is the underscore ('_'
) token.
In this case the right-hand side is evaluated, and then ignored.
Precondition | Statement | Description |
---|---|---|
e: T, T is constructible, a pointer type, a texture type, or a sampler type | _ = e |
Evaluates e.
Note: The resulting value is not stored.
The |
A phony-assignment is useful for:
-
Calling a function that returns a value, but clearly expressing that the resulting value is not needed.
-
Statically accessing a variable, thus establishing it as a part of the shader’s resource interface.
Note: A buffer variable’s store type may not be constructible, e.g. it contains an atomic type, or a runtime-sized array. In these cases, use a pointer to the variable’s contents instead.
var < private> counter : i32; fn increment_and_yield_previous () -> i32{ let previous = counter ; counter = counter + 1 ; return previous ; } fn user () { // Increment the counter, but don’t use the result. _ = increment_and_yield_previous (); }
struct BufferContents { counter : atomic< u32> , data : array< vec4< f32>> } @group ( 0 ) @binding ( 0 ) var < storage> buf : BufferContents ; @group ( 0 ) @binding ( 1 ) var t : texture_2d< f32> ; @group ( 0 ) @binding ( 2 ) var s : sampler; @fragment fn shade_it () -> @location ( 0 ) vec4< f32> { // Declare that buf, t, and s are part of the shader interface, without // using them for anything. _ = & buf ; _ = t ; _ = s ; return vec4< f32> (); }
9.2.3. Compound Assignment
An assignment is a compound assignment when the left-hand side is an expression, and the operator is one of the compound_assignment_operators.
The type requirements, semantics, and behavior of each statement is defined as if the compound assignment expands as in the following table, except that:
-
the reference expression e1 is evaluated only once, and
-
the reference type for e1 must have a read_write access mode.
Statement | Expansion |
---|---|
e1 += e2 | e1 = e1 + (e2) |
e1 -= e2 | e1 = e1 - (e2) |
e1 *= e2 | e1 = e1 * (e2) |
e1 /= e2 | e1 = e1 / (e2) |
e1 %= e2 | e1 = e1 % (e2) |
e1 &= e2 | e1 = e1 & (e2) |
e1 |= e2 | e1 = e1 | (e2) |
e1 ^= e2 | e1 = e1 ^ (e2) |
e1 >>= e2 | e1 = e1 >> (e2) |
e1 <<= e2 | e1 = e1 << (e2) |
Note: The syntax does not allow a compound assignment to also be a phony assignment.
Note: Even though the reference e1 is evaluated once, its underlying memory is accessed twice: first a read access gets the old value, and then a write access stores the updated value.
var < private> next_item : i32= 0 ; fn advance_item () -> i32{ next_item += 1 ; // Adds 1 to next_item. return next_item - 1 ; } fn bump_item () { var data : array< f32, 10 > ; next_item = 0 ; // Adds 5.0 to data[0], calling advance_item() only once. data [ advance_item ()] += 5.0 ; // next_item will be 1 here. } fn precedence_example () { var value = 1 ; // The right-hand side of a compound assignment is its own expression. value *= 2 + 3 ; // Same as value = value * (2 + 3); // 'value' now holds 5. }
For example, when e1 is not a reference to a component inside a vector, then
e1+=
e2;
can be rewritten as
where the identifier{ let p = &(
e1); *p = *p + (
e2); }
p
is chosen to be different from all other identifiers in the program.
When e1 is a reference to a component inside a vector, the above technique needs to be modified because WGSL does not allow taking the address in that case. For example, if ev is a reference to a vector, the statement
evcan be rewritten as[
c] +=
e2;
where identifiers{ let p = &(
ev); let c0 =
c; (*p)[c0] = (*p)[c0] + (
e2); }
c0
and p
are chosen to be different from all other identifiers in the program.
9.3. Increment and Decrement Statements
An increment statement adds 1 to the contents of a variable. A decrement statement subtracts 1 from the contents of a variable.
The expression must evaluate to a reference with a concrete integer scalar store type and read_write access mode.
Precondition | Statement | Description |
---|---|---|
r : ref<AS,T,read_write>, T is a concrete integer scalar | r++
| Adds 1 to the contents of memory referenced by r. Same as r += T(1) |
r : ref<AS,T,read_write>, T is a concrete integer scalar | r--
| Subtracts 1 from the contents of memory referenced by r. Same as r -= T(1) |
fn f () { var a : i32= 20 ; a ++ ; // Now a contains 21 a -- ; // Now a contains 20 }
9.4. Control Flow
Control flow statements may cause the program to execute in non-sequential order.
9.4.1. If Statement
An if statement conditionally executes at most one compound statement based on the evaluation of condition expressions.
An if
statement has an if
clause, followed by zero or more else if
clauses, followed by an optional else
clause.
`'if'` expression compound_statement
`'else'` `'if'` expression compound_statement
`'else'` compound_statement
Type rule precondition:
The expression in each if
and else if
clause must be of bool type.
An if
statement is executed as follows:
-
The condition associated with the
if
clause is evaluated. If the result istrue
, control transfers to the first compound statement (immediately after the condition expression). -
Otherwise, the condition of the next
else if
clause in textual order (if one exists) is evaluated and, if the result istrue
, control transfers to the associated compound statement.-
This behavior is repeated for all
else if
clauses until one of the conditions evaluates totrue
.
-
-
If no condition evaluates to
true
, then control transfers to the compound statement associated with theelse
clause (if it exists).
9.4.2. Switch Statement
A switch statement transfers control to one of a set of case clauses, or to the default clause, depending on the evaluation of a selector expression.
attribute * `'switch'` expression switch_body
`'case'` case_selectors `':'` ? compound_statement
`'default'` `':'` ? compound_statement
case_selector ( `','` case_selector ) * `','` ?
`'default'`
A case clause is the 'case'
token followed by a comma-separated list of case selectors and a
body in the form of a compound statement.
A default-alone clause is the 'default'
token followed by a body in the form of a compound statement.
A default clause is either:
-
a case clause where
'default'
appears as one of its selectors, or
Each switch statement must have exactly one default clause.
The 'default'
token must not appear more than once in a single case_selector list.
Type rule precondition: For a single switch statement, the selector expression and all case selector expressions must be of the same concrete integer scalar type.
The expressions in the case_selectors must be const-expressions.
Two different case selector expressions in the same switch statement must not have the same value.
If the selector value equals the value of an expression in a case_selector list, then control is transferred to the body of that case clause. If the selector value does not equal any of the case selector values, then control is transferred to the body of the default clause.
When control reaches the end of the body of a clause, control transfers to the first statement after the switch statement.
When one of the statements in the body of a clause is a declaration, it follows the normal scope and lifetime rules of a declaration in a compound statement. That is, the body is a sequence of statements, and if one of those is a declaration then the scope of that declaration extends from the start of the next statement in the sequence until the end of the body. The declaration executes when it is reached, creating a new instance of the variable or value, and initializes it.
var a : i32; let x : i32= generateValue (); switch x { case 0 : { // The colon is optional a = 1 ; } default { // The default need not appear last a = 2 ; } case 1 , 2 , { // Multiple selector values can be used a = 3 ; } case 3 , { // The trailing comma is optional a = 4 ; } case 4 { a = 5 ; } }
const c = 2 ; var a : i32; let x : i32= generateValue (); switch x { case 0 : { a = 1 ; } case 1 , c { // Const-expression can be used in case selectors a = 3 ; } case 3 , default { // The default keyword can be used with other clauses a = 4 ; } }
9.4.3. Loop Statement
A loop statement repeatedly executes a loop body; the loop body is specified as a compound statement. Each execution of the loop body is called an iteration.
This repetition can be interrupted by a break, or return statement.
Optionally, the last statement in the loop body may be a continuing statement.
When one of the statements in the loop body is a declaration, it follows the normal scope and lifetime rules of a declaration in a compound statement. That is, the loop body is a sequence of statements, and if one of those is a declaration then the scope of that declaration extends from the start of the next statement in the sequence until the end of the loop body. The declaration executes each time it is reached, so each new iteration creates a new instance of the variable or value, and re-initializes it.
Note: The loop statement is one of the biggest differences from other shader languages.
This design directly expresses loop idioms commonly found in compiled code. In particular, placing the loop update statements at the end of the loop body allows them to naturally use values defined in the loop body.
var a : i32= 2 ; var i : i32= 0 ; // <1> loop { if i >= 4 { break ; } a = a * 2 ; i ++ ; }
- <1> The initialization is listed before the loop.
int a = 2; int step = 1; for (int i = 0; i < 4; i += step) { if (i % 2 == 0) continue; a *= 2; }
var a : i32= 2 ; var i : i32= 0 ; loop { if i >= 4 { break ; } let step : i32= 1 ; i = i + step ; if i % 2 == 0 { continue ; } a = a * 2 ; }
var a : i32= 2 ; var i : i32= 0 ; loop { if i >= 4 { break ; } let step : i32= 1 ; if i % 2 == 0 { continue ; } a = a * 2 ; continuing { // <2> i = i + step ; } }
- <2> The continue construct is placed at the end of the
loop
9.4.4. For Statement
attribute * `'for'` `'('` for_header `')'` compound_statement
for_init ? `';'` expression ? `';'` for_update ?
The for statement takes the form for (initializer; condition; update_part) { body }
and is syntactic sugar on top of a loop statement with the same body
.
Additionally:
-
If
initializer
is non-empty, it is executed inside an additional scope before the first iteration. The scope of a declaration in the initializer extends to the end of the loop body. -
Type rule precondition: If the condition is non-empty, it must be an expression of bool type.
-
If present, the condition is evaluated immediately before executing the loop body. If the condition is false, then a § 9.4.6 Break Statement is executed, finishing execution of the loop. This check is performed at the start of each loop iteration.
-
-
If
update_part
is non-empty, it becomes a continuing statement at the end of the loop body.
The initializer
of a for loop is executed once prior to executing the loop.
When a declaration appears in the initializer, its identifier is in scope until the end of the body
.
Unlike declarations in the body
, the declaration is not re-initialized each iteration.
The condition
, body
and update_part
execute in that order to form a loop iteration.
The body
is a special form of compound statement.
The identifier of a declaration in the body
is in scope from the start of
the next statement until the end of the body
.
The declaration is executed each time it is reached, so each new iteration
creates a new instance of the variable or constant, and re-initializes it.
var a : i32= 2 ; for ( var i : i32= 0 ; i < 4 ; i ++ ) { if a == 0 { continue ; } a = a + 2 ; }
Converts to:
var a : i32= 2 ; { // Introduce new scope for loop variable i var i : i32= 0 ; loop { if ! ( i < 4 ) { break ; } if a == 0 { continue ; } a = a + 2 ; continuing { i ++ ; } } }
9.4.5. While Statement
attribute * `'while'` expression compound_statement
The while statement is a kind of loop parameterized by a condition. At the start of each loop iteration, a boolean condition is evaluated. If the condition is false, the while loop ends execution. Otherwise, the rest of the iteration is executed.
Type rule precondition: The condition must be of bool type.
A while loop can be viewed as syntactic sugar over either a loop or for statement. The following statement forms are equivalent:
-
while
condition{
body_statements}
-
loop { if !
condition{break;}
body_statements}
-
for (;
condition;) {
body_statements}
9.4.6. Break Statement
`'break'`
A break statement transfers control to immediately after the body of the nearest-enclosing loop or switch statement, thus ending execution of the loop or switch statement.
A break
statement must only be used within loop, for, while, and switch statements.
A break
statement must not be placed such that it would exit from a loop’s continuing statement.
Use a break-if statement instead.
var a : i32= 2 ; var i : i32= 0 ; loop { let step : i32= 1 ; if i % 2 == 0 { continue ; } a = a * 2 ; continuing { i = i + step ; if i >= 4 { break ; } // Invalid. Use break-if instead. } }
9.4.7. Break-If Statement
`'break'` `'if'` expression `';'`
A break-if statement evaluates a boolean condition; If the condition is true, control is transferred to immediately after the body of the nearest-enclosing loop statement, ending execution of that loop.
Type rule precondition: The condition must be of bool type.
Note: A break-if statement may only appear as the last statement in the body of a continuing statement.
var a : i32= 2 ; var i : i32= 0 ; loop { let step : i32= 1 ; if i % 2 == 0 { continue ; } a = a * 2 ; continuing { i = i + step ; break if i >= 4 ; } }
9.4.8. Continue Statement
`'continue'`
A continue statement transfers control in the nearest-enclosing loop:
-
forward to the continuing statement at the end of the body of that loop, if it exists.
-
otherwise backward to the first statement in the loop body, starting the next iteration.
A continue
statement must only be used in a loop, for or while statement.
A continue
statement must not be placed such that it would transfer
control to an enclosing continuing statement.
(It is a forward branch when branching to a continuing
statement.)
A continue
statement must not be placed such that it would transfer
control past a declaration used in the targeted continuing statement.
Note: A continue
can only be used in a continuing
statement if it is used for transferring control
flow within another loop nested in the continuing
statement. That is, a continue
cannot be used to transfer control to the start of the currently executing continuing
statement.
var i : i32= 0 ; loop { if i >= 4 { break ; } if i % 2 == 0 { continue ; } // <3> let step : i32= 2 ; continuing { i = i + step ; } }
- <3> The
continue
is invalid because it bypasses the declaration ofstep
used in thecontinuing
construct
9.4.9. Continuing Statement
`'continuing'` continuing_compound_statement
A continuing statement specifies a compound statement to be executed at the end of a loop iteration. The construct is optional.
The compound statement must not contain a return at any compound statement nesting level.
9.4.10. Return Statement
`'return'` expression ?
A return statement ends execution of the current function. If the function is an entry point, then the current shader invocation is terminated. Otherwise, evaluation continues with the next expression or statement after the evaluation of the call site of the current function invocation.
If the function does not have a return type, then the return statement is optional. If the return statement is provided for such a function, it must not supply a value. Otherwise the expression must be present, and is called the return value. In this case the call site of this function invocation evaluates to the return value. The type of the return value must match the return type of the function.
9.4.11. Discard Statement
A discard statement converts the invocation into
a helper invocation and throws away the fragment.
The discard
statement must only be used in a fragment shader stage.
More precisely, executing a discard
statement will:
-
convert the current invocation into a helper invocation, and
-
prevent the current fragment from being processed downstream in the GPURenderPipeline.
Only statements executed prior to the discard
statement will have observable effects.
Note: A discard
statement may be executed by any function in a fragment stage and the effect is the same:
the fragment will be thrown away.
@group ( 0 ) @binding ( 0 ) var < storage, read_write> will_emit_color : u32; fn discard_if_shallow ( pos : vec4< f32> ) { if pos . z < 0.001 { // If this is executed, then the will_emit_color variable will // never be set to 1 because helper invocations will not write // to shared memory. discard ; } will_emit_color = 1 ; } @fragment fn main ( @builtin ( position) coord_in : vec4< f32> ) -> @location ( 0 ) vec4< f32> { discard_if_shallow ( coord_in ); // Set the value to 1 and emit red, but only if the helper function // did not execute the discard statement. will_emit_color = 1 ; return vec4< f32> ( 1.0 , 0.0 , 0.0 , 1.0 ); }
9.5. Function Call Statement
A function call statement executes a function call.
A shader-creation error results if the called function has the must_use attribute.
Note: If the function returns a value, and the function does not have the must_use attribute, that value is ignored.
9.6. Const Assertion Statement
A const assertion statement produces a shader-creation error if the
expression evaluates to false
.
The expression must be a const-expression.
The statement can satisfy static access conditions in
a shader, but otherwise has no effect on the compiled shader.
This statement can be used at module scope and within functions.
Type rule precondition: The expression must be of bool type.
`'const_assert'` expression
const x = 1 ; const y = 2 ; const_assert x < y ; // valid at module-scope. const_assert ( y != 0 ); // parentheses are optional. fn foo () { const z = x + y - 2 ; const_assert z > 0 ; // valid in functions. let a = 3 ; const_assert a != 0 ; // invalid, the expresion must be a const-expression. }
9.7. Statements Grammar Summary
The statement rule matches statements that can be used in most places inside a function body.
| variable_or_value_statement `';'`
| `'discard'` `';'`
Additionally, certain statements may only be used in very specific contexts:
9.8. Statements Behavior Analysis
9.8.1. Rules
Some statements affecting control-flow are only valid in some contexts. For example, continue is invalid outside of a loop, for, or while. Additionally, the uniformity analysis (see § 14.2 Uniformity) needs to know when control flow can exit a statement in multiple different ways.
Both goals are achieved by a system for summarizing execution behaviors of statements. Behavior analysis maps each statement to the set of possible ways execution proceeds after evaluation of the statement completes. As with type analysis for values and expressions, behavior analysis proceeds bottom up: first determine behaviors for certain basic statements, and then determine behavior for higher level constructs by applying combining rules.
A behavior is a set, whose elements may be:
-
Return
-
Break
-
Continue
-
Next
Each of those correspond to a way to exit a compound statement: either through a keyword, or by falling to the next statement ("Next").
We note "s: B" to say that s respects the rules regarding behaviors, and has behavior B.
For each function:
-
Its body must be a valid statement by these rules.
-
If the function has a return type, the behavior of its body must be {Return}.
-
Otherwise, the behavior of its body must be a subset of {Next, Return}.
We assign a behavior to each function: it is its body’s behavior (treating the body as a regular statement), with any "Return" replaced by "Next". As a consequence of the rules above, a function behavior is always one of {}, or {Next}.
Behavior analysis must be able to determine a non-empty behavior for each statement, and function.
Statement | Preconditions | Resulting behavior |
---|---|---|
empty statement | {Next} | |
{s} | s: B | B |
s1 s2
Note: s1 often ends in a semicolon. | s1: B1 Next in B1 s2: B2 | (B1∖{Next}) ∪ B2 |
s1: B1 Next not in B1 s2: B2 | B1 | |
var x:T; | {Next} | |
let x = e; | {Next} | |
var x = e; | {Next} | |
x = e; | {Next} | |
_ = e; | {Next} | |
f(e1, ..., en); | f has behavior B | B |
return; | {Return} | |
return e; | {Return} | |
discard; | {Next} | |
break; | {Break} | |
break if e; | {Break, Next} | |
continue; | {Continue} | |
if e s1 else s2 | s1: B1 s2: B2 | B1 ∪ B2 |
loop {s1 continuing {s2}} | s1: B1 s2: B2 None of {Continue, Return} are in B2 Break is not in (B1 ∪ B2) | (B1 ∪ B2)∖{Continue, Next} |
s1: B1 s2: B2 None of {Continue, Return} are in B2 Break is in (B1 ∪ B2) | (B1 ∪ B2 ∪ {Next})∖{Break, Continue} | |
switch e {case c1: s1 ... case cn: sn} | s1: B1 ... sn: Bn Break is not in (B1 ∪ ... ∪ Bn) | B1 ∪ ... ∪ Bn |
s1: B1 ... sn: Bn Break is in (B1 ∪ ... ∪ Bn) | (B1 ∪ ... ∪ Bn ∪ {Next})∖Break |
Note: ∪ is a set union operation and ∖ is a set difference operation.
Note: The empty statement case occurs when a loop
has an empty body, or when a for
loop lacks an initialization or update statement.
For the purpose of this analysis:
-
for
loops get desugared (see § 9.4.4 For Statement) -
while
loops get desugared (see § 9.4.5 While Statement) -
loop {s}
is treated asloop {s continuing {}}
-
if
statements without anelse
branch are treated as if they had an empty else branch (which adds Next to their behavior) -
if
statements withelse if
branches are treated as if they were nested simpleif/else
statements -
a switch_clause starting with
default
behaves just like a switch_clause starting withcase _:
Each built-in function has a behavior of {Next}. And each operator application not listed in the table above has the same behavior as if it were a function call with the same operands and with a function’s behavior of {Next}.
The behavior of a function must satisfy the rules given above.
Note: It is unnecessary to analyze the behavior of expressions because they will always be {Next} or a previously analyzed function will have produced a error.
9.8.2. Notes
This section is informative, non-normative.
Behavior analysis can cause a program to be rejected in the following ways (restating requirements from above):
-
The body of a function (treated as a regular statement) has a behavior not included in {Next, Return}.
-
The body of a function with a return type has a behavior which is not {Return}.
-
The behavior of a continuing block contains any of Continue, or Return.
-
Some obviously infinite loops have an empty behavior set, and are therefore invalid.
This analysis can be run in linear time, by analyzing the call-graph bottom-up (since the behavior of a function call can depend on the function’s code).
9.8.3. Examples
Here are some examples showing this analysis in action:
fn simple () -> i32{ var a : i32; return 0 ; // Behavior: {Return} a = 1 ; // Valid, statically unreachable code. // Statement behavior: {Next} // Overall behavior (due to sequential statements): {Return} return 2 ; // Valid, statically unreachable code. Behavior: {Return} } // Function behavior: {Return}
fn nested () -> i32{ var a : i32; { // The start of a compound statement. a = 2 ; // Behavior: {Next} return 1 ; // Behavior: {Return} } // The compound statement as a whole has behavior {Return} a = 1 ; // Valid, statically unreachable code. // Statement behavior: {Next} // Overall behavior (due to sequential statements): {Return} return 2 ; // Valid, statically unreachable code. Behavior: {Return} }
fn if_example () { var a : i32= 0 ; loop { if a == 5 { break ; // Behavior: {Break} } // Behavior of the whole if compound statement: {Break, Next}, // as the if has an implicit empty else a = a + 1 ; // Valid, as the previous statement had "Next" in its behavior } }
fn if_example () { var a : i32= 0 ; loop { if a == 5 { break ; // Behavior: {Break} } else { continue ; // Behavior: {Continue} } // Behavior of the whole if compound statement: {Break, Continue} a = a + 1 ; // Valid, statically unreachable code. // Statement behavior: {Next} // Overall behavior: {Break, Continue} } }
fn if_example () { var a : i32= 0 ; loop { // if e1 s1 else if e2 s2 else s3 // is identical to // if e1 else { if e2 s2 else s3 } if a == 5 { break ; // Behavior: {Break} } else if a == 42 { continue ; // Behavior: {Continue} } else { return ; // Behavior {Return} } // Behavior of the whole if compound statement: // {Break, Continue, Return} } // Behavior of the whole loop compound statement {Next, Return} } // Behavior of the whole function {Next}
fn switch_example () { var a : i32= 0 ; switch a { default : { break ; // Behavior: {Break} } } // Behavior: {Next}, as switch replaces Break by Next a = 5 ; // Valid, as the previous statement had Next in its behavior }
fn invalid_infinite_loop () { loop { } // Behavior: { }. Invalid because it’s empty. }
fn invalid_infinite_loop () { loop { discard ; // Behavior { Next }. } // Invalid, behavior of the whole loop is { }. }
fn conditional_continue () { var a : i32; loop { if a == 5 { break ; } // Behavior: {Break, Next} if a % 2 == 1 { // Valid, as the previous statement has Next in its behavior continue ; // Behavior: {Continue} } // Behavior: {Continue, Next} a = a * 2 ; // Valid, as the previous statement has Next in its behavior continuing { // Valid as the continuing statement has behavior {Next} // which does not include any of: // {Break, Continue, Return} a = a + 1 ; } } // The loop as a whole has behavior {Next}, // as it absorbs "Continue" and "Next", // then replaces "Break" with "Next" }
fn redundant_continue_with_continuing () { var a : i32; loop { if a == 5 { break ; } continue ; // Valid. This is redundant, branching to the next statement. continuing { a = a + 1 ; } } }
fn continue_end_of_loop_body () { for ( var i : i32= 0 ; i < 5 ; i ++ ) { continue ; // Valid. This is redundant, // branching to the end of the loop body. } // Behavior: {Next}, // as loops absorb "Continue", // and "for" loops always add "Next" }
for
loops desugar to loop
with a conditional break. As shown in a previous example, the conditional break has behavior {Break, Next}, which leads to adding "Next" to the loop’s behavior.
fn missing_return () -> i32{ var a : i32= 0 ; if a == 42 { return a ; // Behavior: {Return} } // Behavior: {Next, Return} } // Error: Next is invalid in the body of a // function with a return type
fn continue_out_of_loop () { var a : i32= 0 ; if a > 0 { continue ; // Behavior: {Continue} } // Behavior: {Next, Continue} } // Error: Continue is invalid in the body of a function
continue
was replaced by break
.
10. Functions
A function performs computational work when invoked.
A function is invoked in one of the following ways:
-
By evaluating a function call expression. See § 8.10 Function Call Expression.
-
By executing a function call statement. See § 9.5 Function Call Statement.
-
An entry point function is invoked by the WebGPU implementation to perform the work of a shader stage in a pipeline. See § 12 Entry Points
There are two kinds of functions:
-
A built-in function is provided by the WGSL implementation, and is always available to a WGSL module. See § 16 Built-in Functions.
-
A user-defined function is declared in a WGSL module.
10.1. Declaring a User-defined Function
A function declaration creates a user-defined function, by specifying:
-
An optional set of attributes.
-
The name of the function.
-
The formal parameter list: an ordered sequence of zero or more formal parameter declarations, which may have attributes applied, separated by commas, and surrounded by parentheses.
-
An optional return type, which may have attributes applied.
-
The function body. This is the set of statements to be executed when the function is called.
A function declaration must only occur at module scope. A function name is in scope for the entire program.
Note: Each user-defined function only has one overload.
A formal parameter declaration specifies an identifier name and a type for a value that must be provided when invoking the function. A formal parameter may have attributes. See § 10.2 Function Calls. The scope of the identifier is the function body. Two formal parameters for a given function must not have the same name.
Note: Some built-in functions may allow parameters to be abstract numeric types; however, this functionality is not currently supported for user-declared functions.
The return type, if specified, must be constructible.
WGSL defines the following attributes that can be applied to function declarations:
-
the shader stage attributes: vertex, fragment, and compute
WGSL defines the following attributes that can be applied to function parameters and return types:
`'fn'` ident `'('` param_list ? `')'` ( `'->'` attribute * template_elaborated_ident ) ?
// Declare the add_two function. // It has two formal parameters, i and b. // It has a return type of i32. // It has a body with a return statement. fn add_two ( i : i32, b : f32) -> i32{ return i + 2 ; // A formal parameter is available for use in the body. } // A compute shader entry point function, 'main'. // It has no specified return type. // It invokes the add_two function, and captures // the resulting value in the named value 'six'. @compute @workgroup_size ( 1 ) fn main () { let six : i32= add_two ( 4 , 5.0 ); }
10.2. Function Calls
A function call is a statement or expression which invokes a function.
The function containing the function call is the calling function, or caller. The function being invoked is the called function, or callee.
The function call:
-
Names the called function, and
-
Provides a parenthesized, comma-separated list of argument value expressions.
The function call must supply the same number of argument values as there are formal parameters in the called function. Each argument value must evaluate to the same type as the corresponding formal parameter, by position.
In summary, when calling a function:
-
Execution of the calling function is suspended.
-
The called function executes until it returns.
-
Execution of the calling function resumes.
A called function returns as follows:
-
A built-in function returns when its work has completed.
-
A user-defined function with a return type returns when it executes a return statement.
-
A user-defined function with no return type returns when it executes a return statement, or when execution reaches the end of its function body.
In detail, when a function call is executed the following steps occur:
-
Function call argument values are evaluated. The relative order of evaluation is left-to-right.
-
Execution of the calling function is suspended. All function scope variables and constants maintain their current values.
-
If the called function is user-defined, memory is allocated for each function scope variable in the called function.
-
Initialization occurs as described in § 7.3 var Declarations.
-
-
Values for the formal parameters of the called function are determined by matching the function call argument values by position. For example, the first formal parameter of the called function will have the value of the first argument at the call site.
-
Control is transferred to the called function. If the called function is user-defined, execution proceeds starting from the first statement in the body.
-
The called function is executed, until it returns.
-
Control is transferred back to the calling function, and the called function’s execution is unsuspended. If the called function returns a value, that value is supplied for the value of the function call expression.
The location of a function call is referred to as a call site, specifically the location of the first token in the parsed instance of the call_phrase grammar rule. Call sites are a dynamic context. As such, the same textual location may represent multiple call sites.
Note: It is possible that a function call in a fragment shader never returns if all of the invocations in a quad are discarded. In such a case, control will not be tranferred back to the calling function.
10.3. const
Functions
A function declared with a const attribute can be evaluated at shader-creation time. These functions are called const-functions. Calls to these functions can part of const-expressions.
It is a shader-creation error if the function contains any expressions that are not const-expressions, or any declarations that are not const-declarations.
Note: The const attribute cannot be applied to user-declared functions.
const first_one = firstLeadingBit ( 1234 + 4567 ); // Evaluates to 12 // first_one has the type i32, because // firstLeadingBit cannot operate on // AbstractInt @id ( 1 ) override x : i32; override y = firstLeadingBit ( x ); // const-expressions can be // used in override-expressions. // firstLeadingBit(x) is not a // const-expression in this context. fn foo () { var a : array< i32, firstLeadingBit ( 257 ) > ; // const-functions can be used in // const-expressions if all their // parameters are const-expressions. }
10.4. Restrictions on Functions
-
A vertex shader must return the position built-in output value.
-
An entry point must never be the target of a function call.
-
If a function has a return type, it must be a constructible type.
-
A function parameter must one the following types:
-
a constructible type
-
a pointer type
-
a texture type
-
a sampler type
-
-
Each function call argument must evaluate to the type of the corresponding function parameter.
-
In particular, an argument that is a pointer must agree with the formal parameter on address space, store type, and access mode.
-
Note: Recursion is disallowed because cycles are not permitted among any kinds of declarations.
fn bar ( p : ptr< function, f32> ) { } fn baz ( p : ptr< private, i32> ) { } fn bar2 ( p : ptr< function, f32> ) { let a = &*&* ( p ); bar ( p ); // Valid bar ( a ); // Valid } fn baz2 ( p : ptr< storage, f32> ) { } struct S { x : i32} @group ( 0 ) @binding ( 0 ) var < storage> ro_storage : f32; @group ( 0 ) @binding ( 1 ) var < storage, read_write> rw_storage : f32; var usable_priv : i32; var unusable_priv : array< i32, 4 > ; fn foo () { var usable_func : f32; var unusable_func : S ; var i32_func : i32; let a_priv = & usable_priv ; let b_priv = a_priv ; let c_priv = &*& usable_priv ; let d_priv = & ( unusable_priv . x ); let e_priv = d_priv ; let a_func = & usable_func ; let b_func = & unusable_func ; let c_func = & ( * b_func )[ 0 ]; let d_func = c_func ; let e_func = &* a_func ; baz ( & usable_priv ); // Valid, address-of a variable. baz ( a_priv ); // Valid, effectively address-of a variable. baz ( b_priv ); // Valid, effectively address-of a variable. baz ( c_priv ); // Valid, effectively address-of a variable. baz ( d_priv ); // Valid, memory view has changed. baz ( e_priv ); // Valid, memory view has changed. baz ( & i32_func ); // Invalid, address space mismatch. bar ( & usable_func ); // Valid, address-of a variable. bar ( c_func ); // Valid, memory view has changed. bar ( d_func ); // Valid, memory view has changed. bar ( e_func ); // Valid, effectively address-of a variable. baz2 ( & ro_storage ); // Valid, address-of a variable. baz2 ( & rw_storage ); // Invalid, access mode mismatch. }
10.4.1. Alias Analysis
10.4.1.1. Root Identifier
Memory locations can be accessed during the execution of a function using memory views. Within a function, each memory view has a particular root identifier, which names the variable or formal parameter that first provides access to that memory in that function.
Locally derived expressions of reference or pointer type may introduce new names for a particular root identifier, but each expression has a statically determinable root identifier.
Given an expression E of pointer or reference type, the root identifier is the originating variable or formal parameter of pointer type found as follows:
-
If E is an identifier resolving to a variable, then the root identifier is that variable.
-
If E is an identifier resolving to a formal parameter of pointer type, then the root identifier is that formal parameter.
-
If E is an identifier resolving to a let-declaration with initializer E2, then the root identifier is the root identifier of E2.
-
If E is of the form
(
E2)
,&
E2,*
E2, or E2[
Ei]
then the root identifier is the root identifier of E2. -
If E is a vector access expression of the form E2.swiz, where swiz is a swizzle name, then the root identifer is the root identifier of E2.
-
If E is a structure access expression of the form E2.member_name, then the root identifer is the root identifier of E2.
10.4.1.2. Aliasing
While the originating variable of a root identifier is a dynamic concept that depends on the call sites for the function, WGSL modules can be statically analyzed to determine the set of all possible originating variables for each root identifier.
Two root identifiers alias when they have the same originating variable. Execution of a WGSL function must not potentially access memory through aliased root identifiers, where one access is a write and the other is a read or a write. This is determined by analyzing the program from the leaves of the callgraph upwards (i.e. topological order). For each function the analysis records the following sets:
-
Module-scope variables that are written. This includes any module-scope variables that are written in functions called from this function.
-
Module-scope variables that are read. This includes any module-scope variables that are read in functions called from this function.
-
Pointer parameters used as root identifiers of memory views that are written in this function or in called functions.
-
Pointer parameters used as root identifiers of memory views that are read in this function or in called functions.
At each call site of a function, it is a shader-creation error if any of the following occur:
-
Two arguments of pointer type have the same root identifier and either corresponding parameter is in the written parameter set.
-
An argument of pointer type whose root identifier is a module-scope variable where:
-
the corresponding pointer parameter is in the set of written pointer parameters, and
-
the module-scope variable is in the read set for the called function.
-
-
An argument of pointer type whose root identifier is a module-scope variable where:
-
the corresponding pointer parameter is in the set of written pointer parameters, and
-
the module-scope variable is in the written set for the called function.
-
-
An argument of pointer type whose root identifier is a module-scope variable where:
-
the corresponding pointer parameter is in the set of read pointer parameters, and
-
the module-scope variable is in the written set for the called function.
-
var < private> x : i32= 0 ; fn f1 ( p1 : ptr< function, i32> , p2 : ptr< function, i32> ) { * p1 = * p2 ; } fn f2 ( p1 : ptr< function, i32> , p2 : ptr< function, i32> ) { f1 ( p1 , p2 ); } fn f3 () { var a : i32= 0 ; f2 ( & a , & a ); // Invalid. Cannot pass two pointer parameters // with the same root identifier when one or // more are written (even by a subfunction). } fn f4 ( p1 : ptr< function, i32> , p2 : ptr< function, i32> ) -> i32{ return * p1 + * p2 ; } fn f5 () { var a : i32= 0 ; let b = f4 ( & a , & a ); // Valid. p1 and p2 in f4 are both only read. } fn f6 ( p : ptr< private, i32> ) { x = * p ; } fn f7 ( p : ptr< private, i32> ) -> i32{ return x + * p ; } fn f8 () { let a = f6 ( & x ); // Invalid. x is written as a global variable and // read as a parameter. let b = f7 ( & x ); // Valid. x is only read as both a parameter and // a variable. }
11. Attributes
An attribute modifies an object. WGSL provides a unified syntax for applying attributes. Attributes are used for a variety of purposes such as specifying the interface with the API.
Generally speaking, from the language’s point-of-view, attributes can be ignored for the purposes of type and semantic checking. Additionally, the attribute name is a context-dependent name, and some attribute parameters are also context-dependent names.
Unless explicitly permitted below, an attribute must not be specified more than once per object or type.
Attribute | Valid Values | Description |
---|---|---|
align
| Must be a const-expression that resolves to an i32 or u32. Must be positive. |
Must only be applied to a member of a structure type.
Must be a power of 2. Note: This attribute influences how a value of the enclosing structure type can appear in memory: at which byte addresses the structure itself and its component members can appear. In particular, the rules in § 13.4 Memory Layout combine to imply the following constraint: If |
binding
| Must be a const-expression that resolves to an i32 or u32. Must be non-negative. |
Must only be applied to a resource variable.
Specifies the binding number of the resource in a bind group. See § 12.3.2 Resource Interface. |
builtin
| Must be an enumerant for a built-in value. |
Must only be applied to an entry point
function parameter, entry point return type, or member of a structure.
Specifies that the associated object is a built-in value, as denoted by the specified enumerant. See § 12.3.1.1 Built-in Inputs and Outputs. |
const
| None |
Must only be applied to function declarations.
Specifies that the function can be used as a const-function. This attribute must not be applied to a user-defined function. Note: This attribute is used as a notational convention to describe which built-in functions can be used in const-expressions. |
diagnostic
|
Two parameters.
The first parameter is a severity_control_name. The second parameter is a diagnostic_rule_name token specifying a triggering rule. |
Specifies a range diagnostic filter. See § 2.3 Diagnostics.
More than one diagnostic attribute may be specified on a syntactic form, but they must specify different triggering rules. |
group
| Must be a const-expression that resolves to an i32 or u32. Must be non-negative. |
Must only be applied to a resource variable.
Specifies the binding group of the resource. See § 12.3.2 Resource Interface. |
id
| Must be a const-expression that resolves to an i32 or u32. Must be non-negative. |
Must only be applied to an override-declaration of scalar type.
Specifies a numeric identifier as an alternate name for a pipeline-overridable constant. |
interpolate
|
One or two parameters.
The first parameter must be an enumerant for an interpolation type. The second parameter, if present, must be an enumerant for the interpolation sampling. |
Must only be applied to a declaration that
has a location attribute applied.
Specifies how the user-defined IO must be interpolated. The attribute is only significant on user-defined vertex outputs and fragment inputs. See § 12.3.1.4 Interpolation. |
invariant
| None |
Must only be applied to the position built-in value.
When applied to the position built-in output value of a vertex
shader, the computation of the result is invariant across different
programs and different invocations of the same entry point.
That is, if the data and control flow match for two Note: This attribute maps to the |
location
| Must be a const-expression that resolves to an i32 or u32. Must be non-negative. |
Must only be applied to an entry point function parameter, entry point
return type, or member of a structure type. Must only be applied to declarations of objects with numeric scalar or numeric vector type. Must not be used with the compute shader stage.
Specifies a part of the user-defined IO of an entry point. See § 12.3.1.3 Input-output Locations. |
must_use
| None |
Must only be applied to the declaration of a function with a return type.
Specifies that a call to this function must be used as an expression. That is, a call to this function must not be the entirety of a function call statement. Note: Many functions return a value and do not have side effects.
It is often a programming defect to call such a function as the only thing in a function call statement.
Built-in functions with these properties are declared as Note: To deliberately work around the |
size
| Must be a const-expression that resolves to an i32 or u32. Must be positive. |
Must only be applied to a member of a structure type.
The member type must have creation-fixed footprint.
The number of bytes reserved in the struct for this member. This number must be at least the byte-size of the type of the member: If |
workgroup_size
|
One, two or three parameters.
Each parameter must be a const-expression or an override-expression. All parameters must be the same type, either i32 or u32. A shader-creation error results if any specified parameter is a const-expression that evaluates to a non-positive value. A pipeline-creation error results if any specified parameter evaluates to a non-positive value or exceeds an upper bound specified by the WebGPU API, or if the product of the parameter values exceeds the upper bound specified by the WebGPU API (see WebGPU § 3.6.2 Limits). |
Must be applied to a compute shader entry point function. Must not be applied to any other object.
Specifies the x, y, and z dimensions of the workgroup grid for the compute shader. The first parameter specifies the x dimension. The second parameter, if provided, specifies the y dimension, otherwise is assumed to be 1. The third parameter, if provided, specifies the z dimension, otherwise is assumed to be 1. |
The shader stage attributes below designate a function as an entry point for a particular shader stage. These attributes must only be applied to function declarations, and at most one may be present on a given function. They take no parameters.
Attribute | Description |
---|---|
vertex | Declares the function to be an entry point for the vertex shader stage of a render pipeline. |
fragment | Declares the function to be an entry point for the fragment shader stage of a render pipeline. |
compute | Declares the function to be an entry point for the compute shader stage of a compute pipeline. |
`'@'` `'align'` `'('` expression attrib_end
| `'@'` `'binding'` `'('` expression attrib_end
| `'@'` `'builtin'` `'('` expression attrib_end
| `'@'` `'const'`
| `'@'` `'diagnostic'` diagnostic_control
| `'@'` `'group'` `'('` expression attrib_end
| `'@'` `'id'` `'('` expression attrib_end
| `'@'` `'interpolate'` `'('` expression attrib_end
| `'@'` `'interpolate'` `'('` expression `','` expression attrib_end
| `'@'` `'invariant'`
| `'@'` `'location'` `'('` expression attrib_end
| `'@'` `'must_use'`
| `'@'` `'size'` `'('` expression attrib_end
| `'@'` `'workgroup_size'` `'('` expression attrib_end
| `'@'` `'workgroup_size'` `'('` expression `','` expression attrib_end
| `'@'` `'workgroup_size'` `'('` expression `','` expression `','` expression attrib_end
| `'@'` `'vertex'`
| `'@'` `'fragment'`
| `'@'` `'compute'`
`'('` severity_control_name `','` diagnostic_rule_name attrib_end
12. Entry Points
An entry point is a user-defined function that performs the work for a particular shader stage.
12.1. Shader Stages
WebGPU issues work to the GPU in the form of draw or dispatch commands. These commands execute a pipeline in the context of a set of shader stage inputs, outputs, and attached resources.
A pipeline describes the work to be performed on the GPU, as a sequence of stages, some of which are programmable. In WebGPU, a pipeline is created before scheduling a draw or dispatch command for execution. There are two kinds of pipelines: GPUComputePipeline, and GPURenderPipeline.
A dispatch command uses a GPUComputePipeline to run a compute shader stage over a logical grid of points with a controllable amount of parallelism, while reading and possibly updating buffer and image resources.
A draw command uses a GPURenderPipeline to run a multi-stage process with two programmable stages among other fixed-function stages:
-
A vertex shader stage maps input attributes for a single vertex into output attributes for the vertex.
-
Fixed-function stages map vertices into graphic primitives (such as triangles) which are then rasterized to produce fragments.
-
A fragment shader stage processes each fragment, possibly producing a fragment output.
-
Fixed-function stages consume a fragment output, possibly updating external state such as color attachments and depth and stencil buffers.
The WebGPU specification describes pipelines in greater detail.
WGSL defines three shader stages, corresponding to the programmable parts of pipelines:
-
compute
-
vertex
-
fragment
Each shader stage has its own set of features and constraints, described elsewhere.
12.2. Entry Point Declaration
To create an entry point, declare a user-defined function with a shader stage attribute.
When configuring a pipeline in the WebGPU API,
the entry point’s function name maps to the entryPoint
attribute of the
WebGPU GPUProgrammableStage
object.
The entry point’s formal parameters denote the stage’s shader stage inputs. The entry point’s return value, if specified, denotes the stage’s shader stage outputs.
The type of each formal parameter, and the entry point’s return type, must be one of:
-
a structure whose member types are any of bool, numeric scalar, or numeric vector.
A structure type can be used to group user-defined inputs with each other and optionally with built-in inputs. A structure type can be used as the return type to group user-defined outputs with each other and optionally with built-in outputs.
Note: The bool case is forbidden for user-defined inputs and outputs. It is only permitted for the front_facing builtin value.
Note: Compute entry points never have a return type.
@vertex fn vert_main () -> @builtin ( position) vec4< f32> { return vec4< f32> ( 0.0 , 0.0 , 0.0 , 1.0 ); } @fragment fn frag_main ( @builtin ( position) coord_in : vec4< f32> ) -> @location ( 0 ) vec4< f32> { return vec4< f32> ( coord_in . x , coord_in . y , 0.0 , 1.0 ); } @compute @workgroup_size ( 1 ) fn comp_main () { }
The set of functions in a shader stage is the union of:
-
The entry point function for the stage.
-
The targets of function calls from within the body of a function in the shader stage, whether or not that call is executed.
The union is applied repeatedly until it stabilizes. It will stabilize in a finite number of steps.
12.2.1. Function Attributes for Entry Points
WGSL defines the following attributes that can be applied to entry point declarations:
-
the shader stage attributes: vertex, fragment, and compute
@compute @workgroup_size ( 8 , 4 , 1 ) fn sorter () { } @compute @workgroup_size ( 8u ) fn reverser () { } // Using an pipeline-overridable constant. @id ( 42 ) override block_width = 12u ; @compute @workgroup_size ( block_width ) fn shuffler () { } // Error: workgroup_size must be specified on compute shader @compute fn bad_shader () { }
12.3. Shader Interface
The shader interface is the set of objects through which the shader accesses data external to the shader stage, either for reading or writing, and the pipeline-overridable constants used to configure the shader. The interface includes:
-
Attached resources, which include:
A declaration D is statically accessed by a shader when:
-
An identifier resolving to D appears in the declaration of any of the functions in the shader stage.
-
An identifier resolving to D is used to define a type for a statically accessed declaration.
-
An identifier resolving to D is used in the initializer for a statically accessed declaration.
-
An identifier resolving to D is used by an attribute used by a statically accessed declaration.
-
All the parts of a function declaration including attributes, formal parameters, return type, and function body.
-
Any type needed to define the above, including following type aliases.
-
As a particular case of helping to define a type, any override-declaration used in an override-expression that is the element count of an array type for a variable in the workgroup address space, when that variable itself is statically accessed.
-
Any override declarations used to support the evaluation of override-expressions in any of the above.
-
Any attributes on any of the above.
We can now precisely define the interface of a shader as consisting of:
-
The formal parameters of the entry point. These denote the shader stage inputs.
-
The return value of the entry point. This denotes the shader stage outputs.
-
The uniform buffer, storage buffer, texture resource, and sampler resource variables statically accessed by the shader.
-
The override-declarations statically accessed by the shader.
12.3.1. Inter-stage Input and Output Interface
A shader stage input is a datum provided to the shader stage from upstream in the pipeline. Each datum is either a built-in input value, or a user-defined input.
A shader stage output is a datum the shader provides for further processing downstream in the pipeline. Each datum is either a built-in output value, or a user-defined output.
IO attributes are used to establish an object as a shader stage input or a shader stage output, or to further describe the properties of an input or output. The IO attributes are:
12.3.1.1. Built-in Inputs and Outputs
A built-in input value provides access to system-generated control information. An entry point must not contain duplicated built-in inputs.
A built-in input for stage S with name X and type TX is accessed via a formal parameter to an entry point for shader stage S, in one of two ways:
-
The parameter has attribute
builtin(
X)
and is of type TX. -
The parameter has structure type, where one of the structure members has attribute
builtin(
X)
and is of type TX.
Conversely, when a parameter or member of a parameter for an entry point has a builtin attribute, the corresponding builtin must be an input for the entry point’s shader stage.
A built-in output value is used by the shader to convey control information to later processing steps in the pipeline. An entry point must not contain duplicated built-in outputs.
A built-in output for stage S with name Y and type TY is set via the return value for an entry point for shader stage S, in one of two ways:
-
The entry point return type has attribute
builtin(
Y)
and is of type TY. -
The entry point return type has structure type, where one of the structure members has attribute
builtin(
Y)
and is of type TY.
Conversely, when the return type or member of a return type for an entry point has a builtin attribute, the corresponding builtin must be an output for the entry point’s shader stage.
Note: The position built-in is both an output of a vertex shader, and an input to the fragment shader.
Collectively, built-in input and built-in output values are known as built-in values.
The following table summarizes the available built-in values. Each is a predeclared enumerant. Each is described in detail in subsequent sections.
Predeclared Name | Stage | Direction | Type |
---|---|---|---|
vertex_index | vertex | input | u32 |
instance_index | vertex | input | u32 |
position | vertex | output | vec4<f32> |
fragment | input | vec4<f32> | |
front_facing | fragment | input | bool |
frag_depth | fragment | output | f32 |
sample_index | fragment | input | u32 |
sample_mask | fragment | input | u32 |
fragment | output | u32 | |
local_invocation_id | compute | input | vec3<u32> |
local_invocation_index | compute | input | u32 |
global_invocation_id | compute | input | vec3<u32> |
workgroup_id | compute | input | vec3<u32> |
num_workgroups | compute | input | vec3<u32> |
struct VertexOutput { @builtin ( position) my_pos : vec4< f32> } @vertex fn vs_main ( @builtin ( vertex_index) my_index : u32, @builtin ( instance_index) my_inst_index : u32, ) -> VertexOutput {} struct FragmentOutput { @builtin ( frag_depth) depth : f32, @builtin ( sample_mask) mask_out : u32} @fragment fn fs_main ( @builtin ( front_facing) is_front : bool, @builtin ( position) coord : vec4< f32> , @builtin ( sample_index) my_sample_index : u32, @builtin ( sample_mask) mask_in : u32, ) -> FragmentOutput {} @compute @workgroup_size ( 64 ) fn cs_main ( @builtin ( local_invocation_id) local_id : vec3< u32> , @builtin ( local_invocation_index) local_index : u32, @builtin ( global_invocation_id) global_id : vec3< u32> , ) {}
12.3.1.1.1. frag_depth
Name | frag_depth |
Stage | fragment |
Type | f32 |
Direction | Output |
Description | Updated depth of the fragment, in the viewport depth range. |
12.3.1.1.2. front_facing
Name | front_facing |
Stage | fragment |
Type | bool |
Direction | Input |
Description | True when the current fragment is on a front-facing primitive. False otherwise. |
12.3.1.1.3. global_invocation_id
Name | global_invocation_id |
Stage | compute |
Type | vec3<u32> |
Direction | Input |
Description | The current invocation’s global invocation ID, i.e. its position in the compute shader grid. The value of global_invocation_id is equal to workgroup_id * workgroup_size + local_invocation_id. |
12.3.1.1.4. instance_index
Name | instance_index |
Stage | vertex |
Type | u32 |
Direction | Input |
Description |
Instance index of the current vertex within the current API-level draw command.
The first instance has an index equal to the |
12.3.1.1.5. local_invocation_id
Name | local_invocation_id |
Stage | compute |
Type | vec3<u32> |
Direction | Input |
Description | The current invocation’s local invocation ID, i.e. its position in the workgroup grid. |
12.3.1.1.6. local_invocation_index
Name | local_invocation_index |
Stage | compute |
Type | u32 |
Direction | Input |
Description | The current invocation’s local invocation index, a linearized index of the invocation’s position within the workgroup grid. |
12.3.1.1.7. num_workgroups
Name | num_workgroups |
Stage | compute |
Type | vec3<u32> |
Direction | Input |
Description | The dispatch size, vec3<u32>(group_count_x, group_count_y, group_count_z) , of the compute shader dispatched by the API.
|
12.3.1.1.8. position
Name | position |
Stage | vertex |
Type | vec4<f32> |
Direction | Output |
Description |
The clip position of the current vertex,
in clip space coordinates.
An output value (x,y,z,w) will map to (x/w, y/w, z/w) in WebGPU normalized device coordinates. See WebGPU § 3.3 Coordinate Systems and WebGPU § 23.3.4 Primitive Clipping. |
Name | position |
Stage | fragment |
Type | vec4<f32> |
Direction | Input |
Description |
Input position of the current fragment.
Let fp be the input position of the fragment. Then schematically: fp.xy = rp.destination.position In more detail:
See WebGPU § 3.3 Coordinate Systems and WebGPU § 23.3.5 Rasterization. |
12.3.1.1.9. sample_index
Name | sample_index |
Stage | fragment |
Type | u32 |
Direction | Input |
Description |
Sample index for the current fragment. The value is least 0 and at most sampleCount -1, where sampleCount is the MSAA sample count specified for the GPU render pipeline.
|
12.3.1.1.10. sample_mask
Name | sample_mask |
Stage | fragment |
Type | u32 |
Direction | Input |
Description | Sample coverage mask for the current fragment. It contains a bitmask indicating which samples in this fragment are covered by the primitive being rendered. |
Name | sample_mask |
Stage | fragment |
Type | u32 |
Direction | Output |
Description | Sample coverage mask control for the current fragment. The last value written to this variable becomes the shader-output mask. Zero bits in the written value will cause corresponding samples in the color attachments to be discarded. |
12.3.1.1.11. vertex_index
Name | vertex_index |
Stage | vertex |
Type | u32 |
Direction | Input |
Description |
Index of the current vertex within the current API-level draw command,
independent of draw instancing.
For a non-indexed draw, the first vertex has an index equal to the For an indexed draw, the index is equal to the index buffer entry for the
vertex, plus the |
12.3.1.1.12. workgroup_id
Name | workgroup_id |
Stage | compute |
Type | vec3<u32> |
Direction | Input |
Description |
The current invocation’s workgroup ID, i.e. the position of the
workgroup in overall compute shader grid.
All invocations in the same workgroup have the same workgroup ID. Workgroup IDs span from (0,0,0) to (group_count_x - 1, group_count_y - 1, group_count_z - 1). |
12.3.1.2. User-defined Inputs and Outputs
User-defined data can be passed as input to the start of a pipeline, passed between stages of a pipeline or output from the end of a pipeline.
Each user-defined input datum and user-defined output datum must:
-
be of numeric scalar type or numeric vector type.
-
be assigned an IO location. See § 12.3.1.3 Input-output Locations.
A compute shader must not have user-defined inputs or outputs.
12.3.1.3. Input-output Locations
Each input-output location can store a value up to 16 bytes in size. The byte size of a type is defined using the SizeOf column in § 13.4.1 Alignment and Size. For example, a four-component vector of floating-point values occupies a single location.
IO locations are specified via the location attribute.
Each user-defined input and output must have an explicitly specified IO location. Each structure member in the entry point IO must be one of either a built-in value (see § 12.3.1.1 Built-in Inputs and Outputs), or assigned a location.
Locations must not overlap within each of the following sets:
-
Members within a structure type. This applies to any structure, not just those used in shader stage inputs or outputs.
-
An entry point’s shader stage inputs, i.e. locations for its formal parameters, or for the members of its formal parameters of structure type.
Note: Location numbering is distinct between inputs and outputs: Location numbers for an entry point’s shader stage inputs do not conflict with location numbers for the entry point’s shader stage outputs.
Note: No additional rule is required to prevent location overlap within an entry point’s outputs. When the output is a structure, the first rule above prevents overlap. Otherwise, the output is a scalar or a vector, and can have only a single location assigned to it.
Note: The number of available locations for an entry point is defined by the WebGPU API.
User-defined IO can be mixed with built-in values in the same structure. For example,
// Mixed builtins and user-defined inputs. struct MyInputs { @location ( 0 ) x : vec4< f32> , @builtin ( front_facing) y : bool, @location ( 1 ) @interpolate ( flat) z : u32} struct MyOutputs { @builtin ( frag_depth) x : f32, @location ( 0 ) y : vec4< f32> } @fragment fn fragShader ( in1 : MyInputs ) -> MyOutputs { // ... }
struct A { @location ( 0 ) x : f32, // Invalid, x and y cannot share a location. @location ( 0 ) y : f32} struct B { @location ( 0 ) x : f32} struct C { // Invalid, structures with user-defined IO cannot be nested. b : B } struct D { x : vec4< f32> } @fragment // Invalid, location cannot be applied to a structure type. fn fragShader1 ( @location ( 0 ) in1 : D ) { // ... } @fragment // Invalid, in1 and in2 cannot share a location. fn fragShader2 ( @location ( 0 ) in1 : f32, @location ( 0 ) in2 : f32) { // ... } @fragment // Invalid, location cannot be applied to a structure. fn fragShader3 ( @location ( 0 ) in1 : vec4< f32> ) -> @location ( 0 ) D { // ... }
12.3.1.4. Interpolation
Authors can control how user-defined IO data is interpolated through the use of the interpolate attribute. WGSL offers two aspects of interpolation to control: the type of interpolation, and the sampling of the interpolation.
The interpolation type must be one of the following predeclared enumerants:
- perspective
-
Values are interpolated in a perspective correct manner.
- linear
-
Values are interpolated in a linear, non-perspective correct manner.
- flat
-
Values are not interpolated. Interpolation sampling is not used with
flat
interpolation.
The interpolation sampling must be one of the following predeclared enumerants:
- center
-
Interpolation is performed at the center of the pixel.
- centroid
-
Interpolation is performed at a point that lies within all the samples covered by the fragment within the current primitive. This value is the same for all samples in the primitive.
- sample
-
Interpolation is performed per sample. The fragment shader is invoked once per sample when this attribute is applied.
For user-defined IO of scalar or vector floating-point type:
-
If the interpolation attribute is not specified, then
@interpolate(perspective, center)
is assumed. -
If the interpolation attribute is specified with an interpolation type:
-
If the interpolation type is
flat
, then interpolation sampling must not be specified. -
If the interpolation type is
perspective
orlinear
, then:-
Any interpolation sampling is valid.
-
If interpolation sampling is not specified,
center
is assumed.
-
-
User-defined vertex outputs and fragment inputs of scalar or vector
integer type must always be specified as @interpolate(flat)
.
Interpolation attributes must match between vertex outputs and fragment inputs with the same location assignment within the same pipeline.
12.3.2. Resource Interface
A resource is an object which provides access to data external to a shader stage, and which is not an override-declaration and not a shader stage input or output. Resources are shared by all invocations of the shader.
There are four kinds of resources:
The resource interface of a shader is the set of module-scope resource variables statically accessed by functions in the shader stage.
Each resource variable must be declared with both group and binding attributes. Together with the shader’s stage, these identify the binding address of the resource on the shader’s pipeline. See WebGPU § 8.3 GPUPipelineLayout.
Two different resource variables in a shader must not have the same group and binding values, when considered as a pair.
12.3.3. Resource Layout Compatibility
WebGPU requires that a shader’s resource interface match the layout of the pipeline using the shader.
It is a pipeline-creation error if a WGSL variable in a resource interface is bound to an incompatible WebGPU binding resource type or binding type, where compatibility is defined by the following table.
WGSL resource | WebGPU resource type | WebGPU binding member | WebGPU binding type | |
---|---|---|---|---|
uniform buffer | GPUBufferBinding
| buffer
| GPUBufferBindingType | "uniform"
|
storage buffer with read_write access | "storage"
| |||
storage buffer with read access | "read-only-storage"
| |||
sampler | GPUSampler
| sampler
| GPUSamplerBindingType | "filtering"
|
"non-filtering"
| ||||
sampler_comparison | "comparison"
| |||
sampled texture, depth texture, or multisampled texture | GPUTextureView
| texture
| GPUTextureSampleType | "float"
|
"unfilterable-float"
| ||||
"sint"
| ||||
"uint"
| ||||
"depth"
| ||||
write-only storage texture | GPUTextureView
| storageTexture
| GPUStorageTextureAccess
| "write-only"
|
read-write storage texture | "read-write"
| |||
read-only storage texture | "read-only"
| |||
external texture | GPUExternalTexture
| externalTexture
| (not applicable) |
See the WebGPU API specification for interface validation requirements.
12.3.4. Buffer Binding Determines Runtime-Sized Array Element Count
When a storage buffer variable contains a runtime-sized array, then the number of elements in that array
is determined from the size of the corresponding GPUBufferBinding
:
Let T be the store type for a storage buffer variable, where T is a runtime-sized array type or contains a runtime-sized array type.
Let EBS be the effective buffer binding size for the
GPUBufferBinding
bound to the pipeline binding address corresponding to the storage buffer variable.Then NRuntime, i.e. the number of elements in the runtime-sized array, is the largest integer such that SizeOf(T) ≤ EBS.
In more detail, the NRuntime for a runtime-size array of type RAT is:
truncate((EBBS − array_offset) ÷ array_stride), where:
EBBS is the effective buffer binding size associated with the variable,
array_offset is the byte offset of the runtime-sized array within the store type of the variable.
It is zero if the store type is RAT, the runtime-sized array type itself.
Otherwise the store type is a structure, and its last member is the runtime-sized array. In this case array_offset is the byte offset of that member within the structure.
array_stride is the stride of the array type, i.e. StrideOf(RAT).
A shader can compute NRuntime via the arrayLength builtin function.
NRuntime is determined by the size of the corresponding buffer binding, and that can be different for each draw or dispatch command.
WebGPU validation rules ensure that 1 ≤ NRuntime.
-
The
weights
variable is a storage buffer. -
Its store type is the runtime-sized arry type
array<f32>
. -
The array offset is 0.
-
The array stride is StrideOf(array<f32>), which is 4.
@group ( 0 ) @binding ( 1 ) var < storage> weights : array< f32> ;
The following table shows examples of NRuntime for the weights
variable, based on
the corresponding effective buffer binding size.
Effective buffer binding size | NRuntime for weights variable
| Calculation |
---|---|---|
1024 | 256 | truncate( 1024 ÷ 4 ) |
1025 | 256 | truncate( 1025 ÷ 4 ) |
1026 | 256 | truncate( 1026 ÷ 4 ) |
1027 | 256 | truncate( 1027 ÷ 4 ) |
1028 | 257 | truncate( 1028 ÷ 4 ) |
-
The
lights
variable is a storage buffer. -
Its store type is
LightStorage
. -
The
point
member ofLightStorage
is a runtime-sized array of typearray<PointLight>
.
struct PointLight { // align(16) size(32) position: vec3f, // offset(0) align(16) size(12) // -- implicit member alignment padding -- // offset(12) size(4) color : vec3f, // offset(16) align(16) size(12) // -- implicit struct size padding -- // offset(28) size(4) } struct LightStorage { // align(16) pointCount : u32, // offset(0) align(4) size(4) // -- implicit member alignment padding -- // offset(4) size(12) point : array< PointLight > , // offset(16) align(16) elementsize(32) } @group ( 0 ) @binding ( 1 ) var < storage> lights : LightStorage ;
The following table shows examples of NRuntime for the point
member of the lights
variable.
Effective buffer binding size | NRuntime for point member of lights variable
| Calculation |
---|---|---|
1024 | 31 | truncate( ( 1024 - 16 ) ÷ 32) ) |
1025 | 31 | truncate( ( 1025 - 16 ) ÷ 32) ) |
1039 | 31 | truncate( ( 1039 - 16 ) ÷ 32) ) |
1040 | 32 | truncate( ( 1040 - 16 ) ÷ 32) ) |
13. Memory
In WGSL, a value of storable type may be stored in memory, for later retrieval. This section describes the structure of memory, and the semantics of operations accessing memory. See § 6.4 Memory Views for the types of values that can be placed in memory, and the types used to perform memory accesses.
13.1. Memory Locations
Memory consists of a set of distinct memory locations. Each memory location is 8-bits in size. An operation affecting memory interacts with a set of one or more memory locations. Memory operations on composites will not access padding memory locations. Therefore, the set of memory locations accessed by an operation may not be contiguous.
Two sets of memory locations overlap if the intersection of their sets of memory locations is non-empty.
13.2. Memory Access Mode
A memory access is an operation that acts on memory locations.
-
A read access observes the contents of memory locations.
-
A write access sets the contents of memory locations.
A single operation can read, write, or both read and write.
Particular memory locations may support only certain kinds of accesses, expressed as the memory’s access mode.
Access mode | Supported accesses |
---|---|
read | Supports read accesses, but not writes. |
write | Supports write accesses, but not reads. |
read_write | Supports both read and write accesses. |
WGSL predeclares the enumerants read
, write
, and read_write
.
13.3. Address Spaces
Memory locations are partitioned into address spaces. Each address space has unique properties determining mutability, visibility, the values it may contain, and how to use variables with it. See § 7 Variable and Value Declarations for more details.
The access mode of a given memory view is often determined by context:
The storage address spaces supports both read and read_write access modes. Each other address space supports only one access mode. The default access mode for each address space is described in the following table.
Address space | Sharing among invocations | Default access mode | Notes |
---|---|---|---|
function | Same invocation only | read_write | |
private | Same invocation only | read_write | |
workgroup | Invocations in the same compute shader workgroup | read_write | The element count of an outermost array may be a pipeline-overridable constant. |
uniform | Invocations in the same shader stage | read | For uniform buffer variables |
storage | Invocations in the same shader stage | read | For storage buffer variables |
handle | Invocations in the same shader stage | read | For sampler and texture variables. |
WGSL predeclares an enumerant for each address space, except for the handle
address space.
Variables in the workgroup address space must only be statically accessed in a compute shader stage.
Variables in the storage address space (storage
buffers) can only be statically accessed by a vertex shader stage if the access mode is read.
Variables whose store type is a storage texture with a write or read_write access mode cannot be statically accessed by a vertex shader stage.
See WebGPU createBindGroupLayout()
.
Note: Each address space may have different performance characteristics.
When writing a variable declaration or a pointer type in WGSL source:
-
For the storage address space, the access mode is optional, and defaults to read.
-
For other address spaces, the access mode must not be written.
13.4. Memory Layout
The layout of types in WGSL is independent of address space. Strictly speaking, however, that layout can only be observed by host-shareable buffers. Uniform buffer and storage buffer variables are used to share bulk data organized as a sequence of bytes in memory. Buffers are shared between the CPU and the GPU, or between different shader stages in a pipeline, or between different pipelines.
Because buffer data are shared without reformatting or translation, it is a dynamic error if buffer producers and consumers do not agree on the memory layout, which is the description of how the bytes in a buffer are organized into typed WGSL values. These bytes are memory locations of a value relative to a common base location.
The store type of a buffer variable must be host-shareable, with fully elaborated memory layout, as described below.
Each buffer variable must be declared in either the uniform or storage address spaces.
The memory layout of a type is significant only when evaluating an expression with:
An 8-bit byte is the most basic unit of host-shareable memory. The terms defined in this section express counts of 8-bit bytes.
We will use the following notation:
-
AlignOf(T) is the alignment of host-shareable type T.
-
AlignOfMember(S, i) is the alignment of the i’th member of the host-shareable structure S.
-
SizeOf(T) is the byte-size of host-shareable type T.
-
SizeOfMember(S, i) is the size of the i’th member of the host-shareable structure S.
-
OffsetOfMember(S, i) is the offset of the i’th member from the start of the host-shareable structure S.
-
StrideOf(A) is the element stride of host-shareable array type A, defined as the number of bytes from the start of one array element to the start of the next element. It equals the size of the array’s element type, rounded up to the alignment of the element type:
StrideOf(array<E, N>) = roundUp(AlignOf(E), SizeOf(E))
StrideOf(array<E>) = roundUp(AlignOf(E), SizeOf(E))
13.4.1. Alignment and Size
Each host-shareable or fixed footprint data type T has an alignment and size.
The alignment of a type is a constraint on where values of that type may be placed in memory, expressed as an integer: a type’s alignment must evenly divide the byte address of the starting memory location of a value of that type. Alignments enable use of more efficient hardware instructions for accessing the values, or satisfy more restrictive hardware requirements on certain address spaces. (See address space layout constraints).
Note: Each alignment value is always a power of two, by construction.
The byte-size of a type or structure member is the number of contiguous bytes reserved in host-shareable memory for the purpose of storing a value of the type or structure member. The size may include non-addressable padding at the end of the type. Consequently, loads and stores of a value might access fewer memory locations than the value’s size.
Alignment and size of host-shareable types are defined recursively in the following table:
Host-shareable type T | AlignOf(T) | SizeOf(T) |
---|---|---|
i32, u32, or f32 | 4 | 4 |
f16 | 2 | 2 |
atomic<|T|> | 4 | 4 |
vec2<T>, T is i32, u32, or f32 | 8 | 8 |
vec2<f16> | 4 | 4 |
vec3<T>, T is i32, u32, or f32 | 16 | 12 |
vec3<f16> | 8 | 6 |
vec4<T>, T is i32, u32, or f32 | 16 | 16 |
vec4<f16> | 8 | 8 |
matCxR (col-major) (General form) | AlignOf(vecR) | SizeOf(array<vecR, C>) |
mat2x2<f32> | 8 | 16 |
mat2x2<f16> | 4 | 8 |
mat3x2<f32> | 8 | 24 |
mat3x2<f16> | 4 | 12 |
mat4x2<f32> | 8 | 32 |
mat4x2<f16> | 4 | 16 |
mat2x3<f32> | 16 | 32 |
mat2x3<f16> | 8 | 16 |
mat3x3<f32> | 16 | 48 |
mat3x3<f16> | 8 | 24 |
mat4x3<f32> | 16 | 64 |
mat4x3<f16> | 8 | 32 |
mat2x4<f32> | 16 | 32 |
mat2x4<f16> | 8 | 16 |
mat3x4<f32> | 16 | 48 |
mat3x4<f16> | 8 | 24 |
mat4x4<f32> | 16 | 64 |
mat4x4<f16> | 8 | 32 |
struct S with members M1...MN | max(AlignOfMember(S,1), ... , AlignOfMember(S,N)) | roundUp(AlignOf(S), justPastLastMember) where justPastLastMember = OffsetOfMember(S,N) + SizeOfMember(S,N) |
array<E, N> | AlignOf(E) | N × roundUp(AlignOf(E), SizeOf(E)) |
array<E> | AlignOf(E) | NRuntime × roundUp(AlignOf(E),SizeOf(E)) where NRuntime is the runtime-determined number of elements of T |
13.4.2. Structure Member Layout
The internal layout of a structure is computed from the sizes and alignments of its members. By default, the members are arranged tightly, in order, without overlap, while satisfying member alignment requirements.
This default internal layout can be overriden by using layout attributes, which are:
The i’th member of structure type S has a size and alignment, denoted by SizeOfMember(S, i) and AlignOfMember(S, i), respectively. The member sizes and alignments are used to calculate each member’s byte offset from the start of the structure, as described in § 13.4.4 Internal Layout of Values.
SizeOfMember(S, i) is k if the i’th member of S has attribute size(k). Otherwise, it is SizeOf(T) where T is the type of the member.
AlignOfMember(S, i) is k if the i’th member of S has attribute align(k). Otherwise, it is AlignOf(T) where T is the type of the member.
If a structure member has the size attribute applied, the value must be at least as large as the size of the member’s type:
SizeOfMember(S, i) ≥ SizeOf(T)
Where T is the type of the i’th member of S.
The first structure member always has a zero byte offset from the start of the structure:
OffsetOfMember(S, 1) = 0
Each subsequent member is placed at the lowest offset that satisfies the member type alignment, and which avoids overlap with the previous member. For each member index i > 1:
OffsetOfMember(S, i) = roundUp(AlignOfMember(S, i ), OffsetOfMember(S, i-1) + SizeOfMember(S, i-1))
struct A { // align(8) size(24) u : f32, // offset(0) align(4) size(4) v : f32, // offset(4) align(4) size(4) w : vec2< f32> , // offset(8) align(8) size(8) x : f32// offset(16) align(4) size(4) // -- implicit struct size padding -- // offset(20) size(4) } struct B { // align(16) size(160) a : vec2< f32> , // offset(0) align(8) size(8) // -- implicit member alignment padding -- // offset(8) size(8) b : vec3< f32> , // offset(16) align(16) size(12) c : f32, // offset(28) align(4) size(4) d : f32, // offset(32) align(4) size(4) // -- implicit member alignment padding -- // offset(36) size(4) e : A , // offset(40) align(8) size(24) f : vec3< f32> , // offset(64) align(16) size(12) // -- implicit member alignment padding -- // offset(76) size(4) g : array< A , 3 > , // element stride 24 offset(80) align(8) size(72) h : i32// offset(152) align(4) size(4) // -- implicit struct size padding -- // offset(156) size(4) } @group ( 0 ) @binding ( 0 ) var < storage, read_write> storage_buffer : B ;
struct A { // align(8) size(32) u : f32, // offset(0) align(4) size(4) v : f32, // offset(4) align(4) size(4) w : vec2< f32> , // offset(8) align(8) size(8) @size ( 16 ) x : f32// offset(16) align(4) size(16) } struct B { // align(16) size(208) a : vec2< f32> , // offset(0) align(8) size(8) // -- implicit member alignment padding -- // offset(8) size(8) b : vec3< f32> , // offset(16) align(16) size(12) c : f32, // offset(28) align(4) size(4) d : f32, // offset(32) align(4) size(4) // -- implicit member alignment padding -- // offset(36) size(12) @align ( 16 ) e : A , // offset(48) align(16) size(32) f : vec3< f32> , // offset(80) align(16) size(12) // -- implicit member alignment padding -- // offset(92) size(4) g : array< A , 3 > , // element stride 32 offset(96) align(8) size(96) h : i32// offset(192) align(4) size(4) // -- implicit struct size padding -- // offset(196) size(12) } @group ( 0 ) @binding ( 0 ) var < uniform> uniform_buffer : B ;
13.4.3. Array Layout Examples
// Array where: // - alignment is 4 = AlignOf(f32) // - element stride is 4 = roundUp(AlignOf(f32),SizeOf(f32)) = roundUp(4,4) // - size is 32 = stride * number_of_elements = 4 * 8 var small_stride : array< f32, 8 > ; // Array where: // - alignment is 16 = AlignOf(vec3<f32>) = 16 // - element stride is 16 = roundUp(AlignOf(vec3<f32>), SizeOf(vec3<f32>)) // = roundUp(16,12) // - size is 128 = stride * number_of_elements = 16 * 8 var bigger_stride : array< vec3< f32> , 8 > ;
// Array where: // - alignment is 4 = AlignOf(f32) // - element stride is 4 = roundUp(AlignOf(f32),SizeOf(f32)) = 4 // If B is the effective buffer binding size for the binding on the // draw or dispatch command, the number of elements is: // N_runtime = floor(B / element stride) = floor(B / 4) @group ( 0 ) @binding ( 0 ) var < storage> weights : array< f32> ; // Array where: // - alignment is 16 = AlignOf(vec3<f32>) = 16 // - element stride is 16 = roundUp(AlignOf(vec3<f32>), SizeOf(vec3<f32>)) // = roundUp(16,12) // If B is the effective buffer binding size for the binding on the // draw or dispatch command, the number of elements is: // N_runtime = floor(B / element stride) = floor(B / 16) var < storage> directions : array< vec3< f32>> ;
13.4.4. Internal Layout of Values
This section describes how the internals of a value are placed in the byte locations of a buffer, given an assumed placement of the overall value. These layouts depend on the value’s type, and the align and size attributes on structure members.
The buffer byte offset at which a value is placed must satisfy the type alignment requirement: If a value of type T is placed at buffer offset k, then k = c × AlignOf(T), for some non-negative integer c.
The data will appear identically regardless of the address space.
When a value V of type u32 or i32 is placed at byte offset k of a host-shared buffer, then:
-
Byte k contains bits 0 through 7 of V
-
Byte k+1 contains bits 8 through 15 of V
-
Byte k+2 contains bits 16 through 23 of V
-
Byte k+3 contains bits 24 through 31 of V
Note: Recall that i32 uses twos-complement representation, so the sign bit is in bit position 31.
A value V of type f32 is represented in IEEE-754 binary32 format. It has one sign bit, 8 exponent bits, and 23 fraction bits. When V is placed at byte offset k of host-shared buffer, then:
-
Byte k contains bits 0 through 7 of the fraction.
-
Byte k+1 contains bits 8 through 15 of the fraction.
-
Bits 0 through 6 of byte k+2 contain bits 16 through 22 of the fraction.
-
Bit 7 of byte k+2 contains bit 0 of the exponent.
-
Bits 0 through 6 of byte k+3 contain bits 1 through 7 of the exponent.
-
Bit 7 of byte k+3 contains the sign bit.
A value V of type f16 is represented in IEEE-754 binary16 format. It has one sign bit, 5 exponent bits, and 10 fraction bits. When V is placed at byte offset k of host-shared buffer, then:
-
Byte k contains bits 0 through 7 of the fraction.
-
Bits 0 through 1 of byte k+1 contain bits 8 through 9 of the fraction.
-
Bits 2 through 6 of byte k+1 contain bits 0 through 4 of the exponent.
-
Bit 7 of byte k+1 contains the sign bit.
Note: The above rules imply that numeric values in host-shared buffers are stored in little-endian format.
When a value V of atomic type atomic
<T> is placed in a host-shared buffer,
it has the same internal layout as a value of the underlying type T.
When a value V of vector type vecN<T> is placed at byte offset k of a host-shared buffer, then:
-
V.x is placed at byte offset k
-
V.y is placed at byte offset k + SizeOf(T)
-
If N ≥ 3, then V.z is placed at byte offset k + 2 × SizeOf(T)
-
If N ≥ 4, then V.w is placed at byte offset k + 3 × SizeOf(T)
When a value V of matrix type matCxR<T> is placed at byte offset k of a host-shared buffer, then:
-
Column vector i of V is placed at byte offset k + i × AlignOf(vecR<T>)
When a value of array type A is placed at byte offset k of a host-shared memory buffer, then:
-
Element i of the array is placed at byte offset k + i × StrideOf(A)
When a value of structure type S is placed at byte offset k of a host-shared memory buffer, then:
-
The i’th member of the structure value is placed at byte offset k + OffsetOfMember(S,i). See § 13.4.2 Structure Member Layout.
13.4.5. Address Space Layout Constraints
The storage and uniform address spaces have different buffer layout constraints which are described in this section.
Note: All address spaces except uniform have the same constraints as the storage address space.
All structure and array types directly or indirectly referenced by a variable must obey the constraints of the variable’s address space. Violations of an address space constraint results in a shader-creation error.
In this section we define RequiredAlignOf(S, C) as the byte offset alignment requirement of values of host-shareable type S when used in address space C.
Host-shareable type S | RequiredAlignOf(S, storage) | RequiredAlignOf(S, uniform) |
---|---|---|
i32, u32, f32, or f16 | AlignOf(S) | AlignOf(S) |
atomic<T> | AlignOf(S) | AlignOf(S) |
vecN<T> | AlignOf(S) | AlignOf(S) |
matCxR<T> | AlignOf(S) | AlignOf(S) |
array<T, N> | AlignOf(S) | roundUp(16, AlignOf(S)) |
array<T> | AlignOf(S) | roundUp(16, AlignOf(S)) |
struct S | AlignOf(S) | roundUp(16, AlignOf(S)) |
Structure members of type T must have a byte offset from the start of the structure that is a multiple of the RequiredAlignOf(T, C) for the address space C:
OffsetOfMember(S, M) = k × RequiredAlignOf(T, C)
Where k is a positive integer and M is a member of structure S with type T
Arrays of element type T must have an element stride that is a multiple of the RequiredAlignOf(T, C) for the address space C:
StrideOf(array<T, N>) = k × RequiredAlignOf(T, C)
StrideOf(array<T>) = k × RequiredAlignOf(T, C)
Where k is a positive integer
Note: RequiredAlignOf(T, C) does not impose any additional restrictions on the values permitted for an align attribute, nor does it affect the rules of AlignOf(T). Data is laid out with the rules defined in previous sections and then the resulting layout is validated against the RequiredAlignOf(T, C) rules.
The uniform address space also requires that:
-
Array elements are aligned to 16 byte boundaries. That is, StrideOf(array<T,N>) = 16 × k’ for some positive integer k’.
-
If a structure member itself has a structure type
S
, then the number of bytes between the start of that member and the start of any following member must be at least roundUp(16, SizeOf(S)).
Note: The following examples show how to use align and size attributes on structure members to satisfy layout requirements for uniform buffers. In particular, these techniques can be used mechanically transform a GLSL buffer with std140 layout to WGSL.
struct S { x : f32} struct Invalid { a : S , b : f32// invalid: offset between a and b is 4 bytes, but must be at least 16 } @group ( 0 ) @binding ( 0 ) var < uniform> invalid : Invalid ; struct Valid { a : S , @align ( 16 ) b : f32// valid: offset between a and b is 16 bytes } @group ( 0 ) @binding ( 1 ) var < uniform> valid : Valid ;
struct small_stride { a : array< f32, 8 > // stride 4 } // Invalid, stride must be a multiple of 16 @group ( 0 ) @binding ( 0 ) var < uniform> invalid : small_stride ; struct wrapped_f32 { @size ( 16 ) elem : f32} struct big_stride { a : array< wrapped_f32 , 8 > // stride 16 } @group ( 0 ) @binding ( 1 ) var < uniform> valid : big_stride ; // Valid
13.5. Memory Model
In general, WGSL follows the Vulkan Memory Model. The remainder of this section describes how WGSL programs map to the Vulkan Memory Model.
Note: The Vulkan Memory Model is a textual version of a formal Alloy model.
13.5.1. Memory Operation
In WGSL, a read access is equivalent to a memory read operation in the Vulkan Memory Model. In WGSL, a write access is equivalent to a memory write operation in the Vulkan Memory Model.
A read access occurs when an invocation executes one of the following:
-
An evaluation of the Load Rule
-
Any texture builtin function except:
-
Any atomic built-in function except atomicStore
-
A workgroupUniformLoad built-in function
-
A compound assignment statement (for the left-hand side expression)
A write access occurs when an invocation executes one of the following:
-
An assignment statement (simple or compound for the left-hand side expression)
-
A textureStore built-in function
-
Any atomic built-in function except atomicLoad
-
atomicCompareExchangeWeak only performs a write if the
exchanged
member of the returned result istrue
-
Atomic read-modify-write built-in functions perform a single memory operation that is both a read access and a write access.
Read and write accesses do not occur under any other circumstances. Read and write accesses are collectively known as memory operations in the Vulkan Memory Model.
A memory operation accesses exactly the set of locations associated with the particular memory view used in the operation. For example, a memory read that accesses a u32 from a struct containing multiple members, only reads the memory locations associated with that u32 member.
Note: A write access to a component of a vector may access all memory locations associated with that vector.
struct S { a : f32, b : u32, c : f32} @group ( 0 ) @binding ( 0 ) var < storage> v : S ; fn foo () { let x = v . b ; // Does not access memory locations for v.a or v.c. }
13.5.2. Memory Model Reference
Each module-scope resource variable forms a memory model reference for the unique group and binding pair. Each other variable (i.e. variables in the function, private, and workgroup address spaces) forms a unique memory model reference for the lifetime of the variable.
13.5.3. Scoped Operations
When an invocation performs a scoped operation, it will affect one or two sets of invocations. These sets are the memory scope and the execution scope. The memory scope specifies the set of invocations that will see any updates to memory contents affected by the operation. For synchronization built-in functions, this also means that all affected memory operations program ordered before the function are visible to affected operations program ordered after the function. The execution scope specifies the set of invocations which may participate in an operation (see § 14.5 Collective Operations).
Atomic built-in functions map to atomic operations whose memory scope is:
-
Workgroup
if the atomic pointer is in the workgroup address space -
QueueFamily
if the atomic pointer is in the storage address space
Synchronization built-in functions map to control
barriers whose execution and memory scopes are Workgroup
.
Implicit and explicit derivatives have an implicit quad execution scope.
Note: If the Vulkan memory model is not enabled in generated shaders, Device
scope should be used instead of QueueFamily
.
13.5.4. Memory Semantics
All Atomic built-in functions use Relaxed
memory semantics and, thus, no storage class
semantics.
Note: Address space in WGSL is equivalent to storage class in SPIR-V.
workgroupBarrier uses AcquireRelease
memory semantics and WorkgroupMemory
semantics. storageBarrier uses AcquireRelease
memory semantics and UniformMemory
semantics. textureBarrier uses AcquireRelease
memory semantics and ImageMemory
semantics.
Note: A combined workgroupBarrier
and storageBarrier
uses AcquireRelease
ordering semantics and both WorkgroupMemory
and UniformMemory
memory
semantics.
Note: No atomic or synchronization built-in functions use MakeAvailable
or MakeVisible
semantics.
13.5.5. Private vs Non-private
All non-atomic read accesses in the storage or workgroup address spaces are considered non-private and correspond to read operations with NonPrivatePointer | MakePointerVisible
memory operands with the Workgroup
scope.
All non-atomic write accesses in the storage or workgroup address spaces are considered non-private and correspond to write operations
with NonPrivatePointer | MakePointerAvailable
memory operands with the Workgroup
scope.
14. Execution
§ 1.1 Overview describes how a shader is invoked and partitioned into invocations. This section describes further constraints on how invocations execute, individually and collectively.
14.1. Program Order Within an Invocation
Each statement in a WGSL module may be executed zero or more times during execution. For a given invocation, each execution of a given statement represents a unique dynamic statement instance.
When a statement includes an expression, the statement’s semantics determines:
-
Whether the expression is evaluated as part of statement execution.
-
The relative ordering of evaluation between independent expressions in the statement.
Expression nesting defines data dependencies which must be satisfied to
complete evaluation.
That is, a nested expression must be evaluated before the enclosing expression
can be evaluated.
The order of evaluation for operands of an expression is left-to-right in
WGSL.
For example, foo() + bar()
must evaluate foo()
before bar()
.
See § 8 Expressions.
Statements in a WGSL module are executed in control flow order. See § 9 Statements and § 10.2 Function Calls.
14.2. Uniformity
A collective operation (e.g. barrier, derivative, or a texture operation relying on an implicitly computed derivative) requires coordination among different invocations running concurrently on the GPU. The operation executes correctly and portably when all invocations execute it concurrently, i.e. in uniform control flow.
Conversely, incorrect or non-portable behavior occurs when a strict subset of invocations execute the operation, i.e. in non-uniform control flow. Informally, some invocations reach the collective operation, but others do not, or not at the same time, as a result of non-uniform control dependencies. Non-uniform control dependencies arise from control flow statements whose behavior depends on non-uniform values.
For example, a non-uniform control dependency arises when different invocations compute different values for the condition of an if, break-if, while, or for, different values for the selector of a switch, or the left-hand operand of a short-circuiting binary operator (
&&
or||
).
These non-uniform values can often be traced back to certain sources that are not statically proven to be uniform. These sources include, but are not limited to:
-
Mutable module-scope variables
-
Most built-in values, except num_workgroups and workgroup_id
-
Certain built-in functions (see § 14.2.7 Uniformity Rules for Function Calls)
To ensure correct and portable behavior, a WGSL implementation will perform a static uniformity analysis, attempting to prove that each collective operation executes in uniform control flow. Subsequent subsections describe the analysis.
A uniformity failure will be triggered when uniformity analysis cannot prove that a particular collective operation executes in uniform control flow.
-
If a uniformity failure is triggered for a builtin function that computes a derivative, then a derivative_uniformity diagnostic is triggered.
-
The diagnostic’s triggering location is the location of the call site of that builtin.
-
The diagnostic’s severity defaults to an error but can be controlled with a diagnostic filter.
-
-
If a uniformity failure is triggered for a synchronization builtin, an error diagnostic is triggered, which results in a shader-creation error.
14.2.1. Terminology and Concepts
The following definitions are merely informative, trying to give an intuition for what the analysis in the next subsection is computing. The analysis is what actually defines these concepts, and when a program is valid or breaks the uniformity rules.
For a given group of invocations:
-
If all invocations in a given scope execute as if they are executing in lockstep at a given point in the program, that point is said to have uniform control flow.
-
For a compute shader stage, the scope of uniform control flow is all invocations in the same workgroup.
-
For other shader stages, the scope of uniform control flow is all invocations for that entry point in the same draw command.
-
-
If an expression is executed in uniform control flow, and all invocations compute the same value, it is said to be a uniform value.
-
If invocations hold the same value for a local variable at every point where it is live, it is said to be a uniform variable.
14.2.2. Uniformity Analysis Overview
The remaining subsections specify a static analysis that verifies that collective operations are only executed in uniform control flow.
The analysis assumes dynamic errors do not occur. A shader stage with a dynamic error is already non-portable, no matter the outcome of uniformity analysis.
-
Sound, meaning that a uniformity failure will be triggered for a program that would break the uniformity requirements of builtins.
-
Linear time complexity, in the number of tokens in the program.
-
Refactoring a piece of code into a function, or inlining a function, cannot make a shader invalid if it was valid before the transformation.
-
If the analysis refuses a program, it provides a straightforward chain of implications that can be used by the user agent to craft a good error message.
Each function is analyzed, trying to ensure two things:
-
Uniformity requirements are satisfied when it calls other functions, and
-
Uniformity requirements are satisfied whenever it is called.
As part of this work, the analysis computes metadata about the function to help analyze its callers in turn. This means that the call graph must first be built, and functions must be analyzed from the leaves upwards, i.e. from functions that call no function outside the standard library toward the entry point. This way, whenever a function is analyzed, the metadata for all of its callees has already been computed. There is no risk of being trapped in a cycle, as recurrence is forbidden in the language.
Note: Another way of saying the same thing is that we do a topological sort of functions ordered by the "is a (possibly indirect) callee of" partial order, and analyze them in that order.
Additionally, for each function call, the analysis computes and propagates the set of triggering rules, if any, that would be triggered if that call cannot be proven to be in uniform control flow. We call this the potential-trigger-set for the call. The elements of this set are drawn from two possibilites:
-
derivative_uniformity, for functions relying on computing a derivative, or
-
an unnamed triggering rule, for the uniformity requirements that cannot be filtered.
-
This is used for compute shader functions relying on synchronization functions.
-
14.2.3. Analyzing the Uniformity Requirements of a Function
Each function is analyzed in two phases.
The first phase walks over the syntax of the function, building a directed graph along the way based on the rules in the following subsections. The second phase explores that graph, computing the constraints on calling this function, and potentially triggering a uniformity failure.
-
A specific point of the program must be executed in uniform control flow.
-
An expression must be a uniform value.
-
A variable must be a uniform variable.
-
A value stored in memory, that could be loaded via a pointer, must be a uniform value.
An edge can be understood as an implication from the statement corresponding to its source node to the statement corresponding to its target node.
For example, one uniformity requirement is that the workgroupBarrier
builtin function must only be called within uniform control flow.
To express this, we add an edge from RequiredToBeUniform.error to the node corresponding to the workgroupBarrier
call site.
One way to understand this is that RequiredToBeUniform.error corresponds to the proposition True,
so that RequiredToBeUniform.error -> X is the same as saying that X is true.
Reciprocally, to express that we cannot ensure the uniformity of something (e.g. a variable which holds the thread id), we add an edge from the corresponding node to MayBeNonUniform. One way to understand this, is that MayBeNonUniform corresponds to the proposition False, so that X -> MayBeNonUniform is the same as saying that X is false.
A consequence of this interpretation is that every node reachable from RequiredToBeUniform.error corresponds to something which is required to be uniform for the program to be valid, and every node from which MayBeNonUniform is reachable corresponds to something whose uniformity we cannot guarantee. It follows that we have a uniformity violation, triggering a uniformity failure, if there is any path from RequiredToBeUniform.error to MayBeNonUniform.
The nodes RequiredToBeUniform.warning and RequiredToBeUniform.info are used in a similar way, but instead help determine when warning or info diagnostics should be triggered:
-
If there is a path from RequiredToBeUniform.warning to MayBeNonUniform, then a warning diagnostic will be triggered.
-
If there is a path from RequiredToBeUniform.info to MayBeNonUniform, then an info diagnostic will be triggered.
As described in § 2.3 Diagnostics, lower severity diagnostics may be discarded if higher severity diagnostics have also been generated.
For each function, two tags are computed:
-
A call site tag describing the control flow uniformity requirements on the call sites of the function, and
-
A function tag describing the function’s effects on uniformity.
For each formal parameter of a function, one or two tags are computed:
-
A parameter tag describes the uniformity req