wgpu icon indicating copy to clipboard operation
wgpu copied to clipboard

Pointer aliasing with unused parameter

Open hasali19 opened this issue 2 years ago • 9 comments

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?

hasali19 avatar May 23 '22 14:05 hasali19