WGSL file format: language

This page lists what a reader of a .wgsl file meets in the language, taken from the W3C WebGPU Shading Language specification, Candidate Recommendation Draft of 6 November 2026. Built-in functions are on WGSL built-in functions, and the keyword and reserved-word lists are on WGSL keywords.

The file

A .wgsl file holds one WGSL module. Appendix A of the specification registers the text/wgsl media type for it with IANA.

Field	Value
Media type	`text/wgsl`
Required parameters	None
Optional parameters	None
Encoding	Unicode text in UTF-8, with no byte order mark (BOM)
File extension	`.wgsl`
Magic number	None
Macintosh file type code	`TEXT`
Fragment identifiers	None
Change controller	W3C

Textual structure

Blankspace is any run of code points from the Unicode Pattern_White_Space property. Some of these also count as line breaks, which end // comments and count for line numbers in diagnostics.

Code point	Name	Line break
`U+0020`	Space	No
`U+0009`	Horizontal tab	No
`U+000A`	Line feed (LF)	Yes
`U+000B`	Vertical tab (VT)	Yes
`U+000C`	Form feed (FF)	Yes
`U+000D`	Carriage return (CR)	Yes: CR alone, or CR followed by LF as one break
`U+0085`	Next line (NEL)	Yes
`U+200E`	Left-to-right mark	No
`U+200F`	Right-to-left mark	No
`U+2028`	Line separator (LS)	Yes
`U+2029`	Paragraph separator (PS)	Yes

Element	Rule
Line comment	`//` up to the next line break or the end of the module.
Block comment	`/* … */`. Block comments nest, and each one must be closed.
Tokens	Literal, keyword, reserved word, syntactic token (punctuation and operators), identifier and context-dependent name. Blankspace and comments separate tokens.
Identifier	Follows UAX #31: `([_\p{XID_Start}][\p{XID_Continue}]+)\|([\p{XID_Start}])`. It must not be a keyword or reserved word, must not be a single `_`, and must not start with `__`. Identifiers match by code point, with no Unicode normalisation.
Context-dependent name	A word with a fixed meaning only in one place, such as an attribute name, a built-in value name, a diagnostic rule, an extension name, an interpolation name or a swizzle. Elsewhere it is an ordinary identifier.
Template list	A comma-separated list of parameters in `<` `>` after a name, as in `vec3<f32>`. A discovery pass before parsing pairs each `<` with its `>`, so they are not read as comparison or shift operators.

Literals

Kind	Pattern	Type
Boolean	`true`, `false`	`bool`
Decimal integer	`/0[iu]?/` or `/[1-9][0-9]*[iu]?/`	No suffix: AbstractInt; `i`: `i32`; `u`: `u32`
Hexadecimal integer	`/0[xX][0-9a-fA-F]+[iu]?/`	As for decimal integers
Decimal float	Digits with an optional `.`, an optional exponent `e`/`E` with optional sign and decimal digits, and an optional `f` or `h` suffix	No suffix: AbstractFloat; `f`: `f32`; `h`: `f16`
Hexadecimal float	`0x`/`0X`, a hex significand with an optional `.`, an optional exponent `p`/`P` with optional sign and decimal digits, and an optional `f` or `h` suffix	As for decimal floats

A number with no point, no exponent and no float suffix is an integer.
A suffixed literal must fit its type: an i or u integer that does not fit, or an f or h float that overflows or cannot be represented exactly, is an error.
The h suffix needs enable f16;.
Decimal significands are truncated after 20 significant digits, hexadecimal ones after 16.

Directives

Directives come first in a module, before any declaration. Each ends with a semicolon.

Directive	Form	Effect
`enable`	`enable f16, subgroups;`	Turns on enable-extensions. An extension the implementation does not support is a shader-creation error.
`requires`	`requires pointer_composite_access;`	Documents use of language extensions; it does not change behaviour. An unsupported extension is a shader-creation error.
`diagnostic`	`diagnostic(off, derivative_uniformity);`	Global diagnostic filter for the whole module. Same form as the `@diagnostic` attribute, without the `@`.

Enable-extensions

These need hardware support, requested as a WebGPU feature, and must be named in enable.

Enable-extension	WebGPU feature	Allows
`f16`	`"shader-f16"`	The `f16` type
`clip_distances`	`"clip-distances"`	The `clip_distances` built-in value
`dual_source_blending`	`"dual-source-blending"`	The `blend_src` attribute
`subgroups`	`"subgroups"`	Subgroup built-in values, subgroup built-in functions and quad built-in functions
`primitive_index`	`"primitive-index"`	The `primitive_index` built-in value
`subgroup_size_control`	`"subgroup-size-control"`	The `subgroup_size` attribute; also enables `subgroups`
`atomic_vec2u_min_max`	`"atomic-vec2u-min-max"`	`atomicStoreMin` and `atomicStoreMax` on `atomic<vec2<u32>>` in the storage address space

Language extensions

These are available whenever the implementation supports them; requires only records their use.

Language extension	Adds
`readonly_and_readwrite_storage_textures`	`read` and `read_write` access modes for storage textures, and `textureBarrier`
`packed_4x8_integer_dot_product`	Dot products on four 8-bit integers packed in a 32-bit scalar (`dot4U8Packed`, `dot4I8Packed`) and the matching pack and unpack functions
`unrestricted_pointer_parameters`	Lifts the limits on pointer parameters of user-defined functions: address spaces other than `function` and `private`, and arguments that point into part of a variable
`pointer_composite_access`	Member and element access through a pointer, such as `p.m`, giving a reference
`uniform_buffer_standard_layout`	Uniform buffers use the same layout constraints as other address spaces
`subgroup_id`	The `subgroup_id` and `num_subgroups` built-in values (with `subgroups` enabled)
`subgroup_uniformity`	A `subgroup` scope for uniform control flow of subgroup and quad built-in functions
`texture_and_sampler_let`	`let` declarations of texture or sampler type
`texture_formats_tier1`	More storage texel formats (marked in the texel format table)
`linear_indexing`	The `global_invocation_index` and `workgroup_index` built-in values
`immediate_address_space`	The `immediate` address space, for `var<immediate>`
`fragment_depth`	A `depth_mode` parameter (`less` or `greater`) on the `frag_depth` built-in
`buffer_view`	`buffer` types and the buffer-view built-in functions: variables with an opaque store type that can be read as other host-shareable types
`swizzle_assignment`	Swizzle view types, so one assignment can update several components of a vector
`atomic_vec2u_min_max`	`atomicStoreMin` and `atomicStoreMax` on `atomic<vec2<u32>>` (this name is in both tables)

Diagnostics

A diagnostic filter has a severity and a triggering rule. The diagnostic directive covers the whole module; the @diagnostic attribute covers a range of code. A rule name is one token, or two joined by a period; implementations may add two-token rules. An unrecognised one-token rule should give a warning.

Severity	Meaning
`error`	Shader-creation or pipeline-creation error
`warning`	Anomaly worth reporting, not an error
`info`	Notable condition, not an error or warning
`off`	Not reported

Triggering rule	Default	Triggered when
`derivative_uniformity`	`error`	A derivative built-in, `textureSample`, `textureSampleBias` or `textureSampleCompare` is called where uniform control flow cannot be proven
`subgroup_uniformity`	`error`	A subgroup or quad built-in is called where uniform control flow cannot be proven, or the `delta` of `subgroupShuffleUp`/`subgroupShuffleDown` or the `mask` of `subgroupShuffleXor` is not provably uniform

Types

Scalar types

Type	Meaning
`bool`	`true` or `false`
AbstractInt	Type of an unsuffixed integer literal; exists only during compile-time evaluation and cannot be written in source
AbstractFloat	Type of an unsuffixed floating-point literal; compile-time only and cannot be written in source
`i32`	32-bit signed integer
`u32`	32-bit unsigned integer
`f32`	32-bit IEEE-754 binary floating point
`f16`	16-bit IEEE-754 binary floating point; needs `enable f16;`

Vectors, matrices, atomics, arrays and structures

Type	Meaning
`vecN<T>`	Vector of N components (2, 3 or 4) of scalar type T
`matCxR<T>`	Matrix of C columns and R rows (each 2, 3 or 4), stored as C column vectors `vecR<T>`; T is `f32`, `f16` or AbstractFloat
`atomic<T>`	Atomic integer, T is `i32` or `u32`; `atomic<vec2<u32>>` (a 64-bit unsigned value) needs `atomic_vec2u_min_max`. Only in the `workgroup` address space or `read_write` storage buffers
`array<E, N>`	Fixed-size array of N elements of type E; N is an override-expression giving a concrete integer
`array<E>`	Runtime-sized array; the bound buffer sets the length. Allowed only in places such as the last member of a structure in storage
`struct Name { m: T, … }`	Structure declared at module scope; members are scalars, vectors, matrices, atomics, fixed-size arrays or structures, and the last may be a runtime-sized array
`buffer<N>`, `buffer`	Opaque buffer types of the `buffer_view` language extension

Predeclared type aliases

The …h aliases need enable f16;.

Alias	Type	Alias	Type
`vec2i`	`vec2<i32>`	`vec2u`	`vec2<u32>`
`vec3i`	`vec3<i32>`	`vec3u`	`vec3<u32>`
`vec4i`	`vec4<i32>`	`vec4u`	`vec4<u32>`
`vec2f`	`vec2<f32>`	`vec2h`	`vec2<f16>`
`vec3f`	`vec3<f32>`	`vec3h`	`vec3<f16>`
`vec4f`	`vec4<f32>`	`vec4h`	`vec4<f16>`
`mat2x2f`	`mat2x2<f32>`	`mat2x2h`	`mat2x2<f16>`
`mat2x3f`	`mat2x3<f32>`	`mat2x3h`	`mat2x3<f16>`
`mat2x4f`	`mat2x4<f32>`	`mat2x4h`	`mat2x4<f16>`
`mat3x2f`	`mat3x2<f32>`	`mat3x2h`	`mat3x2<f16>`
`mat3x3f`	`mat3x3<f32>`	`mat3x3h`	`mat3x3<f16>`
`mat3x4f`	`mat3x4<f32>`	`mat3x4h`	`mat3x4<f16>`
`mat4x2f`	`mat4x2<f32>`	`mat4x2h`	`mat4x2<f16>`
`mat4x3f`	`mat4x3<f32>`	`mat4x3h`	`mat4x3<f16>`
`mat4x4f`	`mat4x4<f32>`	`mat4x4h`	`mat4x4<f16>`

Predeclared enumerants

New enumerants cannot be declared.

Enumeration	Enumerants
Address space	`function`, `private`, `workgroup`, `uniform`, `storage`, `immediate` (the `handle` address space has no enumerant and is never written)
Access mode	`read`, `write`, `read_write`
Texel format	The 40 formats in the texel format table
Interpolation type	`perspective` (perspective-correct), `linear` (linear, not perspective-correct), `flat` (not interpolated)
Interpolation sampling	`center` (pixel centre), `centroid` (a point inside all covered samples of the primitive), `sample` (per sample; the fragment shader runs once per sample), `first` (value of the first vertex), `either` (first or last vertex, implementation-dependent)
Depth mode	`less`, `greater` (for `@builtin(frag_depth, …)` with `fragment_depth`)
Swizzle names	One to four letters, all from `xyzw` or all from `rgba`

Texel formats for storage textures

Unorm, snorm and float formats are read and written in the shader as vec4<f32>, uint formats as vec4<u32> and sint formats as vec4<i32>.

Texel format	Channel format	Channels in memory order	Needs
`rgba8unorm`	8unorm	r, g, b, a
`rgba8snorm`	8snorm	r, g, b, a
`rgba8uint`	8uint	r, g, b, a
`rgba8sint`	8sint	r, g, b, a
`rgba16unorm`	16unorm	r, g, b, a	`texture_formats_tier1`
`rgba16snorm`	16snorm	r, g, b, a	`texture_formats_tier1`
`rgba16uint`	16uint	r, g, b, a
`rgba16sint`	16sint	r, g, b, a
`rgba16float`	16float	r, g, b, a
`rg8unorm`	8unorm	r, g	`texture_formats_tier1`
`rg8snorm`	8snorm	r, g	`texture_formats_tier1`
`rg8uint`	8uint	r, g	`texture_formats_tier1`
`rg8sint`	8sint	r, g	`texture_formats_tier1`
`rg16unorm`	16unorm	r, g	`texture_formats_tier1`
`rg16snorm`	16snorm	r, g	`texture_formats_tier1`
`rg16uint`	16uint	r, g	`texture_formats_tier1`
`rg16sint`	16sint	r, g	`texture_formats_tier1`
`rg16float`	16float	r, g	`texture_formats_tier1`
`r32uint`	32uint	r
`r32sint`	32sint	r
`r32float`	32float	r
`rg32uint`	32uint	r, g
`rg32sint`	32sint	r, g
`rg32float`	32float	r, g
`rgba32uint`	32uint	r, g, b, a
`rgba32sint`	32sint	r, g, b, a
`rgba32float`	32float	r, g, b, a
`bgra8unorm`	8unorm	b, g, r, a
`r8unorm`	8unorm	r	`texture_formats_tier1`
`r8snorm`	8snorm	r	`texture_formats_tier1`
`r8uint`	8uint	r	`texture_formats_tier1`
`r8sint`	8sint	r	`texture_formats_tier1`
`r16unorm`	16unorm	r	`texture_formats_tier1`
`r16snorm`	16snorm	r	`texture_formats_tier1`
`r16uint`	16uint	r	`texture_formats_tier1`
`r16sint`	16sint	r	`texture_formats_tier1`
`r16float`	16float	r	`texture_formats_tier1`
`rgb10a2unorm`	r, g, b: 10unorm; a: 2unorm	r, g, b, a	`texture_formats_tier1`
`rgb10a2uint`	r, g, b: 10uint; a: 2uint	r, g, b, a	`texture_formats_tier1`
`rg11b10ufloat`	r, g: 11float; b: 10float	r, g, b	`texture_formats_tier1`

Texture and sampler types

In sampled and multisampled types T is the sampled type, f32, i32 or u32. Storage textures take a texel format and an access mode; a bad combination of the two is a pipeline-creation error. Depth textures allow reads only.

Type	Kind	Dimensions	Arrayed
`texture_1d<T>`	Sampled	1D	No
`texture_2d<T>`	Sampled	2D	No
`texture_2d_array<T>`	Sampled	2D	Yes
`texture_3d<T>`	Sampled	3D	No
`texture_cube<T>`	Sampled	Cube	No
`texture_cube_array<T>`	Sampled	Cube	Yes
`texture_multisampled_2d<T>`	Multisampled	2D	No
`texture_depth_multisampled_2d`	Multisampled	2D	No
`texture_external`	External sampled (float, opaque representation)	2D	No
`texture_storage_1d<Format, Access>`	Storage	1D	No
`texture_storage_2d<Format, Access>`	Storage	2D	No
`texture_storage_2d_array<Format, Access>`	Storage	2D	Yes
`texture_storage_3d<Format, Access>`	Storage	3D	No
`texture_depth_2d`	Depth	2D	No
`texture_depth_2d_array`	Depth	2D	Yes
`texture_depth_cube`	Depth	Cube	No
`texture_depth_cube_array`	Depth	Cube	Yes
`sampler`	Sampler for sampled textures	–	–
`sampler_comparison`	Comparison sampler for depth textures	–	–

Memory views

Type	Meaning
`ptr<AS, T, AM>`	Pointer to memory of store type T in address space AS with access mode AM. AS and T are always written; AM only for `storage`, where it defaults to `read`
`ref<AS, T, AM>`	Reference type with the same three parts. Used only when analysing a module; never written in source. The indirection operator `*p` turns a pointer into a reference, and `&r` turns a reference into a pointer
`swizzle<AS, S, N, K>`	Swizzle view type, from the `swizzle_assignment` extension

Declarations and scope

A declaration binds an identifier to a type, value, variable, function or parameter. Predeclared names and module-scope declarations are visible in the whole module, in any order. A function parameter is visible in the function body. Any other declaration is visible from just after it to the end of its enclosing block. Two declarations with the same name must not end their scope at the same place, and module-scope declarations must not be recursive, directly or indirectly.

Declaration	Form	Rules
`const`	`const n = 4;`	Immutable, fixed at shader-creation time. Needs an initial value that is a const-expression. May have an abstract type.
`override`	`@id(0) override scale: f32 = 1.0;`	Immutable, fixed at pipeline-creation time. Module scope only, concrete scalar type, initial value optional.
`let`	`let x = a + b;`	Immutable value, function scope only, initial value required. Type is concrete and constructible, or a pointer.
`var`	`var<private> seed: u32;`	Mutable variable with a store type. See the address-space table below.
`alias`	`alias Colour = vec4f;`	Type alias. Module scope only.
`struct`	`struct Light { pos: vec3f, power: f32 }`	Structure type. Module scope only.
`fn`	`fn area(r: f32) -> f32 { … }`	Function. Parameters are constructible types, pointers, textures or samplers; the return type, after `->`, must be constructible. Entry points carry a stage attribute.
`const_assert`	`const_assert n < 8;`	Shader-creation error if the const-expression is false. Module scope or as a statement; no effect on the compiled shader.

`var` declarations by address space

Form	Scope	Mutable	Initial value in source
`var` or `var<function>`	Function	Yes	Allowed; otherwise zero value
`var<private>`	Module	Yes	Allowed (const- or override-expression); otherwise zero value
`var<workgroup>`	Module	Yes	Not allowed; zero value
`var<uniform>`	Module	No	Not allowed
`var<storage>` or `var<storage, read>`	Module	No	Not allowed
`var<storage, read_write>`	Module	Yes	Not allowed
`var<immediate>`	Module	No	Not allowed
`var` of texture or sampler type (`handle`)	Module	No	Not allowed

The address space must be written for every space except function (optional) and handle (never written).
An access mode may be written only for storage, and only after the address space.
The @group and @binding attributes apply to resource variables and place them in a bind group.

Operator precedence

From strongest to weakest. The binding column says which stronger groups an operand may contain without parentheses: "Unary" means only unary and primary expressions. Shift and relational operators do not chain with themselves, and the bitwise and short-circuit groups do not mix with each other, so a < b < c, a & b | c and a && b || c need parentheses.

Group	Operators	Associativity	Binding
Parenthesised	`(…)`
Primary	`a()`, `a[]`, `a.b`	Left to right
Unary	`-a`, `!a`, `~a`, `*a`, `&a`	Right to left	All above
Multiplicative	`a*b`, `a/b`, `a%b`	Left to right	All above
Additive	`a+b`, `a-b`	Left to right	All above
Shift	`a<<b`, `a>>b`	Requires parentheses	Unary
Relational	`a<b`, `a>b`, `a<=b`, `a>=b`, `a==b`, `a!=b`	Requires parentheses	All above
Binary AND	`a&b`	Left to right	Unary
Binary XOR	`a^b`	Left to right	Unary
Binary OR	`a\|b`	Left to right	Unary
Short-circuit AND	`a&&b`	Left to right	Relational
Short-circuit OR	`a\|\|b`	Left to right	Relational

& and | on bool values do not short-circuit. a[i] indexes vectors, matrix columns and arrays; a.b selects a structure member or a vector swizzle such as v.xy or c.rgb.

Statements

Statement	Form	Rules
Compound	`{ … }`	A block; declarations inside end with it
Assignment	`r = e;`	Left side is a writable reference. Assigning to a multi-letter swizzle needs `swizzle_assignment` and distinct components.
Phony assignment	`_ = e;`	Evaluates `e` and discards it; `e` is constructible, a pointer, a texture or a sampler
Compound assignment	`r += e;` and `-=`, `*=`, `/=`, `%=`, `&=`, `\|=`, `^=`, `>>=`, `<<=`	Same as `r = r op (e)`, with `r` evaluated once
Increment, decrement	`r++;` `r--;`	Concrete integer scalar with `read_write` access; same as `r += T(1)` and `r -= T(1)`
Variable or value declaration	`var`, `let`, `const`	See declarations
If	`if c { … } else if d { … } else { … }`	Conditions are `bool`; at most one block runs
Switch	`switch e { case 1, 2: { … } default: { … } }`	Case selectors are const-expressions of one concrete integer type, with no repeats. Exactly one `default`, alone or mixed into a case list. The colon is optional. No fall-through.
Loop	`loop { … continuing { … } }`	Repeats until `break` or `return`; the optional `continuing` block comes last and runs at the end of each iteration
Break if	`break if c;`	Only as the last statement of a `continuing` block; `bool` condition
For	`for (init; cond; update) { … }`	Same as a `loop` that tests `cond` first and puts `update` in `continuing`
While	`while c { … }`	Same as a `loop` that breaks when `c` is false
Break	`break;`	Inside `loop`, `for`, `while` or `switch`; must not leave a `continuing` block
Continue	`continue;`	Inside `loop`, `for` or `while`; goes to `continuing` if there is one, and must not skip a declaration that block uses
Return	`return;` `return e;`	Value matches the return type; in an entry point it ends the invocation. Not allowed inside `continuing`.
Discard	`discard;`	Fragment shaders only; the invocation becomes a helper invocation and its output is thrown away
Function call	`f(a, b);`	Any result is ignored; an error if `f` has `@must_use`
Const assertion	`const_assert e;`	Shader-creation error if `e` is false

An unbounded loop is a dynamic error.

Attributes

An attribute is written @name or @name(…) before the thing it modifies. Integer parameters are const-expressions resolving to i32 or u32 unless the table says otherwise. The specification defines 18 attributes.

Attribute	Parameters	Applies to	Meaning
`@align(n)`	Positive power of 2	Structure member	Byte alignment of the member. Outside the uniform address space it must be a multiple of the member type's required alignment.
`@binding(n)`	Non-negative	Resource variable	Binding number within the bind group
`@blend_src(n)`	0 or 1	Structure member that also has `@location`, numeric scalar or vector	Which source of dual-source blending the fragment output is; needs `dual_source_blending`. Not on stage inputs, and on outputs only for the fragment stage.
`@builtin(name)`	A built-in value name	Entry point parameter, entry point return type, or structure member	Marks a built-in value. With `fragment_depth`, `frag_depth` takes a second parameter, `less` or `greater`.
`@compute`	None	Function	Entry point for the compute shader stage of a compute pipeline
`@const`	None	Function declaration	Marks a built-in function usable as a const-function. Must not be applied to a user-defined function.
`@diagnostic(severity, rule)`	Severity name, triggering rule name	The code it is written on	Diagnostic filter over a range of code. Several are allowed if their rules differ.
`@fragment`	None	Function	Entry point for the fragment shader stage of a render pipeline
`@group(n)`	Non-negative	Resource variable	Bind group of the resource
`@id(n)`	Non-negative	`override` of scalar type	Numeric identifier used as an alternative name for the pipeline-overridable constant
`@interpolate(type, sampling)`	Interpolation type; optional interpolation sampling	Declaration with `@location`	How user-defined IO is interpolated. Matters only on vertex outputs and fragment inputs.
`@invariant`	None	The `position` built-in	On a vertex `position` output, the result is the same across programs and entry points when data and control flow match. No effect on a `position` input.
`@location(n)`	Non-negative	Entry point parameter, return type or structure member of numeric scalar or vector type	IO location of user-defined IO. Not on compute shader inputs.
`@must_use`	None	Function with a return type	A call must be used as an expression, not as a whole function call statement
`@size(n)`	Positive	Structure member with a creation-fixed footprint	Bytes reserved for the member; at least the size of its type
`@subgroup_size(n)`	Const- or override-expression	Compute entry point	Subgroup size for the invocation; needs `subgroup_size_control`. A value that is not a power of two is an error. The x size of `@workgroup_size` must be a multiple of it, and it must lie between the device's subgroupMinSize and subgroupMaxSize.
`@vertex`	None	Function	Entry point for the vertex shader stage of a render pipeline
`@workgroup_size(x, y, z)`	One to three const- or override-expressions of the same type; y and z default to 1	Compute entry point	Dimensions of the workgroup grid. Each must be positive and within WebGPU limits, as must their product.

At most one of @vertex, @fragment and @compute may be on a function.

Entry points and shader interface

A WebGPU draw command runs a render pipeline with a vertex shader stage, which maps the attributes of each vertex, and a fragment shader stage, which processes each rasterised fragment. A dispatch command runs a compute pipeline's compute shader stage over a grid of invocations. An entry point is a function with @vertex, @fragment or @compute; its parameters are stage inputs and its return value is the stage output.

Built-in values

Name	Stage	Direction	Type	Needs	Meaning
`vertex_index`	vertex	input	`u32`		Index of the vertex within the draw command, independent of instancing
`instance_index`	vertex	input	`u32`		Instance index of the vertex within the draw command
`clip_distances`	vertex	output	`array<f32, N>`, N ≤ 8	`clip_distances`	Distance to each user-defined clip plane: 0 on the plane, positive inside, negative outside
`position`	vertex	output	`vec4<f32>`		Clip-space position of the vertex
`position`	fragment	input	`vec4<f32>`		Position of the fragment's rasterisation point
`front_facing`	fragment	input	`bool`		True when the fragment is on a front-facing primitive
`frag_depth`	fragment	output	`f32`		New depth of the fragment in the viewport depth range
`primitive_index`	fragment	input	`u32`	`primitive_index`	Index of the primitive within the current instance of the draw
`sample_index`	fragment	input	`u32`		Sample index of the fragment's rasterisation point
`sample_mask`	fragment	input	`u32`		Sample coverage bit mask of the fragment
`sample_mask`	fragment	output	`u32`		Sample coverage mask control for the fragment
`local_invocation_id`	compute	input	`vec3<u32>`		Local invocation ID within the workgroup
`local_invocation_index`	compute	input	`u32`		Linearised position within the workgroup grid
`global_invocation_id`	compute	input	`vec3<u32>`		Global invocation ID
`global_invocation_index`	compute	input	`u32`	`linear_indexing`	Global invocation index
`workgroup_id`	compute	input	`vec3<u32>`		Workgroup ID
`workgroup_index`	compute	input	`u32`	`linear_indexing`	Workgroup index
`num_workgroups`	compute	input	`vec3<u32>`		Dispatch size in workgroups
`subgroup_invocation_id`	compute, fragment	input	`u32`	`subgroups`	Invocation's ID within its subgroup
`subgroup_size`	compute, fragment	input	`u32`	`subgroups`	Size of the invocation's subgroup
`subgroup_id`	compute	input	`u32`	`subgroups` and `subgroup_id`	ID of the subgroup within the workgroup
`num_subgroups`	compute	input	`u32`	`subgroups` and `subgroup_id`	Number of subgroups in the workgroup

User-defined inputs and outputs

Each user-defined input or output is a numeric scalar or vector with an explicit @location. Structure members in entry point IO are each either a built-in value or have a location; the two can be mixed in one structure.
One location holds up to 16 bytes, so a vec4<f32> fits in one.
An entry point's inputs must not repeat a location. Within a structure, locations must not repeat, except for dual-source blending: exactly two members, @location(0) @blend_src(0) and @location(0) @blend_src(1), of the same type.
Compute shaders have no user-defined inputs or outputs.

Interpolation

IO type	Rule
Floating-point, no `@interpolate`	`@interpolate(perspective, center)` is assumed
`flat`	Sampling, if given, is `first` or `either`; default `first`
`perspective` or `linear`	Sampling, if given, is `center`, `centroid` or `sample`; default `center`
Integer vertex outputs and fragment inputs	Must be `flat`

A fragment input must match the interpolation of the vertex output at the same location, or pipeline creation fails.

Memory

Address spaces

Only storage supports two access modes (read and read_write); every other address space has one. Access modes are read (reads only), write (writes only) and read_write.

Address space	Shared among	Access mode	Notes
`function`	Same invocation only	`read_write`	Function-scope variables
`private`	Same invocation only	`read_write`	Module-scope, one copy per invocation
`workgroup`	Invocations in the same compute workgroup	`read_write`	Compute stage only. The element count of an outermost array may be an override constant.
`uniform`	Invocations in the same shader stage	`read`	Uniform buffers
`storage`	Invocations in the same shader stage	`read` (default) or `read_write`	Storage buffers. A vertex shader may access them only with `read`.
`immediate`	Invocations in the same shader stage	`read`	Immediate data; each entry point may statically access at most one such variable. Needs `immediate_address_space`.
`handle`	Invocations in the same shader stage	`read`	Textures and samplers; never written in source. A vertex shader must not access a storage texture with `write` or `read_write` access.

Alignment and size

Matrices are column-major. bool takes 4 bytes so that neighbouring values can be written without data races.

Type	AlignOf	SizeOf
`bool`, `i32`, `u32`, `f32`	4	4
`f16`	2	2
`atomic<T>`	4	4
`vec2<T>` (T is bool, i32, u32 or f32)	8	8
`vec2<f16>`	4	4
`vec3<T>` (T is bool, i32, u32 or f32)	16	12
`vec3<f16>`	8	6
`vec4<T>` (T is bool, i32, u32 or f32)	16	16
`vec4<f16>`	8	8
`matCxR` (general)	AlignOf(vecR)	SizeOf(array<vecR, C>)
`mat2x2<f32>`	8	16
`mat2x2<f16>`	4	8
`mat3x2<f32>`	8	24
`mat3x2<f16>`	4	12
`mat4x2<f32>`	8	32
`mat4x2<f16>`	4	16
`mat2x3<f32>`	16	32
`mat2x3<f16>`	8	16
`mat3x3<f32>`	16	48
`mat3x3<f16>`	8	24
`mat4x3<f32>`	16	64
`mat4x3<f16>`	8	32
`mat2x4<f32>`	16	32
`mat2x4<f16>`	8	16
`mat3x4<f32>`	16	48
`mat3x4<f16>`	8	24
`mat4x4<f32>`	16	64
`mat4x4<f16>`	8	32
`struct S` with members M1 … MN	Largest member alignment	roundUp(AlignOf(S), offset of the last member + its size)
`array<E, N>`	AlignOf(E)	N × roundUp(AlignOf(E), SizeOf(E))
`array<E>`	AlignOf(E)	Runtime element count × roundUp(AlignOf(E), SizeOf(E))
`buffer<N>`	–	N
`buffer`	–	Size of the bound GPUBuffer

@align and @size on a structure member change its alignment and the bytes it takes. The extra per-address-space rules in the specification's Address Space Layout Constraints section, which the uniform_buffer_standard_layout extension relaxes for uniform buffers, are not covered on this page.

Sources

WebGPU Shading Language – W3C Candidate Recommendation Draft, 6 November 2026: sections 2.3, 3, 4, 5, 6, 7, 8.20, 9, 10, 11, 12, 13, 14.2 to 14.4 and Appendix A
text/wgsl – IANA media type registration

.WGSL File Format – Language