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.