Pointer aliasing with unused parameter

[hasali19]

The spec says that aliasing is invalid if the aliases access the same memory location and at least one of the accesses is a write.

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

var<private> flag: bool;

fn func(p: ptr<private, bool>) {
    flag = true;
}

@compute
@workgroup_size(1)
fn main() {
    func(&flag);
    if (flag) {
        s_out = 1;
    } else {
        s_out = 2;
    }
}

In this code, a pointer to flag is passed to func but it is never used, and flag is only ever accessed through the originating variable identifier. So based on my understanding (or have I misunderstood the spec?) this should be fine. However, on DirectX this behaves unexpectedly with an output value of 2. On Vulkan it works as expected, outputting 1.

Is this a bug or is it forbidden even if the memory location isn't accessed through the pointer?

[teoxoy]

I think this is an issue with our HLSL backend (or HLSL weirdness); for some reason the generated code for the function (even though it uses inout) writing to the variable has no effect.

[jimblandy]

Yes, the new WGSL aliasing rules permit the program shown.

But HLSL inout is defined to mean that the argument's value is copied into the callee's parameter upon entry to the call, and then the parameter's value is copied back into the argument upon return. HLSL's inout does not mean "pass by reference".

So if Naga translates the WGSL function func into an HLSL function with an inout parameter, that means that the assignment to flag in func will get overwritten when func returns, if main passes &flag.

It might be better for me to think about this after a good night's sleep instead of 3am, but I think this means that Naga needs to use the alias analysis to decide how to translate func into HLSL. If a WGSL function's pointer parameter is never written to, then I think it must not be rendered as an inout parameter in HLSL.

I think this means that passing parameters by pointer cannot be translated well into HLSL: they always become copies.

[jimblandy]

I believe this code is now forbidden by WGSL, per the restrictions in the section Aliasing, and the bug is that Naga fails to validate it out.

[bavalpey]

Correct me if I am wrong, but it seems that Naga still does not fully implement the Alias Analysis as required by the WGSL spec (noted in the comment above).

Doesn't this mean that wgpu, which uses Naga for validation, is not conformant with the WGSL spec?

[ErichDonGubler]

@bavalpey:

Doesn't this mean that wgpu, which uses Naga for validation, is not conformant with the WGSL spec?

Yep! And we're trying to fix it. WebGPU compliance is a major driver for Firefox contributors to WGPU, and you can follow along at https://github.com/orgs/gfx-rs/projects/7/views/5 if you'd like to see the priorities we've identified.

[bavalpey]

Thanks for your response! I do have some stake in this. Sharing thoughts below.

I'm working on an analysis for wgsl and am using Naga to drive it, hooking into it directly after its validation pass. I would like to depend on on the guarantees of Alias Analysis. I suppose I could just continue with the assumption that the aliasing requirements of WGSL aren't violated.

Are there currently plans to implement the alias analysis? I didn't see it explicitly mentioned on that view you shared.

One of the things that would make my life easier is if there was some way, in Naga, to figure out the originating variable for the pointer expression. This is done for a different reason in the uniformity analysis, where we see assignable_global added in ExpressionInfo.

I would presume that, similar to the uniformity requirement, there would be an additional field to expressions in the expression_info field in the ModuleInfo returned by the validation layer that provides this information for all pointer types, not just globals.

I do have one question, and perhaps someone here can shed some light on it. The WGSL spec says that

Execution of a WGSL function must not potentially access memory through aliased root identifiers,

However, in the proceeding discussion, the analysis only discusses prohibitions on aliases of parameters. It doesn't discuss aliases to other variables. I am curious, then, as to the legality of something like this:

@group(0) @binding(0) var<storage, read_write> x: i32;

fn foo() -> i32 {
  let a = &x;
  *a = 4;
  return x;
}

[ErichDonGubler]

@bavalpey: The issue you're looking for is https://github.com/gfx-rs/wgpu/issues/4369.

I don't have anything immediately useful to contribute to the discussion, but @teoxoy or @jimblandy may.

[bavalpey]

Not sure how that issue is related. I don't have questions about uniformity analysis.

[ErichDonGubler]

@bavalpey: Ah, I must be confusing the two, then. I hadn't read the Aliasing section mentioned by @jimblandy, and didn't realize that these are two separate analyses.

This is likely something we should add to our board, then! Sorry for the confusion. 😓