WGSL loop and while that should behave the same behave differently

[suguruwataru]

I'm not familiar with synchronization yet so I'm not sure whether it's WGSL being weird or I'm having wrong expectations.

A friendly warning that running the test code below may make your computer hang.

Consider these compute shaders, where an atomic u32 is used as a mutex with atomicOring with 1 locking it and atomicAnding with 0 unlocking it:

var<workgroup> lock: atomic<u32>;
@group(0) @binding(0)
var<storage, read_write> ctr: u32;

@compute @workgroup_size(2, 1, 1)
fn loop_atomic(@builtin(global_invocation_id) id: vec3u) {
    loop {
        if atomicOr(&lock, 1u) == 1u { continue; }
        ctr++;
        atomicAnd(&lock, 0u);
        break;
    }
}

@compute @workgroup_size(2, 1, 1)
fn while_atomic(@builtin(global_invocation_id) id: vec3u) {
    while atomicOr(&lock, 1u) == 1u { }
    ctr++;
    atomicAnd(&lock, 0u);
}

I expect the 2 shaders to do the same thing: each starts 2 threads, and the threads each acquires the lock, add to the ctr variable, release the lock, and finish. And in CPU code I check whether the ctr variable actually reaches 2. loop_atomic does this. However, while_atomic hangs my computer for several seconds, and in this case ctr is zero when checked in CPU.

I'm using wgpu 25.0.0 and am running on Linux.

I've put my testing code here: https://github.com/suguruwataru/wgsl-loop. I also tried to recreate the logic in Rust in the testing code. They work like the loop case, not the while case. Each case is written as a test, cargo is configured to run all the tests with only one thread, and the WGSL while case is #[ignore], so one can safely run cargo test and see how the other cases go. To run the while case, which may make your computer hang, cargo test -- --ignored.

[cwfitzgerald]

So there's actually a couple problems here. It's actually impossible to write a working mutex in wgsl for two reasons:

• Atomics in wgsl have "Relaxed" memory semantics. This means that the visibility of atomic operations say nothing about the values of the rest of memory. In this case, there is no guarentee that ctr++ will actually do what you intend it to, and may end up reading stale values as there's nothing here forcing the caches to be flushed and memory to be visible to other threads. I think you're seeing a consistent value as it's a single workgroup (which share L1$) doing the addition. Split this across multiple workgroups with multiple L1$, and the output value will likely be garbage.
• There is no guarentee of forward progress in any graphics shading language. If you have one workgroup (group A) where the mutex is locked and one (group B) waiting on it to unlock, there's no guarentee that group A will ever be scheduled on the GPU ever again, the gpu is allowed to spin on B, waiting for it to finish before getting to A. In practice different gpus allow different amounts of forward progress, there was a 2021 paper discussing this in practice.

Then there's control flow issues which I think are causing your problem:

• This is a bit more unclear and may explain why this exact case works. You need to look at how the compiler will compile the code onto the SIMD unit that is the gpu. If this is new to you, I'd recommend looking at this slideshow by one of the WebGPU spec devs for an overview. If the shader compiler compiles the loop into what amounts to an if/else, it will evaluate both sides of the if statement every iteration. The while example however it will try to wait until all values in the warp satisfy the condition before moving on, which due to mutual exclusion can't happen. This is a rather complicated topic, so I am probably missing a bunch in this explanation, and I would be willing to chat about it on the matrix if you want to know more.

[suguruwataru]

Thank you very much for the explanation! Now I get what a "relaxed memory ordering" implicates... In the past the memory orderings are all very vague to me.

I do have some quick follow-up questions:

• If a mutex is not possible, how do GPU programmers usually implement synchronization for the use cases of a mutex? I wonder how a combination of atomics and barriers can achieve that... Or do they just think of algorithms that don't use a mutex?
• If forward progress is not guaranteed across wrokgroups, what about the same workgroup? I have the impression that threads in a single workgroup run in real parallelism, and want to confirm that.

And about the last point, I think I'll need some time to digest it. And I'm really happy to hear that I can keep asking questions on that on Matrix!

[suguruwataru]

So I wrote my current understanding of why the behavior difference occurs into these huge pseudo-markdown tables below. I decided to keep posting here as the huge tables do appear to belong here more than Matrix...

Also the tables are drawn in a way that makes them look clear, and some of my understanding does deviate from them:

• I believe only the enter if/loop and exit if/loop on the same row happen generally at the same time. Things in between can happen in any order, and thing can go more complicated than that, like in the continue case, thread 2 can loop back multiple times when thread 1 is unlocking and waiting for it to exit loop together.
• I'm not sure what the waiting thread actually does. If both sides of the if is evaluated every iteration, does the waiting thread actually break for every loop back the other thead does, instead of doing nothing? What if before break, it does something else, will it also be done as many times as the other thread loops back? I believe it's doing nothing as this seems to be what the slides show, but want to confirm.

And I do have more follow-up questions below the tables...

loop-break case. I checked wgpu source and saw this is what while loops compile to (I believe it's the same for for loops?):

| thread 1                  | thread 2                  | code                                 |
|===========================|===========================|======================================|
| enter loop                | enter loop                |  => loop {                           |
|                           |                           |       if atomicOr(&lock, 1u) == 0u { |
|                           |                           |         break;                       |
|                           |                           |       }                              |
|                           |                           |     }                                |
|                           |                           |     atomicAnd(&lock, 0u);            |
|---------------------------|---------------------------|--------------------------------------|
| enter if                  | enter if                  |     loop {                           |
|                           |                           |  =>   if atomicOr(&lock, 1u) == 0u { |
|                           |                           |         break;                       |
|                           |                           |       }                              |
|                           |                           |     }                                |
|                           |                           |     atomicAnd(&lock, 0u);            |
|---------------------------|---------------------------|--------------------------------------|
| atomicOr(&lock, 1u) == 0u |                           |     loop {                           |
|                           | atomicOr(&lock, 1u) == 1u |  =>   if atomicOr(&lock, 1u) == 0u { |
|                           |                           |         break;                       |
|                           |                           |       }                              |
|                           |                           |     }                                |
|                           |                           |     atomicAnd(&lock, 0u);            |
|---------------------------|---------------------------|--------------------------------------|
| exit if and exit loop     | exit if                   |     loop {                           |
|                           |                           |       if atomicOr(&lock, 1u) == 0u { |
|                           |                           | 1=>     break;                       |
|                           |                           | 2=>   }                              |
|                           |                           |     }                                |
|                           |                           |     atomicAnd(&lock, 0u);            |
|---------------------------|---------------------------|--------------------------------------|
| wait for thread 2 to exit | loop back                 | 2=> loop {                           |
| loop together             |                           |       if atomicOr(&lock, 1u) == 0u { |
|                           |                           |         break;                       |
|                           |                           |       }                              |
|                           |                           | 1=> }                                |
|                           |                           |     atomicAnd(&lock, 0u);            |
|---------------------------|---------------------------|--------------------------------------|
| wait for thread 2 to exit | enter if                  |     loop {                           |
| loop together             |                           | 2=>   if atomicOr(&lock, 1u) == 0u { |
|                           |                           |         break;                       |
|                           |                           |       }                              |
|                           |                           | 1=> }                                |
|                           |                           |     atomicAnd(&lock, 0u);            |
|---------------------------|---------------------------|--------------------------------------|
| wait for thread 2 to exit | atomicOr(&lock, 1u) == 1u |     loop {                           |
| loop together             |                           | 2=>   if atomicOr(&lock, 1u) == 0u { |
|                           |                           |         break;                       |
|                           |                           |       }                              |
|                           |                           | 1=> }                                |
|                           |                           |     atomicAnd(&lock, 0u);            |
|---------------------------|---------------------------|--------------------------------------|
| wait for thread 2 to exit | exit if                   |     loop {                           |
| loop together             |                           |       if atomicOr(&lock, 1u) == 0u { |
|                           |                           |         break;                       |
|                           |                           | 2=>   }                              |
|                           |                           | 1=> }                                |
|                           |                           |     atomicAnd(&lock, 0u);            |
|---------------------------|---------------------------|--------------------------------------|
| wait for thread 2 to exit | loop back                 | 2=> loop {                           |
| loop together             |                           |       if atomicOr(&lock, 1u) == 0u { |
|                           |                           |         break;                       |
|                           |                           |       }                              |
|                           |                           | 1=> }                                |
|                           |                           |     atomicAnd(&lock, 0u);            |
|---------------------------|---------------------------|--------------------------------------|
And infinite loop...

loop-continue case:

| thread 1                  | thread 2                  | code                                 |
|===========================|===========================|======================================|
| enter loop                | enter loop                |  => loop {                           |
|                           |                           |       if atomicOr(&lock, 1u) == 1u { |
|                           |                           |         continue;                    |
|                           |                           |       }                              |
|                           |                           |       atomicAnd(&lock, 0u);          |
|                           |                           |       break;                         |
|                           |                           |     }                                |
|---------------------------|---------------------------|--------------------------------------|
| enter if                  | enter if                  |     loop {                           |
|                           |                           |  =>   if atomicOr(&lock, 1u) == 1u { |
|                           |                           |         continue;                    |
|                           |                           |       }                              |
|                           |                           |       atomicAnd(&lock, 0u);          |
|                           |                           |       break;                         |
|                           |                           |     }                                |
|---------------------------|---------------------------|--------------------------------------|
| atomicOr(&lock, 1u) == 0u |                           |     loop {                           |
|                           | atomicOr(&lock, 1u) == 1u |  =>   if atomicOr(&lock, 1u) == 1u { |
|                           |                           |         continue;                    |
|                           |                           |       }                              |
|                           |                           |       atomicAnd(&lock, 0u);          |
|                           |                           |       break;                         |
|                           |                           |     }                                |
|---------------------------|---------------------------|--------------------------------------|
| exit if                   | exit if                   |     loop {                           |
|                           |                           |       if atomicOr(&lock, 1u) == 1u { |
|                           |                           | 2=>     continue;                    |
|                           |                           | 1=>   }                              |
|                           |                           |       atomicAnd(&lock, 0u);          |
|                           |                           |       break;                         |
|                           |                           |     }                                |
|---------------------------|---------------------------|--------------------------------------|
| atomicAnd(&lock, 0u)      | loop back                 | 2=> loop {                           |
|                           |                           |       if atomicOr(&lock, 1u) == 1u { |
|                           |                           |         continue;                    |
|                           |                           |       }                              |
|                           |                           | 1=>   atomicAnd(&lock, 0u);          |
|                           |                           |       break;                         |
|                           |                           |     }                                |
|---------------------------|---------------------------|--------------------------------------|
| wait for thread 2 to exit | enter if                  |     loop {                           |
| loop together             |                           | 2=>   if atomicOr(&lock, 1u) == 1u { |
|                           |                           |         continue;                    |
|                           |                           |       }                              |
|                           |                           |       atomicAnd(&lock, 0u);          |
|                           |                           | 1=>   break;                         |
|                           |                           |     }                                |
|---------------------------|---------------------------|--------------------------------------|
| wait for thread 2 to exit | atomicOr(&lock, 1u) == 0u |     loop {                           |
| loop together             |                           | 2=>   if atomicOr(&lock, 1u) == 1u { |
|                           |                           |         continue;                    |
|                           |                           |       }                              |
|                           |                           |       atomicAnd(&lock, 0u);          |
|                           |                           | 1=>   break;                         |
|                           |                           |     }                                |
|---------------------------|---------------------------|--------------------------------------|
| wait for thread 2 to exit | exit if                   |     loop {                           |
| loop together             |                           |       if atomicOr(&lock, 1u) == 1u { |
|                           |                           |         continue;                    |
|                           |                           | 2=>   }                              |
|                           |                           |       atomicAnd(&lock, 0u);          |
|                           |                           | 1=>   break;                         |
|                           |                           |     }                                |
|---------------------------|---------------------------|--------------------------------------|
| wait for thread 2 to exit | atomicAnd(&lock, 0u)      |     loop {                           |
| loop together             |                           |       if atomicOr(&lock, 1u) == 1u { |
|                           |                           |         continue;                    |
|                           |                           |       }                              |
|                           |                           | 2=>   atomicAnd(&lock, 0u);          |
|                           |                           | 1=>   break;                         |
|                           |                           |     }                                |
|---------------------------|---------------------------|--------------------------------------|
| exit loop                 | exit loop                 |     loop {                           |
|                           |                           |       if atomicOr(&lock, 1u) == 1u { |
|                           |                           |         continue;                    |
|                           |                           |       }                              |
|                           |                           |       atomicAnd(&lock, 0u);          |
|                           |                           |  =>   break;                         |
|                           |                           |     }                                |
|---------------------------|---------------------------|--------------------------------------|

Follow-up questions:

1. If if and loop are entered and exited at around the same time, can this be used as some sort of synchronization? I think the answer is no and we need barriers, and I guess the reason is caching: after 2 threads exit an if at around the same time, due to caching, thread 1 and thread 2 may see different things when they read a value written by thread 1 inside the if.
2. Even with a relaxed memory ordering, is it possible to use one bit of an atomic as the mutex for the other bits? I believe we can, as access to that memory location is synchronized.
3. While reading wgpu code, I saw there are textureAtomic* functions, but can't seem to find info about them else where. So what's the status of those? They seem incredibly helpful!
4. Can we think of the difference in these loops' behaviors as part of the stable interface of shader languages, instead of implementation details that may change? I am pretty sure the answer is yes, but it's so different from CPU so I'm feeling it's a bit hard to just accept...