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

Variable Value Analysis - Ambiguity in loop rule #5439

@JamesLee-Jones

Description

@JamesLee-Jones

The first rule for loops in the variable value analysis of loops is as follows:

Statement Result Edges from the Result
loop { s1 continuing { s2 } } Vin(s1) Vout(prev), Vout(s2)

This aims to capture the fact that the uniformity of variables at the beginning of the body of a loop depends on both the uniformity of the variable before the loop (indicated by Vout(prev)) and the end of the loop body (indicated by Vout(s)).

In order to avoid being overly conservative, this rule should mean that at the beginning of analysis of the loop body, an new node is created for each variable x that depends on both the node representing the uniformity of x before the loop and at the end of the loop body x. This is how Tint handles this case.

However, since it isn't specified that Vin(...) creates new nodes, I think this rule could be interpreted to mean that the node representing the uniformity of every variable x before the loop body is used to analyse the body and an edge is added to the node representing the uniformity of x at the end of the loop body. If implemented in this way, the following shader is rejected as non-uniform:

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

This could be solved by adding a comment in the Variable Value Analysis section to say that Vin(...) creates new uniformity nodes. Another option would be to add a new 'function' that is used when the variable value analysis has multiple outgoing edges, which would avoid specifying that new nodes should be created when analysing things like sequential statements.

Metadata

Metadata

Assignees

Labels

copyeditingPure editorial stuff (copyediting, *.bs file syntax, etc.)uniformityIssues / discussions around uniformity analysiswgslWebGPU Shading Language Issues

Type

No type

Projects

No projects

Relationships

None yet

Development

No branches or pull requests

Issue actions