[naga] [msl-out] Error: as_type cast from 'long' to 'metal::uint' (aka 'unsigned int') is not allowed

[LegNeato]

Description

When compiling this code with rust-gpu, I get valid spirv asm. I then pass it to naga to translate to metal (as I am on macOS) and I get these errors:

                        float _e11563 = function_119(metal::float2(_e7495 * 7.6543, _e7496 * 7.6543));
                              ^
program_source:16276:67: error: as_type cast from 'long' to 'metal::uint' (aka 'unsigned int') is not allowed
                    int _e13236 = as_type<int>(as_type<uint>(1) + as_type<uint>((_e13226 == _e13226) ? ((_e13226 > as_type<float>(1325400063u)) ? 2147483647 : ((_e13226 < as_type<float>(3472883712u)) ? -2147483648 : naga_f2i32(_e13226))) : 0));
                                                                  ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
program_source:16304:67: error: as_type cast from 'long' to 'metal::uint' (aka 'unsigned int') is not allowed
                    int _e13273 = as_type<int>(as_type<uint>(1) + as_type<uint>((_e13263 == _e13263) ? ((_e13263 > as_type<float>(1325400063u)) ? 2147483647 : ((_e13263 < as_type<float>(3472883712u)) ? -2147483648 : naga_f2i32(_e13263))) : 0));
                                                                  ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~

Repro steps
It's a bit convoluted...

1. Checkout Update deps Rust-GPU/rust-gpu-shadertoys#33
2. Patch in Reduce repo MSRV from 1.85 to 1.84 #7425
3. In build.rs, remove the check that turns it off on macos or in shaders/lib.rs, delete everything in the match except 21 => filtering_procedurals::Inputs

I'll work on getting the raw spirv.

Note I tested with the change from #7422 by @jamienicol and it still was not fixed.

Expected vs observed behavior
The shader to translate to metal code that compiles.

Platform
macOS, latest master of wgpu

[LegNeato]

@jamienicol I assume you are busy, but wanted to point you to this as you have been working in this space and might be able to see something quickly. I tried to find the issue but I am still getting up to speed on the code.

[jamienicol]

From a quick scan I would guess the issue is that like in #7423 -2147483648 is parsed as "unary negation of 2147483648". 2147483648 is too large to fit in a 32 bit int, and is therefore inferred to be a long. Which promotes the rest of the expression to a long. Which then errors when trying to bitcast it to a uint because as_type() must keep the same bitwidth

[jamienicol]

Yep I think that's it. Here's a reduced case:

fn main() {
  var x = bitcast<u32>(-2147483648);
}

Same error message.

[LegNeato]

@jamienicol great, found the fix. PR incoming soon. Thanks!

[LegNeato]

(unless you want to PR, up to you...don't want to step on your toes and you know more what you are doing and I don't know if this affects other backends)

[jamienicol]

Please go ahead 🙂