-
Notifications
You must be signed in to change notification settings - Fork 344
Description
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.