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.

Sign up for GitHub

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

Jump to bottom

[msl-out] BoundsCheckPolicy::ReadZeroSkipWrite produces invalid Metal when calling functions with pointers to array elements #4541

Open
Liamolucko opened this issue Aug 16, 2023 · 4 comments · May be fixed by #7323
Open

[msl-out] BoundsCheckPolicy::ReadZeroSkipWrite produces invalid Metal when calling functions with pointers to array elements #4541

Liamolucko opened this issue Aug 16, 2023 · 4 comments · May be fixed by #7323
Assignees
@andyleiserson
Labels
area: naga back-end Outputs of naga shader conversion lang: Metal Metal Shading Language naga Shader Translator type: bug Something isn't working
Milestone
WebGPU Specificat...

Comments

@Liamolucko
Copy link
Contributor

Say you have a shader like this, which calls a function with a pointer to an array element:

fn takes_ptr(p: ptr<function, i32>) {}

fn foo(i: u32) {
    var arr = array(1, 2, 3, 4);
    takes_ptr(&arr[i]);
}

When BoundsCheckPolicy::ReadZeroSkipWrite is enabled, Naga produces this invalid Metal code:

// language: metal2.0
#include <metal_stdlib>
#include <simd/simd.h>

using metal::uint;
struct DefaultConstructible {
    template<typename T>
    operator T() && {
        return T {};
    }
};

struct type_3 {
    int inner[4];
};

void takes_ptr(
    thread int& p
) {
    return;
}

void foo(
    uint i
) {
    type_3 arr = {};
    arr = type_3 {1, 2, 3, 4};
    takes_ptr(uint(i) < 4 ? arr.inner[i] : DefaultConstructible());
    return;
}

DefaultConstructible() isn't a valid reference, which Metal rightly complains about:

test.metal:28:5: error: no matching function for call to 'takes_ptr'
    takes_ptr(uint(i) < 4 ? arr.inner[i] : DefaultConstructible());
    ^~~~~~~~~
test.metal:17:6: note: candidate function not viable: expects an l-value for 1st argument
void takes_ptr(
     ^
1 error generated.

One solution to this that I could see would be to generate a variable containing the default value of the array element type, and pass a reference to that if the index is out of bounds. Something along the lines of this:

void foo(
    uint i
) {
    type_3 arr = {};
    arr = type_3 {1, 2, 3, 4};
    // I think this is equivalent to `DefaultConstructible()`, looking at its definition, but I don't know much about C++/Metal.
    int oob = {};
    takes_ptr(uint(i) < 4 ? arr.inner[i] : oob);
    return;
}
@teoxoy teoxoy added kind: bug lang: Metal Metal Shading Language area: naga back-end Outputs of naga shader conversion labels Aug 17, 2023
@teoxoy
Copy link
Member

teoxoy commented Aug 17, 2023

I haven't looked at how the code is setup for this right now but I suspect it might be easier to move the whole ternary operation on it own line to solve the issue.

@Liamolucko
Copy link
Contributor Author

I haven't looked at how the code is setup for this right now but I suspect it might be easier to move the whole ternary operation on it own line to solve the issue.

I'm assuming you mean something like this:

int var = uint(i) < 4 ? arr.inner[i] : DefaultConstructible();
takes_ptr(var);

Wouldn't that then write to var instead of the original array element?

I'm willing to make a PR for this, BTW.

@teoxoy
Copy link
Member

teoxoy commented Aug 18, 2023

I was thinking thread int& var = uint(i) < 4 ? arr.inner[i] : DefaultConstructible(); but that won't work for the same reason as it doesn't now (taking a reference to a temporary).

I think we can go with the solution you proposed.

@Liamolucko
Copy link
Contributor Author

Huh, turns out translating that same shader to SPIR-V results in a panic:

thread 'main' panicked at 'internal error: entered unreachable code: Expression [8] is not cached!', /Users/liam/.cargo/registry/src/index.crates.io-6f17d22bba15001f/naga-0.13.0/src/back/spv/block.rs:2066:45
note: run with `RUST_BACKTRACE=1` environment variable to display a backtrace

I'll try and fix that too while I'm at it.

@Liamolucko Liamolucko mentioned this issue Aug 21, 2023
Closed
@cwfitzgerald cwfitzgerald transferred this issue from gfx-rs/naga Oct 25, 2023
@cwfitzgerald cwfitzgerald added naga Shader Translator type: bug Something isn't working and removed kind: bug labels Oct 25, 2023
@teoxoy teoxoy added this to the WebGPU Specification V1 milestone Nov 3, 2023
@andyleiserson andyleiserson self-assigned this Mar 5, 2025
This was referenced Mar 6, 2025
Closed
Open
andyleiserson added a commit to andyleiserson/wgpu that referenced this issue Mar 12, 2025
@andyleiserson
[msl-out] Fix ReadZeroSkipWrite bounds check mode for pointer argum…
43754e1 
…ents

Fixes gfx-rs#4541

--
Co-authored-by: Liam Murphy <liampm32@gmail.com>
@andyleiserson andyleiserson linked a pull request Mar 12, 2025 that will close this issue
Open
6 tasks
@andyleiserson andyleiserson moved this to In Progress in WebGPU for Firefox Mar 21, 2025
@andyleiserson andyleiserson mentioned this issue Mar 24, 2025
Open
andyleiserson added a commit to andyleiserson/wgpu that referenced this issue Mar 26, 2025
@andyleiserson
[msl-out] Fix ReadZeroSkipWrite bounds check mode for pointer argum…
1d40580 
…ents

Fixes gfx-rs#4541

--
Co-authored-by: Liam Murphy <liampm32@gmail.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees

@andyleiserson andyleiserson

Labels
area: naga back-end Outputs of naga shader conversion lang: Metal Metal Shading Language naga Shader Translator type: bug Something isn't working
Projects
WebGPU for Firefox
Status: In Progress
Milestone
WebGPU Specification V1
Development

Successfully merging a pull request may close this issue.

[msl-out] Fix ReadZeroSkipWrite bounds check mode for pointer arguments andyleiserson/wgpu
4 participants
@cwfitzgerald @teoxoy @Liamolucko @andyleiserson