Variable Value Analysis - Ambiguity in loop rule

[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.

[dneto0]

I'm having trouble parsing this:

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.

Is this about using the same node V(in) for both x and y ? That's not the intent.

If that is the issue, then there's a clause in the paragraph before the table saying in part "It is written in terms of an arbitrary variable x." Maybe that doesn't carry through so clearly. Each of the V*(...) nodes has a implicitly subscript x; if we made that explicit in the text, would that solve he problem?

If I got the diagnosis wrong, please explain a bit more about the graph that causes the example to be mis-analyzed.

[JamesLee-Jones]

To clarify, from the spec I don't think it's clear that Vin(x) means you have to create a new node to represent the uniformity of x when analysing the next statement. For some of the rules, like s1;s2, you can get away with using the same node from the exit of s1 at the entrance of s2 for every x.

For the first variable value analysis loop rule to avoid being overly-conservative, Vin(x) has to create a new node to represent the uniformity of every variable in the body. The spec says that:

Vin(next) is the value of x prior to the execution of the next statement.

but it doesn't say that this node has to be created.

The shader I gave is handled properly by tint, but is an example of a shader that would be be conservatively rejected if the loop rule for variable value analysis were to be misread as Vin(x) being the node that represents the uniformity of x prior to the loop body instead of a new node created to be used in the body.

[dneto0]

ok, so conceptually the confusion is something like this:

good case:

   Vout(prev)-for-x    loop  {   Vin(s1)-for-x s1; ... } 
   // with an edge from Vin(s1)-for-x to Vout(prev)-for-x

bad case:

   Vout(prev)-for-x-is-also-Vin(s1)-for-x   loop  { s1; ... }

Yes, that would be bad. So yes, we should clear up any potential confusion. I expect that will lead to a clarification in the spec.

And I think implicitly your assertion is that this is only a problem for loops and not the other control constructs because only loops have the backedge in the CFG. Seems right but worth me looking carefully again.

[dneto0]

I suggest treating this as editorial; committee work not required.