Advanced

This page is about advanced topics. If you haven't read the basics this won't make sense.

Unfinished page!

This page is unfinished, it may contain inaccurate, wrong, or invalid information/code.

Topics covered here:

Atomics and thread safety stuff
Textures
Alignment and padding(unfinished)
Workgroups and threads

Atomics and thread safety

An atomic is a variable that is shared across a workgroup, and when written to guarantees that multiple threads don't write to it at the same time, resulting in data loss. This is referred to as a race condition. Atomics can either be a signed or unsigned 32 bit integer(i32 or u32). Here's what the block to construct them looks like:

    Create atomic of type (type v) :: #4287f5

The type input here is either i32 or u32. You can use this in the declare variable block like this:

    Declare variable (Variable usage (workgroup v) next [] :: #4287f5) variable as [sum] with value []: (Create atomic of type (i32 v) :: #4287f5) :: #4287f5

Let's go through each of these inputs one at a time. The first input where we put the variable usage block says that this variable will be shared across all threads in a workgroup. We named it sum, and we don't initialize it by putting nothing in the input. Then in the next input we say that the type is our i32 atomic. Let's make a simple shader that finds the sum of an array of inputs, and divides it by the thread id.

This is just some random shader I put together in 10 minutes to test atomics, it doesn't actually have any purpose.

Starting from the scratch code, we'll need 3 buffers: one for providing an input to the shader, one for the shader to write to, and a third to read the shader's output. This bind group layout looks like this:

    Create bind group layout called [myBindGroupLayout] {
        Add bind group layout entry with binding [0] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (read-only-storage v) :: #4287f5) :: #4287f5 
        Add bind group layout entry with binding [1] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (storage v) :: #4287f5) :: #4287f5 
    } :: #4287f5

The buffer at binding 0 will be our shader's input buffer, binding 1 will be the output buffer. Next we need to bind our buffers, which look like this:

    Create buffer called [input] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [work] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_SRC v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [output] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (MAP_READ v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5

Our input buffer can be read from via the STORAGE flag, and we can copy data to it using COPY_DST. The work buffer can be written to, and we can copy data from it using COPY_SRC. Finally the output buffer can be read from using MAP_READ, and we can copy data to it. Next we bind the buffers:

    Create bind group called [myBindGroup] using layout [myBindGroupLayout] {
        Add bind group entry with binding [0] of type (buffer v) using resource named [input] :: #4287f5
        Add bind group entry with binding [1] of type (buffer v) using resource named [work] :: #4287f5
    } :: #4287f5

Now we can work on our shader. We start by binding our resources:

    Def shader [myShader] using bind group layout [myBindGroupLayout] :: hat #4287f5
    Bind shader resource # [0] to variable [input] with settings (Variable usage (storage v) next (Variable usage (read v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5) :: #4287f5) :: #4287f5
    Bind shader resource # [1] to variable [output] with settings (Variable usage (storage v) next (Variable usage (read_write v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v), length\(array only!\) [] :: #4287f5) :: #4287f5) :: #4287f5

Now, atomic variables need to be declared in uniform control flow. This means it can't be inside of any functions, if statements or anything. We'll declare our i32 atomic, which we will use to store the sum of the values.

    Def shader [myShader] using bind group layout [myBindGroupLayout] :: hat #4287f5
    Bind shader resource # [0] to variable [input] with settings (Variable usage (storage v) next (Variable usage (read v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v), length\(array only!\) [] :: #4287f5) :: #4287f5) :: #4287f5
    Bind shader resource # [1] to variable [output] with settings (Variable usage (storage v) next (Variable usage (read_write v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5), length\(array only!\) [] :: #4287f5) :: #4287f5
    Declare variable (Variable usage (workgroup v) next [] :: #4287f5) variable as [sum] with value []: (Create atomic of type (i32 v) :: #4287f5) :: #4287f5
    Compute shader with workgroup size [\[1\]] {

    } :: #4287f5

Now, we need to initialize our variable or it will lead to everything being undefined. We could just set it to 0 and go on with our day, but that would cause a race condition. We can fix this using barriers. A barrier is a function that will keep a thread from reading and writing to something until all other threads in the workgroup have reached that point. There's 2 types:

    Barrier (storageBarrier v) :: #4287f5

    Barrier (workgroupBarrier v) :: #4287f5

storageBarrier - A storage barrier pauses threads from reading/writing to a storage buffer until all the threads haved reached the barrier, at which point they continue.

workgroupBarrier - A workgroup barrier pauses threads from reading/writing to a workgroup variable until all the threads have reached the barrier, at which point they continue.

That's cool and all, but first we need to go over how to use atomic variables. Atomics have built in functions to perform operations in a thread safe way. If you try to set the variable without using these functions, you're defeating the whole purpose of atomics.

    Load atomic [atomic] :: reporter #4287f5
    Perform operation (atomicOp v) on atomic [atomic] with value [value] :: #4287f5
    Perform operation (atomicOp v) on atomic [atomic] with value [value] :: reporter #4287f5

Load atomic:

This block returns the value of the atomic specified.

Perform operation:

This performs the operation specified on the atomic using the value. For example, atomicStore sets the value of the atomic to the value. There's a lot of these and I'm lazy, so please read the list of functions.

All of the atomic operations return the original value of the atomic before modification.

Note: all of the "atomic" inputs take a pointer block! You will see a really long and complicated error otherwise.

    Pointer to variable [] :: reporter #4287f5

We'll set the value of the atomic to 0 and then have a workgroup barrier to wait until the variable has been initialized.

    Def shader [myShader] using bind group layout [myBindGroupLayout] :: hat #4287f5
    Bind shader resource # [0] to variable [input] with settings (Variable usage (storage v) next (Variable usage (read v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5), length\(array only!\) [] :: #4287f5) :: #4287f5
    Bind shader resource # [1] to variable [output] with settings (Variable usage (storage v) next (Variable usage (read_write v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5), length\(array only!\) [] :: #4287f5) :: #4287f5
    Declare variable (Variable usage (workgroup v) next [] :: #4287f5) variable as [sum] with value []: (Create atomic of type (i32 v) :: #4287f5) :: #4287f5
    Compute shader with workgroup size [\[1\]] {
        Perform operation (atomicStore v) on atomic (Pointer to variable [sum] :: #4287f5) with value [0] :: #4287f5
        Barrier (workgroupBarrier v) :: #4287f5
    } :: #4287f5

Next we have each thread get their item from the input buffer, and add it to the sum. We'll need to cast it to an i32, and we can get the workgroup id using the (hidden) workgroup_id variable.

    Def shader [myShader] using bind group layout [myBindGroupLayout] :: hat #4287f5
    Bind shader resource # [0] to variable [input] with settings (Variable usage (storage v) next (Variable usage (read v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5), length\(array only!\) [] :: #4287f5) :: #4287f5
    Bind shader resource # [1] to variable [output] with settings (Variable usage (storage v) next (Variable usage (read_write v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5), length\(array only!\) [] :: #4287f5) :: #4287f5
    Declare variable (Variable usage (workgroup v) next [] :: #4287f5) variable as [sum] with value []: (Create atomic of type (i32 v) :: #4287f5) :: #4287f5
    Compute shader with workgroup size [\[1\]] {
        Perform operation (atomicStore v) on atomic (Pointer to variable [sum] :: #4287f5) with value [0] :: #4287f5
        Barrier (workgroupBarrier v) :: #4287f5
        Perform operation (atomicAdd v) on atomic (Pointer to variable [sum] :: #4287f5) with value (WGSL builtin (i32 v) with args (In object [input] get index (In object [workgroup_id] get property [x] :: #4287f5) :: #4287f5) :: #4287f5) :: #4287f5
    } :: #4287f5

After we've done that, we can add a storage barrier(I think a workgroup barrier would work here too because we're both reading from a buffer and writing to an atomic.) so we wait until the sum is finished, then set the respective item in the output buffer to the sum divided by the index(plus one so we don't get zero division issues). We'll need to cast everything to floats.

    Def shader [myShader] using bind group layout [myBindGroupLayout] :: hat #4287f5
    Bind shader resource # [0] to variable [input] with settings (Variable usage (storage v) next (Variable usage (read v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5), length\(array only!\) [] :: #4287f5) :: #4287f5
    Bind shader resource # [1] to variable [output] with settings (Variable usage (storage v) next (Variable usage (read_write v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5), length\(array only!\) [] :: #4287f5) :: #4287f5
    Declare variable (Variable usage (workgroup v) next [] :: #4287f5) variable as [sum] with value []: (Create atomic of type (i32 v) :: #4287f5) :: #4287f5
    Compute shader with workgroup size [\[1\]] {
        Perform operation (atomicStore v) on atomic (Pointer to variable [sum] :: #4287f5) with value [0] :: #4287f5
        Barrier (workgroupBarrier v) :: #4287f5
        Perform operation (atomicAdd v) on atomic (Pointer to variable [sum] :: #4287f5) with value (WGSL builtin (i32 v) with args (In object [input] get index (In object [workgroup_id] get property [x] :: #4287f5) :: #4287f5) :: #4287f5) :: #4287f5
        Barrier (storageBarrier v) :: #4287f5
        Variable (In object [output] get index (In object [workgroup_id] get property [x] :: #4287f5) :: #4287f5) (= v) (  (WGSL builtin (f32 v) with args (Load atomic (Pointer to variable [sum] :: #4287f5) :: #4287f5) :: #4287f5) / (( WGSL builtin (f32 v) with args (In object [workgroup_id] get property [x] :: #4287f5) :: #4287f5) + [1.0])  ) :: #4287f5
    } :: #4287f5

That's our shader finished. Now we need to finish the data management side, which isn't too complicated. Right now we have:

    Create buffer called [input] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [work] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_SRC v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [output] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (MAP_READ v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create bind group layout called [myBindGroupLayout] {
        Add bind group layout entry with binding [0] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (read-only-storage v) :: #4287f5) :: #4287f5 
        Add bind group layout entry with binding [1] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (storage v) :: #4287f5) :: #4287f5 
    } :: #4287f5
    Create bind group called [myBindGroup] using layout [myBindGroupLayout] {
        Add bind group entry with binding [0] of type (buffer v) using resource named [input] :: #4287f5
        Add bind group entry with binding [1] of type (buffer v) using resource named [work] :: #4287f5
    } :: #4287f5

Our data is going to be an array of random numbers, which we can make using the turbowarp JSON extension:

I highly recommend using a different extension, this one is really slow because it parses/stringifies the object in every block!

    set [input data v] to [\[\]]
    repeat [32] {
        set [input data v] to (add (pick random [-15] to [15]) to array (input data) :: #3271D0)
    } :: control

Next we can create an arraybuffer and view from this data and copy it to our input buffer:

    set [input data v] to [\[\]]
    repeat [32] {
        set [input data v] to (add (pick random [-15] to [15]) to array (input data) :: #3271D0)
    } :: control
    Create arraybuffer and view called [data] from array (input data) of type (Float32Array v) :: #4287f5
    Write [32] elements of data from arraybuffer [data] to buffer [input] from offset [0] to offset [0] :: #4287f5

We can add that to our code, along with compiling and running the shader:

    Create buffer called [input] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [work] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_SRC v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [output] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (MAP_READ v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create bind group layout called [myBindGroupLayout] {
        Add bind group layout entry with binding [0] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (read-only-storage v) :: #4287f5) :: #4287f5 
        Add bind group layout entry with binding [1] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (storage v) :: #4287f5) :: #4287f5 
    } :: #4287f5
    Create bind group called [myBindGroup] using layout [myBindGroupLayout] {
        Add bind group entry with binding [0] of type (buffer v) using resource named [input] :: #4287f5
        Add bind group entry with binding [1] of type (buffer v) using resource named [work] :: #4287f5
    } :: #4287f5
    set [input data v] to [\[\]]
    repeat [32] {
        set [input data v] to (add (pick random [-15] to [15]) to array (input data) :: #3271D0)
    } :: control
    Create arraybuffer and view called [data] from array (input data) of type (Float32Array v) :: #4287f5
    Write [32] elements of data from arraybuffer [data] to buffer [input] from offset [0] to offset [0] :: #4287f5
    compile shaders :: #4287f5
    Run shader [myShader] using bind group [myBindGroup] dimensions x: [32] y: [1] z: [1] :: #4287f5

You don't need to write the data to a buffer before binding it, hence why our number generation code is right before compiling and running the shader.

Finally we can do the same thing we did in the basic doubling shader where we copy data from the work buffer to the read buffer, and read that:

    Copy ((32) * (4)) bytes of data from buffer [work] from position [0] to buffer [output] at position [0] :: #4287f5
    Read buffer [output] to arraybuffer [outputArrayBuffer] :: #4287f5
    View arraybuffer (outputArrayBuffer v) as (Float32Array v) called [outputView] :: #4287f5
    set [out v] to (Get view [outputView] as array :: #4287f5)

That's everything finished! If you run this you should see a bunch of weird numbers, here's all the code put together:

    Def shader [myShader] using bind group layout [myBindGroupLayout] :: hat #4287f5
    Bind shader resource # [0] to variable [input] with settings (Variable usage (storage v) next (Variable usage (read v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5), length\(array only!\) [] :: #4287f5) :: #4287f5
    Bind shader resource # [1] to variable [output] with settings (Variable usage (storage v) next (Variable usage (read_write v) :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5), length\(array only!\) [] :: #4287f5) :: #4287f5
    Declare variable (Variable usage (workgroup v) next [] :: #4287f5) variable as [sum] with value []: (Create atomic of type (i32 v) :: #4287f5) :: #4287f5
    Compute shader with workgroup size [\[1\]] {
        Perform operation (atomicStore v) on atomic (Pointer to variable [sum] :: #4287f5) with value [0] :: #4287f5
        Barrier (workgroupBarrier v) :: #4287f5
        Perform operation (atomicAdd v) on atomic (Pointer to variable [sum] :: #4287f5) with value (WGSL builtin (i32 v) with args (In object [input] get index (In object [workgroup_id] get property [x] :: #4287f5) :: #4287f5) :: #4287f5) :: #4287f5
        Barrier (storageBarrier v) :: #4287f5
        Variable (In object [output] get index (In object [workgroup_id] get property [x] :: #4287f5) :: #4287f5) (= v) (  (WGSL builtin (f32 v) with args (Load atomic (Pointer to variable [sum] :: #4287f5) :: #4287f5) :: #4287f5) / (( WGSL builtin (f32 v) with args (In object [workgroup_id] get property [x] :: #4287f5) :: #4287f5) + [1.0])  ) :: #4287f5
    } :: #4287f5

    when flag clicked
    Create buffer called [input] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [work] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_SRC v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [output] with size\(in bytes\) ([32] * [4]) and usage flags (Usage (Buffer usage (MAP_READ v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create bind group layout called [myBindGroupLayout] {
        Add bind group layout entry with binding [0] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (read-only-storage v) :: #4287f5) :: #4287f5 
        Add bind group layout entry with binding [1] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (storage v) :: #4287f5) :: #4287f5 
    } :: #4287f5
    Create bind group called [myBindGroup] using layout [myBindGroupLayout] {
        Add bind group entry with binding [0] of type (buffer v) using resource named [input] :: #4287f5
        Add bind group entry with binding [1] of type (buffer v) using resource named [work] :: #4287f5
    } :: #4287f5
    set [input data v] to [\[\]]
    repeat [32] {
        set [input data v] to (add (pick random [-15] to [15]) to array (input data) :: #3271D0)
    } :: control
    Create arraybuffer and view called [data] from array (input data) of type (Float32Array v) :: #4287f5
    Write [32] elements of data from arraybuffer [data] to buffer [input] from offset [0] to offset [0] :: #4287f5
    compile shaders :: #4287f5
    Run shader [myShader] using bind group [myBindGroup] dimensions x: [32] y: [1] z: [1] :: #4287f5
    Copy ((32) * (4)) bytes of data from buffer [work] from position [0] to buffer [output] at position [0] :: #4287f5
    Read buffer [output] to arraybuffer [outputArrayBuffer] :: #4287f5
    View arraybuffer (outputArrayBuffer v) as (Float32Array v) called [outputView] :: #4287f5
    set [out v] to (Get view [outputView] as array :: #4287f5)

SB3: here

Textures

Part of the appeal of compute shaders is the ability to make fast post processing effects. You could copy texture data to a buffer and read from that, but it's faster to use built in texture which I'll show here by making a simple shader to read a pixel from a costume. To start off, we need to create the bind group layout and stuff.

    Add bind group layout entry with binding [0] for type (storageTexture v) and descriptor (Texture layout entry descriptor with usage type (type v) and format (format v) :: reporter #4287f5) :: #4287f5

This code does a couple of things. As with buffer entries, it has its own descriptor for a specific type. The texture descriptor has the following inputs:

type - This describes how the texture will be used. Possible values:

read-only. We're only reading, so we'll use this.
write-only
read-write

format - The color format to use. I have arbitrarily chosen rgba8unorm, meaning that there's 8 bits per channel, and each channel is unsigned(positive only) and normalized(0-1). For information on these different settings, either google it or look at the webgpu spec.

We'll also need a buffer, so our bind group layout looks like this:

    Create bind group layout called [myBindGroupLayout] {
        Add bind group layout entry with binding [0] for type (storageTexture v) and descriptor (Texture layout entry descriptor with usage type (read-only v) and format (rgba8unorm v) :: reporter #4287f5) :: #4287f5
        Add bind group layout entry with binding [1] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (storage v) :: #4287f5) :: #4287f5
    } :: #4287f5

Next we need to actually create our data. We can use the create texture block for this:

    Create texture called [texture] with dimensions [width] [height], color format (colorFormat v) and usage [usage] :: #4287f5
    Texture usage (usage v) :: reporter #4287f5

Create texture block - Creates a new texture

texture - The name of the texture to create
width/height - The size of the texture, in pixels
colorFormat - The number of channels and bits per channel. There's a lot, see the webgpu spec.
usage - Similar to the usage input for the create buffer block, this takes in usage flags.

Texture usage flag block - A flag describing how a texture will be used

usage - This block's flag. Possible values:

COPY_SRC - This can be the source texture in the copy texture to buffer block.
COPY_DST - This can be the destination texture in the write texture data block.
TEXTURE_BINDING - I don't think this is usable in compute shaders, but it's here just in case.
STORAGE_BINDING - This lets you bind this texture in a bind group.

We can use the looks+ extension to get the size of the texture we're using. Since we just want to be able to copy data to the texture and bind it to the shader, we will have a usage of STORAGE_BINDING | COPY_DST. We'll also need a buffer to write our texture data into, and to read from:

    Create texture called [myTexture] with dimensions ([width v] of (costume2 v) :: #9966FF) ([height v] of (costume2 v) :: #9966FF), color format (rgba8unorm v) and usage (Usage (Texture usage (STORAGE_BINDING v) :: #4287f5) | (Texture usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [workBuffer] with size\(in bytes\) (16) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_SRC v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [readBuffer] with size\(in bytes\) (16) and usage flags (Usage (Buffer usage (MAP_READ v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5

And we can create a bind group:

    Create bind group layout called [myBindGroupLayout] {
        Add bind group layout entry with binding [0] for type (storageTexture v) and descriptor (Texture layout entry descriptor with usage type (read-only v) and format (rgba8unorm v) :: reporter #4287f5) :: #4287f5
        Add bind group layout entry with binding [1] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (storage v) :: #4287f5) :: #4287f5
    } :: #4287f5
    Create texture called [myTexture] with dimensions ([width v] of (costume2 v) :: #9966FF) ([height v] of (costume2 v) :: #9966FF), color format (rgba8unorm v) and usage (Usage (Texture usage (STORAGE_BINDING v) :: #4287f5) | (Texture usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [workBuffer] with size\(in bytes\) (16) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_SRC v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [readBuffer] with size\(in bytes\) (16) and usage flags (Usage (Buffer usage (MAP_READ v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create bind group called [myBindGroup] using layout [myBindGroupLayout] {
        Add bind group entry with binding [0] of type (storageTexture v) using resource named [myTexture] :: #4287f5
        Add bind group entry with binding [1] of type (buffer v) using resource named [workBuffer] :: #4287f5
    } :: #4287f5

Next we can write the data from costume2 to our texture:

    Create bind group layout called [myBindGroupLayout] {
        Add bind group layout entry with binding [0] for type (storageTexture v) and descriptor (Texture layout entry descriptor with usage type (read-only v) and format (rgba8unorm v) :: reporter #4287f5) :: #4287f5
        Add bind group layout entry with binding [1] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (storage v) :: #4287f5) :: #4287f5
    } :: #4287f5
    Create texture called [myTexture] with dimensions ([width v] of (costume2 v) :: #9966FF) ([height v] of (costume2 v) :: #9966FF), color format (rgba8unorm v) and usage (Usage (Texture usage (STORAGE_BINDING v) :: #4287f5) | (Texture usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [workBuffer] with size\(in bytes\) (16) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_SRC v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [readBuffer] with size\(in bytes\) (16) and usage flags (Usage (Buffer usage (MAP_READ v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create bind group called [myBindGroup] using layout [myBindGroupLayout] {
        Add bind group entry with binding [0] of type (storageTexture v) using resource named [myTexture] :: #4287f5
        Add bind group entry with binding [1] of type (buffer v) using resource named [workBuffer] :: #4287f5
    } :: #4287f5
    Write texture data from (costume2 v) to texture [myTexture] :: #4287f5

And we can add our code to compile/run the shader and get the output:

    Create bind group layout called [myBindGroupLayout] {
        Add bind group layout entry with binding [0] for type (storageTexture v) and descriptor (Texture layout entry descriptor with usage type (read-only v) and format (rgba8unorm v) :: reporter #4287f5) :: #4287f5
        Add bind group layout entry with binding [1] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (storage v) :: #4287f5) :: #4287f5
    } :: #4287f5
    Create texture called [myTexture] with dimensions ([width v] of (costume2 v) :: #9966FF) ([height v] of (costume2 v) :: #9966FF), color format (rgba8unorm v) and usage (Usage (Texture usage (STORAGE_BINDING v) :: #4287f5) | (Texture usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [workBuffer] with size\(in bytes\) (16) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_SRC v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [readBuffer] with size\(in bytes\) (16) and usage flags (Usage (Buffer usage (MAP_READ v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create bind group called [myBindGroup] using layout [myBindGroupLayout] {
        Add bind group entry with binding [0] of type (storageTexture v) using resource named [myTexture] :: #4287f5
        Add bind group entry with binding [1] of type (buffer v) using resource named [workBuffer] :: #4287f5
    } :: #4287f5
    Write texture data from (costume2 v) to texture [myTexture] :: #4287f5
    compile shaders :: #4287f5
    Run shader [myShader] using bind group [myBindGroup] dimensions x: [1] y: [1] z: [1] :: #4287f5
    Copy [16] bytes of data from buffer [workBuffer] from position [0] to buffer [readBuffer] at position [0] :: #4287f5
    Read buffer [readBuffer] to arraybuffer [myArrayBuffer] :: #4287f5
    View arraybuffer (myArrayBuffer v) as (Float32Array v) called [myView] :: #4287f5
    set [my variable v] to (Get view [myView] as array :: #4287f5)

Now we can start on our shader. We can start by binding the texture in slot 1 to a variable, which we'll call myTex

    Bind shader resource #[0] to variable [myTex] with settings [] type (Texture type of (rgba8unorm v) with access (read v) :: #4287f5) :: #4287f5

Small note here: The settings input is blank because that makes the binding become var instead of the usual var<usage settings>. Next we bind the buffer:

    Bind shader resource #[0] to variable [myTex] with settings [] type (Texture type of (rgba8unorm v) with access (read v) :: #4287f5) :: #4287f5
    Bind shader resource #[1] to variable [workBuffer] with settings (Variable usage (storage v) next (Variable usage (read_write v) next [] :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5), length\(array only!\) [] :: #4287f5) :: #4287f5

Now in our compute shader, we can declare a variable and set its value to the wgsl builtin function textureLoad. This returns the color data differently depending on the color format we set earlier, which in this case is rgba8unorm so it returns a vec4 of floats. It takes an argument of a vec2<i32> or vec2<u32>, and it ends up looking like this:

    Compute shader with workgroup size [\[1\]] {
        declare (let v) variable with value (WGSL builtin (textureLoad v) with args (Function arg input [someText], next (Function arg input (Construct type (Create type (vec2 v) of (Base type (i32 v) :: #4287f5) :: #4287f5) with values [0] :: #4287f5), next [] :: #4287f5) :: #4287f5) :: #4287f5) :: #4287f5
    } :: #4287f5

Next we can set the data in our output buffer:

    Compute shader with workgroup size [\[1\]] {
        declare (let v) variable as [c] with value (WGSL builtin (textureLoad v) with args (Function arg input [someText], next (Function arg input (Construct type (Create type (vec2 v) of (Base type (i32 v) :: #4287f5) :: #4287f5) with values [0] :: #4287f5), next [] :: #4287f5) :: #4287f5) :: #4287f5) :: #4287f5
        Variable (In object [workBuffer] get index [0] :: #4287f5) (= v) (In object [c] get property [r] :: #4287f5) :: #4287f5
        Variable (In object [workBuffer] get index [1] :: #4287f5) (= v) (In object [c] get property [g] :: #4287f5) :: #4287f5
        Variable (In object [workBuffer] get index [2] :: #4287f5) (= v) (In object [c] get property [b] :: #4287f5) :: #4287f5
        Variable (In object [workBuffer] get index [3] :: #4287f5) (= v) (In object [c] get property [a] :: #4287f5) :: #4287f5
    } :: #4287f5

That's our code finished! If we put it all together it looks like this:

    Bind shader resource #[0] to variable [myTex] with settings [] type (Texture type of (rgba8unorm v) with access (read v) :: #4287f5) :: #4287f5
    Bind shader resource #[1] to variable [workBuffer] with settings (Variable usage (storage v) next (Variable usage (read_write v) next [] :: #4287f5) :: #4287f5) type (Create type (array v) of (Base type (f32 v) :: #4287f5), length\(array only!\) [] :: #4287f5) :: #4287f5
    Compute shader with workgroup size [\[1\]] {
        declare (let v) variable as [c] with value (WGSL builtin (textureLoad v) with args (Function arg input [someText], next (Function arg input (Construct type (Create type (vec2 v) of (Base type (i32 v) :: #4287f5) :: #4287f5) with values [0] :: #4287f5), next [] :: #4287f5) :: #4287f5) :: #4287f5) :: #4287f5
        Variable (In object [workBuffer] get index [0] :: #4287f5) (= v) (In object [c] get property [r] :: #4287f5) :: #4287f5
        Variable (In object [workBuffer] get index [1] :: #4287f5) (= v) (In object [c] get property [g] :: #4287f5) :: #4287f5
        Variable (In object [workBuffer] get index [2] :: #4287f5) (= v) (In object [c] get property [b] :: #4287f5) :: #4287f5
        Variable (In object [workBuffer] get index [3] :: #4287f5) (= v) (In object [c] get property [a] :: #4287f5) :: #4287f5
    } :: #4287f5

    when flag clicked
    Create bind group layout called [myBindGroupLayout] {
        Add bind group layout entry with binding [0] for type (storageTexture v) and descriptor (Texture layout entry descriptor with usage type (read-only v) and format (rgba8unorm v) :: reporter #4287f5) :: #4287f5
        Add bind group layout entry with binding [1] for type (buffer v) and descriptor (Buffer layout entry descriptor with usage type (storage v) :: #4287f5) :: #4287f5
    } :: #4287f5
    Create texture called [myTexture] with dimensions ([width v] of (costume2 v) :: #9966FF) ([height v] of (costume2 v) :: #9966FF), color format (rgba8unorm v) and usage (Usage (Texture usage (STORAGE_BINDING v) :: #4287f5) | (Texture usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [workBuffer] with size\(in bytes\) (16) and usage flags (Usage (Buffer usage (STORAGE v) :: #4287f5) | (Buffer usage (COPY_SRC v) :: #4287f5) :: #4287f5) :: #4287f5
    Create buffer called [readBuffer] with size\(in bytes\) (16) and usage flags (Usage (Buffer usage (MAP_READ v) :: #4287f5) | (Buffer usage (COPY_DST v) :: #4287f5) :: #4287f5) :: #4287f5
    Create bind group called [myBindGroup] using layout [myBindGroupLayout] {
        Add bind group entry with binding [0] of type (storageTexture v) using resource named [myTexture] :: #4287f5
        Add bind group entry with binding [1] of type (buffer v) using resource named [workBuffer] :: #4287f5
    } :: #4287f5
    Write texture data from (costume2 v) to texture [myTexture] :: #4287f5
    compile shaders :: #4287f5
    Run shader [myShader] using bind group [myBindGroup] dimensions x: [1] y: [1] z: [1] :: #4287f5
    Copy [16] bytes of data from buffer [workBuffer] from position [0] to buffer [readBuffer] at position [0] :: #4287f5
    Read buffer [readBuffer] to arraybuffer [myArrayBuffer] :: #4287f5
    View arraybuffer (myArrayBuffer v) as (Float32Array v) called [myView] :: #4287f5
    set [my variable v] to (Get view [myView] as array :: #4287f5)

SB3 here

You might be wondering why you can't just set the type of the output buffer to an array of vec4s. This brings us into our next topic:

Alignment and padding

This is complicated and I can't be bothered explaining it right now. Go read the webgpufundamentals page

Workgroups and Threads

In the compute shader block there's an input for the workgroup size. This is the number of threads in a workgroup. A workgroup is a set of threads that is able to share data via workgroup variables, which we talked about earlier in the atomic variables section. Here's a visual aid I stole from webgpufundamentals:

this is supposed to be a visual aid but you're kinda screwed cuz you can't see it

As you can see the size of the workgroup is described in three dimensions. Now, in the run shader block you can describe the dispatch dimensions. This is the number of workgroups to create on each dimension. I'm not sure how to describe this in any other way, so here's another stolen visual aid

Each unit on the dispatch is made up of the previously described workgroup threads. Your compute shader comes with some builtin positions in the dispatch/workgroup:

vec3<u32> local_invocation_id - The position of the thread within the workgroup

vec3<u32> workgroup_id - The position of the workgroup within the dispatch

u32 local_invocation_index - local_invocation_id linearized

vec3<u32> global_invocation_id - Freaky thing that is apparently equal to workgroup_id * workgroup_size + local_invocation_id

vec3<u32> num_workgroups - The number of workgroups on each axis