[hlsl-out] error X4010: Unsigned integer divide by zero

[hasali19]

I've found what seems to be another hlsl compilation failure caused by FXC. This one incorrectly reports a divide by zero error. Interestingly it's only rejected when the divide call is in a nested loop.

fn divide(a: i32, b: i32) -> i32 {
    if (b == 0) {
        return a / 2;
    } else {
        return a / b;
    }
}

struct Out {
    value: u32,
}

@group(0)
@binding(0)
var<storage, read_write> s_out: Out;

@compute
@workgroup_size(1)
fn main() {
    loop {
        loop {
            var x = divide(0, 0);
            break;
        }
        break;
    }
    s_out.value = 0u;
}

Output:

The application panicked (crashed).
Message:  wgpu error: Validation Error

Caused by:
    In Device::create_compute_pipeline
    Internal error: D3DCompile error (0x80004005): \\wsl.localhost\Ubuntu\home\hasan\dev\wgslsmith\Shader@0x0000025A33C8DBA0(19,17-21): error X4010: Unsigned integer divide by zero
\\wsl.localhost\Ubuntu\home\hasan\dev\wgslsmith\Shader@0x0000025A33C8DBA0(28,5-15): warning X3557: loop only executes for 0 iteration(s), forcing loop to unroll

Tint seems to be able to satisfy FXC by adding an extra check on b in the divide function: https://shader-playground.timjones.io/ac86c91d9f0f9b09c7668fa99534fe3f.

[teoxoy]

The WGSL spec requires us to add some checks (akin to tint's) to handle undefined behavior of some operators and built-in functions which we aren't currently doing.

[jimblandy]

If implementing WGSL's division semantics happens to work around this FXC bug, that'd be great. The Naga HLSL back end code implementing the workaround should mention this FXC bug, so future developers will know we have unexpected constraints we're trying to meet.