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

MATERIAL TO BE MOVED TO A NEW HOME OR DELETED

1. Introduction

[[ location ( 0 )]] var < out > gl_FragColor : vec4 < f32 > ; [[ stage ( fragment )]] fn main () -> void { gl_FragColor = vec4 < f32 > ( 0.4 , 0.4 , 0.8 , 1.0 ); return ; }

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

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.

Comments begin with a # 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 FLOAT_LITERAL (-?[0-9]*.[0-9]+ | -?[0-9]+.[0-9]*)((e|E)(+|-)?[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 STRING_LITERAL "[^"]*"

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.

const_literal : INT_LITERAL | UINT_LITERAL | FLOAT_LITERAL | TRUE | FALSE

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 .

(dneto) also lifetime.

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.

(dneto) complete

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:

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

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 type is a numeric vector type if its component type is a numeric scalar.

EXAMPLE: Vector vec2 < f32 > # is a vector of two f32s .

3.3.6. Matrix Types

Type Description matNxM<T> Matrix of N columns and M rows, where N and M are both in {2, 3, 4}. T must be f32.

EXAMPLE: Matrix mat2x3 < f32 > # is a 2 column , 3 row matrix of 32 - bit floats .

3.3.7. Array Types

Type Description array<E,N> An N-element array of elements of type E.

array<E> A runtime-sized array of elements of type E, also known as a runtime array. These may only appear in specific contexts.



(dneto): Complete description of Array<E,N>

(dneto): Runtime-sized array is only usable as the last element of a struct defining the contents of a storage buffer.

3.3.8. Structure Types

Type Description struct<T1,...,Tn> An ordered tuple of N members of types T1 through Tn, with N being an integer greater than 0.

EXAMPLE: Structure struct Data { a : i32 ; b : vec2 < f32 > ; }

struct_decl : struct_decoration_decl* STRUCT IDENT struct_body_decl struct_decoration_decl : ATTR_LEFT struct_decoration ATTR_RIGHT struct_decoration : BLOCK struct_body_decl : BRACE_LEFT struct_member* BRACE_RIGHT struct_member : struct_member_decoration_decl+ variable_ident_decl SEMICOLON struct_member_decoration_decl : | ATTR_LEFT (struct_member_decoration COMMA)* struct_member_decoration ATTR_RIGHT struct_member_decoration : OFFSET PAREN_LEFT INT_LITERAL PAREN_RIGHT

Note: Layout decorations are required if the struct is used in an SSBO, UBO or Push Constant. Otherwise, the layout will be ignored.

(dneto): MatrixStride, RowMajor, ColMajor layout decorations are needed for matrices.

EXAMPLE: Structure 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 type RTArr = [[stride 16]] array<vec4<f32>>; [[block]] struct S { [[offset(0)]] a : f32; [[offset(4)]] b : f32; [[offset(16)]] data : RTArr; };

3.4. Memory TODO

TODO: This section is a stub.

In WGSL, a value of § 3.4.2 Storable Types may be stored in memory, for later retrieval.

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.4.3. Host-shareable Types

The following types are host-shareable :

TODO: This is a stub: Collectively, the stride and offset attributes are called layout attributes.

3.4.4. Storage Classes TODO

storage_class : INPUT | OUTPUT | UNIFORM | WORKGROUP | UNIFORM_CONSTANT | STORAGE_BUFFER | IMAGE | PRIVATE | FUNCTION

Name SPIR-V Storage Class input Input output Output uniform Uniform workgroup Workgroup uniform_constant UniformConstant storage_buffer StorageBuffer image Image private Private function Function

3.4.5. Memory Layout Rules TODO

TODO: The following is a stub

Variables in certain storage classes must have host-shareable store type with fully elaborated memory layout.

The memory layout of a type is significant only when referring to a value in those storage classes. This affects evaluation of a variable in one of those storage classes, or a pointer into one of those storage classes.

3.5. Pointer Types TODO

Type Description ptr<SC,T> Pointer (or reference) to storage in § 3.4.4 Storage Classes TODO SC which can hold a value of the § 3.4.2 Storable Types 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.

EXAMPLE: Pointer ptr < storage_buffer , i32 > ptr < private , array < i32 , 12 >>

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 E.g. a[12] If a is a pointer to an array, this evaluates to a.Subaccess(12)

E.g. s.foo If s is a pointer to a structure of type S, k is the index of the foo element of S, this evaluates to s.Subaccess(k)

Assigning (L-Value) On the left hand side of an assignment operation, and the right hand side matches the pointee type of the pointer. E.g. v = 12; assuming prior declaration var v : i32 Copying On the right hand side of a const-declaration, and the type of the const-declaration matches the pointer type. E.g. const v2 : ptr<private,i32> = v; assuming prior declaration var<private> v:i32 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 Types TODO

3.6.1. Sampled Texture Types

texture_sampled_1d<type> %1 = OpTypeImage %type 1D 0 0 0 1 Unknown texture_sampled_1d_array<type> %1 = OpTypeImage %type 1D 0 1 0 1 Unknown texture_sampled_2d<type> %1 = OpTypeImage %type 2D 0 0 0 1 Unknown texture_sampled_2d_array<type> %1 = OpTypeImage %type 2D 0 1 0 1 Unknown texture_sampled_3d<type> %1 = OpTypeImage %type 3D 0 0 0 1 Unknown texture_sampled_cube<type> %1 = OpTypeImage %type Cube 0 0 0 1 Unknown texture_sampled_cube_array<type> %1 = OpTypeImage %type Cube 0 1 0 1 Unknown

type must be f32 , i32 or u32

The parameterized type for the images is the type after conversion from sampling. E.g. you can have an image with texels with 8bit unorm components, but when you sample them you get a 32-bit float result (or vec-of-f32).

3.6.2. Multisampled Texture Types

texture_multisampled_2d<type> %1 = OpTypeImage %type 2D 0 0 1 1 Unknown

type must be f32 , i32 or u32

3.6.3. Read-only Storage Texture Types

texture_ro_1d<type, image_storage_type> %1 = OpTypeImage %type 1D 0 0 0 2 image_storage_type ReadOnly texture_ro_1d_array<type, image_storage_type> %1 = OpTypeImage %type 1D 0 1 0 2 image_storage_type ReadOnly texture_ro_2d<type, image_storage_type> %1 = OpTypeImage %type 2D 0 0 0 2 image_storage_type ReadOnly texture_ro_2d_array<type, image_storage_type> %1 = OpTypeImage %type 2D 0 1 0 2 image_storage_type ReadOnly texture_ro_3d<type, image_storage_type> %1 = OpTypeImage %type 3D 0 0 0 2 image_storage_type ReadOnly

type must be f32 , i32 or u32

The parameterized type for the images is the type after conversion from reading. E.g. you can have an image with texels with 8bit unorm components, but when you read them you get a 32-bit float result (or vec-of-f32).

3.6.4. Write-only Storage Texture Types

texture_wo_1d<image_storage_type> %1 = OpTypeImage %void 1D 0 0 0 2 image_storage_type WriteOnly texture_wo_1d_array<image_storage_type> %1 = OpTypeImage %void 1D 0 1 0 2 image_storage_type WriteOnly texture_wo_2d<image_storage_type> %1 = OpTypeImage %void 2D 0 0 0 2 image_storage_type WriteOnly texture_wo_2d_array<image_storage_type> %1 = OpTypeImage %void 2D 0 1 0 2 image_storage_type WriteOnly texture_wo_3d<image_storage_type> %1 = OpTypeImage %void 3D 0 0 0 2 image_storage_type WriteOnly

3.6.5. Depth Texture Types

texture_depth_2d %1 = OpTypeImage %f32 2D 1 0 0 1 Unknown texture_depth_2d_array %1 = OpTypeImage %f32 2D 1 1 0 1 Unknown texture_depth_cube %1 = OpTypeImage %f32 Cube 1 0 0 1 Unknown texture_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

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 image_storage_type GREATER_THAN sampler_type : SAMPLER | SAMPLER_COMPARISON sampled_texture_type : TEXTURE_SAMPLED_1D | TEXTURE_SAMPLED_1D_ARRAY | TEXTURE_SAMPLED_2D | TEXTURE_SAMPLED_2D_ARRAY | TEXTURE_SAMPLED_3D | TEXTURE_SAMPLED_CUBE | TEXTURE_SAMPLED_CUBE_ARRAY multisampled_texture_type : TEXTURE_MULTISAMPLED_2D storage_texture_type : TEXTURE_RO_1D | TEXTURE_RO_1D_ARRAY | TEXTURE_RO_2D | TEXTURE_RO_2D_ARRAY | TEXTURE_RO_3D | TEXTURE_WO_1D | TEXTURE_WO_1D_ARRAY | TEXTURE_WO_2D | TEXTURE_WO_2D_ARRAY | TEXTURE_WO_3D depth_texture_type : TEXTURE_DEPTH_2D | TEXTURE_DEPTH_2D_ARRAY | TEXTURE_DEPTH_CUBE | TEXTURE_DEPTH_CUBE_ARRAY image_storage_type : 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

EXAMPLE: Type Alias type Arr = array<i32, 5>; type RTArr = [[stride(16)]] array<vec4<f32>>;

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 | array_decoration_list* ARRAY LESS_THAN type_decl COMMA INT_LITERAL GREATER_THAN | array_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_list : ATTR_LEFT (array_decoration COMMA)* array_decoration ATTR_RIGHT array_decoration : STRIDE PAREN_LEFT INT_LITERAL PAREN_RIGHT

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

4. Variable and const

TODO: Stub (describe what a constant is): A constant is a name for a value, declared via a const declaration.

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.

A variable declaration :

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 , Function , or Output § 3.4.4 Storage Classes TODO. If present, the intiailizer’s type must match the store type of the variable.

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 type_decl variable_storage_decoration : LESS_THAN storage_class GREATER_THAN

Two variables with overlapping lifetimes must not have overlapping storage.

When a variable is created, its storage contains an initial value as follows:

For variables in the Private , Function , or Output 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.

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

var x : f32 = 1.0 ; const y = x * x + x + 1 ; Consider the following snippet of WGSL: Because x 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 However, it is expected that either the browser or the driver optimizes this intermediate representation such that the redundant loads are eliminated.

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.

global_variable_decl : variable_decoration_list* variable_decl | variable_decoration_list* sampler_or_texture_decl | variable_decoration_list* variable_decl EQUAL const_expr variable_decoration_list : ATTR_LEFT (variable_decoration COMMA)* variable_decoration ATTR_RIGHT variable_decoration : LOCATION PAREN_LEFT INT_LITERAL PAREN_RIGHT | BUILTIN PAREN_LEFT IDENT PAREN_RIGHT | BINDING PAREN_LEFT INT_LITERAL PAREN_RIGHT | SET PAREN_LEFT INT_LITERAL PAREN_RIGHT

EXAMPLE: Variable Decorations [[location(2)]] OpDecorate %gl_FragColor Location 2 [[binding(3), set(4)]] OpDecorate %gl_FragColor Binding 3 OpDecorate %gl_FragColor DescriptorSet 4

See § 14 Built-in variables TODO for the decorations for specifying built-in variables.

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.

EXAMPLE: Module constants 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 is the application supplies a constant ID that is not in the program? Proposal: pipeline creation fails with an error.

EXAMPLE: Module constants, pipeline-overrideable [[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 : global_const_decoration_list* CONST variable_ident_decl global_const_initializer? global_const_decoration_list : ATTR_LEFT global_const_decoration ATTR_RIGHT global_const_decoration : CONSTANT_ID PAREN_LEFT INT_LITERAL PAREN_RIGHT global_const_initializer : EQUAL const_expr const_expr : type_decl PAREN_LEFT (const_expr COMMA)* const_expr PAREN_RIGHT | const_literal

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

A variable or constant declared in a declaration statement in a function body is in function scope. The name is available for use immedately after its declaration statement, and until the end of the brace-delimited list of statements immediately enclosing the declaration.

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

Scalar literal type rules 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

Scalar constructor type rules Precondition Conclusion Notes e : bool bool(e) : bool Pass-through (OpCopyObject) e : i32 i32(e) : i32 Pass-through (OpCopyObject) e : u32 u32(e) : u32 Pass-through (OpCopyObject) e : f32 f32(e) : f32 Pass-through (OpCopyObject)

Vector constructor type rules, where T is a scalar type Precondition Conclusion Notes e1 : T

e2 : T vec2<T>(e1,e2) : vec2<T> OpCompositeConstruct e1 : T

e2 : T

e3 : T vec3<T>(e1,e2,e3) : vec3<T> OpCompositeConstruct e1 : T

e2 : T

e3 : T

e4 : T vec4<T>(e1,e2,e3,e4) : vec4<T> OpCompositeConstruct e1 : T

e2 : vec2<T> vec3<T>(e1,e2) : vec3<T>

vec3<T>(e2,e1) : vec3<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

Matrix constructor type rules 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

Array constructor type rules Precondition Conclusion Notes e1 : T

...

eN : T

array< T,N >(e1,...,eN) : array<T, N> Construction of an array from elements

Structure constructor type rules 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() is false

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.

Scalar zero value type rules 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)

Vector zero type rules, where T is a scalar type Precondition Conclusion Notes vec2<T>() : vec2<T> Zero value (OpConstantNull) vec3<T>() : vec3<T> Zero value (OpConstantNull) vec4<T>() : vec4<T> Zero value (OpConstantNull)

EXAMPLE: Zero-valued vectors 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 .

Matrix zero type rules 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)

Array zero type rules Precondition Conclusion Notes T is storable array< T,N >() : array<T, N> Zero-valued array (OpConstantNull)

EXAMPLE: Zero-valued arrays array < bool , 2 > () # The zero - valued array of two booleans . array < bool , 2 > ( false , false ) # The same value , written explicitly .

Structure zero type rules Precondition Conclusion Notes S is a 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)

EXAMPLE: Zero-valued structures struct Student { grade : i32 ; GPA : f32 ; attendance : array < bool , 4 > ; }; # The zero value for Student Student () # The same value , written explicitly . Student ( 0 , 0.0 , array < bool , 4 > ( false , false , false , false )) # The same value , written with zero - valued members . Student ( i32 (), f32 (), array < bool , 4 > ())

5.4. Conversion Expressions

Scalar conversion type rules Precondition Conclusion Notes 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.

Vector conversion type rules Precondition Conclusion Notes e : vec N <u32> vec N < i32 > ( e ) : vec N <i32> Component-wise reinterpretation of bits.

Component i of the result is i32( e [ i ])

(OpBitcast) e : vec N <f32> vec N < i32 > ( e ) : vec N <i32> Component-wise value conversion to signed integer, including invalid cases.

Component i of the result is i32( e [ i ])

(OpConvertFToS) e : vec N <i32> vec N < u32 > ( e ) : vec N <u32> Component-wise reinterpretation of bits.

Component i of the result is u32( e [ i ])

(OpBitcast) e : vec N <f32> vec N < u32 > ( e ) : vec N <u32> Component-wise value conversion to unsigned integer, including invalid cases.

Component i of the result is u32( e [ i ])

(OpConvertFToU) e : vec N <i32> vec N < f32 > ( e ) : vec N <f32> Component-wise value conversion to floating point, including invalid cases.

Component i of the result is f32( e [ i ])

(OpConvertSToF) e : vec N <u32> vec N < f32 > ( e ) : vec N <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.

Scalar bitcast type rules Precondition Conclusion Notes e : T ,

T is one of i32, u32, f32 bitcast< T >( e ) : T Identity transform.

The result is e . (OpCopyObject) 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)

Vector bitcast type rules Precondition Conclusion Notes e : vec< N > T >,

T is one of i32, u32, f32 bitcast<vec N < T >>( e ) : T Identity transform.

The result is e . (OpCopyObject) e : vec< N > T >,

T is one of u32, f32 bitcast<vec N <i32>>( e ) : vec N <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<vec N <u32>>( e ) : vec N <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<vec N <f32>>( e ) : vec N <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.

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

TODO: Type rules for vector access

5.6.2. Matrix Access Expression TODO

5.6.3. Array Access Expression TODO

5.6.4. Structure Access Expression TODO

5.7. Logical Expressions TODO

Unary logical operations Precondition Conclusion Notes e : bool !e : bool OpLogicalNot

Binary logical expressions Precondition Conclusion Notes e1 : bool

e2 : bool e1 || e2 : bool Short-circuiting "or". Yields true if either e1 or e2 are true; evaluates e2 only if e1 is false. e1 : bool

e2 : bool e1 && e2 : bool Short-circuiting "and". Yields true if both e1 and e2 are true; evaluates e2 only if e1 is true. e1 : 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

Unary arithmetic expressions Precondition Conclusion Notes e : T, T is Integral -e : T OpSNegate e : T, T is Floating -e : T OpFNegate

Binary arithmetic expressions over scalars 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 (OpFAdd) 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)

Binary arithmetic expressions over vectors 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)

Binary arithmetic expressions with mixed scalar, vector, and matrix operands Precondition Conclusion Notes e1 : f32

e2 : T

T is FloatVec e1 * e2 : T

e2 * e1 : T Multiplication of a vector and a scalar (OpVectorTimesScalar) e1 : f32

e2 : T

T is matNxM<f32> e1 * e2 : T

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

Comparisons over scalars 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)

Comparisons over vectors 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

(dneto): Bitwise-complement is under discussion. https://github.com/gpuweb/gpuweb/pull/727

Binary bitwise operations 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

Bit shift expressions 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 : vec N < T >

e2 : vec N <u32>

T is Int e1 << e2 : vec N < 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 : vec N <u32>

e2 : u32

e1 >> e2 : vec N <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 : vec N <i32>

e2 : vec N <u32>

e1 >> e2 : vec N <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 | type_decl argument_expression_list | const_literal | paren_rhs_statement | BITCAST LESS_THAN type_decl GREATER_THAN paren_rhs_statement OpBitcast postfix_expression : | BRACKET_LEFT short_circuit_or_expression BRACKET_RIGHT postfix_expression | argument_expression_list postfix_expression | PERIOD IDENT postfix_expression argument_expression_list : PAREN_LEFT ((short_circuit_or_expression COMMA)* short_circuit_or_expression)? PAREN_RIGHT unary_expression : singular_expression | MINUS unary_expression OpSNegate OpFNegate | BANG 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.

EXAMPLE: GLSL Loop int a = 2; for (int i = 0; i < 4; i++) { a *= 2; }

EXAMPLE: WGSL Loop 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.

EXAMPLE: GLSL Loop with continue int a = 2; const int step = 1; for (int i = 0; i < 4; i += step) { if (i % 2 == 0) continue; a *= 2; }

EXAMPLE: WGSL Loop with continue const 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; }

EXAMPLE: WGSL Loop with continue and continuing const 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.

EXAMPLE: For to Loop transformation 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; 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 empty else clause no elseif clauses The only statement in the else clause of an if that has an empty true-branch clause and no elseif clauses.

That if statement must appear last in the continuing clause.

EXAMPLE: WGSL Valid loop if-break from a continuing clause const 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; } } }

EXAMPLE: WGSL Valid loop if-else-break from a continuing clause const 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; } } }

EXAMPLE: WGSL Invalid breaks from a continuing clause const 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.

EXAMPLE: Invalid continue bypasses declaration 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 of step used in the continuing 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 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.

Functions must end with a return statement. The return may be given with a value to be returned.

Function names must be unique over all functions and all variables in the module.

function_decl : function_decoration_decl? function_header body_statement function_decoration_decl : ATTR_LEFT (function_decoration COMMA)* function_decoration ATTR_RIGHT 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

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

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:

compute

vertex

fragment

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.

An entry point function must have no parameters, and its return type must be void.

EXAMPLE: Entry Point [[stage(vertex)]] fn vtx_main() -> void { gl_Position = vec4<f32>(); } # OpEntryPoint Vertex %vtx_main "vtx_main" %gl_Position [[stage(fragment)]] fn frag_main -> void { gl_FragColor = vec4<f32>(); } # OpEntryPoint Fragment %frag_main "frag_main" %gl_FragColor [[stage(compute)]] fn main() -> void { } # OpEntryPoint GLCompute %main "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

function_decoration : STAGE PAREN_LEFT pipeline_stage PAREN_RIGHT | WORKGROUP_SIZE PAREN_LEFT INT_LITERAL ( COMMA INT_LITERAL ( COMMA INT_LITERAL )? )? PAREN_RIGHT pipeline_stage : COMPUTE | VERTEX | FRAGMENT

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?

EXAMPLE: workgroup_size Attribute [[ 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 TODO

8.3.1. Built-in variables TODO

8.3.2. Pipeline Input and Output Interface TODO

TODO(dneto): The following sentence was moved from elsewhere. Expand this.

The input and output parameters to the entry point are determined by which global variables are used in the function and any called functions.

8.3.2.1. Built-in inputs and outputs TODO

8.3.2.2. User Data TODO

8.3.2.3. Input-output Locations TODO

8.3.3. Resource interface TODO

8.4. Pipeline compatibility TODO

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 varibales 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 a thread) 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.

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

( ⌊ 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 Requireing 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 adjact 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, 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.



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?

11. Memory Model TODO

12. Keyword and Token Summary

12.1. Keyword Summary

Type-defining keywords 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_SAMPLED_1D texture_sampled_1d TEXTURE_SAMPLED_1D_ARRAY texture_sampled_1d_array TEXTURE_SAMPLED_2D texture_sampled_2d TEXTURE_SAMPLED_2D_ARRAY texture_sampled_2d_array TEXTURE_SAMPLED_3D texture_sampled_3d TEXTURE_SAMPLED_CUBE texture_sampled_cube TEXTURE_SAMPLED_CUBE_ARRAY texture_sampled_cube_array TEXTURE_MULTISAMPLED_2D texture_multisampled_2d TEXTURE_RO_1D texture_ro_1d TEXTURE_RO_1D_ARRAY texture_ro_1d_array TEXTURE_RO_2D texture_ro_2d TEXTURE_RO_2D_ARRAY texture_ro_2d_array TEXTURE_RO_3D texture_ro_3d TEXTURE_WO_1D texture_wo_1d TEXTURE_WO_1D_ARRAY texture_wo_1d_array TEXTURE_WO_2D texture_wo_2d TEXTURE_WO_2D_ARRAY texture_wo_2d_array TEXTURE_WO_3D texture_wo_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

Other keywords Token Definition AS as BINDING binding BITCAST bitcast BLOCK block BREAK break BUILTIN builtin CASE case COMPUTE compute CONST const CONSTANT_ID constant_id CONTINUE continue CONTINUING continuing DEFAULT default DISCARD discard ELSE else ELSE_IF elseif FALLTHROUGH fallthrough FALSE false FN fn FOR for FRAGMENT fragment FUNCTION function IF if IMAGE image IN in INPUT input LOCATION location LOOP loop OFFSET offset OUT out OUTPUT output PRIVATE private RETURN return SET set STAGE stage STORAGE_BUFFER storage_buffer STRIDE stride SWITCH switch TRUE true TYPE type UNIFORM uniform UNIFORM_CONSTANT uniform_constant VAR var VERTEX vertex WORKGROUP workgroup WORKGROUP_SIZE workgroup_size

Image format keywords 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 bgraunorm 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

12.2. Reserved Keywords

asm bf16 do enum f16 f64 i8 i16 i64 let typedef u8 u16 u64 unless using while regardless premerge

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 * 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-0002: 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-0005: Functions must be declared before use.

v-0006: Variables must be defined before use.

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 only permitted in loop

v-0011: Global variable names must be unique

v-0012: Structure names must be unique

v-0013: Variables declared in a function must be unique between that function and any global variables.

v-0014: Variables declared in a function must have unique names

v-0015: Runtime arrays may only appear as the last member of a struct

v-0016: Function names must be unique

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-0019: Functions used in entry points must exist

v-0020: The pair of <entry point name, pipeline stage> must be unique in the module

v-0021: Can not 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

14. Built-in variables TODO

EXAMPLE: Valid Built-in Decoration Identifiers [[builtin(position)]] OpDecorate %gl_Position BuiltIn Position [[builtin(vertex_idx)]] OpDecorate %gl_VertexIdx BuiltIn VertexIndex [[builtin(instance_idx)]] OpDecorate %gl_InstanceId BuiltIn InstanceIndex [[builtin(front_facing)]] OpDecorate %gl_FrontFacing BuiltIn FrontFacing [[builtin(frag_coord)]] OpDecorate %gl_FragCoord BuiltIn FragCoord [[builtin(frag_depth)]] OpDecorate %gl_FragDepth BuiltIn FragDepth [[builtin(local_invocation_id)]] OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId [[builtin(local_invocation_idx)]] OpDecorate %gl_LocalInvocationIndex BuiltIn LocalInvocationIndex [[builtin(global_invocation_id)]] OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId

The usages of the variable builtin decorations is further restricted in the type, function decorations and storage class.

TODO: list storage class and shader stage restrictions.

Name Type Restrictions position vec4<f32> Vertex Output vertex_idx i32 Vertex Input instance_idx i32 Vertex Input front_facing bool Fragment Input frag_coord vec4<f32> Fragment Input frag_depth f32 Fragment Output local_invocation_id vec3<u32> Compute Input global_invocation_id vec3<u32> Compute Input local_invocation_idx u32 Compute Input

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

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

15.3. Float built-in functions

Precondition Built-in Description T is f32 abs( e : T ) -> T (GLSLstd450FAbs) T is f32 abs( e : vec N < T > ) -> vec N < T > (GLSLstd450FAbs) T is f32 acos( e : T ) -> T (GLSLstd450Acos) T is f32 acos( e : vec N < T > ) -> vec N < T > (GLSLstd450Acos) T is f32 asin( e : T ) -> T (GLSLstd450Asin) T is f32 asin( e : vec N < T > ) -> vec N < T > (GLSLstd450Asin) T is f32 atan( e : T ) -> T (GLSLstd450Atan) T is f32 atan( e : vec N < T > ) -> vec N < T > (GLSLstd450Atan) T is f32 atan2( e1 : T , e2 : T ) -> T (GLSLstd450Atan2) T is f32 atan2( e1 : vec N < T > , e2 : vec N < T > ) -> vec N < T > (GLSLstd450Atan2) T is f32 ceil( e : T ) -> T (GLSLstd450Ceil) T is f32 ceil( e : vec N < T > ) -> vec N < T > (GLSLstd450Ceil) T is f32 clamp( e1 : T , e2 : T , e3 : T ) -> T (GLSLstd450NClamp) T is f32 clamp( e1 : vec N < T > , e2 : vec N < T > , e3 : vec N < T > ) -> vec N < T > (GLSLstd450NClamp) T is f32 cos( e : T ) -> T (GLSLstd450Cos) T is f32 cos( e : vec N < T > ) -> vec N < T > (GLSLstd450Cos) T is f32 cosh( e : T ) -> T (GLSLstd450Cosh) T is f32 cosh( e : vec N < T > ) -> vec N < 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 : vec N < T > , e2 : vec N < T > ) -> vec N < T > (GLSLstd450Distance) T is f32 exp( e : T ) -> T (GLSLstd450Exp) T is f32 exp( e : vec N < T > ) -> vec N < T > (GLSLstd450Exp) T is f32 exp2( e : T ) -> T (GLSLstd450Exp2) T is f32 exp2( e : vec N < T > ) -> vec N < T > (GLSLstd450Exp2) T is f32 faceForward( e1 : T , e2 : T , e3 : T ) -> T (GLSLstd450FaceForward) T is f32 faceForward( e1 : vec N < T > , e2 : vec N < T > , e3 : vec N < T > ) -> vec N < T > (GLSLstd450FaceForward) T is f32 floor( e : T ) -> T (GLSLstd450Floor) T is f32 floor( e : vec N < T > ) -> vec N < T > (GLSLstd450Floor) T is f32 fma( e1 : T , e2 : T , e3 : T ) -> T (GLSLstd450Fma) T is f32 fma( e1 : vec N < T > , e2 : vec N < T > , e3 : vec N < T > ) -> vec N < T > (GLSLstd450Fma) T is f32 fract( e : T ) -> T (GLSLstd450Fract) T is f32 fract( e : vec N < T > ) -> vec N < 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 : vec N < T > , e2 : ptr<vec N < I >> ) -> vec N < T > (GLSLstd450Frexp) T is f32 inverseSqrt( e : T ) -> T (GLSLstd450InverseSqrt) T is f32 inverseSqrt( e : vec N < T > ) -> vec N < T > (GLSLstd450InverseSqrt) T is f32

I is i32 or u32 ldexp( e1 : T , e2 : ptr< I > ) -> T (GLSLstd450Ldexp) T is f32

I is i32 or u32 ldexp( e1 : vec N < T > , e2 : ptr<vec N < I >> ) -> vec N < T > (GLSLstd450Ldexp) T is f32 length( e : T ) -> T (GLSLstd450Length) T is f32 length( e : vec N < T > ) -> vec N < T > (GLSLstd450Length) T is f32 log( e : T ) -> T (GLSLstd450Log) T is f32 log( e : vec N < T > ) -> vec N < T > (GLSLstd450Log) T is f32 log2( e : T ) -> T (GLSLstd450Log2) T is f32 log2( e : vec N < T > ) -> vec N < T > (GLSLstd450Log2) T is f32 max( e1 : T , e2 : T ) -> T (GLSLstd450NMax) T is f32 max( e1 : vec N < T > , e2 : vec N < T > ) -> vec N < T > (GLSLstd450NMax) T is f32 min( e1 : T , e2 : T ) -> T (GLSLstd450NMin) T is f32 min( e1 : vec N < T > , e2 : vec N < T > ) -> vec N < T > (GLSLstd450NMin) T is f32 mix( e1 : T , e2 : T , e3 : T ) -> T (GLSLstd450FMix) T is f32 mix( e1 : vec N < T > , e2 : vec N < T > , e3 : vec N < T > ) -> vec N < T > (GLSLstd450Modf) T is f32

modf( e1 : T , e2 : ptr< T > ) -> T (GLSLstd450Modf) T is f32 modf( e1 : vec N < T > , e2 : ptr<vec N < T >> ) -> vec N < T > (GLSLstd450Modf) T is f32 normalize( e : vec N < T > ) -> vec N < T > (GLSLstd450Normalize) T is f32 pow( e1 : T , e2 : T ) -> T (GLSLstd450Pow) T is f32 pow( e1 : vec N < T > , e2 : vec N < T > ) -> vec N < T > (GLSLstd450Pow) T is f32 reflect( e1 : T , e2 : T ) -> T (GLSLstd450Reflect) T is f32 reflect( e1 : vec N < T > , e2 : vec N < T > ) -> vec N < T > (GLSLstd450Reflect) T is f32 round( e : T ) -> T (GLSLstd450Round) T is f32 round( e : vec N < T > ) -> vec N < T > (GLSLstd450Round) T is f32 sign( e : T ) -> T (GLSLstd450FSign) T is f32 sign( e : vec N < T > ) -> vec N < T > (GLSLstd450FSign) T is f32 sin( e : T ) -> T (GLSLstd450Sin) T is f32 sin( e : vec N < T > ) -> vec N < T > (GLSLstd450Sin) T is f32 sinh( e : T ) -> T (GLSLstd450Sinh) T is f32 sinh( e : vec N < T > ) -> vec N < T > (GLSLstd450Sinh) T is f32 smoothStep( e1 : T , e2 : T , e3 : T ) -> T (GLSLstd450SmoothStep) T is f32 smoothStep( e1 : vec N < T > , e2 : vec N < T > , e3 : vec N < T > ) -> vec N < T > (GLSLstd450SmoothStep) T is f32 sqrt( e : T ) -> T (GLSLstd450Sqrt) T is f32 sqrt( e : vec N < T > ) -> vec N < T > (GLSLstd450Sqrt) T is f32 step( e1 : T , e2 : T ) -> T (GLSLstd450Step) T is f32 step( e1 : vec N < T > , e2 : vec N < T > ) -> vec N < T > (GLSLstd450Step) T is f32 tan( e : T ) -> T (GLSLstd450Tan) T is f32 tan( e : vec N < T > ) -> vec N < T > (GLSLstd450Tan) T is f32 tanh( e : T ) -> T (GLSLstd450Tanh) T is f32 tanh( e : vec N < T > ) -> vec N < T > (GLSLstd450Tanh) T is f32 trunc( e : T ) -> T (GLSLstd450Trunc) T is f32 trunc( e : vec N < T > ) -> vec N < T > (GLSLstd450Trunc)

15.4. Integer built-in functions

Precondition Built-in Description T is u32 or i32 abs( e : T ) -> T (GLSLstd450SAbs) T is u32 or i32 abs( e : vec N < T > ) -> vec N < T > (GLSLstd450SAbs) T is u32 clamp( e1 : T , e2 : T , e3 : T ) -> T (GLSLstd450UClamp) T is i32 clamp( e1 : vec N < T > , e2 : vec N < T > , e3 : vec N < T > ) -> vec N < T > (GLSLstd450UClamp) T is i32 clamp( e1 : T , e2 : T , e3 : T ) -> T (GLSLstd450SClamp) T is i32 clamp( e1 : vec N < T > , e2 : vec N < T > , e3 : vec N < T > ) -> vec N < 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 : vec N < T > ) -> vec N < 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 i32 max( e1 : vec N < T > , e2 : vec N < T > ) -> vec N < T > (GLSLstd450UMax) T is i32 max( e1 : T , e2 : T ) -> T (GLSLstd450SMax) T is i32 max( e1 : vec N < T > , e2 : vec N < T > ) -> vec N < T > (GLSLstd450SMax) T is u32 min( e1 : T , e2 : T ) -> T (GLSLstd450UMin) T is i32 min( e1 : vec N < T > , e2 : vec N < T > ) -> vec N < T > (GLSLstd450UMin) T is i32 min( e1 : T , e2 : T ) -> T (GLSLstd450SMin) T is i32 min( e1 : vec N < T > , e2 : vec N < T > ) -> vec N < 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 : vec N < T > ) -> vec N < 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 : mat N x N < T > ) -> vec N < T > (GLSLstd450Determinant)

15.6. Vector built-in functions

Vector built-in functions SPIR-V dot(vecN<f32>, vecN<f32>) -> float OpDot outerProduct(vecN<f32>, vecM<f32>) -> matNxM<f32> OpOuterProduct

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

vec4<type> textureLoad(texture_ro_1d , i32 coords) vec4<type> textureLoad(texture_ro_2d , vec2<i32> coords) vec4<type> textureLoad(texture_ro_1d_array, vec2<i32> coords) vec4<type> textureLoad(texture_ro_3d , vec3<i32> coords) vec4<type> textureLoad(texture_ro_2d_array, vec3<i32> coords) %32 = OpImageRead %v4float %texture_ro %coords vec4<type> textureLoad(texture_sampled_1d , i32 coords, i32 level_of_detail) vec4<type> textureLoad(texture_sampled_2d , vec2<i32> coords, i32 level_of_detail) vec4<type> textureLoad(texture_sampled_1d_array, vec2<i32> coords, i32 level_of_detail) vec4<type> textureLoad(texture_sampled_3d , vec3<i32> coords, i32 level_of_detail) vec4<type> textureLoad(texture_sampled_2d_array, vec3<i32> coords, i32 level_of_detail) %32 = OpImageFetch %v4float %texture_sampled %coords Lod %level_of_detail vec4<type> textureLoad(texture_multisampled_2d, vec2<i32> coords, i32 sample_index) %32 = OpImageFetch %v4float %texture_multisampled %coords Sample %sample_index TODO(dsinclair): Add textureWrite method for texture_wo with integral coords TODO(dsinclair): Allow a small constant offset on the coordinate? May not be portable. vec4<type> textureSample(texture_sampled_1d , sampler, f32 coords) vec4<type> textureSample(texture_sampled_2d , sampler, vec2<f32> coords) vec4<type> textureSample(texture_sampled_1d_array , sampler, vec2<f32> coords) vec4<type> textureSample(texture_sampled_3d , sampler, vec3<f32> coords) vec4<type> textureSample(texture_sampled_2d_array , sampler, vec3<f32> coords) vec4<type> textureSample(texture_sampled_cube , sampler, vec3<f32> coords) vec4<type> textureSample(texture_sampled_cube_array, sampler, vec4<f32> coords) %24 = OpImageSampleImplicitLod %v4float %sampled_image %coords vec4<type> textureSampleLevel(texture_sampled_1d , sampler, f32 coords, f32 lod) vec4<type> textureSampleLevel(texture_sampled_2d , sampler, vec2<f32> coords, f32 lod) vec4<type> textureSampleLevel(texture_sampled_1d_array , sampler, vec2<f32> coords, f32 lod) vec4<type> textureSampleLevel(texture_sampled_3d , sampler, vec3<f32> coords, f32 lod) vec4<type> textureSampleLevel(texture_sampled_2d_array , sampler, vec3<f32> coords, f32 lod) vec4<type> textureSampleLevel(texture_sampled_cube , sampler, vec3<f32> coords, f32 lod) vec4<type> textureSampleLevel(texture_sampled_cube_array, sampler, vec4<f32> coords, f32 lod) %25 = OpImageSampleExplicitLod %v4float %sampled_image %coords Lod %lod vec4<type> textureSampleBias(texture_sampled_1d , sampler, f32 coords, f32 bias) vec4<type> textureSampleBias(texture_sampled_2d , sampler, vec2<f32> coords, f32 bias) vec4<type> textureSampleBias(texture_sampled_1d_array , sampler, vec2<f32> coords, f32 bias) vec4<type> textureSampleBias(texture_sampled_3d , sampler, vec3<f32> coords, f32 bias) vec4<type> textureSampleBias(texture_sampled_2d_array , sampler, vec3<f32> coords, f32 bias) vec4<type> textureSampleBias(texture_sampled_cube , sampler, vec3<f32> coords, f32 bias) vec4<type> textureSampleBias(texture_sampled_cube_array, sampler, vec4<f32> coords, f32 bias) %19 = OpImageSampleImplicitLod %v4float %sampled_image %coords Bias %bias f32 textureSampleCompare(texture_depth_2d , sampler_comparison, vec2<f32> coords, f32 depth_reference) f32 textureSampleCompare(texture_depth_2d_array , sampler_comparison, vec3<f32> coords, f32 depth_reference) f32 textureSampleCompare(texture_depth_cube , sampler_comparison, vec3<f32> coords, f32 depth_reference) f32 textureSampleCompare(texture_depth_cube_array, sampler_comparison, vec4<f32> coords, f32 depth_reference) %65 = OpImageSampleDrefExplicitLod %float %sampled_image %coord %depth_reference Lod %float_0 TODO(dsinclair): Add Level-of-Detail via explicit gradient. "Grad" image operand in SPIR-V TODO(dsinclair): Need gather operations

15.9. Atomic built-in functions

16. Glossary

TODO: Remove terms unused in the rest of the specification.

Term Definition Dominates Basic block A dominates basic block B if: A and B are both in the same function F

Every control flow path in F that goes to B must also to through A 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

var e : f32 = 3 ; # error : literal is the wrong type var f : f32 = 1.0 ; var t : i32 = i32 ( f );

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

17.4. Preamble

TODO(dneto): Forcing the Vulkan memory model is obsolete and should be removed.

EXAMPLE: Preamble .... OpCapability Shader OpCapability VulkanMemoryModel OpMemoryModel Logical VulkanKHR ....

While we recognize that most Vulkan devices will not support VulkanMemoryModel we expect the SPIR-V generated to be converted by SPIRV-Tools after the fact to make the shader compatible.