Sitelet https://github.com/gpuweb/gpuweb/issues/3602
Skip to content

wgsl: full reference to module-scope variable is uniform, the value stored there may not be #3602

Description

@dneto0

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

Activity

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    wgslWebGPU Shading Language Issues

    Type

    No type

    Projects

    No projects

      Milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions