Skip to content

Latest commit

 

History

History
132 lines (96 loc) · 5.07 KB

WGSL-forDummies.md

File metadata and controls

132 lines (96 loc) · 5.07 KB

Intro to WGSL

Most of the following document is written based on this (the January 12, 2023 draft).

WebGPU issues a unit of work to the GPU in the form of a GPU command. WGSL is concerned with two kinds of GPU commands:

  • a draw command executes a render pipeline in the context of inputs, outputs, and attached resources.
  • a dispatch command executes a compute pipeline in the context of inputs and attached resources.

Both kinds of pipelines use shaders written in WGSL.

General Notes

  • WGSL is an imperative language: behavior is specified as a sequence of statements to execute.
  • WGSL is statically typed: each value computed by a particular expression is in a specific type, determined only by examining the program source.
  • WGSL does not have implicit conversions or promotions from concrete types, but does provide implicit conversions and promotions from abstract types. Converting a value from one concrete numeric or boolean type to another requires an explicit conversion, construction, or reinterpretation of bits.
  • In a compute shader, invocations in the same workgroup share variables in the workgroup address space. Invocations in different workgroups do not share those variables.
  • Directives are optional. If present, all directives must appear before any declarations or static assertions.

Shader Life Cycle

  1. Shader module creation - This occurs when the WebGPU createShaderModule method is called. The source text for a WGSL program is provided at this time.

  2. Pipeline creation - This occurs when the WebGPU createComputePipeline method or the WebGPU createRenderPipeline method is invoked. These methods use one or more previously created shader modules, together with other configuration information.

  3. Shader execution start - This occurs when a draw or dispatch command is issued to the GPU, begins executing the pipeline, and invokes the shader stage entry point function.

  4. Shader execution end - This occurs when all work in the shader completes.

Variables and Declarations

For reference, see here.

Some examples on valid and invalid variable declaration:

// Valid, i_1 is in scope until the end of the for loop
for ( var i: i32 = 0; i < 10; i++ ) { // i_1
  // Invalid, i_2 has the same end scope as i_1.
  var i: i32 = 1; // i_2.
}

Variables out of functions should have scope:

var<private> bar: u32 = 1u;

fn my_func(foo: f32) {

}

Type inference

Variable definition should containt type, however it can be inferred by the initilizer litterals. For instance let x = 1 + 2.5; is the same as let x : f32 = 1.0f + 2.5f;

See more examples.

Attributes

Attribute descriptions.

attribute :
| '@' 'align' '(' expression attrib_end
| '@' 'binding' '(' expression attrib_end
| '@' 'builtin' '(' builtin_value_name attrib_end
| '@' 'const'
| '@' 'group' '(' expression attrib_end
| '@' 'id' '(' expression attrib_end
| '@' 'interpolate' '(' interpolation_type_name attrib_end
| '@' 'interpolate' '(' interpolation_type_name ',' interpolation_sample_name attrib_end
| '@' 'invariant'
| '@' 'location' '(' expression attrib_end
| '@' 'size' '(' expression attrib_end
| '@' 'workgroup_size' '(' expression attrib_end
| '@' 'workgroup_size' '(' expression ',' expression attrib_end
| '@' 'workgroup_size' '(' expression ',' expression ',' expression attrib_end
| '@' 'vertex'
| '@' 'fragment'
| '@' 'compute'

Bind group layout and bind group

A bind group layout defines the input/output interface expected by a shader, while a bind group represents the actual input/output data for a shader.

Shader module creation with WGSL

const shaderModule = device.createShaderModule({
  code: `
    struct Matrix {
      size : vec2<f32>,
      numbers: array<f32>,
    }

    @group(0) @binding(0) var<storage, read> firstMatrix : Matrix;
    @group(0) @binding(1) var<storage, read> secondMatrix : Matrix;
    @group(0) @binding(2) var<storage, read_write> resultMatrix : Matrix;

    @compute @workgroup_size(8, 8)
    fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
      // Guard against out-of-bounds work group sizes
      if (global_id.x >= u32(firstMatrix.size.x) || global_id.y >= u32(secondMatrix.size.y)) {
        return;
      }

      resultMatrix.size = vec2(firstMatrix.size.x, secondMatrix.size.y);

      let resultCell = vec2(global_id.x, global_id.y);
      var result = 0.0;
      for (var i = 0u; i < u32(firstMatrix.size.y); i = i + 1u) {
        let a = i + resultCell.x * u32(firstMatrix.size.y);
        let b = resultCell.y + i * u32(secondMatrix.size.y);
        result = result + firstMatrix.numbers[a] * secondMatrix.numbers[b];
      }

      let index = resultCell.y + resultCell.x * u32(secondMatrix.size.y);
      resultMatrix.numbers[index] = result;
    }
  `
});

References