WebGPU

https://gpuweb.github.io/gpuweb/
Admin
WebGPU is a W3C standard API that exposes graphics processing unit capabilities for performing...
Tokens:150,286
Snippets:434
Trust Score:8.3
Update:11 months ago
Show doc for...
Tokens:
Raw
### WGSL Vector Zero Value Construction Example

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates the construction of zero-valued vectors in WGSL using the zero value constructor. It shows examples for `vec2<f32>` and `vec3<i32>`, illustrating both the direct constructor call and explicit value initialization.

```WGSL
vec2<f32>()                 // The zero-valued vector of two f32 components.
vec2<f32>(0.0, 0.0)         // The same value, written explicitly.

vec3<i32>()                 // The zero-valued vector of three i32 components.
vec3<i32>(0, 0, 0)          // The same value, written explicitly.


```

--------------------------------

### WGSL Structure Zero Value Construction Example

Source: https://gpuweb.github.io/gpuweb/wgsl

Provides an example of constructing a zero-valued structure in WGSL. It defines a `Student` struct and demonstrates assigning its zero value using the struct constructor, both with explicit component values and with zero-valued members.

```WGSL
struct Student {
  grade: i32,
  GPA: f32,
  attendance: array<bool,4>
}

fn func() {
  var s: Student;

  // The zero value for Student
  s = Student();

  // The same value, written explicitly.
  s = Student(0, 0.0, array<bool,4>(false, false, false, false));

  // The same value, written with zero-valued members.
  s = Student(i32(), f32(), array<bool,4>());
}


```

--------------------------------

### WGSL Compound Assignment Examples

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates the usage of compound assignment operators like += in WGSL. It shows how these operators modify variable values and provides examples of their application in functions.

```wgsl
var<private> next_item: i32 = 0;

fn advance_item() -> i32 {
   next_item += 1;   // Adds 1 to next_item.
   return next_item - 1;
}

fn bump_item() {
  var data: array<f32,10>;
  next_item = 0;
  // Adds 5.0 to data[0], calling advance_item() only once.
  data[advance_item()] += 5.0;
  // next_item will be 1 here.
}

fn precedence_example() {
  var value = 1;
  // The right-hand side of a compound assignment is its own expression.
  value *= 2 + 3; // Same as value = value * (2 + 3);
  // 'value' now holds 5.
}

```

--------------------------------

### WGSL Array Zero Value Construction Example

Source: https://gpuweb.github.io/gpuweb/wgsl

Illustrates how to construct zero-valued arrays in WGSL using the `array` constructor. This example shows creating a two-element boolean array and compares the constructor call with explicit value assignment.

```WGSL
array<bool, 2>()               // The zero-valued array of two booleans.
array<bool, 2>(false, false)   // The same value, written explicitly.


```

--------------------------------

### WGSL Example: Enabling Extensions

Source: https://gpuweb.github.io/gpuweb/wgsl

Illustrates how to use the 'enable' directive in WGSL to activate hypothetical extensions like 'arbitrary_precision_float' and 'rounding_mode'. It shows enabling extensions and using them with custom types and attributes.

```WGSL
// Enable a hypothetical extension for arbitrary precision floating point types.
enable arbitrary_precision_float;
enable arbitrary_precision_float; // A redundant enable directive is ok.

// Enable a hypothetical extension to control the rounding mode.
enable rounding_mode;

// Assuming arbitrary_precision_float enables use of:
//    - a type f<E,M>
//    - as a type in function return, formal parameters and let-declarations
//    - as a value constructor from AbstractFloat
//    - operands to division operator: /
// Assuming @rounding_mode attribute is enabled by the rounding_mode enable directive.
@rounding_mode(round_to_even)
fn halve_it(x : f<8, 7>) -> f<8, 7> {
  let two = f<8, 7>(2);
  return x / 2; // uses round to even rounding mode.
}

```

--------------------------------

### WGSL Pointer Example: Function Parameter

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates using a pointer as a formal parameter in a WGSL function. The example shows how to dereference the pointer using `*` to modify the value it points to within the function.

```wgsl
fn add_one(x: ptr<function,i32>) {
  /* Update the locations for 'x' to contain the next higher integer value,
     (or to wrap around to the largest negative i32 value).
     On the left-hand side, unary '*' converts the pointer to a reference that
     can then be assigned to. It has a read_write access mode, by default.
     /* On the right-hand side:
        - Unary '*' converts the pointer to a reference, with a read_write
          access mode.
        - The only matching type rule is for addition (+) and requires '*x' to

```

--------------------------------

### WGSL: User-Defined Function Declaration Examples

Source: https://gpuweb.github.io/gpuweb/wgsl

Illustrates the declaration and usage of user-defined functions in WGSL. It includes a simple function `add_two` and a compute shader entry point `main` that calls `add_two`.

```wgsl
// Declare the add_two function.
// It has two formal parameters, i and b.
// It has a return type of i32.
// It has a body with a return statement.
fn add_two(i: i32, b: f32) -> i32 {
  return i + 2;  // A formal parameter is available for use in the body.
}

// A compute shader entry point function, 'main'.
// It has no specified return type.
// It invokes the add_two function, and captures
// the resulting value in the named value 'six'.
@compute @workgroup_size(1)
fn main() {
   let six: i32 = add_two(4, 5.0);
}
```

--------------------------------

### WGSL Runtime-Sized Array Layout Examples

Source: https://gpuweb.github.io/gpuweb/wgsl

Illustrates runtime-sized array declarations in WGSL for storage buffers, showing how the number of elements is calculated based on the buffer binding size and element stride for f32 and vec3<f32>.

```wgsl
// Array where:
//   - alignment is 4 = AlignOf(f32)
//   - element stride is 4 = roundUp(AlignOf(f32),SizeOf(f32)) = 4
// If B is the effective buffer binding size for the binding on the
// draw or dispatch command, the number of elements is:
//   N_runtime = floor(B / element stride) = floor(B / 4)
@group(0) @binding(0)
var<storage> weights: array<f32>;

// Array where:
//   - alignment is 16 = AlignOf(vec3<f32>) = 16
//   - element stride is 16 = roundUp(AlignOf(vec3<f32>), SizeOf(vec3<f32>))
//                          = roundUp(16,12)
// If B is the effective buffer binding size for the binding on the
// draw or dispatch command, the number of elements is:
//   N_runtime = floor(B / element stride) = floor(B / 16)
var<storage> directions: array<vec3<f32>>;

```

--------------------------------

### WGSL Switch Statement Basic Example

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates a basic WGSL switch statement with various case clauses and a default clause. It illustrates optional colons, the placement of the default clause, and the use of multiple selector values.

```WGSL
var a : i32;
let x : i32 = generateValue();
switch x {
  case 0: {      // The colon is optional
    a = 1;
  }
  default {      // The default need not appear last
    a = 2;
  }
  case 1, 2, {   // Multiple selector values can be used
    a = 3;
  }
  case 3, {      // The trailing comma is optional
    a = 4;
  }
  case 4 {
    a = 5;
  }
}
```

--------------------------------

### Pointer Desugaring Example in WGSL

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates the desugaring of pointer types in a WGSL function for uniformity analysis. It shows the original function and its equivalent version after pointer desugaring, simplifying pointer usage for analysis.

```wgsl
fn foo(p : ptr<function, array<f32, 4>>, i : i32) -> f32 {
  let p1 = p;
  var x = i;
  let p2 = &((*p1)[x]);
  x = 0;
  *p2 = 5;
  return (*p1)[x];
}
```

```wgsl
fn foo_for_analysis(p : ptr<function, array<f32, 4>>, i : i32) -> f32 {
  var p_var = *p;            // Introduce variable for p.
  let p1 = &p_var;           // Use the variable for p1
  var x = i;
  let x_tmp1 = x;            // Capture value of x
  let p2 = &(p_var[x_tmp1]); // Substitute p1's initializer
  x = 0;
  *(&(p_var[x_tmp1])) = 5;   // Substitute p2's initializer
  return (*(&p_var))[x];     // Substitute p1's initializer
}
```

--------------------------------

### WGSL Example: Gather Depth Comparison

Source: https://gpuweb.github.io/gpuweb/wgsl

An example demonstrating how to use the textureGatherCompare function in WGSL for gathering depth comparison results. It defines a function that takes texture coordinates and a depth reference, and returns the comparison vector.

```WGSL
@group(0) @binding(0) var dt: texture_depth_2d;
@group(0) @binding(1) var s: sampler;

fn gather_depth_compare(c: vec2<f32>, depth_ref: f32) -> vec4<f32> {
  return textureGatherCompare(dt,s,c,depth_ref);
}

```

--------------------------------

### WGSL Pointer Type Example

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates the declaration and usage of pointer types in WGSL functions. It shows how to specify the address space, storable type, and access mode for pointers to scalars and arrays.

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

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

```

--------------------------------

### WGSL Fixed-Size Array Layout Examples

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates fixed-size array declarations in WGSL, illustrating how alignment, element stride, and total size are determined based on data types like f32 and vec3<f32>.

```wgsl
// Array where:
//   - alignment is 4 = AlignOf(f32)
//   - element stride is 4 = roundUp(AlignOf(f32),SizeOf(f32)) = roundUp(4,4)
//   - size is 32 = stride * number_of_elements = 4 * 8
var small_stride: array<f32, 8>;

// Array where:
//   - alignment is 16 = AlignOf(vec3<f32>) = 16
//   - element stride is 16 = roundUp(AlignOf(vec3<f32>), SizeOf(vec3<f32>))
//                          = roundUp(16,12)
//   - size is 128 = stride * number_of_elements = 16 * 8
var bigger_stride: array<vec3<f32>, 8>;

```

--------------------------------

### Const-Expression Examples in WebGPU Shading Language

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates various scenarios of const-expressions, including simple integer literals, negative numbers, and their analysis in const and let declarations. It also highlights the behavior of short-circuiting operators like &&.

```wgsl
const minint = -2147483648;
let minint = -2147483648;
false && (10i < i32(5 * 1000 * 1000 * 1000))
false && array<u32, 1 + 2>(0, 1, 2)[0] == 0
```

--------------------------------

### WGSL Alias Analysis Example

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates WGSL's rules for aliasing pointer parameters to prevent memory safety violations, covering scenarios where writes and reads occur through aliased identifiers.

```wgsl
var<private> x : i32 = 0;

fn f1(p1 : ptr<function, i32>, p2 : ptr<function, i32>) {
  *p1 = *p2;
}

fn f2(p1 : ptr<function, i32>, p2 : ptr<function, i32>) {
  f1(p1, p2);
}

fn f3() {
  var a : i32 = 0;
  f2(&a, &a);  // Invalid. Cannot pass two pointer parameters
               // with the same root identifier when one or
               // more are written (even by a subfunction).
}

fn f4(p1 : ptr<function, i32>, p2 : ptr<function, i32>) -> i32 {
  return *p1 + *p2;
}

fn f5() {
  var a : i32 = 0;
  let b = f4(&a, &a); // Valid. p1 and p2 in f4 are both only read.
}

fn f6(p : ptr<private, i32>) {
  x = *p;
}

fn f7(p : ptr<private, i32>) -> i32 {
  return x + *p;
}

fn f8() {
  let a = f6(&x); // Invalid. x is written as a global variable and
                  // read as a parameter.
  let b = f7(&x); // Valid. x is only read as both a parameter and
                  // a variable.
}

```

--------------------------------

### Const Function Usage Example (WGSL)

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates the usage of 'const' functions within WGSL, including their evaluation at shader-creation time and their role in const-expressions. It highlights type inference and the use of const functions in array declarations.

```wgsl
const first_one = firstLeadingBit(1234 + 4567); // Evaluates to 12
                                                // first_one has the type i32, because
                                                // firstLeadingBit cannot operate on
                                                // AbstractInt

@id(1) override x : i32;
override y = firstLeadingBit(x); // const-expressions can be
                                 // used in override-expressions.
                                 // firstLeadingBit(x) is not a
                                 // const-expression in this context.

fn foo() {
  var a : array<i32, firstLeadingBit(257)>; // const-functions can be used in
                                            // const-expressions if all their
                                            // parameters are const-expressions.
}
```

--------------------------------

### WGSL Return Statement

Source: https://gpuweb.github.io/gpuweb/wgsl

Provides an example of a 'return' statement in WGSL. A 'return' statement ends the execution of the current function. If the function has a return type, a value must be provided. If not, the 'return' statement is optional and should not supply a value.

```wgsl
return expression;
// or for functions without a return type:
return;
```

--------------------------------

### WGSL Reference Example: Variable Operations

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates the use of reference types in WGSL for variable declaration, arithmetic operations, and assignments. The 'Load Rule' is shown where references are automatically dereferenced to satisfy type requirements.

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

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

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

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

```

--------------------------------

### WGSL Structure Layout with Implicit Member Sizes and Alignments

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates the default calculation of structure member sizes, alignments, and offsets in WGSL when explicit attributes are not used. This example shows how padding is implicitly added to satisfy alignment requirements.

```WGSL
struct A {
    u: f32,
    v: f32,
    w: vec2<f32>,
    x: f32
}

struct B {
    a: vec2<f32>,
    b: vec3<f32>,
    c: f32,
    d: f32,
    e: A,
    f: vec3<f32>,
    g: array<A, 3>,
    h: i32
}

@group(0) @binding(0)
var<storage,read_write> storage_buffer: B;
```

--------------------------------

### Built-in Values Overview

Source: https://gpuweb.github.io/gpuweb/wgsl

This section provides an overview of the available built-in input and output values, summarizing their stage, direction, type, and any relevant extensions.

```APIDOC
## Built-in Values Overview

This section summarizes the available built-in values, detailing their name, stage, direction, type, and any applicable extensions.

### Built-in Input and Output Values Table

| Name                 | Stage   | Direction | Type                       | Extension       |
|----------------------|---------|-----------|----------------------------|-----------------|
| `vertex_index`       | vertex  | input     | u32                        |                 |
| `instance_index`     | vertex  | input     | u32                        |                 |
| `clip_distances`     | vertex  | output    | array<f32, N> (`N` ≤ `8`) | clip_distances  |
| `position`           | vertex  | output    | vec4<f32>                  |                 |
|                      | fragment| input     | vec4<f32>                  |                 |
| `front_facing`       | fragment| input     | bool                       |                 |
| `frag_depth`         | fragment| output    | f32                        |                 |
| `primitive_index`    | fragment| input     | u32                        | primitive_index |
| `sample_index`       | fragment| input     | u32                        |                 |
| `sample_mask`        | fragment| input     | u32                        |                 |
|                      | fragment| output    | u32                        |                 |
| `local_invocation_id`| compute | input     | vec3<u32>                  |                 |
| `local_invocation_index`| compute | input     | u32                        |                 |
| `global_invocation_id`| compute | input     | vec3<u32>                  |                 |
| `workgroup_id`       | compute | input     | vec3<u32>                  |                 |
| `num_workgroups`     | compute | input     | vec3<u32>                  |                 |
| `subgroup_invocation_id`| compute | input     | u32                        | subgroups      |
|                      | fragment|           |                            |                 |
| `subgroup_size`      | compute | input     | u32                        | subgroups      |
|                      | fragment|           |                            |                 |

### Example: Declaring Built-in Values

```wgsl
struct VertexOutput {
  @builtin(position) my_pos: vec4<f32>,
  @builtin(clip_distances) my_clip_distances: array<f32, 8>,
}

@vertex
fn vs_main(
  @builtin(vertex_index) my_index: u32,
  @builtin(instance_index) my_inst_index: u32,
) -> VertexOutput {}

struct FragmentOutput {
  @builtin(frag_depth) depth: f32,
  @builtin(sample_mask) mask_out: u32
}

@fragment
fn fs_main(
  @builtin(front_facing) is_front: bool,
  @builtin(position) coord: vec4<f32>,
  @builtin(sample_index) my_sample_index: u32,
  @builtin(sample_mask) mask_in: u32,
) -> FragmentOutput {}

@compute @workgroup_size(64)
fn cs_main(
  @builtin(local_invocation_id) local_id: vec3<u32>,
  @builtin(local_invocation_index) local_index: u32,
  @builtin(global_invocation_id) global_id: vec3<u32>,
) {}
```
```

--------------------------------

### WGSL Comments: Line-Ending and Block Examples

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates the syntax for both line-ending and block comments in WGSL. Line-ending comments start with '//' and extend to the end of the line. Block comments start with '/*' and end with '*/', and can span multiple lines and be nested.

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

--------------------------------

### WGSL workgroup_size Attribute Examples

Source: https://gpuweb.github.io/gpuweb/wgsl

Illustrates the usage of the @workgroup_size attribute in WGSL for compute shaders. Examples include fixed-size workgroups, and workgroups defined by pipeline-overridable constants. An error case for incorrect placement is also shown.

```WGSL
@compute @workgroup_size(8,4,1)
fn sorter() { }

@compute @workgroup_size(8u)
fn reverser() { }

// Using an pipeline-overridable constant.
@id(42) override block_width = 12u;
@compute @workgroup_size(block_width)
fn shuffler() { }

// Error: workgroup_size must be specified on compute shader
@compute
fn bad_shader() { }

```

--------------------------------

### WGSL: User-Defined Function Call Uniformity Example

Source: https://gpuweb.github.io/gpuweb/wgsl

Shows how user-defined function calls can affect uniformity analysis. In this example, the `scale` function's parameters are marked as `ParameterReturnContentsRequiredToBeUniform`, creating an invalid path in the analysis graph from the function's return value to a built-in value.

```WGSL
fn scale(in1 : f32, in2 : f32) -> f32 {
  let v = in1 / in2;
  return v;
}

@group(0) @binding(0) var t : texture_2d<f32>;
@group(0) @binding(1) var s : sampler;

@fragment
fn main(@builtin(position) pos : vec4<f32>) {
  let tmp = scale(pos.x, 0.5);
  if tmp > 1.0 {
    _ = textureSample(t, s, pos.xy);
  }
}
```

--------------------------------

### WGSL Discard Statement Example

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates the usage of the discard statement in a WGSL fragment shader. The discard statement is used conditionally to prevent a fragment from being processed further. Helper invocations, like those resulting from 'discard', do not write to shared memory. This example shows how 'discard' can prevent a color from being emitted.

```wgsl
@group(0) @binding(0)
var<storage, read_write> will_emit_color : u32;

fn discard_if_shallow(pos: vec4<f32>) {
  if pos.z < 0.001 {
    // If this is executed, then the will_emit_color variable will
    // never be set to 1 because helper invocations will not write
    // to shared memory.
    discard;
  }
  will_emit_color = 1;
}

@fragment
fn main(@builtin(position) coord_in: vec4<f32>)
  -> @location(0) vec4<f32>
{
  discard_if_shallow(coord_in);

  // Set the value to 1 and emit red, but only if the helper function
  // did not execute the discard statement.
  will_emit_color = 1;
  return vec4<f32>(1.0, 0.0, 0.0, 1.0);
}

```

--------------------------------

### WGSL Invalid textureSample Function Call Example

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates an invalid call to the textureSample built-in function within a conditional block that depends on a non-uniform value (fragment shader position). This highlights how non-uniform control flow can invalidate certain operations. The example also touches upon control flow returning to a uniform state after the conditional block.

```wgsl
@group(0) @binding(0) var t : texture_2d<f32>;
@group(0) @binding(1) var s : sampler;

@fragment
fn main(@builtin(position) pos : vec4<f32>) {
  if (pos.x < 0.5) {
    // Invalid textureSample function call.
    _ = textureSample(t, s, pos.xy);
  }
}

```

--------------------------------

### Behavior Analysis Rules for Statements

Source: https://gpuweb.github.io/gpuweb/wgsl

This table outlines the rules for analyzing and validating the behaviors of various statements in the GPUWeb specification. It details preconditions and the resulting behavior sets (e.g., {Next}, {Return}, {Break, Continue}) for each statement type.

```markdown
Statement | Preconditions | Resulting behavior   
---
---
---
_empty statement_ |  | {Next}
{s} | s: B | B
s1 s2 Note: s1 often ends in a semicolon. | s1: B1  
Next in B1  
s2: B2 | (B1\{Next}) ∪ B2
s1: B1  
Next not in B1  
s2: B2 | B1
var x:T; |  | {Next}
let x = e; |  | {Next}
var x = e; |  | {Next}
x = e; |  | {Next}
_ = e; |  | {Next}
f(e1, ..., en); | f has behavior B | B
return; |  | {Return}
return e; |  | {Return}
discard; |  | {Next}
break; |  | {Break}
break if e; |  | {Break, Next}
continue; |  | {Continue}
const_assert e; |  | {Next}
if e s1 else s2 |  s1: B1  
s2: B2 | B1 ∪ B2
loop {s1 continuing {s2}} |  s1: B1  
s2: B2  
B1 = {Return}  
None of {Continue, Return} are in B2  
| {Return}
s1: B1  
s2: B2  
B1 ≠ {Return}  
None of {Continue, Return} are in B2  
Break is not in (B1 ∪ B2) | (B1 ∪ B2)∖{Continue, Next}
s1: B1  
s2: B2  
B1 ≠ {Return}  
None of {Continue, Return} are in B2  
Break is in (B1 ∪ B2) | (B1 ∪ B2 ∪ {Next})∖{Break, Continue}
switch e {case c1: s1 ... case cn: sn} |  s1: B1  
...  
sn: Bn  
Break is not in (B1 ∪ ... ∪ Bn) | B1 ∪ ... ∪ Bn
s1: B1  
...  
sn: Bn  
Break is in (B1 ∪ ... ∪ Bn) | (B1 ∪ ... ∪ Bn ∪ {Next})∖Break
```

--------------------------------

### WGSL Function-scope Variable Uniformity Example

Source: https://gpuweb.github.io/gpuweb/wgsl

Illustrates the difference between valid and invalid barrier function calls based on the uniformity of function-scope variables. It shows that a workgroupBarrier call is invalid when 'x' is derived from a mutable module-scope variable ('a'), while a storageBarrier call is valid when 'x' is derived from an immutable module-scope variable ('b'). This example emphasizes the analysis's ability to track uniformity across a variable's lifetime and control flow.

```wgsl
@group(0) @binding(0) var<storage, read_write> a : i32;
@group(0) @binding(1) var<uniform> b : i32;

@compute @workgroup_size(16,1,1)
fn main() {
  var x : i32;
  x = a;
  if x > 0 {
    // Invalid barrier function call.
    workgroupBarrier();
  }
  x = b;
  if x < 0 {
    // Valid barrier function call.
    storageBarrier();
  }
}

```

--------------------------------

### WGSL Variable Declaration Syntax

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates the basic syntax for declaring variables in WGSL, including specifying address space and access mode.

```wgsl
var<private> my_variable: f32;
var<uniform, read> uniform_buffer: mat4x4<f32>;
var<storage, read_write> storage_buffer: array<u32>;
var<workgroup> workgroup_var: vec4<i32>;
var<function> function_var: bool = true;

```

--------------------------------

### Get Multisample Count (WGSL)

Source: https://gpuweb.github.io/gpuweb/wgsl

Returns the number of samples per texel in a multisampled texture. This is applicable to `texture_multisampled_2d` and `texture_depth_multisampled_2d` types.

```wgsl
@must_use fn textureNumSamples(t: T) -> u32
```

--------------------------------

### Rust: Invalid Infinite Loop

Source: https://gpuweb.github.io/gpuweb/wgsl

Shows an example of an invalid, empty infinite loop in Rust. Such loops are disallowed because they offer no path to termination.

```rust
fn invalid_infinite_loop() {
  loop { }     // Behavior: { }.  Invalid because it's empty.
}

```

--------------------------------

### WGSL Vector Type Declaration

Source: https://gpuweb.github.io/gpuweb/wgsl

Declares a vector type in WGSL, specifying the number of components and the component type. The example shows a vector of two 32-bit floating-point numbers.

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

```

--------------------------------

### Sample Texture 3D with Offset and Mip Level

Source: https://gpuweb.github.io/gpuweb/wgsl

Samples a 3D texture with an optional texel offset at a specified mip level. Requires a texture, sampler, 3D texture coordinates, mip level, and a 3D integer offset. Returns a 4-component floating-point vector.

```wgsl
@must_use fn textureSampleLevel(t: texture_3d<f32>,
                                s: sampler,
                                coords: vec3<f32>,
                                level: f32,
                                offset: vec3<i32>) -> vec4<f32>
```

--------------------------------

### WGSL subgroupBallot: Get active invocations mask

Source: https://gpuweb.github.io/gpuweb/wgsl

Returns a bitmask indicating which invocations in the subgroup have a true predicate. The result is a vec4<u32> where each component covers a range of 32 invocations.

```WGSL
@must_use fn subgroupBallot(pred : bool) -> vec4<u32>
```

--------------------------------

### WGSL: Struct with location attributes

Source: https://gpuweb.github.io/gpuweb/wgsl

Defines a structure with members explicitly assigned to different IO locations using the `@location` attribute. This example demonstrates how to structure data for shader inputs or outputs.

```WGSL
struct A {
  @location(0) x: f32,
  // Despite locations being 16-bytes, x and y cannot share a location
  @location(1) y: f32
}

// in1 occupies locations 0 and 1.
// in2 occupies location 2.
// The return value occupies location 0.
@fragment
fn fragShader(in1: A, @location(2) in2: f32) -> @location(0) vec4<f32> {
 // ...
}
```

--------------------------------

### Construct Structure (WGSL)

Source: https://gpuweb.github.io/gpuweb/wgsl

Constructs an instance of a structure `S` by providing its members in order. The types of the provided members must match the declared types of the structure's members.

```wgsl
@const @must_use fn S(e1 : T1, ..., eN : TN) -> S
```

--------------------------------

### Sample Texture 2D with Mip Level

Source: https://gpuweb.github.io/gpuweb/wgsl

Samples a 2D texture at a specified mip level. Requires a texture, sampler, 2D texture coordinates, and the mip level. Returns a 4-component floating-point vector.

```wgsl
@must_use fn textureSampleLevel(t: texture_2d<f32>,
                                s: sampler,
                                coords: vec2<f32>,
                                level: f32) -> vec4<f32>
```

--------------------------------

### WGSL: Const Assertion Statement Examples

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates the usage of const_assert in WGSL for compile-time checks. It shows valid and invalid uses at module scope and within functions, highlighting the requirement for const-expressions.

```wgsl
const x = 1;
const y = 2;
const_assert x < y; // valid at module-scope.
const_assert(y != 0); // parentheses are optional.

fn foo() {
  const z = x + y - 2;
  const_assert z > 0; // valid in functions.
  let a  = 3;
  const_assert a != 0; // invalid, the expression must be a const-expression.
}
```

--------------------------------

### WGSL Shader-Creation Error Example with Override Expressions

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates a shader-creation error in WGSL that occurs due to division by zero within an override expression. This error is independent of any API-provided override values.

```wgsl
override a : i32 = 0;
override b = a / 0; // shader-creation error,
                    // regardless of attempting to override c

```

--------------------------------

### Sample Texture 1D with Mip Level

Source: https://gpuweb.github.io/gpuweb/wgsl

Samples a 1D texture at a specified mip level. Requires a texture, sampler, texture coordinates, and the mip level. Returns a 4-component floating-point vector.

```wgsl
@must_use fn textureSampleLevel(t: texture_1d<f32>,
                                s: sampler,
                                coords: f32,
                                level: f32) -> vec4<f32>
```

--------------------------------

### Sample Texture 2D with Offset and Mip Level

Source: https://gpuweb.github.io/gpuweb/wgsl

Samples a 2D texture with an optional texel offset at a specified mip level. Requires a texture, sampler, 2D texture coordinates, mip level, and a 2D integer offset. Returns a 4-component floating-point vector.

```wgsl
@must_use fn textureSampleLevel(t: texture_2d<f32>,
                                s: sampler,
                                coords: vec2<f32>,
                                level: f32,
                                offset: vec2<i32>) -> vec4<f32>
```

--------------------------------

### Sample Texture 3D or Cube with Mip Level

Source: https://gpuweb.github.io/gpuweb/wgsl

Samples a 3D texture or cube texture at a specified mip level. Requires a texture, sampler, 3D texture coordinates, and mip level. Returns a 4-component floating-point vector.

```wgsl
@must_use fn textureSampleLevel(t: T,
                                s: sampler,
                                coords: vec3<f32>,
                                level: f32) -> vec4<f32>
```

--------------------------------

### WGSL Operator Precedence Corner Cases

Source: https://gpuweb.github.io/gpuweb/wgsl

Demonstrates invalid WGSL expressions due to incorrect operator precedence and associativity. These examples highlight situations where parentheses are required to enforce the intended order of operations.

```wgsl
let a = x & (y ^ (z | w)); // Invalid: x & y ^ z | w
let b = (x + y) << (z >= w); // Invalid: x + y << z >= w
let c = x < (y > z); // Invalid: x < y > z
let d = x && (y || z); // Invalid: x && y || z
```

--------------------------------

### Sample Texture Cube Array with Mip Level

Source: https://gpuweb.github.io/gpuweb/wgsl

Samples a cube texture array at a specified mip level. Requires a texture, sampler, 3D texture coordinates, array index, and mip level. Returns a 4-component floating-point vector.

```wgsl
@must_use fn textureSampleLevel(t: texture_cube_array<f32>,
                                s: sampler,
                                coords: vec3<f32>,
                                array_index: A,
                                level: f32) -> vec4<f32>
```

--------------------------------

### Get Texture Mip Levels (WGSL)

Source: https://gpuweb.github.io/gpuweb/wgsl

Retrieves the number of mipmap levels for a given texture. This function supports various texture types like 1D, 2D, 3D, cube, and their array/depth variants.

```wgsl
@must_use fn textureNumLevels(t: T) -> u32
```
Tokens:
Raw