wgsl: full reference to module-scope variable is uniform, the value stored there may not be #3602
Description
Spawned from the discussion at: #2586 (comment)
Example:
var<workgroup> w: array<i32,10>;
@compute @workgroup_size(10)
fn foo(@builtin(local_invocation_index) i:u32) {
// The *reference* w is uniform.
let pw = &w; // should be considered uniform
let v0 = w[0]; // should be considered non-uniform because of mutability.
let vi = w[i]; // should be considered non-uniform because of nonuniform indexing *and* mutability
let p0 = &w[0]; // should be considered uniform because it's an address combining uniform things.
let pi = &w[i]; // should be considered non-uniform
let vp0 = *p0; // should be considered non-uniform because of mutability.
let vp1 = *pi; // non-uniform
}In short:
- full reference to a module-scope variable is uniform
- address calculation combining only uniform things should also be uniform
- applying LoadRule to a mutable module-scope reference yields MayBeNonUniform
- applying LoadRule to immutable thing yields Uniform
However, uniformity analysis currently says
- full reference to module-scope variable in "lvalue position" is MayBeNonUniform.
This matters because with the current formulation we can't ever say that a pointer to a workgroup variable is uniform.
And that's what need for workgroupUniformLoad. See #2586