1. Introduction
WebGPU Shader 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.
[[ stage ( fragment )]] fn main () ->[[ location ( 0 )]] vec4 < f32 > { return vec4 < f32 > ( 0.4 , 0.4 , 0.8 , 1.0 ); }
1.1. Goals
-
Trivially convertable to SPIR-V
-
Constructs are defined as normative references to their SPIR-V counterparts
-
All features in WGSL are directly translatable to SPIR-V. (No polymorphism, no general pointers, no overloads, etc)
-
Features and semantics are exactly the ones of SPIR-V
-
Each item in this spec must provide the mapping to SPIR-V for the construct
1.2. Technical Overview TODO
1.3. Notation
The floor expression is defined over real numbers x:
-
⌊x⌋ = k, where k is the unique integer such that k ≤ x < k+1
The ceiling expression is defined over real numbers x:
-
⌈x⌉ = k, where k is the unique integer such that k-1 < x ≤ k
The roundUp function is defined for positive integers k and n as:
-
roundUp(k, n) = ⌈n ÷ k⌉ × k
2. Textual structure TODO
TODO: This is a stub.
A WGSL program is text. This specification does not prescribe a particular encoding for that text.
2.1. Comments
Comments begin with //
and continue to the end of the current line. There are no multi-line comments.
TODO: What indicates the end of a line? (E.g. A line ends at the next linefeed or at the end of the program)
2.2. Tokens TODO
2.3. Literals TODO
Token | Definition |
---|---|
DECIMAL_FLOAT_LITERAL
| (-?[0-9]*.[0-9]+ | -?[0-9]+.[0-9]*)((e|E)(+|-)?[0-9]+)?
|
HEX_FLOAT_LITERAL
| -?0x([0-9a-fA-F]*.?[0-9a-fA-F]+ | [0-9a-fA-F]+.[0-9a-fA-F]*)(p|P)(+|-)?[0-9]+
|
INT_LITERAL
| -?0x[0-9a-fA-F]+ | 0 | -?[1-9][0-9]*
|
UINT_LITERAL
| 0x[0-9a-fA-F]+u | 0u | [1-9][0-9]*u
|
Note: literals are parsed greedily. This means that for statements like a -5
this will not parse as a
minus
5
but instead as a
-5
which
may be unexpected. A space must be inserted after the -
if the first
expression is desired.
const_literal : INT_LITERAL | UINT_LITERAL | FLOAT_LITERAL | TRUE | FALSE
FLOAT_LITERAL : DECIMAL_FLOAT_LITERAL | HEX_FLOAT_LITERAL
2.4. Keywords TODO
TODO: Stub
See § 12.1 Keyword Summary for a list of keywords.
2.5. Identifiers TODO
Token | Definition |
---|---|
IDENT
| [a-zA-Z][0-9a-zA-Z_]*
|
Note: literals are parsed greedy. This means that for statements like a -5
this will not parse as a
minus
5
but instead as a
-5
which
may be unexpected. A space must be inserted after the -
if the first
expression is desired.
2.6. Attributes TODO
2.7. Declarations TODO
TODO: This is a stub.
(Forward Reference) A name can denote a value, a type, a function, or a variable.
2.7.1. Scoping
A declaration introduces a name, given by an identifier token. Scoping is the set of rules determining where that name may be used, in relation to the position of the declaration in the program. If a name may be used at a particular point in the program, then we say it is in scope.
There are multiple levels of scoping depending on how and where things are declared.
A declaration must not introduce a name when that name is already in scope at the start of the declaration. That is, shadow names are not allowed in WGSL.
3. Types
Note: For the syntax of declaring types in WGSL please see the § 12 Keyword and Token Summary. TODO(dneto): This note is probably editorially obsolete.
Programs calculate values. Each value in WGSL belongs to exactly one type. A type is a set of (mathematical) values.
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 document is the same as its WGSL syntax. The spelling is different for structure types, or types containing structures.
3.1. Type Checking
Type checking is the process of mapping terms in the WGSL source language to § 3 Types.
Generally, we start by determining types for the smallest WGSL source phrases, and then build up via combining rules.
If we can derive a type for the whole WGSL source program via the type rules, then we say the program is well-typed. Otherwise there is a type error and is not a valid WGSL program.
3.1.1. Explanation for those familiar with formal type checking
Much of it can be bottom-up, like usual.
The interesting bit is that the type of a pointer expression is either straightforward pointer type itself, or the pointee type, depending on its § 3.5.2 Pointer Evaluation TODO context:
-
In Indexing, Assigning (LValue), and Copying contexts, the pointer expression denotes a pointer value.
-
In a Parameter context:
-
If the parameter type matches the pointer expression’s straightforward pointer type, then the expression denotes that pointer type.
-
Otherwise the pointer expression denotes a value of the pointee type, being the value loaded (at that time) from the referenced storage.
-
In a Reading (RValue) context, the pointer expression denotes a value of the pointee type.
3.1.2. How to read type-checking rules
A type assertion is a mapping from some WGSL source expression to an WGSL type. When this specification has
e : T
we are saying the WGSL expression e is of type T. In the type rules, the WGSL source expression will often have placeholders in italics that represent sub-expressions in the grammar.
In the type checking tables, each row represents a type deduction rule: If the conditions in the precondition column are satisfied, then the type assertion in the conclusion column is also satisfied.
For convenience, we will use the following shorthands:
Scalar | scalar types: one of bool, i32, u32, f32 |
BoolVec | § 3.3.5 Vector Types with bool component |
Int | i32 or u32 |
IntVec | § 3.3.5 Vector Types with an Int component |
Integral | Int or § 3.3.5 Vector Types with an Int component |
SignedIntegral | i32 or § 3.3.5 Vector Types with an i32 component |
FloatVec | § 3.3.5 Vector Types with f32 component |
Floating | f32 or FloatVec |
Arity(T) | number of components in § 3.3.5 Vector Types T |
3.2. Void Type
The void type contains no values.
It is used where a type is required by the language but where no values are produced or consumed. For example, it is used for the return type of a function which does not produce a value.
3.3. Value Types
3.3.1. Boolean Type
The bool type contains the values true
and false
.
3.3.2. Integer Types
The u32 type is the set of 32-bit unsigned integers.
The i32 type is the set of 32-bit signed integers. It uses a two’s complementation representation, with the sign bit in the most significant bit position.
3.3.3. Floating Point Type
The f32 type is the set of 32-bit floating point values of the IEEE 754 binary32 (single precision) format. See § 10.5 Floating Point Evaluation TODO for details.
3.3.4. Scalar Types
The scalar types are bool, i32, u32, and f32.
The numeric scalar types are i32, u32, and f32.
3.3.5. Vector Types
A vector is a grouped sequence of 2, 3, or 4 scalar components.
Type | Description |
---|---|
vecN<T> | Vector of N elements 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.
3.3.6. Matrix Types
A matrix is a grouped sequence of 2, 3, or 4 floating point vectors.
Type | Description |
---|---|
matNxM<f32> | Matrix of N columns and M rows, where N and M are both in {2, 3, 4}. Equivalently, it can be viewed as N column vectors of type vecM<f32>. |
The key use case for a matrix is to embody a linear transformation. In this interpretation, the vectors of a matrix are treated as column vectors.
The product operator (*
) is used to either:
-
scale the transformation by a scalar magnitude.
-
apply the transformation to a vector.
-
combine the transformation with another matrix.
See § 5.8 Arithmetic Expressions TODO.
mat2x3 < f32 > // This is a 2 column, 3 row matrix of 32-bit floats. // Equivalently, it is 2 column vectors of type vec3<f32>.
3.3.7. Array Types
An array is an indexable grouping of element values.
Type | Description |
---|---|
array<E,N> | An N-element array of elements of type E. N must be 1 or larger. |
array<E> | A runtime-sized array of elements of type E,
also known as a runtime array.
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 § 5.6.3 Array Access Expression TODO.
An array element type must be one of:
Restrictions on runtime-sized arrays:
-
The last member of the structure type defining the store type for a variable in the storage storage class may be a runtime-sized array.
-
A runtime-sized array must not be used as the store type or contained within a store type in any other cases.
-
The type of an expression must not be a runtime-sized array type.
(dneto): Complete description of Array<E,N>
3.3.8. Structure Types
A structure is a grouping of named member values.
Type | Description |
---|---|
struct<T1,...,TN> | An ordered tuple of N members of types Tn through TN, with N being an integer greater than 0. A structure type declaration specifies an identifier name for each member. Two members of the same structure type must not have the same name. |
A structure member type must be one of:
Note: The structure member type restriction and the array element type restriction are mutually reinforcing. Combined, they imply that a pointer may not appear in any level of nesting within either an array or structure. Similarly, the same limitations apply to textures and samplers.
Attribute | Description |
---|---|
block
| Applies to a structure type. Indicates this structure type represents the contents of a buffer resource occupying a single binding slot in the shader’s resource interface. The block attribute must be applied to a structure type used
as the store type of a uniform buffer or storage buffer variable.
|
A structure type with the block attribute must not be:
-
the element type of an array type.
-
the member type in another structure.
struct_decl : decoration_list* STRUCT IDENT struct_body_decl
Struct decoration keys | Valid values | Note |
---|---|---|
block
| The block decoration takes no parameters |
struct_body_decl : BRACE_LEFT struct_member* BRACE_RIGHT struct_member : decoration_list* variable_ident_decl SEMICOLON
Struct member decoration keys | Valid values | Note |
---|---|---|
offset
| non-negative i32 literal |
Note: Layout attributes are required if the structure type is used to define a uniform buffer or a storage buffer. See § 3.4.6 Memory Layout.
// Offset decorations struct my_struct { [[offset(0)]] a : f32; [[offset(4)]] b : vec4<f32>; };
OpName %my_struct "my_struct" OpMemberName %my_struct 0 "a" OpMemberDecorate %my_struct 0 Offset 0 OpMemberName %my_struct 1 "b" OpMemberDecorate %my_struct 1 Offset 4 %my_struct = OpTypeStruct %float %v4float
// Runtime Array type RTArr = [[stride(16)]] array<vec4<f32>>; [[block]] struct S { [[offset(0)]] a : f32; [[offset(4)]] b : f32; [[offset(16)]] data : RTArr; };
OpName %my_struct "my_struct" OpMemberName %my_struct 0 "a" OpMemberDecorate %my_struct 0 Offset 0 OpMemberName %my_struct 1 "b" OpMemberDecorate %my_struct 1 Offset 4 OpMemberName %my_struct 2 "data" OpMemberDecorate %my_struct 2 Offset 16 OpDecorate %rt_arr ArrayStride 16 %rt_arr = OpTypeRuntimeArray %v4float %my_struct = OpTypeStruct %float %v4float %rt_arr
3.4. Memory TODO
TODO: This section is a stub.
In WGSL, a value of storable type may be stored in memory, for later retrieval.
In general WGSL follows the Vulkan Memory Model with the following exceptions
-
No Acquire/Release semantics
3.4.1. Memory Locations TODO
TODO: This is a stub
Memory consists of distinct locations.
3.4.2. Storable Types
The following types are storable:
-
§ 3.3.7 Array Types if its element type is storable.
-
§ 3.3.8 Structure Types if all its members are storable.
3.4.3. IO-shareable Types
The following types are IO-shareable:
-
scalar types
-
numeric vector types
-
§ 3.3.7 Array Types if its element type is IO-shareable, and the array is not runtime-sized
-
§ 3.3.8 Structure Types if all its members are IO-shareable
The following kinds of values must be of IO-shareable type:
-
Values read from or written to built-in variables.
-
Values accepted as inputs from an upstream pipeline stage.
-
Values written as output for downstream processing in the pipeline, or to an output attachment.
Note: Only built-in pipeline inputs may have a boolean type. A user input or output data attribute must not be of bool type or contain a bool type. See § 8.3.1 Pipeline Input and Output Interface.
3.4.4. Host-shareable Types
Host-shareable types are used to describe the contents of buffers which are shared between the host and the GPU, or copied between host and GPU without format translation. When used for this purpose, the type must be additionally decorated with layout attributes as described in § 3.4.6 Memory Layout. We will see in § 4.1 Module Scope Variables that the store type of uniform buffer and storage buffer variables must be host-shareable.
The following types are host-shareable:
-
numeric scalar types
-
numeric vector types
-
§ 3.3.7 Array Types if the array type has a stride attribute and its element type is host-shareable
-
§ 3.3.8 Structure Types if each member is host-shareable and has an offset attribute
Decoraton | Operand | Description |
---|---|---|
stride
| positive i32 literal | Applied to an array type. The number of bytes from the start of one element of the array to the start of the next element. |
offset
| non-negative i32 literal | Applied to a member of a structure type. The number of bytes between the start of the structure and the location of this member. |
Note: An IO-shareable type T would also be host-shareable if T and its subtypes have appropriate stride and offset attributes, and if T is not bool and does not contain a bool. Additionally, a runtime-sized array is host-shareable but is not IO-shareable.
Note: Both IO-shareable and host-shareable types have concrete sizes, but counted differently. IO-shareable types are sized by a location-count metric, see § 8.3.1.3 Input-output Locations TODO. Host-shareable types are sized by a byte-count metric, see § 3.4.6 Memory Layout.
3.4.5. Storage Classes
Memory locations are partitioned into storage classes. Each storage class has unique properties determining mutability, visibility, the values it may contain, and how to use variables with it.
Storage class | Readable by shader? Writable by shader? | Sharing among invocations | Variable scope | Restrictions on stored values | Notes |
---|---|---|---|---|---|
in | Read-only | Same invocation only | Module scope | IO-shareable | Input from an upstream pipeline stage, or from the implementation. |
out | Read-write | Same invocation only | Module scope | IO-shareable | Output to a downstream pipeline stage. |
function | Read-write | Same invocation only | Function scope | Storable | |
private | Read-write | Same invocation only | Module scope | Storable | |
workgroup | Read-write | Invocations in the same compute shader workgroup | Module scope | Storable | |
uniform | Read-only | Invocations in the same shader stage | Module scope | Host-shareable | For uniform buffer variables |
storage | Readable. Also writable if the variable is not read-only. | Invocations in the same shader stage | Module scope | Host-shareable | For storage buffer variables |
handle | Read-only | Invocations in the same shader stage | Module scope | Opaque representation of handle to a sampler or texture | Used for sampler and texture variables The token handle is reserved: it is never used in a WGSL program.
|
The note about read-only storage variables may change depending on the outcome of https://github.com/gpuweb/gpuweb/issues/935
can we remove the "in" and "out" classes entirely?
storage_class : IN | OUT | FUNCTION | PRIVATE | WORKGROUP | UNIFORM | STORAGE
WGSL storage class | SPIR-V storage class |
---|---|
in | Input |
out | Output |
uniform | Uniform |
workgroup | Workgroup |
handle | UniformConstant |
storage | StorageBuffer |
private | Private |
function | Function |
3.4.6. Memory Layout
Uniform buffer and storage buffer variables are used to share bulk data organized as a sequence of bytes in memory. Buffers are shared between the CPU and the GPU, or between different shader stages in a pipeline, or between different pipelines.
Because buffer data are shared without reformatting or translation, buffer producers and consumers must agree on the memory layout, which is the description of how the bytes in a buffer are organized into typed WGSL values.
The store type of a buffer variable must be host-shareable, with fully elaborated memory layout, as described below.
Each buffer variable must be declared in either the uniform or storage storage classes.
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:
-
Stride(A) is the value of the stride attribute of array type A.
-
Offset(S,i) is the value of the offset attribute of the i’th member of structure type S.
The remainder of this section is structured as follows:
-
Section § 3.4.6.1 Memory Layout Intent describes the intent of the layout rules. It is not normative.
-
Section § 3.4.6.2 Internal Layout of Values says how the internals of a value are placed in byte memory locations starting from a given byte offset in a buffer. For composite types the layout depends on the properties of the type itself, the storage class, and the user-specified stride and offset attributes.
-
Section § 3.4.6.3 Layout Constraints and Standard Buffer Layout describes the contraints on array strides and structure member offsets, ultimately yielding definitions for uniform buffer layout and storage buffer layout.
3.4.6.1. Memory Layout Intent
This section is informative, not normative.
The layout rules describe two sets of constraints, one for uniform buffers and one for storage buffers. They are similar in many respects, but the uniform buffer layout is more restrictive.
In particular:
-
Scalar data values are aligned to their own size.
-
There is no padding between components of a vector.
-
Two-element and four-element vectors are aligned to their size.
-
Three-element vectors are aligned as if they were four-element vectors.
-
An array’s element alignment is a multiple of the element type’s alignment.
-
In a uniform buffer, the array element alignment is also a multiple of 16.
-
-
An array’s alignment is the same as its element alignment.
-
A matrix with N columns is aligned as N column vectors without additional padding.
-
If the columns are 3-element vector, then each vector has its own internal padding to the size of a 4-element vector, as noted above.
-
-
A structure inherits the worst-case alignment of any of its members.
-
In a uniform buffer, the structure alignment is also a multiple of 16.
-
-
The members of a structure are laid out in order: earlier members are appear earlier in the buffer
Additionally we define a value’s allocation extent, or memory footprint, which determines how many memory locations must be reserved to store that value in host-shareable memory. Allocation extent is a determining factor of the minimum size of a buffer that can be bound to a uniform buffer variable or to a storage buffer variable. See § 8.3.3 Resource layout compatibility.
Compared to OpenGL:
-
For any type except column major
mat2x2
and types that contain column majormat2x2
, OpenGLstd140
layout is the same as using the tightest offset and stride assignments in WGSL uniform buffer layout. -
OpenGL
std140
layout ofmat2x2
has extra padding between column vectors that is not present in amat2x2
type in WGSL.-
The OpenGL
std140
layout of themat2x2
type has the second column vector starting 16 bytes after the first column vector. But in WGSL the second column vector ofmat2x2<f32>
starts 8 bytes after the first column vector.
-
-
For any type, OpenGL
std430
layout is the same as using the tightest offset and stride assignments in WGSL storage buffer layout. -
OpenGL supports row-major matrices, but WGSL does not.
Compared to Vulkan § 15.6.4 Offset and Stride Assignment:
-
Vulkan standard buffer layout maps to WGSL standard buffer layout rules with the following qualifications:
-
Vulkan allows a vector to be aligned to the size of its scalar component, but WGSL requires a more constrained alignment.
-
The Vulkan
scalarBlockLayout
anduniformBufferStandardLayout
features do not apply to WGSL. -
The Vulkan concept of scalar alignment does not correspond to a concept in WGSL;
-
-
The Vulkan base alignment for a type S corresponds to the WGSL alignment requirement for S in the storage storage class: Align(S,
storage
). -
The Vulkan extended alignment for a type S corresponds to the WGSL alignment requirement for S in the uniform storage class: Align(S,
uniform
). -
The Vulkan concept of improperly straddle is not permitted in WGSL, because WGSL requires vectors to be aligned to their whole size.
-
Vulkan supports non-32-bit scalar types and vector types with non-32-bit components, but WGSL does not.
-
Vulkan supports row-major matrices, but WGSL does not.
-
Vulkan allows offsets to be non-monotonic, but WGSL does not.
3.4.6.2. 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, the storage class of the buffer, the stride attribute on array types, and the offset attribute on structure type members.
A type can be used for values in both uniform and storage storage classes. This is valid as long as the layout constraints are satisifed for both storage classes. The data will appear identically in both storage classes, except for the case of matrices noted above.
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 23 of the fraction.
-
Bit 7 of byte k+2 contains bit 0 bit of the exponent.
-
Bits 0 through 6 of byte k+3 contain bits 1 through 7 of the exponent.
-
Bit 7 of byte k+3 contains the sign bit.
Note: The above rules imply that numeric values in host-shared buffers are stored in little-endian format.
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+4
-
If N ≥ 3, then V.z is placed at byte offset k+8
-
If N ≥ 4, then V.w is placed at byte offset k+12
When a matrix value M is placed at byte offset k of a host-shared memory buffer, then:
-
If M has 2 rows, then:
-
Column vector i of M is placed at byte offset k + 8 × i
-
-
If M has 3 or 4 rows, then:
-
Column vector i of M is placed at byte offset k + 16 × i
-
When a value of array type A is placed at byte offset k of a host-shared memory buffer, then:
-
Element i of the array is placed at byte offset k + i × Stride(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 + Offset(S,i)
3.4.6.3. Layout Constraints and Standard Buffer Layout
This section defines a standard buffer layout, parameterized on storage class, and the associated constraints on array strides and structure member offsets. It also provides a way to compute the number of bytes occupied by a buffer variable and by its internal components.
The alignment of a type constrains the byte index at which a value of that type may be placed relative to the start of the host-shareable buffer. The constraint is expressed below, after other necessary terms are also defined. Alignment is a function of both the type and the storage class of the buffer.
We write Align(S,C) for the alignment of host-shareable type S in storage class C, where C is either storage or storage. It is defined recursively in the following table:
Host-shareable type S | Align(S,storage )
| Align(S,uniform )
|
---|---|---|
i32, u32, or f32 | 4 | 4 |
vec2<T>, where T is one of i32, u32, or f32 | 8 | 8 |
vec3<T>, where T is one of i32, u32, or f32 | 16 | 16 |
vec4<T>, where T is one of i32, u32, or f32 | 16 | 16 |
matNx2<f32> | 8 | 8 |
matNx3<f32> | 16 | 16 |
matNx4<f32> | 16 | 16 |
array<T,N> | Align(T,storage )
| roundUp(16, Align(T,uniform ))
|
array<T> | Align(T,storage )
| roundUp(16, Align(T,uniform ))
|
struct<T1,...,Tn> | max(Align(T1,storage ),..., Align(Tn,storage ))
| roundUp(16, A), where A = max(Align(T1, uniform ),..., Align(Tn,uniform )))
|
The allocation extent of a value V is the number of contiguous bytes reserved in host-shareable memory for the purpose of storing V. It is a function of the type of V, the size of any runtime-sized array that V may contain, and the storage class of the buffer.
Note: The allocation extent may include padding inserted to satisfy alignment rules. Consequently, loads and stores of a value might access fewer memory locations than value’s allocation extent.
We write Extent(V,C) for the allocation extent of value V of host-shareable type S in storage class C, where C is either storage or storage. It is defined recursively in the following table:
Host-shareable type S | Extent(V,storage )where V is of type S | Extent(V,uniform )where V is of type S |
---|---|---|
i32, u32, or f32 | 4 | 4 |
vecN<T>, where T is one of i32, u32, or f32 | N × 4 | N × 4 |
matNx2<f32> | N × 8 | N × 8 |
matNx3<f32> | N × 16 | N × 16 |
matNx4<f32> | N × 16 | N × 16 |
array<T,N> | N × Stride(S) | N × Stride(S) |
array<T> | Nruntime × Stride(S), where Nruntime is the runtime-determined number of elements of V | Not applicable: runtime-sized arrays cannot appear in storage storage |
struct<T1,...,Tn> | roundUp(Align(S,storage ),L),where L = Offset(S,n) + Extent(Vn, storage )),and Vn is the last member of V | roundUp(Align(S,uniform ),L),where L = Offset(S,n) + Extent(Vn, uniform )),and Vn is the last member of V |
When a type S is not a runtime-sized array and it does not contain a runtime-sized array, then all values V of type S will have the same allocation extent for a storage class C. In these cases we define the allocation extent of the type S as that common value: Extent(S,C) = Extent(V,C), for any V of type S.
Note: When underlying the target is a Vulkan device, we assume the device does
not support the scalarBlockLayout
feature.
Therefore, a data value must not be placed in the padding at the end of a structure or matrix,
nor in the padding at the last element of an array.
Counting such padding as part of the allocation extent allows WGSL to capture this constraint.
Host-shareable type S satisfies standard buffer layout rules for storage class C when:
-
If S is a structure type struct<T1,...,Tn>, then it satisfies standard buffer layout rules for C when all the following are satisifed:
-
Each member type Ti satisfies standard buffer layout rules for C
-
Members do not overlap, and are laid out in declaration order:
Offset(S,i) + Extent(Ti,C) ≤ Offset(S,i+1), for 1 ≤ i < n
-
If the structure is aligned, then members will also be aligned:
Offset(S,i) = k × Align(Ti,C), for some non-negative integer k
-
-
If S is an array type
array
<E,N> orarray
<E>, then it satisfies standard buffer layout rules for C when all the following are satisifed:-
Element type E satisfies standard buffer layout rules for C
-
The stride ensures elements don’t overlap:
Stride(S) ≥ Extent(E,C)
-
If the array is aligned, then each array element is aligned:
Stride(S) = k × Align(E,C), for some positive integer k
-
For the uniform storage class, array elements are aligned to 16 byte boundaries:
If C is uniform, then Stride(S) = k × 16 for some non-negative integer k
-
-
Other host-shareable types S are not futher constrained. They always satisfy standard buffer layout rules.
Note: The consistency and completeness of these rules rely on the fact that a runtime-sized array may only appear as the last element of a structure that is the store type for a buffer variable in the storage storage class.
Host-shareable type S satisfies uniform buffer layout when S satisfies standard buffer layout rules for storage class uniform.
Host-shareable type S satisfies storage buffer layout when S satisfies standard buffer layout rules for storage class storage.
3.5. Pointer Types TODO
Type | Description |
---|---|
ptr<SC,T> | Pointer (or reference) to storage in storage class SC which can hold a value of the storable T. Here, T is the known as the pointee type. |
Note: We’ve described a SPIR-V logical pointer type.
Note: Pointers are not storable.
3.5.1. Abstract Operations on Pointers TODO
A pointer value P supports the following operations:
P.Write(V) | Place a value V into the referenced storage. V’s type must match P’s pointee type. |
P.Read() | An evaluation yielding the value currently in the P’s referenced storage. The result type is P’s pointee type. |
P.Subaccess(K) | Valid for pointers with a composite pointee type where K must evaluate to an integer between 0 and one
less than the number of components in P’s pointee type.
The subaccess evaluation yields a pointer to the storage for
the K’th component within P’s referenced storage,
using zero-based indexing. If P’s storage class is SC, and
the K’th member of P’s pointee type is of type T, then
the result type is ptr<SC,T> .
|
Note: Assignment of swizzled values is not permitted (SubaccessSwizzle).
e.g. vec4<i32> v; v.xz = vec2<i32>(0, 1);
is not allowed.
3.5.2. Pointer Evaluation TODO
TODO: This is a stub: Using pointers in context. Disambiguating which abstract operation occurs based on context: pointer semantics vs. dereferenced value semantics.
A pointer may appear in exactly the following contexts
Indexing |
A subaccessing evaluation
|
Assigning (L-Value) |
On the left hand side of an assignment operation, and the right hand side
matches the pointee type of the pointer.
|
Copying |
On the right hand side of a const-declaration, and the type of the
const-declaration matches the pointer type.
|
Parameter | Used in a function call, where the function’s parameter type matches the pointer type. |
Reading (R-Value) | Any other context. Evaluates to P.Read(), yielding a value of P’s pointee type. |
3.6. Texture and Sampler Types
A texel is a scalar or vector used as the smallest independently accessible element of a texture. The word texel is short for texture element.
A texture is a collection of texels supporting special operations useful for rendering. In WGSL, those operations are invoked via texture builtin functions. See § 15.8 Texture built-in functions for a complete list.
A WGSL texture corresponds to a WebGPU GPUTexture.
A texture is either arrayed, or non-arrayed:
-
A non-arrayed texture is a grid of texels. Each texel has a unique grid coordinate.
-
An arrayed texture is a homegeneous array of grids of texels. In an arrayed texture, each texel is identified with its unique combination of array index and grid coordinate.
A texture has the following features:
- texel format
-
The data in each texel. See § 3.6.1 Texel formats
- dimensionality
-
The number of dimensions in the grid coordinates, and how the coordinates are interpreted. The number of dimensions is 1, 2, or 3. In some cases the third coordinate is decomposed so as to specify a cube face and a layer index.
- size
-
The extent of grid coordinates along each dimension
- mipmap levels
-
The mipmap level count is at least 1 for sampled textures, and equal to 1 for storage textures.
Mip level 0 contains a full size version of the texture. Each successive mip level contains a filtered version of the previous mip level at half the size (within rounding) of the previous mip level.
When sampling a texture, an explicit or implicitly-computed level-of-detail is used to select the mip levels from which to read texel data. These are then combined via filtering to produce the sampled value. - arrayed
-
whether the texture is arrayed
- array size
-
the number of homogeneous grids, if the texture is arrayed
A texture’s representation is typically optimized for rendering operations. To achieve this, many details are hidden from the programmer, including data layouts, data types, and internal operations that cannot be expressed directly in the shader language.
As a consequence, a shader does not have direct access to the texel storage within a texture variable. Instead, use texture builtin functions as follows:
-
Within the shader:
-
Declare a module-scope variable in the handle storage class, where the store type is one of the texture types described in later sections.
-
Inside a function, call one of the texture builtin functions, and provide the texture variable as the first parameter.
-
-
When constructing the WebGPU pipeline, the texture variable’s store type and binding must be compatible with the corresponding bind group layout entry.
In this way, the set of supported operations for a texture type is determined by the availability of texture builtin functions accepting that texture type as the first parameter.
3.6.1. Texel formats
In WGSL, certain texture types are parameterized by texel format.
A texel format is characterized by:
- channels
-
Each channel contains a scalar. A texel format has up to four channels:
r
,g
,b
, 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 specfies the conversion from the stored channel bits to the value used in the shader. This is also known as the channel transfer function, or CTF.
Channel format | Number of stored bits | Interpetation 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 ÷ 255 |
8sint | 8 | signed integer v ∈ {-128,...,127} | i32 | max(-1, v ÷ 127) |
16uint | 16 | unsigned integer v ∈ {0,...,65535} | u32 | v |
16sint | 16 | signed integer v ∈ {-32768,...,32767} | i32 | v |
16float | 16 | IEEE 754 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 32-bit floating point value v | f32 | v |
The texel formats listed in the Texel Formats for Storage Textures table correspond to the WebGPU plain color formats which support the WebGPU STORAGE usage. These texel formats are used to parameterize the storage texture types defined in § 3.6.4 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)) |
The following table lists the correspondence between WGSL texel formats and SPIR-V image formats.
Texel format | SPIR-V Image Format | SPIR-V Enabling Capability |
---|---|---|
rgba8unorm | Rgba8 | Shader |
rgba8snorm | Rgba8Snorm | Shader |
rgba8uint | Rgba8ui | Shader |
rgba8sint | Rgba8i | Shader |
rgba16uint | Rgba16ui | Shader |
rgba16sint | Rgba16i | Shader |
rgba16float | Rgba16f | Shader |
r32uint | R32ui | Shader |
r32sint | R32i | Shader |
r32float | R32f | Shader |
rg32uint | Rg32ui | StorageImageExtendedFormats |
rg32sint | Rg32i | StorageImageExtendedFormats |
rg32float | Rg32f | StorageImageExtendedFormats |
rgba32uint | Rgba32ui | Shader |
rgba32sint | Rgba32i | Shader |
rgba32float | Rgba32f | Shader |
3.6.2. Sampled Texture Types
texture_1d<type>
%1 = OpTypeImage %type 1D 0 0 0 1 Unknowntexture_2d<type>
%1 = OpTypeImage %type 2D 0 0 0 1 Unknowntexture_2d_array<type>
%1 = OpTypeImage %type 2D 0 1 0 1 Unknowntexture_3d<type>
%1 = OpTypeImage %type 3D 0 0 0 1 Unknowntexture_cube<type>
%1 = OpTypeImage %type Cube 0 0 0 1 Unknowntexture_cube_array<type>
%1 = OpTypeImage %type Cube 0 1 0 1 Unknown
-
type must be
f32
,i32
oru32
-
The parameterized type for the images is the type after conversion from sampling. E.g. you can have an image with texels with 8bit unorm components, but when you sample them you get a 32-bit float result (or vec-of-f32).
3.6.3. Multisampled Texture Types
texture_multisampled_2d<type>
%1 = OpTypeImage %type 2D 0 0 1 1 Unknown
-
type must be
f32
,i32
oru32
3.6.4. Storage Texture Types
A read-only storage texture supports reading a single texel without the use of a sampler, with automatic conversion of the stored texel value to a usable shader value. A write-only storage texture supports writing a single texel, with automatic conversion of the shader value to a stored texel value. See § 15.8 Texture built-in functions.
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 § 3.6.1 Texel formats.
For a write-only storage texture the inverse of the conversion function is used to convert the shader value to the stored texel.
TODO(dneto): Move description of the conversion to the builtin function that actually does the reading.
texture_storage_1d<texel_format>
// %1 = OpTypeImage sampled_type 1D 0 0 0 2 image_formattexture_storage_2d<texel_format>
// %1 = OpTypeImage sampled_type 2D 0 0 0 2 image_formattexture_storage_2d_array<texel_format>
// %1 = OpTypeImage sampled_type 2D 0 1 0 2 image_formattexture_storage_3d<texel_format>
// %1 = OpTypeImage sampled_type 3D 0 0 0 2 texel_format
In the SPIR-V mapping:
-
The Image Format parameter of the image type declaration is as specified by the SPIR-V texel format correspondence table in § 3.6.1 Texel formats.
-
The Sampled Type parameter of the image type declaration is the SPIR-V scalar type corresponding to the channel format for the texel format.
When mapping to SPIR-V, a read-only storage texture variable must have a NonWritable
decoration and
a write-only storage texture variable must have a NonReadable
decoration.
For example:
var tbuf : [[access(read)]] texture_storage_1d<rgba8unorm>; // Maps to the following SPIR-V: // OpDecorate %tbuf NonWritable // ... // %float = OpTypeFloat 32 // %image_type = OpTypeImage %float 1D 0 0 0 2 Rgba8 // %image_ptr_type = OpTypePointer UniformConstant %image_type // %tbuf = OpVariable %image_ptr_type UniformConstant
var tbuf : [[access(write)]] texture_storage_1d<rgba8unorm>; // Maps to the following SPIR-V: // OpDecorate %tbuf NonReadable // ... // %float = OpTypeFloat 32 // %image_type = OpTypeImage %float 1D 0 0 0 2 Rgba8 // %image_ptr_type = OpTypePointer UniformConstant %image_type // %tbuf = OpVariable %image_ptr_type UniformConstant
3.6.5. Depth Texture Types
texture_depth_2d
%1 = OpTypeImage %f32 2D 1 0 0 1 Unknowntexture_depth_2d_array
%1 = OpTypeImage %f32 2D 1 1 0 1 Unknowntexture_depth_cube
%1 = OpTypeImage %f32 Cube 1 0 0 1 Unknowntexture_depth_cube_array
%1 = OpTypeImage %f32 Cube 1 1 0 1 Unknown
3.6.6. Sampler Type
sampler OpTypeSampler sampler_comparison OpTypeSampler
3.6.7. Texture Types Grammar
TODO: Add texture usage validation rules.texture_sampler_types : sampler_type | depth_texture_type | sampled_texture_type LESS_THAN type_decl GREATER_THAN | multisampled_texture_type LESS_THAN type_decl GREATER_THAN | storage_texture_type LESS_THAN texel_format GREATER_THAN sampler_type : SAMPLER | SAMPLER_COMPARISON sampled_texture_type : TEXTURE_1D | TEXTURE_2D | TEXTURE_2D_ARRAY | TEXTURE_3D | TEXTURE_CUBE | TEXTURE_CUBE_ARRAY multisampled_texture_type : TEXTURE_MULTISAMPLED_2D storage_texture_type : TEXTURE_STORAGE_1D | TEXTURE_STORAGE_2D | TEXTURE_STORAGE_2D_ARRAY | TEXTURE_STORAGE_3D depth_texture_type : TEXTURE_DEPTH_2D | TEXTURE_DEPTH_2D_ARRAY | TEXTURE_DEPTH_CUBE | TEXTURE_DEPTH_CUBE_ARRAY texel_format : R8UNORM R8 -- Capability: StorageImageExtendedFormats | R8SNORM R8Snorm -- Capability: StorageImageExtendedFormats | R8UINT R8ui -- Capability: StorageImageExtendedFormats | R8SINT R8i -- Capability: StorageImageExtendedFormats | R16UINT R16ui -- Capability: StorageImageExtendedFormats | R16SINT R16i -- Capability: StorageImageExtendedFormats | R16FLOAT R16f -- Capability: StorageImageExtendedFormats | RG8UNORM Rg8 -- Capability: StorageImageExtendedFormats | RG8SNORM Rg8Snorm -- Capability: StorageImageExtendedFormats | RG8UINT Rg8ui -- Capability: StorageImageExtendedFormats | RG8SINT Rg8i -- Capability: StorageImageExtendedFormats | R32UINT R32ui | R32SINT R32i | R32FLOAT R32f | RG16UINT Rg16ui -- Capability: StorageImageExtendedFormats | RG16SINT Rg16i -- Capability: StorageImageExtendedFormats | RG16FLOAT Rg16f -- Capability: StorageImageExtendedFormats | RGBA8UNORM Rgba8 | RGBA8UNORM-SRGB ??? | RGBA8SNORM Rgba8Snorm | RGBA8UINT Rgba8ui | RGBA8SINT Rgba8i | BGRA8UNORM Rgba8 ??? | BGRA8UNORM-SRGB ??? | RGB10A2UNORM Rgb10A2 -- Capability: StorageImageExtendedFormats | RG11B10FLOAT R11fG11fB10f -- Capability: StorageImageExtendedFormats | RG32UINT Rg32ui -- Capability: StorageImageExtendedFormats | RG32SINT Rg32i -- Capability: StorageImageExtendedFormats | RG32FLOAT Rg32f -- Capability: StorageImageExtendedFormats | RGBA16UINT Rgba16ui | RGBA16SINT Rgba16i | RGBA16FLOAT Rgba16f | RGBA32UINT Rgba32ui | RGBA32SINT Rgba32i | RGBA32FLOAT Rgba32f
3.7. Type Aliases TODO
type_alias : TYPE IDENT EQUAL type_decl
3.8. Type Declaration Grammar
type_decl : IDENT | BOOL | FLOAT32 | INT32 | UINT32 | VEC2 LESS_THAN type_decl GREATER_THAN | VEC3 LESS_THAN type_decl GREATER_THAN | VEC4 LESS_THAN type_decl GREATER_THAN | POINTER LESS_THAN storage_class COMMA type_decl GREATER_THAN | decoration_list* ARRAY LESS_THAN type_decl COMMA INT_LITERAL GREATER_THAN | decoration_list* ARRAY LESS_THAN type_decl GREATER_THAN | MAT2x2 LESS_THAN type_decl GREATER_THAN | MAT2x3 LESS_THAN type_decl GREATER_THAN | MAT2x4 LESS_THAN type_decl GREATER_THAN | MAT3x2 LESS_THAN type_decl GREATER_THAN | MAT3x3 LESS_THAN type_decl GREATER_THAN | MAT3x4 LESS_THAN type_decl GREATER_THAN | MAT4x2 LESS_THAN type_decl GREATER_THAN | MAT4x3 LESS_THAN type_decl GREATER_THAN | MAT4x4 LESS_THAN type_decl GREATER_THAN | texture_sampler_types
When the type declaration is an identifer, then the expression must be in scope of a declaration of the identifier as a type alias or structure type.
Array decoration keys | Valid values | Note |
---|---|---|
stride
| greater than zero i32 literal |
identifier Allows to specify types created by the type command bool %1 = OpTypeBool f32 %2 = OpTypeFloat 32 i32 %3 = OpTypeInt 32 1 u32 %4 = OpTypeInt 32 0 vec2<f32> %7 = OpTypeVector %float 2 array<f32, 4> %uint_4 = OpConstant %uint 4 %9 = OpTypeArray %float %uint_4 [[stride(32)]] array<f32, 4> OpDecorate %9 ArrayStride 32 %uint_4 = OpConstant %uint 4 %9 = OpTypeArray %float %uint_4 array<f32> %rtarr = OpTypeRuntimeArray %float mat2x3<f32> %vec = OpTypeVector %float 3 %6 = OpTypeMatrix %vec 2
// Storage buffers var<storage> buf1 : [[access(read)]] Buffer; // Can read, cannot write. var<storage> buf2 : [[access(read_write)]] Buffer; // Can both read and write // Uniform buffer. Always read-only, and has more restrictive layout rules. struct ParamsTable {}; var<uniform> params : ParamsTable;
4. Variable and const
TODO: Stub (describe what a constant is): A constant is a name for a value, declared via a const
declaration.
What types are permitted? Storable, plus pointer to store type.
TODO(dneto): A const may not be of type pointer-to-handle. A function parameter may not have type pointer-to-handle. Otherwise we’d have a need to make a pointer-to-handle type expression. But we’ve reserved the handle keyword. When translating from SPIR-V, you must trace through the OpCopyObject (or no-index OpAccessChain) instructions that might be between the pointer-to-array and the pointer-to-struct.
A variable is a named reference to storage 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 storage) and its reference type (the type of the variable itself). If a variable has store type T and storage class S, then its reference type is pointer-to-T-in-S.
-
Determines the variable’s name, storage class, and store type (and hence its reference type).
-
Ensures the execution environment allocates storage for a value of the store type, for the lifetime of the variable.
-
Optionally have an initializer expression, if the variable is in the private or function storage classes. If present, the initializer’s type must match the store type of the variable.
See § 4.1 Module Scope Variables and § 4.3 Function Scope Variables and Constants for rules about where a variable in a particular storage class can be declared, and when the storage class decoration is required, optional, or forbidden.
variable_statement : variable_decl | variable_decl EQUAL short_circuit_or_expression | CONST variable_ident_decl EQUAL short_circuit_or_expression variable_decl : VAR variable_storage_decoration? variable_ident_decl variable_ident_decl : IDENT COLON decoration_list* type_decl variable_storage_decoration : LESS_THAN storage_class GREATER_THAN
Variable declaration decoration keys | Valid values | Note |
---|---|---|
access
| read , write or read_write
|
The access decoration must only appear on a type used as the store type for a variable in the storage storage class. The access decoration must not appear on a type of const declaration nor as the store type for variable with a storage class other than storage. The access decoration is required for variables in the storage storage class.
Two variables with overlapping lifetimes will not have overlapping storage.
When a variable is created, its storage contains an initial value as follows:
-
For variables in the private or function storage classes:
-
The zero value for the store type, if the variable declaration has no initializer.
-
Otherwise, it is the result of evaluating the initializer expression at that point in the program execution.
-
-
For variables in other storage classes, the execution environment provides the initial value.
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 = i + 1 ; break if ( i == 5 ); }
i
will take on values 0, 1, 2, 3, 4, 5, and variable twice
will take on values 0, 2, 4, 6, 8.
Consider the following snippet of WGSL:
Becausex
is a variable, all accesses to it turn into load and store operations.
If this snippet was compiled to SPIR-V, it would be represented as
%temp_1 =OpLoad %float %x %temp_2 =OpLoad %float %x %temp_3 =OpFMul %float %temp_1 %temp_2 %temp_4 =OpLoad %float %x %temp_5 =OpFAdd %float %temp_3 %temp_4 %y =OpFAdd %float %temp_5 %one
4.1. Module Scope Variables
A variable or constant declared outside a function is at module scope. The name is available for use immediately after its declaration statement, until the end of the program.
Variables at module scope are restricted as follows:
-
The variable must not be in the function, in, or out storage classes.
-
A variable in the private, workgroup, uniform, or storage storage classes:
-
Must be declared with an explicit storage class decoration.
-
Must use a store type as described in § 3.4.5 Storage Classes.
-
-
If the store type is a texture type or a sampler type, then the variable declaration must not have a storage class decoration. The storage class will always be handle.
A variable in the uniform storage class is a uniform buffer variable. Its store type must be a host-shareable structure type with block attribute, satisfying the uniform buffer layout rules.
A variable in the storage storage class is a storage buffer variable. Its store type must be a host-shareable structure type with block attribute, satisfying the storage buffer layout rules.
As described in § 8.3.2 Resource interface, uniform buffers, storage buffers, textures, and samplers form the resource interface of a shader. Such variables are declared with group and binding decorations.
var<private> decibels: f32; var<workgroup> worklist: array<i32,10>; [[block]] struct Params { [[offset(0)]] specular: f32; [[offset(4)]] count: i32; }; var<uniform> param: Params; // A uniform buffer [[block]] struct PositionsBuffer { [[offset(0)]] pos: [[stride(8)]] array<vec2<f32>>; }; [[group(0), binding(0)]] var<storage> pbuf: PositionsBuffer; // A storage buffer [[group(0), binding(1)]] var filter_params: sampler; // Textures and samplers are always in "handle" storage.
global_variable_decl : decoration_list* variable_decl | decoration_list* variable_decl EQUAL const_expr decoration_list : ATTR_LEFT (decoration COMMA)* decoration ATTR_RIGHT decoration : IDENT PAREN_LEFT literal_or_ident PAREN_RIGHT | IDENT literal_or_ident : FLOAT_LITERAL | INT_LITERAL | UINT_LITERAL | IDENT
[[location(2)]] OpDecorate %variable Location 2 [[group(4), binding(3)]] OpDecorate %variable DescriptorSet 4 OpDecorate %variable Binding 3
Global variable decoration keys | Valid values | Note |
---|---|---|
binding
| non-negative i32 literal | See § 8.3.2 Resource interface |
builtin
| a builtin variable identifier | See § 14 Built-in variables |
group
| non-negative i32 literal | See § 8.3.2 Resource interface |
location
| non-negative i32 literal | See TBD |
4.2. Module Constants
A module constant declares a name for a value, outside of all function declarations. The name is available for use after the end of the declaration, until the end of the WGSL program.
When the declaration has no attributes, an initializer expression must be present, and the name denotes the value of that expression.
const golden : f32 = 1.61803398875; // The golden ratio const e2 : vec3<i32> = vec3<i32>(0,1,0); // The second unit vector for three dimensions.
When the declaration uses the constant_id
attribute,
the constant is pipeline-overridable. In this case:
-
The type must one of the scalar types.
-
The initializer expression is optional.
-
The attribute’s literal operand is known as the pipeline constant ID, and must be a non-negative integer value representable in 32 bits.
-
Pipeline constant IDs must be unique within the WGSL program: Two module constants must not use the same pipeline constant ID.
-
The application can specify its own value for the name at pipeline-creation time. The pipeline creation API accepts a mapping from the pipeline constant ID to a value of the constant’s type. If the mapping has an entry for the ID, the value in the mapping is used. Otherwise, the initializer expression must be present, and its value is used.
What happens if the application supplies a constant ID that is not in the program? Proposal: pipeline creation fails with an error.
[[constant_id(0)]] const has_point_light : bool = true; // Algorithmic control [[constant_id(1200)]] const specular_param : f32 = 2.3; // Numeric control [[constant_id(1300)]] const gain : f32; // Must be overridden
When a variable or feature is used within control flow that depends on the value of a constant, then that variable or feature is considered to be used by the program. This is true regardless of the value of the constant, whether that value is the one from the constant’s declaration or from a pipeline override.
global_constant_decl : decoration_list* CONST variable_ident_decl global_const_initializer? global_const_initializer : EQUAL const_expr const_expr : type_decl PAREN_LEFT (const_expr COMMA)* const_expr PAREN_RIGHT | const_literal
Global const decoration keys | Valid values | Note |
---|---|---|
constant_id
| non-negative i32 literal |
-1 %a = OpConstant %int -1 2 %b = OpConstant %uint 2 3.2 %c = OpConstant %float 3.2 true %d = OpConstantTrue false %e = OpConstant False vec4<f32>(1.2, 2.3, 3.4, 2.3) %f0 = OpConstant %float 1.2 %f1 = OpConstant %float 2.3 %f2 = OpConstant %float 3.4 %f = OpConstantComposite %v4float %f0 %f1 %f2 %f1
The WebGPU pipeline creation API must specify how API-supplied values are mapped to
shader scalar values. For booleans, I suggest using a 32-bit integer, where only 0 maps to false
.
If WGSL gains non-32-bit numeric scalars, I recommend overridable constants continue being 32-bit
numeric types.
4.3. Function Scope Variables and Constants
A variable or constant declared in a declaration statement in a function body is in function scope. The name is available for use immediately after its declaration statement, and until the end of the brace-delimited list of statements immediately enclosing the declaration.
A variable declared in function scope is always in the function storage class. The variable storage decoration is optional. The variable’s store type must be storable.
fn f () ->void { var < function > count :u32 ; // A variable in function storage class. var delta :i32 ; // Another variable in the function storage class. var sum :f32 = 0.0 ; // A function storage class variable with initializer. const unit :i32 = 1 ; // A constant. Const declarations don’t use a storage class. }
A variable or constant declared in the first clause of a for
statement is available for use in the second
and third clauses and in the body of the for
statement.
4.4. Never-alias assumption TODO
5. Expressions TODO
5.1. Literal Expressions TODO
Precondition | Conclusion | Notes |
---|---|---|
true : bool
| OpConstantTrue %bool | |
false : bool
| OpConstantFalse %bool | |
INT_LITERAL : i32 | OpConstant %int literal | |
UINT_LITERAL : u32 | OpConstant %uint literal | |
FLOAT_LITERAL : f32 | OpConstant %float literal |
5.2. Type Constructor Expressions TODO
Precondition | Conclusion | Notes |
---|---|---|
e : bool | bool(e) : bool
| Identity. In the SPIR-V translation, the ID of this expression reuses the ID of the operand. |
e : i32 | i32(e) : i32
| Identity. In the SPIR-V translation, the ID of this expression reuses the ID of the operand. |
e : u32 | u32(e) : u32
| Identity. In the SPIR-V translation, the ID of this expression reuses the ID of the operand. |
e : f32 | f32(e) : f32
| Identity. In the SPIR-V translation, the ID of this expression reuses the ID of the operand. |
Precondition | Conclusion | Notes |
---|---|---|
e1 : T e2 : T | vec2<T>(e1,e2) : vec2<T>
| OpCompositeConstruct |
e : vec2<T> | vec2<T>(e) : vec2<T>
| Identity. The result is e. |
e1 : T e2 : T e3 : T | vec3<T>(e1,e2,e3) : vec3<T>
| OpCompositeConstruct |
e1 : T e2 : vec2<T> | vec3<T>(e1,e2) : vec3<T>vec3<T>(e2,e1) : vec3<T>
| OpCompositeConstruct |
e : vec3<T> | vec3<T>(e) : vec3<T>
| Identity. The result is e. |
e1 : T e2 : T e3 : T e4 : T | vec4<T>(e1,e2,e3,e4) : vec4<T>
| OpCompositeConstruct |
e1 : T e2 : T e3 : vec2<T> | vec4<T>(e1,e2,e3) : vec4<T>vec4<T>(e1,e3,e2) : vec4<T>vec4<T>(e3,e1,e2) : vec4<T>
| OpCompositeConstruct |
e1 : vec2<T> e2 : vec2<T> | vec4<T>(e1,e2) : vec4<T>
| OpCompositeConstruct |
e1 : T e2 : vec3<T> | vec4<T>(e1,e2) : vec4<T>vec4<T>(e2,e1) : vec4<T>
| OpCompositeConstruct |
e : vec4<T> | vec4<T>(e) : vec4<T>
| Identity. The result is e. |
Precondition | Conclusion | Notes |
---|---|---|
e1 : vec2
e2 : vec2 e3 : vec2 e4 : vec2 |
mat2x2<f32>(e1,e2) : mat2x2
mat3x2<f32>(e1,e2,e3) : mat3x2
mat4x2<f32>(e1,e2,e3,e4) : mat4x2
| Column by column construction. OpCompositeConstruct |
e1 : vec3
e2 : vec3 e3 : vec3 e4 : vec3 |
mat2x3<f32>(e1,e2) : mat2x3
mat3x3<f32>(e1,e2,e3) : mat3x3
mat4x3<f32>(e1,e2,e3,e4) : mat4x3
| Column by column construction. OpCompositeConstruct |
e1 : vec4
e2 : vec4 e3 : vec4 e4 : vec4 |
mat2x4<f32>(e1,e2) : mat2x4
mat3x4<f32>(e1,e2,e3) : mat3x4
mat4x4<f32>(e1,e2,e3,e4) : mat4x4
| Column by column construction. OpCompositeConstruct |
Precondition | Conclusion | Notes |
---|---|---|
e1 : T ... eN : T | array< T,N>(e1,...,eN) : array<T, N>
| Construction of an array from elements |
Precondition | Conclusion | Notes |
---|---|---|
e1 : T1 ... eN : TN T1 is storable ... TN is storable S is a structure type with members having types T1 ... TN. The expression is in the scope of declaration of S. | S(e1,...,eN) : S
| Construction of a structure from members |
5.3. Zero Value Expressions
Each storable type T has a unique zero value, written in WGSL as the type followed by an empty pair of parentheses: T ()
.
We should exclude being able to write the zero value for an runtime-sized array. https://github.com/gpuweb/gpuweb/issues/981
The zero values are as follows:
-
bool()
isfalse
-
i32()
is 0 -
u32()
is 0 -
f32()
is 0.0 -
The zero value for an N-element vector of type T is the N-element vector of the zero value for T.
-
The zero value for an N-column M-row matrix of
f32
is the matrix of those dimensions filled with 0.0 entries. -
The zero value for an N-element array with storable element type E is an array of N elements of the zero value for E.
-
The zero value for a storable structure type S is the structure value S with zero-valued members.
Precondition | Conclusion | Notes |
---|---|---|
bool() : bool
| false Zero value (OpConstantNull for bool) | |
i32() : i32
| 0 Zero value (OpConstantNull for i32) | |
u32() : u32
| 0u Zero value (OpConstantNull for u32) | |
f32() : f32
| 0.0 Zero value (OpConstantNull for f32) |
Precondition | Conclusion | Notes |
---|---|---|
vec2<T>() : vec2<T>
| Zero value (OpConstantNull) | |
vec3<T>() : vec3<T>
| Zero value (OpConstantNull) | |
vec4<T>() : vec4<T>
| Zero value (OpConstantNull) | |
vec2 < f32 > () // The zero-valued vector of two f32 elements. vec2 < f32 > ( 0.0 , 0.0 ) // The same value, written explicitly. vec3 < i32 > () // The zero-valued vector of four i32 elements. vec3 < i32 > ( 0 , 0 , 0 ) // The same value, written explicitly.
Precondition | Conclusion | Notes |
---|---|---|
mat2x2<f32>() : mat2x2
mat3x2<f32>() : mat3x2
mat4x2<f32>() : mat4x2
| Zero value (OpConstantNull) | |
mat2x3<f32>() : mat2x3
mat3x3<f32>() : mat3x3
mat4x3<f32>() : mat4x3
| Zero value (OpConstantNull) | |
mat2x4<f32>() : mat2x4
mat3x4<f32>() : mat3x4
mat4x4<f32>() : mat4x4
| Zero value (OpConstantNull) |
Precondition | Conclusion | Notes |
---|---|---|
T is storable | array< T,N>() : array<T, N>
| Zero-valued array (OpConstantNull) |
array < bool , 2 > () // The zero-valued array of two booleans. array < bool , 2 > ( false , false ) // The same value, written explicitly.
Precondition | Conclusion | Notes |
---|---|---|
S is a storable structure type.The expression is in the scope of declaration of S. | S() : S
| Zero-valued structure: a structure of type S where each member is the zero value for its member type. (OpConstantNull) |
struct Student { grade :i32 ; GPA :f32 ; attendance :array < bool , 4 > ; }; fn func () ->void { var s :Student ; // The zero value for Student s = Student (); // The same value, written explicitly. s = Student ( 0 , 0.0 , array < bool , 4 > ( false , false , false , false )); // The same value, written with zero-valued members. s = Student ( i32 (), f32 (), array < bool , 4 > ()); }
5.4. Conversion Expressions
Precondition | Conclusion | Notes |
---|---|---|
e : u32 | bool( e) : bool
| Coercion to boolean. The result is false if e is 0, and true otherwise. (Use OpINotEqual to compare e against 0.) |
e : i32 | bool( e) : bool
| Coercion to boolean. The result is false if e is 0, and true otherwise. (Use OpINotEqual to compare e against 0.) |
e : f32 | bool( e) : bool
| Coercion to boolean. The result is false if e is 0.0 or -0.0, and true otherwise. In particular NaN and infinity values map to true. (Use OpFUnordNotEqual to compare e against 0.0 .)
|
e : u32 | i32( e) : i32
| Reinterpretation of bits. The result is the unique value in i32 that is equal to (e mod 232). (OpBitcast) |
e : f32 | i32( e) : i32
| Value conversion, including invalid cases. (OpConvertFToS) |
e : i32 | u32( e) : u32
| Reinterpretation of bits. The result is the unique value in u32 that is equal to (e mod 232). (OpBitcast) |
e : f32 | u32( e) : u32
| Value conversion, including invalid cases. (OpConvertFToU) |
e : i32 | f32( e) : f32
| Value conversion, including invalid cases. (OpConvertSToF) |
e : u32 | f32( e) : f32
| Value conversion, including invalid cases. (OpConvertUToF) |
Details of conversion to and from floating point are explained in § 10.5.1 Floating point conversion.
Precondition | Conclusion | Notes |
---|---|---|
e : vecN<u32> | vec N<bool >( e) : vecN<bool>
| Component-wise coercion of a unsigned integer vector to a boolean vector. Component i of the result is bool( e[ i]) (OpINotEqual to compare e against a zero vector.) |
e : vecN<i32> | vec N<bool >( e) : vecN<bool>
| Component-wise coercion of a signed integer vector to a boolean vector. Component i of the result is bool( e[ i]) (OpINotEqual to compare e against a zero vector.) |
e : vecN<f32> | vec N<bool >( e) : vecN<bool>
| Component-wise coercion of a floating point vector to a boolean vector. Component i of the result is bool( e[ i]) (OpFUnordNotEqual to compare e against a zero vector.) |
e : vecN<u32> | vec N<i32 >( e) : vecN<i32>
| Component-wise reinterpretation of bits. Component i of the result is i32( e[ i]) (OpBitcast) |
e : vecN<f32> | vec N<i32 >( e) : vecN<i32>
| Component-wise value conversion to signed integer, including invalid cases. Component i of the result is i32( e[ i]) (OpConvertFToS) |
e : vecN<i32> | vec N<u32 >( e) : vecN<u32>
| Component-wise reinterpretation of bits. Component i of the result is u32( e[ i]) (OpBitcast) |
e : vecN<f32> | vec N<u32 >( e) : vecN<u32>
| Component-wise value conversion to unsigned integer, including invalid cases. Component i of the result is u32( e[ i]) (OpConvertFToU) |
e : vecN<i32> | vec N<f32 >( e) : vecN<f32>
| Component-wise value conversion to floating point, including invalid cases. Component i of the result is f32( e[ i]) (OpConvertSToF) |
e : vecN<u32> | vec N<f32 >( e) : vecN<f32>
| Component-wise value conversion to floating point, including invalid cases. Component i of the result is f32( e[ i]) (ConvertUToF) |
5.5. Reinterpretation of Representation Expressions
A bitcast
expression is used to reinterpet the bit representation of a
value in one type as a value in another type.
Precondition | Conclusion | Notes |
---|---|---|
e : T, T is one of i32, u32, f32 | bitcast<T>(e) : T | Identity transform. The result is e. In the SPIR-V translation, the ID of this expression reuses the ID of the operand. |
e : T, T is one of u32, f32 | bitcast<i32>(e) : i32 | Reinterpretation of bits as a signed integer. The result is the reinterpretation of the 32 bits in the representation of e as a i32 value. (OpBitcast) |
e : T, T is one of i32, f32 | bitcast<u32>(e) : u32 | Reinterpretation of bits as an unsigned integer. The result is the reinterpretation of the 32 bits in the representation of e as a u32 value. (OpBitcast) |
e : T, T is one of i32, u32 | bitcast<f32>(e) : f32 | Reinterpretation of bits as a floating point value. The result is the reinterpretation of the 32 bits in the representation of e as a f32 value. (OpBitcast) |
Precondition | Conclusion | Notes |
---|---|---|
e : vec<N>T>, T is one of i32, u32, f32 | bitcast<vecN<T>>(e) : T | Identity transform. The result is e. In the SPIR-V translation, the ID of this expression reuses the ID of the operand. |
e : vec<N>T>, T is one of u32, f32 | bitcast<vecN<i32>>(e) : vecN<i32> | Component-wise reinterpretation of bits. Component i of the result is bitcast<i32>( e[ i]) (OpBitcast) |
e : vec<N>T>, T is one of i32, f32 | bitcast<vecN<u32>>(e) : vecN<u32> | Component-wise reinterpretation of bits. Component i of the result is bitcast<u32>( e[ i]) (OpBitcast) |
e : vec<N>T>, T is one of i32, u32 | bitcast<vecN<f32>>(e) : vecN<f32> | Component-wise Reinterpretation of bits. Component i of the result is bitcast<f32>( e[ i]) (OpBitcast) |
5.6. Composite Value Expressions TODO
5.6.1. Vector Access Expression
Accessing members of a vector can be done either using array subscripting (e.g. a[2]
) or using a sequence of convenience names, each mapping to an element of the source vector.
- The colour set of convenience names:
r
,g
,b
,a
for vector elements 0, 1, 2, and 3 respectively. - The dimensional set of convenience names:
x
,y
,z
,w
for vector elements 0, 1, 2, and 3, respectively.
The convenience names are accessed using the .
notation. (e.g. color.bgra
).
NOTE: the convenience letterings can not be mixed. (i.e. you can not use rybw
).
Using a convenience letter, or array subscript, which accesses an element past the end of the vector is an error.
The convenience letterings can be applied in any order, including duplicating letters as needed. You can provide 1 to 4 letters when extracting components from a vector. Providing more then 4 letters is an error.
The result type depends on the number of letters provided. Assuming a vec4<f32>
Accessor | Result type |
---|---|
r | f32
|
rg | vec2<f32>
|
rgb | vec3<f32>
|
rgba | vec4<f32>
|
var a :vec3 < f32 > = vec3 < f32 > ( 1. , 2. , 3. ); var b :f32 = a . y ; // b = 2.0 var c :vec2 < f32 > = a . bb ; // c = (3.0, 3.0) var d :vec3 < f32 > = a . zyx ; // d = (3.0, 2.0, 1.0) var e :f32 = a [ 1 ]; // e = 2.0
5.6.1.1. Vector single component selection
Precondition | Conclusion | Description |
---|---|---|
e : vecN<T> | e.x : Te .r : T
| Select the first component of e (OpCompositeExtract with selection index 0) |
e : vecN<T> | e.y : Te .g : T
| Select the second component of e (OpCompositeExtract with selection index 1) |
e : vecN<T> N is 3 or 4 | e.z : Te .b : T
| Select the third component of e (OpCompositeExtract with selection index 2) |
e : vec4<T> | e.w : Te .a : T
| Select the fourth component of e (OpCompositeExtract with selection index 3) |
e : vecN<T> i : Int | 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], then an index in the range [0, N-1] is used instead. (OpVectorExtractDynamic) |
Which index is used when it’s out of bounds?
5.6.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-element 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.(OpVectorShuffle) |
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-element 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.(OpVectorShuffle) |
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-element 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.(OpVectorShuffle) |
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-element 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.(OpVectorShuffle) |
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-element 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.(OpVectorShuffle) |
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-element 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.(OpVectorShuffle) |
5.6.2. Matrix Access Expression TODO
Precondition | Conclusion | Description |
---|---|---|
e : matNxM<T> i : Int | e[i] : vecM<T> | The result is the i’th column vector of e. The first column vector is at index i=0. If i is outside the range [0,N-1], then an index in the range [0, N-1] is used instead. (OpCompositeExtract) |
Which index is used when it’s out of bounds?
5.6.3. Array Access Expression TODO
Precondition | Conclusion | Description |
---|---|---|
e : array<T,N> i : Int | e[i] : T | The result is the value of the i’th element of the array value e. The first element is at index i=0. If i is outside the range [0,N-1], then an index in the range [0, N-1] is used instead. (OpCompositeExtract) |
Which index is used when it’s out of bounds?
5.6.4. Structure Access Expression TODO
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. (OpCompositeExtract, using the member index) |
5.7. Logical Expressions TODO
Precondition | Conclusion | Notes |
---|---|---|
e : bool | ! e : bool
| Logical negation. Yields true when e is false, and false when e is true. (OpLogicalNot) |
e : vecN<bool> | ! e : vecN<bool>
| Component-wise logical negation. Component i of the result is !( e[ i]) .(OpLogicalNot) |
Precondition | Conclusion | Notes |
---|---|---|
e1 : bool e2 : bool | e1 || e2 : bool
| Short-circuiting "or". Yields true if either e1 or e2 are true; evaluates e2 only if e1 is false.
|
e1 : bool e2 : bool | e1 && e2 : bool
| Short-circuiting "and". Yields true if both e1 and e2 are true; evaluates e2 only if e1 is true.
|
e1 : bool e2 : bool | e1 | e2 : bool
| Logical "or". Evaluates both e1 and e2 ; yields true if either are true .
|
e1 : bool e2 : bool | e1 & e2 : bool
| Logical "and". Evaluates both e1 and e2 ; yields true if both are true .
|
e1 : T e2 : T T is BoolVec | e1 | e2 : T
| Component-wise logical "or" |
e1 : T e2 : T T is BoolVec | e1 & e2 : T
| Component-wise logical "and" |
5.8. Arithmetic Expressions TODO
Precondition | Conclusion | Notes |
---|---|---|
e : T, T is SignedIntegral | -e : T
| Signed integer negation. OpSNegate |
e : T, T is Floating | -e : T
| Floating point negation. OpFNegate |
Precondition | Conclusion | Notes |
---|---|---|
e1 : u32 e2 : u32 | e1 + e2 : u32
| Integer addition, modulo 232 (OpIAdd) |
e1 : i32 e2 : i32 | e1 + e2 : i32
| Integer addition, modulo 232 (OpIAdd) |
e1 : f32 e2 : f32 | e1 + e2 : f32
| Floating point addition (OpFAdd) |
e1 : u32 e2 : u32 | e1 - e2 : u32
| Integer subtraction, modulo 232 (OpISub) |
e1 : i32 e2 : i32 | e1 - e2 : i32
| Integer subtraction, modulo 232 (OpISub) |
e1 : f32 e2 : f32 | e1 - e2 : f32
| Floating point subtraction (OpFSub) |
e1 : u32 e2 : u32 | e1 * e2 : u32
| Integer multiplication, modulo 232 (OpIMul) |
e1 : i32 e2 : i32 | e1 * e2 : i32
| Integer multiplication, modulo 232 (OpIMul) |
e1 : f32 e2 : f32 | e1 * e2 : f32
| Floating point multiplication (OpFMul) |
e1 : u32 e2 : u32 | e1 / e2 : u32
| Unsigned integer division (OpUDiv) |
e1 : i32 e2 : i32 | e1 / e2 : i32
| Signed integer division (OpSDiv) |
e1 : f32 e2 : f32 | e1 / e2 : f32
| Floating point division (OpFDiv) |
e1 : u32 e2 : u32 | e1 % e2 : u32
| Unsigned integer modulus (OpUMod) |
e1 : i32 e2 : i32 | e1 % e2 : i32
| Signed integer remainder, where sign of non-zero result matches sign of e2 (OpSMod) |
e1 : f32 e2 : f32 | e1 % e2 : f32
| Floating point modulus, where sign of non-zero result matches sign of e2 (OpFMod) |
Precondition | Conclusion | Notes |
---|---|---|
e1 : T e2 : T T is IntVec | e1 + e2 : T
| Component-wise integer addition (OpIAdd) |
e1 : T e2 : T T is FloatVec | e1 + e2 : T
| Component-wise floating point addition (OpIAdd) |
e1 : T e2 : T T is IntVec | e1 - e2 : T
| Component-wise integer subtraction (OpISub) |
e1 : T e2 : T T is FloatVec | e1 - e2 : T
| Component-wise floating point subtraction (OpISub) |
e1 : T e2 : T T is IntVec | e1 * e2 : T
| Component-wise integer multiplication (OpIMul) |
e1 : T e2 : T T is FloatVec | e1 * e2 : T
| Component-wise floating point multiplication (OpIMul) |
e1 : T e2 : T T is IntVec with unsigned component | e1 / e2 : T
| Component-wise unsigned integer division (OpUDiv) |
e1 : T e2 : T T is IntVec with signed component | e1 / e2 : T
| Component-wise signed integer division (OpSDiv) |
e1 : T e2 : T T is FloatVec | e1 / e2 : T
| Component-wise floating point division (OpFDiv) |
e1 : T e2 : T T is IntVec with unsigned component | e1 % e2 : T
| Component-wise unsigned integer modulus (OpUMod) |
e1 : T e2 : T T is IntVec with signed component | e1 % e2 : T
| Component-wise signed integer remainder (OpSMod) |
e1 : T e2 : T T is FloatVec | e1 % e2 : T
| Component-wise floating point modulus (OpFMod) |
Precondition | Conclusion | Notes |
---|---|---|
e1 : f32 e2 : T T is FloatVec | e1 * e2 : Te2 * e1 : T
| Multiplication of a vector and a scalar (OpVectorTimesScalar) |
e1 : f32 e2 : T T is matNxM<f32> | e1 * e2 : Te2 * e1 : T
| Multiplication of a matrix and a scalar (OpMatrixTimesScalar) |
e1 : vecM<f32> e2 : matNxM<f32> | e1 * e2 : vecN<f32> | Vector times matrix (OpVectorTimesMatrix) |
e1 : matNxM<f32> e2 : vecN<f32> | e1 * e2 : vecM<f32> | Matrix times vector (OpMatrixTimesVector) |
e1 : matKxN<f32> e2 : matMxK<f32> | e1 * e2 : matMxN<f32> | Matrix times matrix (OpMatrixTimesMatrix) |
5.9. Comparison Expressions TODO
Precondition | Conclusion | Notes |
---|---|---|
e1 : bool e2 : bool | e1 == e2 : bool
| Equality (OpLogicalEqual) |
e1 : bool e2 : bool | e1 != e2 : bool
| Inequality (OpLogicalNotEqual) |
e1 : i32 e2 : i32 | e1 == e2 : bool
| Equality (OpIEqual) |
e1 : i32 e2 : i32 | e1 != e2 : bool
| Inequality (OpINotEqual) |
e1 : i32 e2 : i32 | e1 < e2 : bool
| Less than (OpSLessThan) |
e1 : i32 e2 : i32 | e1 <= e2 : bool
| Less than or equal (OpSLessThanEqual) |
e1 : i32 e2 : i32 | e1 >= e2 : bool
| Greater than or equal (OpSGreaterThanEqual) |
e1 : i32 e2 : i32 | e1 > e2 : bool
| Greater than or equal (OpSGreaterThan) |
e1 : u32 e2 : u32 | e1 == e2 : bool
| Equality (OpIEqual) |
e1 : u32 e2 : u32 | e1 != e2 : bool
| Inequality (OpINotEqual) |
e1 : u32 e2 : u32 | e1 < e2 : bool
| Less than (OpULessThan) |
e1 : u32 e2 : u32 | e1 <= e2 : bool
| Less than or equal (OpULessThanEqual) |
e1 : u32 e2 : u32 | e1 >= e2 : bool
| Greater than or equal (OpUGreaterThanEqual) |
e1 : u32 e2 : u32 | e1 > e2 : bool
| Greater than or equal (OpUGreaterThan) |
e1 : f32 e2 : f32 | e1 == e2 : bool
| Equality (OpFOrdEqual) |
e1 : f32 e2 : f32 | e1 != e2 : bool
| Equality (OpFOrdNotEqual) |
e1 : f32 e2 : f32 | e1 < e2 : bool
| Less than (OpFOrdLessThan) |
e1 : f32 e2 : f32 | e1 <= e2 : bool
| Less than or equal (OpFOrdLessThanEqual) |
e1 : f32 e2 : f32 | e1 >= e2 : bool
| Greater than or equal (OpFOrdGreaterThanEqual) |
e1 : f32 e2 : f32 | e1 > e2 : bool
| Greater than or equal (OpFOrdGreaterThan) |
Precondition | Conclusion | Notes |
---|---|---|
e1 : T e2 : T T is vecN<bool> | e1 == e2 : vecN<bool>
| Component-wise equality Component i of the result is ( e1[ i] == e2[ i]) (OpLogicalEqual) |
e1 : T e2 : T T is vecN<bool> | e1 != e2 : vecN<bool>
| Component-wise inequality Component i of the result is ( e1[ i] != e2[ i]) (OpLogicalNotEqual) |
e1 : T e2 : T T is vecN<i32> | e1 == e2 : vecN<bool>
| Component-wise equality (OpIEqual) |
e1 : T e2 : T T is vecN<i32> | e1 != e2 : vecN<bool>
| Component-wise inequality (OpINotEqual) |
e1 : T e2 : T T is vecN<i32> | e1 < e2 : vecN<bool>
| Component-wise less than (OpSLessThan) |
e1 : T e2 : T T is vecN<i32> | e1 <= e2 : vecN<bool>
| Component-wise less than or equal (OpSLessThanEqual) |
e1 : T e2 : T T is vecN<i32> | e1 >= e2 : vecN<bool>
| Component-wise greater than or equal (OpSGreaterThanEqual) |
e1 : T e2 : T T is vecN<i32> | e1 > e2 : vecN<bool>
| Component-wise greater than or equal (OpSGreaterThan) |
e1 : T e2 : T T is vecN<u32> | e1 == e2 : vecN<bool>
| Component-wise equality (OpIEqual) |
e1 : T e2 : T T is vecN<u32> | e1 != e2 : vecN<bool>
| Component-wise inequality (OpINotEqual) |
e1 : T e2 : T T is vecN<u32> | e1 < e2 : vecN<bool>
| Component-wise less than (OpULessThan) |
e1 : T e2 : T T is vecN<u32> | e1 <= e2 : vecN<bool>
| Component-wise less than or equal (OpULessThanEqual) |
e1 : T e2 : T T is vecN<u32> | e1 >= e2 : vecN<bool>
| Component-wise greater than or equal (OpUGreaterThanEqual) |
e1 : T e2 : T T is vecN<u32> | e1 > e2 : vecN<bool>
| Component-wise greater than or equal (OpUGreaterThan) T is vecN<u32> |
e1 : T e2 : T T is vecN<f32> | e1 == e2 : vecN<bool>
| Component-wise equality (OpFOrdEqual) |
e1 : T e2 : T T is vecN<f32> | e1 != e2 : vecN<bool>
| Component-wise inequality (OpFOrdNotEqual) |
e1 : T e2 : T T is vecN<f32> | e1 < e2 : vecN<bool>
| Component-wise less than (OpFOrdLessThan) |
e1 : T e2 : T T is vecN<f32> | e1 <= e2 : vecN<bool>
| Component-wise less than or equal (OpFOrdLessThanEqual) |
e1 : T e2 : T T is vecN<f32> | e1 >= e2 : vecN<bool>
| Component-wise greater than or equal (OpFOrdGreaterThanEqual) |
e1 : T e2 : T T is vecN<f32> | e1 > e2 : vecN<bool>
| Component-wise greater than or equal (OpFOrdGreaterThan) |
5.10. Bit Expressions TODO
Precondition | Conclusion | Notes |
---|---|---|
e : u32 | ~ e : u32
| Bitwise complement on unsigned integers. Result is the mathematical value (232 - 1 - e). OpNot |
e : vecN<u32> | ~ e : vecN<u32>
| Component-wise unsigned complement. Component i of the result is ~( e[ i]) . OpNot |
e : i32 | ~ e : i32
| Bitwise complement on signed integers. Result is i32(~u32(e)). OpNot |
e : vecN<i32> | ~ e : vecN<i32>
| Component-wise signed complement. Component i of the result is ~( e[ i]) . OpNot |
Precondition | Conclusion | Notes |
---|---|---|
e1 : T e2 : T T is Integral | e1 | e2 : T
| Bitwise-or |
e1 : T e2 : T T is Integral | e1 & e2 : T
| Bitwise-and |
e1 : T e2 : T T is Integral | e1 ^ e2 : T
| Bitwise-exclusive-or |
Precondition | Conclusion | Notes |
---|---|---|
e1 : T e2 : u32 T is Int | e1 << e2 : T
| Shift left: Shift e1 left, inserting zero bits at the least significant positions, and discarding the most significant bits. The number of bits to shift is the value of e2 modulo the bit width of e1. (OpShiftLeftLogical) |
e1 : vecN<T> e2 : vecN<u32> T is Int | e1 << e2 : vecN<T>
| Component-wise shift left: Component i of the result is ( e1[ i] << e2[ i]) (OpShiftLeftLogical) |
e1 : u32 e2 : u32 | e1 >> e2 : u32
| Logical shift right: Shift e1 right, inserting zero bits at the most significant positions, and discarding the least significant bits. The number of bits to shift is the value of e2 modulo the bit width of e1. (OpShiftRightLogical) |
e1 : vecN<u32> e2 : u32 | e1 >> e2 : vecN<u32>
| Component-wise logical shift right: Component i of the result is ( e1[ i] >> e2[ i]) (OpShiftRightLogical)
|
e1 : i32 e2 : u32 | e1 >> e2 : i32
| Arithmetic shift right: Shift e1 right, copying the sign bit of e1 into the most significant positions, and discarding the least significant bits. The number of bits to shift is the value of e2 modulo the bit width of e1. (OpShiftRightArithmetic) |
e1 : vecN<i32> e2 : vecN<u32> | e1 >> e2 : vecN<i32>
| Component-wise arithmetic shift right: Component i of the result is ( e1[ i] >> e2[ i]) (OpShiftRightArithmetic)
|
5.11. Function Call Expression TODO
TODO: Stub. Call to function returning non-void, is an expression.
5.12. Variable or const reference TODO
5.13. Pointer Expressions TODO
TODO: Stub: how to write each of the abstract pointer operations
5.14. Expression Grammar Summary
primary_expression : IDENT argument_expression_list? | type_decl argument_expression_list | const_literal | paren_rhs_statement | BITCAST LESS_THAN type_decl GREATER_THAN paren_rhs_statement OpBitcast argument_expression_list : PAREN_LEFT ((short_circuit_or_expression COMMA)* short_circuit_or_expression)? PAREN_RIGHT postfix_expression : | BRACKET_LEFT short_circuit_or_expression BRACKET_RIGHT postfix_expression | PERIOD IDENT postfix_expression unary_expression : singular_expression | MINUS unary_expression OpSNegate OpFNegate | BANG unary_expression OpLogicalNot | TILDE unary_expression OpNot singular_expression : primary_expression postfix_expression multiplicative_expression : unary_expression | multiplicative_expression STAR unary_expression OpVectorTimesScalar OpMatrixTimesScalar OpVectorTimesMatrix OpMatrixTimesVector OpMatrixTimesMatrix OpIMul OpFMul | multiplicative_expression FORWARD_SLASH unary_expression OpUDiv OpSDiv OpFDiv | multiplicative_expression MODULO unary_expression OpUMOd OpSMod OpFMod additive_expression : multiplicative_expression | additive_expression PLUS multiplicative_expression OpIAdd OpFAdd | additive_expression MINUS multiplicative_expression OpFSub OpISub shift_expression : additive_expression | shift_expression SHIFT_LEFT additive_expression OpShiftLeftLogical | shift_expression SHIFT_RIGHT additive_expression OpShiftRightLogical or OpShiftRightArithmetic relational_expression : shift_expression | relational_expression LESS_THAN shift_expression OpULessThan OpFOrdLessThan | relational_expression GREATER_THAN shift_expression OpUGreaterThan OpFOrdGreaterThan | relational_expression LESS_THAN_EQUAL shift_expression OpULessThanEqual OpFOrdLessThanEqual | relational_expression GREATER_THAN_EQUAL shift_expression OpUGreaterThanEqual OpFOrdGreaterThanEqual equality_expression : relational_expression | relational_expression EQUAL_EQUAL relational_expression OpIEqual OpFOrdEqual | relational_expression NOT_EQUAL relational_expression OpINotEqual OpFOrdNotEqual and_expression : equality_expression | and_expression AND equality_expression exclusive_or_expression : and_expression | exclusive_or_expression XOR and_expression inclusive_or_expression : exclusive_or_expression | inclusive_or_expression OR exclusive_or_expression short_circuit_and_expression : inclusive_or_expression | short_circuit_and_expression AND_AND inclusive_or_expression short_circuit_or_expression : short_circuit_and_expression | short_circuit_or_expression OR_OR short_circuit_and_expression
6. Statements TODO
6.1. Assignment TODO
assignment_statement : singular_expression EQUAL short_circuit_or_expression If singular_expression is a variable, this maps to OpStore to the variable. Otherwise, singular expression is a pointer expression in an Assigning (L-value) context which maps to OpAccessChain followed by OpStore
6.1.1. Writing to a variable TODO
6.1.2. Writing to a part of a composite variable TODO
6.2. Control flow TODO
6.2.1. Sequence TODO
6.2.2. If/elseif/else Statement TODO
if_statement : IF paren_rhs_statement body_statement elseif_statement? else_statement? elseif_statement : ELSE_IF paren_rhs_statement body_statement elseif_statement? else_statement : ELSE body_statement
6.2.3. Switch Statement
switch_statement : SWITCH paren_rhs_statement BRACE_LEFT switch_body+ BRACE_RIGHT switch_body : CASE case_selectors COLON BRACE_LEFT case_body BRACE_RIGHT | DEFAULT COLON BRACE_LEFT case_body BRACE_RIGHT case_selectors : const_literal (COMMA const_literal)* case_body : | statement case_body | FALLTHROUGH SEMICOLON
A switch statement transfers control to one of a set of case clauses, or to the default
clause,
depending on the evaluation of a selector expression.
The selector expression must be of a scalar integer type.
If the selector value equals a value in a case selector list, then control is transferred to
the body of that case clause.
If the selector value does not equal any of the case selector values, then control is
transferred to the default
clause.
Each switch statement must have exactly one default clause.
The case selector values must have the same type as the selector expression.
A literal value must not appear more than once in the case selectors for a switch statement.
Note: The value of the literal is what matters, not the spelling.
For example 0
, 00
, and 0x0000
all denote the zero value.
When control reaches the end of a case body, control normally transfers to the first statement
after the switch statement.
Alternately, executing a fallthrough
statement transfers control to the body of the next case clause or
default clause, whichever appears next in the switch body.
A fallthrough
statement must not appear as the last statement in the last clause of a switch.
6.2.4. Loop Statement
loop_statement : LOOP BRACE_LEFT statements continuing_statement? BRACE_RIGHT
The loop construct causes a block of statements, the loop body, to execute repeatedly.
This repetition can be interrupted by a § 6.2.6 Break, return
, or discard
.
Optionally, the last statement in the loop body may be a § 6.2.8 Continuing Statement.
Note: The loop statement is one of the biggest differences from other shader languages.
This design directly expresses loop idioms commonly found in compiled code. In particular, placing the loop update statements at the end of the loop body allows them to naturally use values defined in the loop body.
const a : i32 = 2; var i : i32 = 0; // <1> loop { if (i >= 4) { break; } a = a * 2; i = i + 1; }
- <1> The initialization is listed before the loop.
int a = 2; const 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; } const step : i32 = 1; i = i + 1; if (i % 2 == 0) { continue; } a = a * 2; }
var a : i32 = 2; var i : i32 = 0; loop { if (i >= 4) { break; } const 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
6.2.5. For Statement
for_statement : FOR PAREN_LEFT for_header PAREN_RIGHT body_statement for_header : (variable_statement | assignment_statement | func_call_statement)? SEMICOLON short_circuit_or_expression? SEMICOLON (assignment_statement | func_call_statement)?
The for(initializer; condition; continuing) { body }
statement is syntactic sugar on top of a § 6.2.4 Loop Statement with the same body
. Additionally:
-
If
initializer
is non-empty, it is executed inside an additional scope before the first iteration. -
If
condition
is non-empty, it is checked at the beginning of the loop body and if unsatisfied then a § 6.2.6 Break is executed. -
If
continuing
is non-empty, it becomes a § 6.2.8 Continuing Statement at the end of the loop body.
for(var i : i32 = 0; i < 4; i = i + 1) { if (a == 0) { continue; } a = a + 2; }
Converts to:
{ // Introduce new scope for loop variable i var i : i32 = 0; var a : i32 = 0; loop { if (!(i < 4)) { break; } if (a == 0) { continue; } a = a + 2; continuing { i = i + 1; } } }
6.2.6. Break
break_statement : BREAK
Use a break
statement to transfer control to the first statement
after the body of the nearest-enclosing § 6.2.4 Loop Statement or § 6.2.3 Switch Statement.
When a break
statement is placed such that it would exit from a loop’s § 6.2.8 Continuing Statement,
then:
-
The
break
statement must appear as either:-
The only statement in the true-branch clause of an
if
that has:-
no
else
clause or an emptyelse
clause -
no
elseif
clauses
-
-
The only statement in the
else
clause of anif
that has an empty true-branch clause and noelseif
clauses.
-
-
That
if
statement must appear last in thecontinuing
clause.
var a : i32 = 2; var i : i32 = 0; loop { const step : i32 = 1; if (i % 2 == 0) { continue; } a = a * 2; continuing { i = i + step; if (i >= 4) { break; } } }
var a : i32 = 2; var i : i32 = 0; loop { const step : i32 = 1; if (i % 2 == 0) { continue; } a = a * 2; continuing { i = i + step; if (i < 4) {} else { break; } } }
var a : i32 = 2; var i : i32 = 0; loop { const step : i32 = 1; if (i % 2 == 0) { continue; } a = a * 2; continuing { i = i + step; break; // Invalid: too early if (i < 4) { i = i + 1; } else { break; } // Invalid: if is too complex, and too early if (i >= 4) { break; } else { i = i + 1; } // Invalid: if is too complex } }
6.2.7. Continue
continue_statement : CONTINUE
Use a continue
statement to transfer control in the nearest-enclosing § 6.2.4 Loop Statement:
-
forward to the § 6.2.8 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 not be placed such that it would transfer
control to an enclosing § 6.2.8 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 construct.
var i : i32 = 0; loop { if (i >= 4) { break; } if (i % 2 == 0) { continue; } // <3> const step : i32 = 2; continuing { i = i + step; } }
- <3> The
continue
is invalid because it bypasses the declaration ofstep
used in thecontinuing
construct
6.2.8. Continuing Statement
continuing_statement : CONTINUING body_statement
A continuing construct is a block of statements to be executed at the end of a loop iteration. The construct is optional.
The block of statements must not contain a return or discard statement.
6.2.9. Return
return_statement : RETURN short_circuit_or_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 return type of the function is the void type, then the return statement is optional. If the return statement is provided for a void function it must not have an expression. 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.
6.2.10. Discard TODO
The discard
statement must only be used in a fragment shader stage.
6.3. Function Call Statement TODO
func_call_statement : IDENT argument_expression_list
6.4. Statements Grammar Summary
body_statement : BRACE_LEFT statements BRACE_RIGHT paren_rhs_statement : PAREN_LEFT short_circuit_or_expression PAREN_RIGHT statements : statement* statement : SEMICOLON | return_statement SEMICOLON | if_statement | switch_statement | loop_statement | for_statement | func_call_statement SEMICOLON | variable_statement SEMICOLON | break_statement SEMICOLON | continue_statement SEMICOLON | DISCARD SEMICOLON | assignment_statement SEMICOLON | body_statement
7. Functions TODO
A function declaration may only occur at module scope. The function name is available for use after its declaration, until the end of the program.
If the return type of the function is not the void type, then the last statement in the function body must be a return statement.
Function names must be unique over all functions and all variables in the module.
function_decl : decoration_list* function_header body_statement function_type_decl : type_decl | VOID function_header : FN IDENT PAREN_LEFT param_list PAREN_RIGHT ARROW function_type_decl param_list : | (variable_ident_decl COMMA)* variable_ident_decl
Function decoration keys | Valid values | Note |
---|---|---|
stage
| compute or vertex or fragment
| |
workgroup_size
| non-negative i32 literals | The workgroup_size accepts a comma separated list of up to 3 values. The values provide the x, y and z dimensions. |
void %6 = OpTypeVoid fn my_func(i : i32, b : f32) -> i32 { return 2; } OpName %my_func "my_func" OpName %a "a" OpName %b "b" %my_func = OpFunction %int None %10 %a = OpFunctionParameter %_ptr_Function_int %b = OpFunctionParameter %_ptr_Function_float %14 = OpLabel OpReturnValue %int_2 OpFunctionEnd
7.1. Function declaration TODO
TODO: Stub
The names in the parameter list of a function definition are available for use in the body of the function. During a particular function evaluation, the parameter names denote the values specified to the function call expression or statement which initiated the function evaluation; the names and values are associated by position.
7.2. Function calls TODO
7.3. Restrictions TODO
TODO: This is a stub-
Recursion is not permitted. (No cycle in the call graph.)
-
Function call parameters
-
Match type and number
-
Restrictions on pointers
-
Aliasing (?)
8. Entry Points TODO
8.1. Shader Stages
In WebGPU, a pipeline is a unit of work executed on the GPU. There are two kinds of pipelines: GPUComputePipeline, and GPURenderPipeline.
A GPUComputePipeline runs 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 GPURenderPipeline is 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:
Each shader stage has its own set of features and constraints, described elsewhere.
8.2. Entry point declaration
An entry point is a function that is invoked to perform the work for a particular shader stage.
Specify a stage
attribute on a function declaration to declare that function
as an entry point.
When configuring the stage in the pipeline, the entry point is specified by providing the WGSL module and the entry point’s function name.
The parameters of an entry point have to be within Entry point IO types. They have in storage class. The return type of an entry point has to be of an Entry point IO type, or void.
[[stage(vertex)]] fn vert_main() -> [[builtin(position)]] vec4<f32> { return vec4<f32>(0.0, 0.0, 0.0, 1.0); } // OpEntryPoint Vertex %vert_main "vert_main" %return_value // OpDecorate %return_value BuiltIn Position // %float = OpTypeFloat 32 // %v4float = OpTypeVector %float 4 // %ptr = OpTypePointer Output %v4float // %return_value = OpVariable %ptr Output [[stage(fragment)]] fn frag_main([[builtin(frag_coord)]] coord_in: vec4<f32>) -> [[location(0)]] vec4<f32> { return vec4<f32>(coord_in.x, coord_in.y, 0.0, 1.0); } // OpEntryPoint Fragment %frag_main "frag_main" %return_value %coord_in // OpDecorate %return_value Location 0 // %float = OpTypeFloat 32 // %v4float = OpTypeVector %float 4 // %ptr = OpTypePointer Output %v4float // %return_value = OpVariable %ptr Output [[stage(compute)]] fn comp_main() -> void { } // OpEntryPoint GLCompute %comp_main "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.
8.2.1. Function attributes for entry points
- stage
-
The
stage
attribute declares that a function is an entry point for particular pipeline stage. - workgroup_size
-
The
workgroup_size
attribute specifies the x, y, and z dimensions of the workgroup grid for a compute shader. The size in the x dimension is provided by the first literal. The size in the y dimension is provided by the second literal, when present, and otherwise is assumed to be 1. The size in the z dimension is provided by the third literal, when present, and otherwise is assumed to be 1. Each dimension size must be at least 1 and at most an upper bound specified by the WebGPU API. This attribute must only be used with a compute shader stage entry point.
Can we query upper bounds on workgroup size dimensions? Is it independent of the shader, or a property to be queried after creating the shader module?
[[ stage(compute), workgroup_size(8,1,1) ]] fn sorter() -> void { } // OpEntryPoint GLCompute %sorter "sorter" // OpExecutionMode %sorter LocalSize 8 1 1 [[ stage(compute), workgroup_size(8) ]] fn reverser() -> void { } // OpEntryPoint GLCompute %reverser "reverser" // OpExecutionMode %reverser LocalSize 8 1 1 [[ stage(compute) ]] fn do_nothing() -> void { } // OpEntryPoint GLCompute %do_nothing "do_nothing" // OpExecutionMode %do_nothing LocalSize 1 1 1
8.3. Shader Interface
The shader interface is the set of objects through which the shader accesses data external to the shader stage, either for reading or writing. The interface includes:
-
Pipeline inputs and outputs
-
Buffer resources
-
Texture resources
These objects are represented by module-scope variables in certain storage classes.
We say a variable is statically accessed by a function if any subexpression in the body of the function uses the variable’s identifier, and that subexpression is in scope of the variable’s declaration. Note that being statically accessed is independent of whether an execution of the shader will actually evaluate the subexpression, or even execute the enclosing statement.
More precisely, the interface of a shader stage consists of:
-
all parameters of the entry point (as in)
-
the result value of the entry point
-
all module scope variables that are statically accessed by functions in the shader stage, and which are in storage classes uniform, storage, or handle.
8.3.1. Pipeline Input and Output Interface
The Entry point IO types include the following:
-
[[builtin(x)]]
-decorated pipeline built-ins. -
[[location(x)]]
-decorated IO-shareable types. -
structures containing only Entry point IO types.
A pipeline input is data provided to the shader stage from upstream in the pipeline. A pipeline input is denoted by the arguments of the entry point.
A pipeline output is data the shader provides for further processing downstream in the pipeline. A pipeline output is denoted by the return type of the entry point.
Each pipeline input or output is one of:
-
A built-in variable. See § 8.3.1.1 Built-in inputs and outputs.
-
A user data attribute. See § 8.3.1.2 User Data Attribute TODO.
8.3.1.1. Built-in inputs and outputs
A built-in input variable provides access to system-generated control information. The set of built-in inputs are listed in § 14 Built-in variables.
To declare a variable for accessing a particular input built-in X from an entry point:
-
Declare a parameter of the entry point function, where the store type is the listed store type for X.
-
Apply a
builtin(
X)
attribute to the parameter.
A built-in output variable is used by the shader to convey control information to later processing steps in the pipeline. The set of built-in outputs are listed in § 14 Built-in variables.
To declare a variable for accessing a particular output built-in Y from an entry point:
-
Add a variable to the result of the entry point, where store type is the listed store type for Y:
-
If the result type was void, change it to the variable type.
-
Otherwise, make the result type to be a structure, where one of the fields is the new variable.
-
Apply a
builtin(
Y)
attribute to the result variable.
The builtin
attribute must not be applied to a variables in module scope,
or the local variables in the function scope.
A variable must not have more than one builtin
attribute.
Each built-in variable has an associated shader stage, as described in § 14 Built-in variables. If a built-in variable has stage S and is statically accessed by a function F, then F must be a function in a shader for stage S.
-
The statement makes it clear that in/out storage classes for builtins are redundant.
-
On the other hand, in Vulkan, builtin variables occoupy I/O location slots (counting toward limits),
8.3.1.2. User Data Attribute TODO
TODO: User data attributes must not be of bool type or contain a bool type.
8.3.1.3. Input-output Locations TODO
TODO: Stub. Location-sizing of types, non-overlap among variables referenced within an entry point static call tree.8.3.2. Resource interface
A resource is an object, other than a pipeline input or output, which provides access to data external to a shader stage. Resources are shared by all invocations of the shader.
There are four kinds of resources:
-
textures
-
samplers
The resource interface of a shader is the set of module-scope resource variables statically accessed by functions in the shader stage.
Each resource variable must be declared with both group
and binding
attributes.
Together with the shader’s stage, these identify the binding address
of the resource on the shader’s pipeline.
See WebGPU § GPUPipelineLayout.
Bindings must not alias within a shader stage: two different variables in the resource interface of a given shader must not have the same group and binding values, when considered as a pair of values.
Decoraton | Operand | Description |
---|---|---|
group
| non-negative i32 literal | Bind group index |
binding
| non-negative i32 literal | Binding number index |
8.3.3. Resource layout compatibility
WebGPU requires that a shader’s resource interface match the layout of the pipeline using the shader.
Each WGSL variable in a resource interface must be bound to a WebGPU resource with a compatible GPUBindingType, where compatibility is defined by the following table.
WGSL resource | WebGPU GPUBindingType |
---|---|
uniform buffer | uniform-buffer |
read-write storage buffer | storage-buffer |
read-only storage buffer | readonly-storage-buffer |
sampler | sampler |
sampler_comparison | comparison-sampler |
sampled texture | sampled-texture or multisampled-texture |
read-only storage texture | readonly-storage-texture |
write-only storage texture | writeonly-storage-texture |
TODO: Rewrite the phrases 'read-only storage buffer' and 'read-write storage buffer' after we settle on how to express those concepts. See https://github.com/gpuweb/gpuweb/pull/1183
If B is a uniform buffer variable in a resource interface, and WB is the WebGPU GPUBuffer bound to B, then:
-
The size of WB must be at least as large as the allocation extent of the store type of B in the storage storage class.
If B is a storage buffer variable in a resource interface, and WB is the WebGPU GPUBuffer bound to B, then:
-
If the store type S of B does not contain a runtime-sized array, then the size of WB must be at least as large as the allocation extent of S in the storage storage class.
-
If the store type S of B contains a runtime-sized array as its last member, then:
-
The runtime-determined array length of that member must be at least 1.
-
The size of WB must be at least as large as the allocation extent in storage class storage of the value stored in B.
-
Note: Recall that a runtime-sized array may only appear as the last element in the structure type that is the store type of a storage buffer variable.
TODO: Describe other interface matching requirements, e.g. for images?
8.4. Pipeline compatibility TODO
TODO: match flat attribute
TODO: user data inputs of fragment stage must be subset of user data outputs of vertex stage
8.4.1. Input-output matching rules TODO
9. WGSL program TODO
TODO: Stub A WGSL program is a sequence of module-scope declarations.
translation_unit : global_decl* EOF
global_decl : SEMICOLON | global_variable_decl SEMICOLON | global_constant_decl SEMICOLON | type_alias SEMICOLON | struct_decl SEMICOLON | function_decl
10. Execution TODO
10.1. Invocation of an entry point TODO
10.1.1. Before an entry point begins TODO
TODO: Stub
-
Setting values of builtin variables
-
External-interface variables have initialized backing storage
-
Internal module-scope variables have backing storage
-
Initializers evaluated in textual order
-
No two variables have overlapping storage (might already be covered earlier?)
10.1.2. Program order (within an invocation) TODO
10.1.2.1. Function-scope variable lifetime and initialization TODO
10.1.2.2. Statement order TODO
10.1.2.3. Intra-statement order (or lack) TODO
TODO: Stub: Expression evaluation
10.2. Uniformity TODO
10.2.1. Uniform control flow TODO
10.2.2. Divergence and reconvergence TODO
10.2.3. Uniformity restrictions TODO
10.3. Compute Shaders and Workgroups
A workgroup is a set of invocations which concurrently execute a compute shader stage entry point, and share access to shader variables in the workgroup storage class.
The workgroup grid for a compute shader is the set of points with integer coordinates (i,j,k) with:
-
0 ≤ i < workgroup_size_x
-
0 ≤ j < workgroup_size_y
-
0 ≤ k < workgroup_size_z
where (workgroup_size_x, workgroup_size_y, workgroup_size_z) is the value specified for the workgroup_size attribute of the entry point, or (1,1,1) if the entry point has no such attribute.
There is exactly one invocation in a workgroup for each point in the workgroup grid.
An invocation’s local invocation ID is the coordinate triple for the invocation’s corresponding workgroup grid point.
When an invocation has local invocation ID (i,j,k), then its local invocation index is
i + (j * workgroup_size_x) + (k * workgroup_size_x * workgroup_size_y)
Note that if a workgroup has W invocations, then each invocation I the workgroup has a unique local invocation index L(I) such that 0 ≤ L(I) < W, and that entire range is covered.
A compute shader begins execution when a WebGPU implementation removes a dispatch command from a queue and begins the specified work on the GPU. The dispatch command specifies a dispatch size, which is an integer triple (group_count_x, group_count_y, group_count_z) indicating the number of workgroups to be executed, as described in the following.
The compute shader grid for a particular dispatch is the set of points with integer coordinates (CSi,CSj,CSk) with:
-
0 ≤ CSi ≤ workgroup_size_x × group_count_x
-
0 ≤ CSj ≤ workgroup_size_y × group_count_y
-
0 ≤ CSk ≤ workgroup_size_z × group_count_z
where workgroup_size_x, workgroup_size_y, and workgroup_size_z are as above for the compute shader entry point.
The work to be performed by a compute shader dispatch is to execute exactly one invocation of the entry point for each point in the compute shader grid.
An invocation’s global invocation ID is the coordinate triple for the invocation’s corresponding compute shader grid point.
The invocations are organized into workgroups, so that each invocation (CSi, CSj, CSk) is identified with the workgroup grid point
( CSi mod workgroup_size_x , CSj mod workgroup_size_y , CSk mod workgroup_size_z )
in workgroup ID
( ⌊ CSi ÷ workgroup_size_x ⌋, ⌊ CSj ÷ workgroup_size_y ⌋, ⌊ CSk ÷ workgroup_size_z ⌋).
WebGPU provides no guarantees about:
-
Whether invocations from different workgroups execute concurrently. That is, you cannot assume more than one workgroup executes at a time.
-
Whether, once invocations from a workgroup begin executing, that other workgroups are blocked from execution. That is, you cannot assume that only one workgroup executes at a time. While a workgroup is executing, the implementation may choose to concurrently execute other workgroups as well, or other queued but unblocked work.
-
Whether invocations from one particular workgroup begin executing before the invocations of another workgroup. That is, you cannot assume that workgroups are launched in a particular order.
WebGPU issue 1045: Dispatch group counts must be positive. However, how do we handle an indirect dispatch that specifies a group count of zero.
10.4. Collective operations TODO
10.4.1. Barrier TODO
10.4.2. Image Operations Requiring Uniformity TODO
10.4.3. Derivatives TODO
10.4.4. Arrayed resource access TODO
10.5. Floating Point Evaluation TODO
TODO: Stub
-
Infinities, NaNs, negative zeros
-
Denorms, flushing
-
fast-math rules: e.g. reassociation, fusing
-
Invariance (or is this more general than floating point)
-
Rounding
-
Error bounds on basic operations
10.5.1. Floating point conversion
When converting a floating point scalar value to an integral type:
-
If the original value is exactly representable in the destination type, then the result is that value.
-
If the original value has a fractional component, then it cannot be represented exactly in the destination type, and the result is TODO
-
If the original value is out of range of the destination type, then TODO.
When converting a value to a floating point type:
-
If the original value is exactly representable in the destination type, then the result is that value.
-
If the original value is zero and of integral type, then the resulting value has a zero sign bit.
-
-
Otherwise, the original value is not exactly representable.
-
If the original value is different from but lies between two adjacent values representable in the destination type, then the result is one of those two values. WGSL does not specify whether the larger or smaller representable value is chosen, and different instances of such a conversion may choose differently.
-
Otherwise, if the original value lies outside the range of the destination type.
-
This does not occur when the original types is one of i32 or u32 and the destination type is f32.
-
This does not occur when the source type is a floating point type with fewer exponent and mantissa bits.
-
If the source type is a floating point type with more mantissa bits than the destination type, then:
-
The extra mantissa bits of the source value may be discarded (treated as if they are 0).
-
If the resulting value is the maximum normal value of the destination type, then that is the result.
-
-
Otherwise the result is the infinity value with the same sign as the source value.
-
-
-
Otherwise, if the original value is a NaN for the source type, then the result is a NaN in the destination type.
-
NOTE: An integer value may lie between two adjacent representable floating point values. In particular, the f32 type uses 23 explicit fractional bits. Additionally, when the floating point value is in the normal range (the exponent is neither extreme value), then the mantissa is the set of fractional bits together with an extra 1-bit at the most significant position at bit position 23. Then, for example, integers 228 and 1+228 both map to the same floating point value: the difference in the least significant 1 bit is not representable by the floating point format. This kind of collision occurs for pairs of adjacent integers with a magnitude of at least 225.
(dneto) Default rounding mode is an implementation choice. Is that what we want?
Check behaviour of the f32 to f16 conversion for numbers just beyond the max normal f16 values. I’ve written what an NVIDIA GPU does. See https://github.com/google/amber/pull/918 for an executable test case.
11. Memory Model TODO
12. Keyword and Token Summary
12.1. Keyword Summary
Token | Definition |
---|---|
ARRAY
| array |
BOOL
| bool |
FLOAT32
| f32 |
INT32
| i32 |
MAT2x2
| mat2x2 // 2 column x 2 row |
MAT2x3
| mat2x3 // 2 column x 3 row |
MAT2x4
| mat2x4 // 2 column x 4 row |
MAT3x2
| mat3x2 // 3 column x 2 row |
MAT3x3
| mat3x3 // 3 column x 3 row |
MAT3x4
| mat3x4 // 3 column x 4 row |
MAT4x2
| mat4x2 // 4 column x 2 row |
MAT4x3
| mat4x3 // 4 column x 3 row |
MAT4x4
| mat4x4 // 4 column x 4 row |
POINTER
| ptr |
SAMPLER
| sampler |
SAMPLER_COMPARISON
| sampler_comparison |
STRUCT
| struct |
TEXTURE_1D
| texture_1d |
TEXTURE_2D
| texture_2d |
TEXTURE_2D_ARRAY
| texture_2d_array |
TEXTURE_3D
| texture_3d |
TEXTURE_CUBE
| texture_cube |
TEXTURE_CUBE_ARRAY
| texture_cube_array |
TEXTURE_MULTISAMPLED_2D
| texture_multisampled_2d |
TEXTURE_STORAGE_1D
| texture_storage_1d |
TEXTURE_STORAGE_2D
| texture_storage_2d |
TEXTURE_STORAGE_2D_ARRAY
| texture_storage_2d_array |
TEXTURE_STORAGE_3D
| texture_storage_3d |
TEXTURE_DEPTH_2D
| texture_depth_2d |
TEXTURE_DEPTH_2D_ARRAY
| texture_depth_2d_array |
TEXTURE_DEPTH_CUBE
| texture_depth_cube |
TEXTURE_DEPTH_CUBE_ARRAY
| texture_depth_cube_array |
UINT32
| u32 |
VEC2
| vec2 |
VEC3
| vec3 |
VEC4
| vec4 |
VOID
| void |
Token | Definition |
BITCAST
| bitcast |
BLOCK
| block |
BREAK
| break |
CASE
| case |
CONST
| const |
CONTINUE
| continue |
CONTINUING
| continuing |
DEFAULT
| default |
DISCARD
| discard |
ELSE
| else |
ELSE_IF
| elseif |
FALLTHROUGH
| fallthrough |
FALSE
| false |
FN
| fn |
FOR
| for |
FUNCTION
| function |
IF
| if |
IN
| in |
LOOP
| loop |
OUT
| out |
PRIVATE
| private |
RETURN
| return |
STORAGE
| storage |
SWITCH
| switch |
TRUE
| true |
TYPE
| type |
UNIFORM
| uniform |
VAR
| var |
WORKGROUP
| workgroup |
Token | Definition |
R8UNORM
| r8unorm |
R8SNORM
| r8snorm |
R8UINT
| r8uint |
R8SINT
| r8sint |
R16UINT
| r16uint |
R16SINT
| r16sint |
R16FLOAT
| r16float |
RG8UNORM
| rg8unorm |
RG8SNORM
| rg8snorm |
RG8UINT
| rg8uint |
RG8SINT
| rg8sint |
R32UINT
| r32uint |
R32SINT
| r32sint |
R32FLOAT
| r32float |
RG16UINT
| rg16uint |
RG16SINT
| rg16sint |
RG16FLOAT
| rg16float |
RGBA8UNORM
| rgba8unorm |
RGBA8UNORM-SRGB
| rgba8unorm_srgb |
RGBA8SNORM
| rgba8snorm |
RGBA8UINT
| rgba8uint |
RGBA8SINT
| rgba8sint |
BGRA8UNORM
| bgra8unorm |
BGRA8UNORM-SRGB
| bgra8unorm_srgb |
RGB10A2UNORM
| rgb10a2unorm |
RG11B10FLOAT
| rg11b10float |
RG32UINT
| rg32uint |
RG32SINT
| rg32sint |
RG32FLOAT
| rg32float |
RGBA16UINT
| rgba16uint |
RGBA16SINT
| rgba16sint |
RGBA16FLOAT
| rgba16float |
RGBA32UINT
| rgba32uint |
RGBA32SINT
| rgba32sint |
RGBA32FLOAT
| rgba32float |
TODO(dneto): Eliminate the image formats that are not used in storage images. For example SRGB formats (bgra8unorm_srgb), mixed channel widths (rg11b10float), out-of-order channels (bgra8unorm)
12.2. Reserved Keywords
The following is a list of keywords which are reserved for future expansion.asm | bf16 | do | enum | f16 |
f64 | i8 | i16 | i64 | let |
typedef | u8 | u16 | u64 | unless |
using | while | regardless | premerge | handle |
12.3. Syntactic Tokens
AND
| &
|
AND_AND
| &&
|
ARROW
| ->
|
ATTR_LEFT
| [[
|
ATTR_RIGHT
| ]]
|
FORWARD_SLASH
| /
|
BANG
| !
|
BRACKET_LEFT
| [
|
BRACKET_RIGHT
| ]
|
BRACE_LEFT
| {
|
BRACE_RIGHT
| }
|
COLON
| :
|
COMMA
| ,
|
EQUAL
| =
|
EQUAL_EQUAL
| ==
|
NOT_EQUAL
| !=
|
GREATER_THAN
| >
|
GREATER_THAN_EQUAL
| >=
|
SHIFT_RIGHT
| >>
|
LESS_THAN
| <
|
LESS_THAN_EQUAL
| <=
|
SHIFT_LEFT
| <<
|
MODULO
| %
|
MINUS
| -
|
PERIOD
| .
|
PLUS
| +
|
OR
| |
|
OR_OR
| ||
|
PAREN_LEFT
| (
|
PAREN_RIGHT
| )
|
SEMICOLON
| ;
|
STAR
| *
|
TILDE
| ~
|
XOR
| ^
|
13. Validation
TODO: Move these to the subject-matter sections.
Each validation item will be given a unique ID and a test must be provided when the validation is added. The tests will reference the validation ID in the test name.
-
v-0001: A declaration must not introduce a name when that name is already in scope at the start of the declaration.
-
v-0002: Non-void functions must end with a return statement.
-
v-0003: At least one of vertex, fragment or compute shader must be present.
-
v-0004: Recursion is not allowed.
-
v-0007: Structures must be defined before use.
-
v-0008: switch statements must have exactly one default clause.
-
v-0009: Break is only permitted in loop and switch constructs.
-
v-0010: continue is only permitted in loop.
-
v-0015: The last member of the structure type defining the "store type" for variable in the storage storage class may be a runtime-sized array.
-
v-0017: Builtin decorations must have the correct types.
-
v-0018: Builtin decorations must be used with the correct shader type and storage class.
-
v-0020: The pair of
<entry point name, pipeline stage>
must be unique in the module. -
v-0021: Cannot re-assign a constant.
-
v-0022: Global variables must have a storage class.
-
v-0023: Entry point functions accept no parameters.
-
v-0024: Entry point functions return void.
-
v-0025: Switch statement selector expression must be of a scalar integer type.
-
v-0026: The case selector values must have the same type as the selector expression.
-
v-0027: A literal value must not appear more than once in the case selectors for a switch statement.
-
v-0028: A fallthrough statement must not appear as the last statement in last clause of a switch.
-
v-0029: Return must come last in its block.
-
v-0030: A runtime-sized array must not be used as the store type or contained within a store type except as allowed by v-0015.
-
v-0031: The type of an expression must not be a runtime-sized array type.
-
v-0032: A runtime-sized array must have a stride attribute.
14. Built-in variables
See § 8.3.1.1 Built-in inputs and outputs for how to declare a built-in variable.
Built-in | Stage | Input or Output | Store type | Description |
---|---|---|---|---|
vertex_index
| vertex | in | u32 |
Index of the current vertex within the current API-level draw command,
independent of draw instancing.
For a non-indexed draw, the first vertex has an index equal to the For an indexed draw, the index is equal to the index buffer entry for
vertex, plus the |
instance_index
| vertex | in | u32 |
Instance index of the current vertex within the current API-level draw command.
The first instance has an index equal to the |
position
| vertex | out | vec4<f32> | Output position of the current vertex, using homogeneous coordinates. After homogeneous normalization (where each of the x, y, and z components are divided by the w component), the position is in the WebGPU normalized device coordinate space. See WebGPU § Coordinate Systems. |
frag_coord
| fragment | in | vec4<f32> | Framebuffer position of the current fragment, using normalized homogeneous coordinates. (The x, y, and z components have already been scaled such that w is now 1.) See WebGPU § Coordinate Systems. |
front_facing
| fragment | in | bool | True when the current fragment is on a front-facing primitive. False otherwise. See WebGPU § Rasterization State. |
frag_depth
| fragment | out | f32 | Updated depth of the fragment, in the viewport depth range. See WebGPU § Coordinate Systems. |
local_invocation_id
| compute | in | vec3<u32> | The current invocation’s local invocation ID, i.e. its position in the workgroup grid. |
local_invocation_index
| compute | in | u32 | The current invocation’s local invocation index, a linearized index of the invocation’s position within the workgroup grid. |
global_invocation_id
| compute | in | vec3<u32> | The current invocation’s global invocation ID, i.e. its position in the compute shader grid. |
workgroup_id
| compute | in | vec3<u32> | The current invocation’s workgroup ID, i.e. the position of the workgroup in the the workgroup grid. |
workgroup_size
| compute | in | vec3<u32> | The workgroup_size of the current entry point. |
subgroup_size
| compute | in | u32 | The subgroup size of the current entry point. |
subgroup_invocation_index
| compute | in | u32 | The current invocation’s subgroup invocation index.
Must be in range [0, subgroup_size -1].
|
sample_index
| fragment | in | u32 | Sample index for the current fragment.
The value is least 0 and at most sampleCount -1, where sampleCount is the number of MSAA samples specified for the GPU render pipeline. See WebGPU § GPURenderPipeline. |
sample_mask_in
| fragment | in | u32 | Sample coverage mask for the current fragment.
It contains a bitmask indicating which samples in this fragment are covered
by the primitive being rendered. See WebGPU § Sample Masking. |
sample_mask_out
| fragment | out | u32 | Sample coverage mask control for the current fragment.
The last value written to this variable becomes the shader-output mask.
Zero bits in the written value will cause corresponding samples in
the color attachments to be discarded. The value in the variable is significant only if the sample_mask_out variable is statically accessed by the fragment shader. If the variable is not statically accessed,
then other factors determine sample coverage. See WebGPU § Sample Masking. |
struct VertexOutput { [[builtin(position)]] my_pos: vec4<f32>; // OpDecorate %my_pos BuiltIn Position // %float = OpTypeFloat 32 // %v4float = OpTypeVector %float 4 // %ptr = OpTypePointer Output %v4float // %my_pos = OpVariable %ptr Output }; [[stage(vertex)]] fn vs_main( [[builtin(vertex_index)]] my_index: u32, // OpDecorate %my_index BuiltIn VertexIndex // %uint = OpTypeInt 32 0 // %ptr = OpTypePointer Input %uint // %my_index = OpVariable %ptr Input [[builtin(instance_index)]] my_inst_index : u32, // OpDecorate %my_inst_index BuiltIn InstanceIndex ) -> VertexOutput; struct FragmentOutput { [[builtin(frag_depth)]] depth: f32; // OpDecorate %depth BuiltIn FragDepth [[builtin(sample_mask_out)]] mask_out : u32; // OpDecorate %mask_out BuiltIn SampleMask ; an output variable }; [[stage(fragment)]] fn fs_main( [[builtin(front_facing)]] is_front : u32, // OpDecorate %is_front BuiltIn FrontFacing [[builtin(frag_coord)]] coord : vec4<f32>, // OpDecorate %coord BuiltIn FragCoord [[builtin(sample_index)]] my_sample_index : u32, // OpDecorate %my_sample_index BuiltIn SampleId [[builtin(sample_mask_in)]] mask_in : u32, // OpDecorate %mask_in BuiltIn SampleMask ; an input variable // OpDecorate %mask_in Flat ) -> FragmentOutput; [[stage(compute)]] fn cs_main( [[builtin(local_invocation_id)]] local_id : vec3<u32>, // OpDecorate %local_id BuiltIn LocalInvocationId [[builtin(local_invocation_index)]] local_index : u32, // OpDecorate %local_index BuiltIn LocalInvocationIndex [[builtin(global_invocation_id)]] global_id : vec3<u32>, // OpDecorate %global_id BuiltIn GlobalInvocationId ) -> void;
15. Built-in functions
Certain functions are always available in a WGSL program, and are provided by the implementation. These are called built-in functions.
Since a built-in function is always in scope, it is an error to attempt to redefine one or to use the name of a built-in function as an identifier for any other kind of declaration.
Unlike ordinary functions defined in a WGSL program, a built-in function may use the same function name with different sets of parameters. In other words, a built-in function may have more than one overload, but ordinary function definitions in WGSL may not.
When calling a built-in function, all arguments to the function are evaluated before function evaulation begins.
TODO(dneto): Elaborate the descriptions of the built-in functions. So far I’ve only reorganized the contents of the existing table.
15.1. Logical built-in functions
Logical built-in functions | SPIR-V |
---|---|
all(BoolVec) -> bool | OpAll |
any(BoolVec) -> bool | OpAny |
select(T,T,bool) -> T | For scalar or vector type T. select(a,b,c) evaluates to a when c is true, and b otherwise.OpSelect |
select(vecN<T>,vecN<T>,vecN<bool>) -> vecN<T> | For scalar type T. select(a,b,c) evaluates to a vector with component i being select(a[i], b[i], c[i]) .OpSelect |
15.2. Value-testing built-in functions
Value-testing built-in functions | SPIR-V |
---|---|
isFinite(float) -> bool | OpIsFinite |
isInf(float) -> bool | OpIsInf |
isNan(float) -> bool | OpIsNan |
isNormal(float) -> bool | OpIsNormal |
TODO: deduplicate these tables
Precondition | Conclusion | Notes |
---|---|---|
e : f32 | isNan(e) : bool
| OpIsNan |
e : T, T is FloatVec | isNan(e) : bool<N>, where N = Arity(T)
| OpIsNan |
e : f32 | isInf(e) : bool
| OpIsInf |
e : T, T is FloatVec | isInf(e) : bool<N>, where N = Arity(T)
| OpIsInf |
e : f32 | isFinite(e) : bool
| OpIsFinite |
e : T, T is FloatVec | isFinite(e) : bool<N>, where N = Arity(T)
| OpIsFinite, or emulate |
e : f32 | isNormal(e) : bool
| OpIsNormal |
e : T, T is FloatVec | isNormal(e) : bool<N>, where N = Arity(T)
| OpIsNormal, or emulate |
e : array<E> | arrayLength(e) : u32
| OpArrayLength |
15.3. Float built-in functions
Precondition | Built-in | Description |
---|---|---|
T is f32 | abs( e: T ) -> T
| (GLSLstd450FAbs) |
T is f32 | abs( e: vecN<T> ) -> vecN<T>
| (GLSLstd450FAbs) |
T is f32 | acos( e: T ) -> T
| (GLSLstd450Acos) |
T is f32 | acos( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Acos) |
T is f32 | asin( e: T ) -> T
| (GLSLstd450Asin) |
T is f32 | asin( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Asin) |
T is f32 | atan( e: T ) -> T
| (GLSLstd450Atan) |
T is f32 | atan( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Atan) |
T is f32 | atan2( e1: T , e2: T ) -> T
| (GLSLstd450Atan2) |
T is f32 | atan2( e1: vecN<T> , e2: vecN<T> ) -> vecN<T>
| (GLSLstd450Atan2) |
T is f32 | ceil( e: T ) -> T
| (GLSLstd450Ceil) |
T is f32 | ceil( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Ceil) |
T is f32 | clamp( e1: T , e2: T , e3: T) -> T
| (GLSLstd450NClamp) |
T is f32 | clamp( e1: vecN<T> , e2: vecN<T>, e3: vecN<T>) -> vecN<T>
| (GLSLstd450NClamp) |
T is f32 | cos( e: T ) -> T
| (GLSLstd450Cos) |
T is f32 | cos( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Cos) |
T is f32 | cosh( e: T ) -> T
| (GLSLstd450Cosh) |
T is f32 | cosh( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Cosh) |
T is f32 | cross( e1: vec3<T> , e2: vec3<T>) -> vec3<T>
| (GLSLstd450Cross) |
T is f32 | distance( e1: T , e2: T ) -> T
| (GLSLstd450Distance) |
T is f32 | distance( e1: vecN<T> , e2: vecN<T>) -> T
| (GLSLstd450Distance) |
T is f32 | exp( e: T ) -> T
| (GLSLstd450Exp) |
T is f32 | exp( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Exp) |
T is f32 | exp2( e: T ) -> T
| (GLSLstd450Exp2) |
T is f32 | exp2( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Exp2) |
T is f32 | faceForward( e1: T , e2: T , e3: T ) -> T
| (GLSLstd450FaceForward) |
T is f32 | faceForward( e1: vecN<T> , e2: vecN<T>, e3: vecN<T>) -> vecN<T>
| (GLSLstd450FaceForward) |
T is f32 | floor( e: T ) -> T
| (GLSLstd450Floor) |
T is f32 | floor( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Floor) |
T is f32 | fma( e1: T , e2: T , e3: T ) -> T
| (GLSLstd450Fma) |
T is f32 | fma( e1: vecN<T> , e2: vecN<T>, e3: vecN<T>) -> vecN<T>
| (GLSLstd450Fma) |
T is f32 | fract( e: T ) -> T
| (GLSLstd450Fract) |
T is f32 | fract( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Fract) |
T is f32 I is i32 or u32 | frexp( e1: T , e2: ptr<I> ) -> T
| (GLSLstd450Frexp) |
T is f32 I is i32 or u32 | frexp( e1: vecN<T> , e2: ptr<vecN<I>>) -> vecN<T>
| (GLSLstd450Frexp) |
T is f32 | inverseSqrt( e: T ) -> T
| (GLSLstd450InverseSqrt) |
T is f32 | inverseSqrt( e: vecN<T> ) -> vecN<T>
| (GLSLstd450InverseSqrt) |
T is f32 I is i32 or u32 | ldexp( e1: T , e2: I ) -> T
| (GLSLstd450Ldexp) |
T is f32 I is i32 or u32 | ldexp( e1: vecN<T> , e2: vecN<I>) -> vecN<T>
| (GLSLstd450Ldexp) |
T is f32 | length( e: T ) -> T
| (GLSLstd450Length) |
T is f32 | length( e: vecN<T> ) -> T
| (GLSLstd450Length) |
T is f32 | log( e: T ) -> T
| (GLSLstd450Log) |
T is f32 | log( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Log) |
T is f32 | log2( e: T ) -> T
| (GLSLstd450Log2) |
T is f32 | log2( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Log2) |
T is f32 | max( e1: T , e2: T ) -> T
| (GLSLstd450NMax) |
T is f32 | max( e1: vecN<T> , e2: vecN<T>) -> vecN<T>
| (GLSLstd450NMax) |
T is f32 | min( e1: T , e2: T ) -> T
| (GLSLstd450NMin) |
T is f32 | min( e1: vecN<T> , e2: vecN<T>) -> vecN<T>
| (GLSLstd450NMin) |
T is f32 | mix( e1: T , e2: T , e3: T) -> T
| (GLSLstd450FMix) |
T is f32 | mix( e1: vecN<T> , e2: vecN<T>, e3: vecN<T>) -> vecN<T>
| (GLSLstd450FMix) |
T is f32 | modf( e1: T , e2: ptr<T> ) -> T
| (GLSLstd450Modf) |
T is f32 | modf( e1: vecN<T> , e2: ptr<vecN<T>>) -> vecN<T>
| (GLSLstd450Modf) |
T is f32 | normalize( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Normalize) |
T is f32 | pow( e1: T , e2: T ) -> T
| (GLSLstd450Pow) |
T is f32 | pow( e1: vecN<T> , e2: vecN<T> ) -> vecN<T>
| (GLSLstd450Pow) |
T is f32 | reflect( e1: T , e2: T ) -> T
| (GLSLstd450Reflect) |
T is f32 | reflect( e1: vecN<T> , e2: vecN<T>) -> vecN<T>
| (GLSLstd450Reflect) |
T is f32 | round( e: T ) -> T
| Result is the integer k nearest to e, as a floating point value. When e lies halfway between integers k and k+1, the result is k when k is even, and k+1 when k is odd. (GLSLstd450RoundEven) |
T is f32 | round( e: vecN<T> ) -> vecN<T>
| Component-wise rounding. Component i of the result is round (e[i])(GLSLstd450RoundEven) |
T is f32 | sign( e: T ) -> T
| (GLSLstd450FSign) |
T is f32 | sign( e: vecN<T> ) -> vecN<T>
| (GLSLstd450FSign) |
T is f32 | sin( e: T ) -> T
| (GLSLstd450Sin) |
T is f32 | sin( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Sin) |
T is f32 | sinh( e: T ) -> T
| (GLSLstd450Sinh) |
T is f32 | sinh( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Sinh) |
T is f32 | smoothStep( e1: T , e2: T , e3: T ) -> T
| (GLSLstd450SmoothStep) |
T is f32 | smoothStep( e1: vecN<T> , e2: vecN<T>, e3: vecN<T>) -> vecN<T>
| (GLSLstd450SmoothStep) |
T is f32 | sqrt( e: T ) -> T
| (GLSLstd450Sqrt) |
T is f32 | sqrt( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Sqrt) |
T is f32 | step( e1: T , e2: T ) -> T
| (GLSLstd450Step) |
T is f32 | step( e1: vecN<T> , e2: vecN<T>) -> vecN<T>
| (GLSLstd450Step) |
T is f32 | tan( e: T ) -> T
| (GLSLstd450Tan) |
T is f32 | tan( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Tan) |
T is f32 | tanh( e: T ) -> T
| (GLSLstd450Tanh) |
T is f32 | tanh( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Tanh) |
T is f32 | trunc( e: T ) -> T
| (GLSLstd450Trunc) |
T is f32 | trunc( e: vecN<T> ) -> vecN<T>
| (GLSLstd450Trunc) |
15.4. Integer built-in functions
Precondition | Built-in | Description |
---|---|---|
abs (e: i32 ) -> i32
| The absolute value of e. (GLSLstd450SAbs) | |
abs (e : vecN<i32> ) -> vecN<i32>
| Component-wise absolute value:
Component i of the result is abs( e[ i]) (GLSLstd450SAbs) | |
abs (e : u32 ) -> u32
| Result is e. This is provided for symmetry with abs for signed integers.
| |
abs( e: vecN<u32> ) -> vecN<u32>
| Result is e. This is provided for symmetry with abs for signed integer vectors.
| |
T is u32 | clamp( e1: T , e2: T, e3: T) -> T
| (GLSLstd450UClamp) |
T is u32 | clamp( e1: vecN<T> , e2: vecN<T>, e3: vecN<T> ) -> vecN<T>
| (GLSLstd450UClamp) |
T is i32 | clamp( e1: T , e2: T, e3: T) -> T
| (GLSLstd450SClamp) |
T is i32 | clamp( e1: vecN<T> , e2: vecN<T>, e3: vecN<T> ) -> vecN<T>
| (GLSLstd450SClamp) |
T is u32 or i32 | countOneBits( e: T ) -> T
| The number of 1 bits in the representation of e. Also known as "population count". (SPIR-V OpBitCount) |
T is u32 or i32 | countOneBits( e: vecN<T>) -> vecN<T> | Component-wise population count:
Component i of the result is countOneBits( e[ i]) (SPIR-V OpBitCount) |
T is u32 | max( e1: T , e2: T) -> T
| (GLSLstd450UMax) |
T is u32 | max( e1: vecN<T> , e2: vecN<T>) -> vecN<T>
| (GLSLstd450UMax) |
T is i32 | max( e1: T , e2: T) -> T
| (GLSLstd450SMax) |
T is i32 | max( e1: vecN<T> , e2: vecN<T>) -> vecN<T>
| (GLSLstd450SMax) |
T is u32 | min( e1: T , e2: T) -> T
| (GLSLstd450UMin) |
T is u32 | min( e1: vecN<T> , e2: vecN<T>) -> vecN<T>
| (GLSLstd450UMin) |
T is i32 | min( e1: T , e2: T) -> T
| (GLSLstd450SMin) |
T is i32 | min( e1: vecN<T> , e2: vecN<T>) -> vecN<T>
| (GLSLstd450SMin) |
T is u32 or i32 | reverseBits( e: T ) -> T
| Reverses the bits in e: The bit at position k of the result equals the
bit at position 31-k of e. (SPIR-V OpBitReverse) |
T is u32 or i32 | reverseBits( e: vecN<T> ) -> vecN<T> | Component-wise bit reversal:
Component i of the result is reverseBits( e[ i]) (SPIR-V OpBitReverse) |
15.5. Matrix built-in functions
Precondition | Built-in | Description |
---|---|---|
T is f32 | determinant( e: matNxN<T> ) -> T
| (GLSLstd450Determinant) |
15.6. Vector built-in functions
Vector built-in functions | SPIR-V |
---|---|
dot(vecN<f32>, vecN<f32>) -> float | OpDot |
15.7. Derivative built-in functions
Derivative built-in functions | SPIR-V |
---|---|
dpdx(IDENT) -> float | OpDPdx |
dpdxCoarse(IDENT) -> float | OpDPdxCoarse |
dpdxFine(IDENT) -> float | OpDPdxFine |
dpdy(IDENT) -> float | OpDPdy |
dpdyCoarse(IDENT) -> float | OpDPdyCoarse |
dpdyFine(IDENT) -> float | OpDPdyFine |
fwidth(IDENT) -> float | OpFwidth |
fwidthCoarse(IDENT) -> float | OpFwidthCoarse |
fwidthFine(IDENT) -> float | OpFwidthFine |
15.8. Texture built-in functions
15.8.1. textureDimensions
Returns the dimensions of a texture, or texture’s mip level in texels.
textureDimensions ( t :texture_1d < T > ) ->i32 textureDimensions ( t :texture_2d < T > ) ->vec2 < i32 > textureDimensions ( t :texture_2d < T > , level :i32 ) ->vec2 < i32 > textureDimensions ( t :texture_2d_array < T > ) ->vec2 < i32 > textureDimensions ( t :texture_2d_array < T > , level :i32 ) ->vec2 < i32 > textureDimensions ( t :texture_3d < T > ) ->vec3 < i32 > textureDimensions ( t :texture_3d < T > , level :i32 ) ->vec3 < i32 > textureDimensions ( t :texture_cube < T > ) ->vec3 < i32 > textureDimensions ( t :texture_cube < T > , level :i32 ) ->vec3 < i32 > textureDimensions ( t :texture_cube_array < T > ) ->vec3 < i32 > textureDimensions ( t :texture_cube_array < T > , level :i32 ) ->vec3 < i32 > textureDimensions ( t :texture_multisampled_2d < T > ) ->vec2 < i32 > textureDimensions ( t :texture_multisampled_2d_array < T > ) ->vec2 < i32 > textureDimensions ( t :texture_depth_2d ) ->vec2 < i32 > textureDimensions ( t :texture_depth_2d , level :i32 ) ->vec2 < i32 > textureDimensions ( t :texture_depth_2d_array ) ->vec2 < i32 > textureDimensions ( t :texture_depth_2d_array , level :i32 ) ->vec2 < i32 > textureDimensions ( t :texture_depth_cube ) ->vec3 < i32 > textureDimensions ( t :texture_depth_cube , level :i32 ) ->vec3 < i32 > textureDimensions ( t :texture_depth_cube_array ) ->vec3 < i32 > textureDimensions ( t :texture_depth_cube_array , level :i32 ) ->vec3 < i32 > textureDimensions ( t :texture_storage_1d < F > ) ->i32 textureDimensions ( t :texture_storage_2d < F > ) ->vec2 < i32 > textureDimensions ( t :texture_storage_2d_array < F > ) ->vec2 < i32 > textureDimensions ( t :texture_storage_3d < F > ) ->vec3 < i32 >
Parameters:
t
| The sampled, multisampled, depth, or storage texture. |
level
| The mip level, with level 0 containing a full size version of the texture. If omitted, the dimensions of level 0 are returned. |
Returns:
The dimensions of the texture in texels.
15.8.2. textureLoad
Reads a single texel from a texture without sampling or filtering.
textureLoad ( t :texture_1d < T > , coords :i32 , level :i32 ) ->vec4 < T > textureLoad ( t :texture_2d < T > , coords :vec2 < i32 > , level :i32 ) ->vec4 < T > textureLoad ( t :texture_2d_array < T > , coords :vec2 < i32 > , array_index :i32 , level :i32 ) ->vec4 < T > textureLoad ( t :texture_3d < T > , coords :vec3 < i32 > , level :i32 ) ->vec4 < T > textureLoad ( t :texture_multisampled_2d < T > , coords :vec2 < i32 > , sample_index :i32 ) ->vec4 < T > textureLoad ( t :texture_multisampled_2d_array < T > , coords :vec2 < i32 > , array_index :i32 , sample_index :i32 ) ->vec4 < T > textureLoad ( t :texture_depth_2d , coords :vec2 < i32 > , level :i32 ) ->f32 textureLoad ( t :texture_depth_2d_array , coords :vec2 < i32 > , array_index :i32 , level :i32 ) ->f32 textureLoad ( t :[[ access ( read )]] texture_storage_1d < F > , coords :i32 ) ->vec4 < T > textureLoad ( t :[[ access ( read )]] texture_storage_2d < F > , coords :vec2 < i32 > ) ->vec4 < T > textureLoad ( t :[[ access ( read )]] texture_storage_2d_array < F > , coords :vec2 < i32 > , array_index :i32 ) ->vec4 < T > textureLoad ( t :[[ access ( read )]] texture_storage_3d < F > , coords :vec3 < i32 > ) ->vec4 < T >
For read-only storage textures the returned channel format T
depends on the texel format F
. See the texel format table for the mapping of texel
format to channel format.
Parameters:
t
| The sampled, multisampled, depth or read-only storage texture. |
coords
| The 0-based texel coordinate. |
array_index
| The 0-based texture array index. |
level
| The mip level, with level 0 containing a full size version of the texture. |
sample_index
| The 0-based sample index of the multisampled texture. |
Returns:
If all the parameters are within bounds, the unfiltered texel data.
If any of the parameters are out of bounds, then zero in all components.
15.8.3. textureNumLayers
Returns the number of layers (elements) of an array texture.
textureNumLayers ( t :texture_2d_array < T > ) ->i32 textureNumLayers ( t :texture_cube_array < T > ) ->i32 textureNumLayers ( t :texture_multisampled_2d_array < T > ) ->i32 textureNumLayers ( t :texture_depth_2d_array ) ->i32 textureNumLayers ( t :texture_depth_cube_array ) ->i32 textureNumLayers ( t :texture_storage_2d_array < F > ) ->i32
Parameters:
t
| The sampled, multisampled, depth or storage array texture. |
Returns:
If the number of layers (elements) of the array texture.
15.8.4. textureNumLevels
Returns the number of mip levels of a texture.
textureNumLevels ( t :texture_2d < T > ) ->i32 textureNumLevels ( t :texture_2d_array < T > ) ->i32 textureNumLevels ( t :texture_3d < T > ) ->i32 textureNumLevels ( t :texture_cube < T > ) ->i32 textureNumLevels ( t :texture_cube_array < T > ) ->i32 textureNumLevels ( t :texture_depth_2d ) ->i32 textureNumLevels ( t :texture_depth_2d_array ) ->i32 textureNumLevels ( t :texture_depth_cube ) ->i32 textureNumLevels ( t :texture_depth_cube_array ) ->i32
Parameters:
t
| The sampled or depth texture. |
Returns:
If the number of mip levels for the texture.
15.8.5. textureNumSamples
Returns the number samples per texel in a multisampled texture.
textureNumSamples ( t :texture_multisampled_2d < T > ) ->i32 textureNumSamples ( t :texture_multisampled_2d_array < T > ) ->i32
Parameters:
t
| The multisampled texture. |
Returns:
If the number of samples per texel in the multisampled texture.
15.8.6. textureSample
Samples a texture.
textureSample ( t :texture_1d < f32 > , s :sampler , coords :f32 ) ->vec4 < f32 > textureSample ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > ) ->vec4 < f32 > textureSample ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , offset :vec2 < i32 > ) ->vec4 < f32 > textureSample ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 ) ->vec4 < f32 > textureSample ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , offset :vec2 < i32 > ) ->vec4 < f32 > textureSample ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > ) ->vec4 < f32 > textureSample ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , offset :vec3 < i32 > ) ->vec4 < f32 > textureSample ( t :texture_cube < f32 > , s :sampler , coords :vec3 < f32 > ) ->vec4 < f32 > textureSample ( t :texture_cube_array < f32 > , s :sampler , coords :vec3 < f32 > , array_index :i32 ) ->vec4 < f32 > textureSample ( t :texture_depth_2d , s :sampler , coords :vec2 < f32 > ) ->f32 textureSample ( t :texture_depth_2d , s :sampler , coords :vec2 < f32 > , offset :vec2 < i32 > ) ->f32 textureSample ( t :texture_depth_2d_array , s :sampler , coords :vec2 < f32 > , array_index :i32 ) ->f32 textureSample ( t :texture_depth_2d_array , s :sampler , coords :vec2 < f32 > , array_index :i32 , offset :vec2 < i32 > ) ->f32 textureSample ( t :texture_depth_cube , s :sampler , coords :vec3 < f32 > ) ->f32 textureSample ( t :texture_depth_cube_array , s :sampler , coords :vec3 < f32 > , array_index :i32 ) ->f32
Parameters:
t
| The sampled or depth texture to sample. |
s
| The sampler type. |
coords
| The texture coordinates used for sampling. |
array_index
| The 0-based texture array index to sample. |
offset
| The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes.offset must be compile time constant, and may only be provided as a literal or const_expr expression (e.g. vec2<i32>(1, 2) ).Each offset component must be at least -8 and at most 7 . Values outside
of this range will be treated as a compile time error.
|
Returns:
The sampled value.
15.8.7. textureSampleBias
Samples a texture with a bias to the mip level.
textureSampleBias ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , bias :f32 ) ->vec4 < f32 > textureSampleBias ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , bias :f32 , offset :vec2 < i32 > ) ->vec4 < f32 > textureSampleBias ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , bias :f32 ) ->vec4 < f32 > textureSampleBias ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , bias :f32 , offset :vec2 < i32 > ) ->vec4 < f32 > textureSampleBias ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , bias :f32 ) ->vec4 < f32 > textureSampleBias ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , bias :f32 , offset :vec3 < i32 > ) ->vec4 < f32 > textureSampleBias ( t :texture_cube < f32 > , s :sampler , coords :vec3 < f32 > , bias :f32 ) ->vec4 < f32 > textureSampleBias ( t :texture_cube_array < f32 > , s :sampler , coords :vec3 < f32 > , array_index :i32 , bias :f32 ) ->vec4 < f32 >
Parameters:
t
| The texture to sample. |
s
| The sampler type. |
coords
| The texture coordinates used for sampling. |
array_index
| The 0-based texture array index to sample. |
bias
| The bias to apply to the mip level before sampling. bias must be between -16.0 and 15.99 .
|
offset
| The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes.offset must be compile time constant, and may only be provided as a literal or const_expr expression (e.g. vec2<i32>(1, 2) ).Each offset component must be at least -8 and at most 7 . Values outside
of this range will be treated as a compile time error.
|
Returns:
The sampled value.
15.8.8. textureSampleCompare
Samples a depth texture and compares the sampled depth values against a reference value.
textureSampleCompare ( t :texture_depth_2d , s :sampler_comparison , coords :vec2 < f32 > , depth_ref :f32 ) ->f32 textureSampleCompare ( t :texture_depth_2d , s :sampler_comparison , coords :vec2 < f32 > , depth_ref :f32 , offset :vec2 < i32 > ) ->f32 textureSampleCompare ( t :texture_depth_2d_array , s :sampler_comparison , coords :vec2 < f32 > , array_index :i32 , depth_ref :f32 ) ->f32 textureSampleCompare ( t :texture_depth_2d_array , s :sampler_comparison , coords :vec2 < f32 > , array_index :i32 , depth_ref :f32 , offset :vec2 < i32 > ) ->f32 textureSampleCompare ( t :texture_depth_cube , s :sampler_comparison , coords :vec3 < f32 > , depth_ref :f32 ) ->f32 textureSampleCompare ( t :texture_depth_cube_array , s :sampler_comparison , coords :vec3 < f32 > , array_index :i32 , depth_ref :f32 ) ->f32
Parameters:
t
| The depth texture to sample. |
s
| The sampler comparision type. |
coords
| The texture coordinates used for sampling. |
array_index
| The 0-based texture array index to sample. |
depth_ref
| The reference value to compare the sampled depth value against. |
offset
| The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes.offset must be compile time constant, and may only be provided as a literal or const_expr expression (e.g. vec2<i32>(1, 2) ).Each offset component must be at least -8 and at most 7 . Values outside
of this range will be treated as a compile time error.
|
Returns:
A value in the range [0.0..1.0]
.
Each sampled texel is compared against the reference value using the comparision
operator defined by the sampler_comparison
, resulting in either a 0
or 1
value for each texel.
If the sampler_comparison
uses bilinear filtering then the returned value is
the filtered average of these values, otherwise the comparision result of a
single texel is returned.
15.8.9. textureSampleGrad
Samples a texture using explicit gradients.
textureSampleGrad ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , ddx :vec2 < f32 > , ddy :vec2 < f32 > ) ->vec4 < f32 > textureSampleGrad ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , ddx :vec2 < f32 > , ddy :vec2 < f32 > , offset :vec2 < i32 > ) ->vec4 < f32 > textureSampleGrad ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , ddx :vec2 < f32 > , ddy :vec2 < f32 > ) ->vec4 < f32 > textureSampleGrad ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , ddx :vec2 < f32 > , ddy :vec2 < f32 > , offset :vec2 < i32 > ) ->vec4 < f32 > textureSampleGrad ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , ddx :vec3 < f32 > , ddy :vec3 < f32 > ) ->vec4 < f32 > textureSampleGrad ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , ddx :vec3 < f32 > , ddy :vec3 < f32 > , offset :vec3 < i32 > ) ->vec4 < f32 > textureSampleGrad ( t :texture_cube < f32 > , s :sampler , coords :vec3 < f32 > , ddx :vec3 < f32 > , ddy :vec3 < f32 > ) ->vec4 < f32 > textureSampleGrad ( t :texture_cube_array < f32 > , s :sampler , coords :vec3 < f32 > , array_index :i32 , ddx :vec3 < f32 > , ddy :vec3 < f32 > ) ->vec4 < f32 >
Parameters:
t
| The texture to sample. |
s
| The sampler type. |
coords
| The texture coordinates used for sampling. |
array_index
| The 0-based texture array index to sample. |
ddx
| The x direction derivative vector used to compute the sampling locations. |
ddy
| The y direction derivative vector used to compute the sampling locations. |
offset
| The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes.offset must be compile time constant, and may only be provided as a literal or const_expr expression (e.g. vec2<i32>(1, 2) ).Each offset component must be at least -8 and at most 7 . Values outside
of this range will be treated as a compile time error.
|
Returns:
The sampled value.
15.8.10. textureSampleLevel
Samples a texture using an explicit mip level.
textureSampleLevel ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , level :f32 ) ->vec4 < f32 > textureSampleLevel ( t :texture_2d < f32 > , s :sampler , coords :vec2 < f32 > , level :f32 , offset :vec2 < i32 > ) ->vec4 < f32 > textureSampleLevel ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , level :f32 ) ->vec4 < f32 > textureSampleLevel ( t :texture_2d_array < f32 > , s :sampler , coords :vec2 < f32 > , array_index :i32 , level :f32 , offset :vec2 < i32 > ) ->vec4 < f32 > textureSampleLevel ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , level :f32 ) ->vec4 < f32 > textureSampleLevel ( t :texture_3d < f32 > , s :sampler , coords :vec3 < f32 > , level :f32 , offset :vec3 < i32 > ) ->vec4 < f32 > textureSampleLevel ( t :texture_cube < f32 > , s :sampler , coords :vec3 < f32 > , level :f32 ) ->vec4 < f32 > textureSampleLevel ( t :texture_cube_array < f32 > , s :sampler , coords :vec3 < f32 > , array_index :i32 , level :f32 ) ->vec4 < f32 > textureSampleLevel ( t :texture_depth_2d , s :sampler , coords :vec2 < f32 > , level :i32 ) ->f32 textureSampleLevel ( t :texture_depth_2d , s :sampler , coords :vec2 < f32 > , level :i32 , offset :vec2 < i32 > ) ->f32 textureSampleLevel ( t :texture_depth_2d_array , s :sampler , coords :vec2 < f32 > , array_index :i32 , level :i32 ) ->f32 textureSampleLevel ( t :texture_depth_2d_array , s :sampler , coords :vec2 < f32 > , array_index :i32 , level :i32 , offset :vec2 < i32 > ) ->f32 textureSampleLevel ( t :texture_depth_cube , s :sampler , coords :vec3 < f32 > , level :i32 ) ->f32 textureSampleLevel ( t :texture_depth_cube_array , s :sampler , coords :vec3 < f32 > , array_index :i32 , level :i32 ) ->f32
Parameters:
t
| The sampled or depth texture to sample. |
s
| The sampler type. |
coords
| The texture coordinates used for sampling. |
array_index
| The 0-based texture array index to sample. |
level
| The mip level, with level 0 containing a full size version of the texture.
For the functions where level is a f32 , fractional values may interpolate
between two levels if the format is filterable according to the Texture Format Capabilities.
|
offset
| The optional texel offset applied to the unnormalized texture coordinate
before sampling the texture. This offset is applied before applying any
texture wrapping modes.offset must be compile time constant, and may only be provided as a literal or const_expr expression (e.g. vec2<i32>(1, 2) ).Each offset component must be at least -8 and at most 7 . Values outside
of this range will be treated as a compile time error.
|
Returns:
The sampled value.
15.8.11. textureStore
Writes a single texel to a texture.
textureStore ( t :[[ access ( write )]] texture_storage_1d < F > , coords :i32 , value :vec4 < T > ) ->void textureStore ( t :[[ access ( write )]] texture_storage_2d < F > , coords :vec2 < i32 > , value :vec4 < T > ) ->void textureStore ( t :[[ access ( write )]] texture_storage_2d_array < F > , coords :vec2 < i32 > , array_index :i32 , value :vec4 < T > ) ->void textureStore ( t :[[ access ( write )]] texture_storage_3d < F > , coords :vec3 < i32 > , value :vec4 < T > ) ->void
The channel format T
depends on the storage texel format F
. See the texel format table for the mapping of texel
format to channel format.
Parameters:
t
| The write-only storage texture. |
coords
| The 0-based texel coordinate. |
array_index
| The 0-based texture array index. |
value
| The new texel value. |
Note:
If any of the parameters are out of bounds, then the call to textureStore()
does nothing.
TODO:
TODO(dsinclair): Need gather operations
15.9. Atomic built-in functions
15.10. Data packing built-in functions
Data packing builtin functions can be used to encode values using data formats that do not correspond directly to types in WGSL. This enables a program to write many densely packed values to memory, which can reduce a shader’s memory bandwidth demand.
Built-in | Description |
---|---|
pack4x8snorm (e: vec4<f32>) -> u32
| Converts four normalized floating point values to 8-bit signed integers, and then combines them
into one u32 value.Component e[i] of the input is converted to an 8-bit twos complement integer value ⌊ 0.5 + 127 × min(1, max(-1, e[i])) ⌋ which is then placed in bits 8 × i through 8 × i + 7 of the result. |
pack4x8unorm (e: vec4<f32>) -> u32
| Converts four normalized floating point values to 8-bit unsigned integers, and then combines them
into one u32 value.Component e[i] of the input is converted to an 8-bit unsigned integer value ⌊ 0.5 + 255 × min(1, max(0, e[i])) ⌋ which is then placed in bits 8 × i through 8 × i + 7 of the result. |
pack2x16snorm (e: vec2<f32>) -> u32
| Converts two normalized floating point values to 16-bit signed integers, and then combines them
into one u32 value.Component e[i] of the input is converted to a 16-bit twos complement integer value ⌊ 0.5 + 32767 × min(1, max(-1, e[i])) ⌋ which is then placed in bits 16 × i through 16 × i + 15 of the result. |
pack2x16unorm (e: vec2<f32>) -> u32
| Converts two normalized floating point values to 16-bit unsigned integers, and then combines them
into one u32 value.Component e[i] of the input is converted to a 16-bit unsigned integer value ⌊ 0.5 + 65535 × min(1, max(0, e[i])) ⌋ which is then placed in bits 16 × i through 16 × i + 15 of the result. |
pack2x16float (e: vec2<f32>) -> u32
| Converts two floating point values to half-precision floating point numbers, and then combines
them into one one u32 value.Component e[i] of the input is converted to a IEEE 754 binary16 value, which is then placed in bits 16 × i through 16 × i + 15 of the result. See § 10.5.1 Floating point conversion for edge case behaviour. |
15.11. Data unpacking built-in functions
Data unpacking builtin functions can be used to decode values in data formats that do not correspond directly to types in WGSL. This enables a program to read many densely packed values from memory, which can reduce a shader’s memory bandwidth demand.
Built-in | Description |
---|---|
unpack4x8snorm (e: u32) -> vec4<f32>
| Decomposes a 32-bit value into four 8-bit chunks, then reinterprets
each chunk as a signed normalized floating point value. Component i of the result is max(v ÷ 127, -1), where v is the interpretation of bits 8×i through 8×i+7 of e as a twos-complement signed integer. |
unpack4x8unorm (e: u32) -> vec4<f32>
| Decomposes a 32-bit value into four 8-bit chunks, then reinterprets
each chunk as an unsigned normalized floating point value. Component i of the result is v ÷ 255, where v is the interpretation of bits 8×i through 8×i+7 of e as an unsigned integer. |
unpack2x16snorm (e: u32) -> vec2<f32>
| Decomposes a 32-bit value into two 16-bit chunks, then reinterprets
each chunk as a signed normalized floating point value. Component i of the result is max(v ÷ 32767, -1), where v is the interpretation of bits 16×i through 16×i+15 of e as a twos-complement signed integer. |
unpack2x16unorm (e: u32) -> vec2<f32>
| Decomposes a 32-bit value into two 16-bit chunks, then reinterprets
each chunk as an unsigned normalized floating point value. Component i of the result is v ÷ 65535, where v is the interpretation of bits 16×i through 16×i+15 of e as an unsigned integer. |
unpack2x16float (e: u32) -> vec2<f32>
| Decomposes a 32-bit value into two 16-bit chunks, and reinterpets each chunk
as a floating point value. Component i of the result is the f32 representation of v, where v is the interpretation of bits 16×i through 16×i+15 of e as an IEEE 754 binary16 value. See § 10.5.1 Floating point conversion for edge case behaviour. |
15.12. Subgroup built-in functions
Subgroup built-in functions | SPIR-V |
subgroupIsFirst() -> bool | OpGroupNonUniformElect |
subgroupAll(bool) -> bool | OpGroupNonUniformAll |
subgroupAny(bool) -> bool | OpGroupNonUniformAny |
subgroupBallot(bool) -> vec4<u32> | OpGroupNonUniformBallot |
subgroupBroadcastFirst(Integral) -> Integral | OpGroupNonUniformBroadcastFirst |
subgroupBroadcastFirst(Floating) -> Floating | OpGroupNonUniformBroadcastFirst |
subgroupAdd(Integral) -> Integral | OpGroupNonUniformIAdd with Reduce |
subgroupAdd(Floating) -> Floating | OpGroupNonUniformFAdd with Reduce |
subgroupMul(Integral) -> Integral | OpGroupNonUniformIMul with Reduce |
subgroupMul(Floating) -> Floating | OpGroupNonUniformFMul with Reduce |
subgroupMin(Integral) -> Integral | OpGroupNonUniformUMin or OpGroupNonUniformSMin with Reduce |
subgroupMin(Floating) -> Floating | OpGroupNonUniformFMin with Reduce |
subgroupMax(Integral) -> Integral | OpGroupNonUniformUMax or OpGroupNonUniformSMax with Reduce |
subgroupMax(Floating) -> Floating | OpGroupNonUniformFMax with Reduce |
subgroupAnd(Integral) -> Integral | OpGroupNonUniformBitwiseAnd |
subgroupOr(Integral) -> Integral | OpGroupNonUniformBitwiseOr |
subgroupXor(Integral) -> Integral | OpGroupNonUniformBitwiseXor |
subgroupPrefixAdd(Integral) -> Integral | OpGroupNonUniformIAdd with ExclusiveScan |
subgroupPrefixAdd(Floating) -> Floating | OpGroupNonUniformFAdd with ExclusiveScan |
subgroupPrefixMul(Integral) -> Integral | OpGroupNonUniformIMul with ExclusiveScan |
subgroupPrefixMul(Floating) -> Floating | OpGroupNonUniformFMul with ExclusiveScan |
Note: Subgroup built-in functions exist if "subgroup-operations" is enabled in requestDevice.
16. Glossary
TODO: Remove terms unused in the rest of the specification.
Term | Definition |
---|---|
Dominates |
Basic block A dominates basic block B if:
|
Strictly dominates | A strictly dominates B if A dominates B and A != B
|
DomBy(A) | The basic blocks dominated by A
|
17. MATERIAL TO BE MOVED TO A NEW HOME OR DELETED
17.1. Composite types
A type is composite if its values have a well-defined internal structure of typed components.
The following types are composite types:
WGSL has operations for:
-
extracting one of the components of a composite value
-
creating a new composite value from an old one by replacing one of its components
-
creating a new composite value from components
17.2. Type Promotions
There are no implicit type promotions in WGSL. If you want to convert between types you must use the cast syntax to do it.The non-promotion extends to vector classes as well. There are no overrides to
shorten vector declarations based on the type or number of elements provided.
If you want vec4<f32>
you must provide 4 float values in the constructor.
17.3. Precedence
(dsinclair) Write out precedence rules. Matches c and glsl rules ....