gpu: WGSL incorrect control flow

[miura-hiroka]

OS: Windows 11
CPU: Ryzen 5 5600
GPU: Radeon RX 6600
Zig: 2024.3.0-mach

This problem seems to be related to Dawn but I don't know much about it.

output:

expected: 0 0 0 0
actual: 1 1 1 1

compute.wgsl:

@group(0) @binding(0) var<storage, read_write> output: array<u32>;

@compute @workgroup_size(4)
fn main(@builtin(global_invocation_id) coords: vec3<u32>) {
  if arrayLength(&output) <= coords.x {
    return;
  }

  var input = array<u32, 4>(0, 0, 0, 0);
  input[coords.x] = 1;

  output[coords.x] = bug(input);
}

// This function should allways return 0.
fn bug(input: array<u32, 4>) -> u32 {
  var a = array<bool, 4>(false, false, false, false);
  var i: u32 = 0;
  loop {
    if input[i] == 1 {
      a[i] = true;
      // Uncommenting this change the result even though i == 4 is impossible.
      // if (i == 4) {
      //   return 1;
      // }
      break; // This is the only way to reach the end of the loop.
    }
    if i == 3 {
      return 0;
    }
    i++;
  }
  // Here, a[i] must be true.
  if !a[i] {
    return 1;
  }
  return 0;
}

source code: mach-bug.zip

[miura-hiroka]

I reported this issue to Google.
https://issues.chromium.org/issues/346386733