这是indexloc提供的服务,不要输入任何密码
Skip to content

Uniformity Analysis - conservative treatement of assignments #5440

@JamesLee-Jones

Description

@JamesLee-Jones

Under the analysis in the spec, the following shader is considered non-uniform:

@compute @workgroup_size(16,1,1)
fn main(@builtin(local_invocation_index) lid: u32, ) {
    var x: u32;
    loop {
        workgroupBarrier();
        continuing {
            loop {
                x = 1;
                continuing {
                    break if (lid == 0);
                }
            }
            break if (x == 0);
        }
    }
}

While this shader (with the assignment in the inner loop removed) is considered uniform:

@compute @workgroup_size(16,1,1)
fn main(@builtin(local_invocation_index) lid: u32, ) {
    var x: u32;
    loop {
        workgroupBarrier();
        continuing {
            loop {
                continuing {
                    break if (lid == 0);
                }
            }
            break if (x == 0);
        }
    }
}

The non-uniform break condition of the inner loop in the first shader taints the inner loop header as non-uniform, which in turn taints the variable x as non-uniform. This leads the break if condition in the outer loop to be considered non-uniform and thus the shader is rejected.

Because the inner loop will always do at least one iteration and is assigned a literal, it will always have the value 1 regardless of when the loop breaks.

The conservative treatment seems to be a side effect of the fact that the analysis of an expression often produces a value node that depends on CF which is needed for the correct handling of value merging at the end of if statements, loops, etc.

Metadata

Metadata

Assignees

No one assigned

    Labels

    uniformityIssues / discussions around uniformity analysiswgslWebGPU Shading Language Issues

    Type

    No type

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions