Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[naga msl-out] invalid code emitted for minimum i32/i64 value literal. #7437

Open
wants to merge 2 commits into
base: trunk
Choose a base branch
from

Conversation

LegNeato
Copy link
Contributor

@LegNeato LegNeato commented Mar 26, 2025

Connections
Fixes #7435.
Similar to #7423

Description
-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

Testing
I tried my repro from #7435 and saw it now works.

Squash or Rebase?

Rebase

Checklist

  • Run cargo fmt.
  • Run taplo format.
  • Run cargo clippy --tests. If applicable, add:
    • --target wasm32-unknown-unknown
  • Run cargo xtask test to run tests.
  • If this contains user-facing changes, add a CHANGELOG.md entry.

@LegNeato
Copy link
Contributor Author

The real change is in the 2nd commit...the first is the dep.

@LegNeato LegNeato changed the title Fix msl [naga msl-out] invalid code emitted for minimum i32/i64 value literal. Mar 26, 2025
Copy link
Contributor

@jamienicol jamienicol left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks good, thanks! Couple nits but I can make those changes for you.

I'm also going to rebase it to not depend on #7422 as that review might take slightly longer for someone to get to.

var<storage, read_write> output_arrays: StorageCompatible;

fn int32_function(x: i32) -> i32 {
var min_val = bitcast<u32>(-2147483648);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't think we need a whole int32.wgsl testcase. These should all be covered by the rest of the test suite.

(The reason we have an int64.wgsl / f64.wgsl etc is because those types are extensions)

We just want to make sure we test the i32::MIN and i64::MIN literals.

We can use the bitcast trick as you've done here for i32. AFAICT that's the only way to ensure we generate a compiler error rather than the accidentally-long value being implicitly converted to an int. Metal doesn't even appear to emit a compiler warning. The question is where to put it if not in a new int32.wgsl test. There's not really a great place but perhaps under arithmetic -> unary in operators.wgsl?

We should test the min int64 literal too. Unfortunately we cannot use the bitcast trick as the value is still a long. However, the metal compiler does emit a constant-conversion warning. So if we run our validation with -Werror=constant-conversion we will catch it.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Oh I though of another issue with relying on the bitcast to catch this bug: currently it works fine as we haven't implemented bitcasts during constant evaluation.

But when we do get round to that, the bitcast would be evaluated away, meaning the output MSL wouldn't contain the bitcast and the test would no longer do what it's meant to. So we need to ensure we add a non-const expression to the bitcast. Or we can even just use signed integer addition, as we wrap those in bitcasts to avoid signed overflow (again we'd have to prevent const evaluation removing the operation)

@jamienicol jamienicol force-pushed the fix-msl branch 2 times, most recently from 22cf7da to 7129ad5 Compare March 27, 2025 10:52
@jamienicol
Copy link
Contributor

r+ from me for @LegNeato's changes, but would @ErichDonGubler mind having a quick look at my changes since I've now co-authored?

I've pushed an extra commit which improves the fix for #7424 so that we actually test the min i64 literal. And then I think the only controversial change I made to the patch was to pass -Werror=constant-conversion to metal so that validation would catch the bug

@jamienicol
Copy link
Contributor

ha, #7422 got merged quicker than I expected. I'll rebase this again 😆

In gfx-rs#7424 we fixed a bug where the representation of the minimum int64
literal generated by naga was invalid WGSL. It changed us from
expressing it as `-9223372036854775808` which was invalid, to
`-9223372036854775807li - 1li`.

This is valid WGSL. However, as the values are concrete i64 types if the
shader is parsed again by naga the expression does not get const
evaluated away, leading to suboptimal code generated by the
backends. This patch makes us perform the subtraction using abstract
integers before casting to i64, solving this problem.

Additionally the input WGSL test is updated to use the same construct.
The literal `-2147483648` is parsed by Metal as negation of positive
2147483648. As 2147483648 is too large for a int, the expression is
silently promoted to a long. Sometimes this does not matter as it will
often be implicitly converted back to an int after the negation.
However, if the expression is used in a bitcast then we hit a compiler
error due to mismatched bitwidths.

Similarily for `-9223372036854775808`, as 9223372036854775808 is too
large for a long, metal emits a `-Wconstant-conversion` warning and
changes the value to -9223372036854775808. This would then be negated
again, possibly causing undefined behaviour.

In both cases we can avoid the issue by expressing the literals as the
second most negative value expressible by the type, minus one.

eg `-2147483647 - 1` and `-9223372036854775807L - 1L`.

We have added a test which uses the most negative i32 literal in an
addition. Because we bitcast addition operands to unsigned in metal,
this would cause a validation error without this fix. For the i64 case
existing tests already make use of the minimum literal value. Passing
the flag `-Werror=constant-conversion` to Metal during validation will
therefore catch this issue.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

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