WGSL Bug

Code

 1: @group(0)
 2: @binding(0)
 3: var<storage, read_write> flow: array<u32, 2>;
 4: 
 5: @compute
 6: @workgroup_size(1)
 7: fn main() {
 8:     var LOOP_COUNTER: u32 = 0u;
 9:     loop {
10:         (flow)[0u] += 1u;
11:         if (LOOP_COUNTER >= 1u) {
12:             (flow)[1u] += 1u;
13:             break;
14:         }
15:         LOOP_COUNTER += 1u;
16:     }
17: }

Expected output is [2, 1] after executing break statement in second iteration of loop.

Results

Vendor GPU Dawn Vulkan Linux wGPU Vulkan Linux Dawn Vulkan Windows wGPU Vulkan Windows Dawn DX12 Windows wGPU DX12 Windows
Nvidia RTX 2060 Max-Q Timeout ND1 Zeros2 Zeros2 Correct Correct
Nvidia RTX 2060 Super Timeout ND1 Zeros2   Correct Correct
Nvidia GTX 1050 Ti Timeout Correct Zeros2 Correct Correct Correct
Nvidia Quadro 4000 Correct Correct        
AMD Ryzen 9 4900HS (iGPU) Correct Correct        
Intel i7-9700K (iGPU) Correct Inconsistent3        
Intel i5-7400 (iGPU) Correct Correct        

Footnotes:

1

Non Deterministic results where the first element is large and the second element is zero; produces correct result if either of the flow increments are removed.

2

Output is [0, 0], presumed to be as a result of a timeout.

3

Sometimes produces correct output, other times produces [0, 0] without timing out.

Author: Kabir Kwatra

Created: 2022-09-13 Tue 13:15