Uniformity Analysis - conservative treatement of assignments

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

[dneto0]

Yes, that does look conservative.

An easy way to see it: Start with the first shader, and pull the assignment x = 1; to before the loop, and it now passes uniformity:

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

[dneto0]

Regarding

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.

Yes, that looks right. In the statement break if (x == 0);, the uniformity of x as an RHSValue depends on the Vout(prev) for x. and then the graph traces a path back to the assignment x=1; which then applies the LHSValue rule for the assignment:

expression	new node	recursive analysis	resulting control flow node, value node	new edges
identifier resolving to function-scope variable "x"	Result	X is the node corresponding to the value of "x" at the output to the statement containing this expression	CF, Result	Result -> {CF, X}

That introduces the dependence on the control flow uniformity node CF for the assignment statement, in the inner loop. And that CF traces back to MayBeNonUniform.

That dependence is meant to capture the nonuniformity introduced in a case like this:

   x = 0;
   loop  {
       if lid == 0 { break; }   // nonuniform
       x = 1;
   }
   // here x must be seen to be MayBeNonUniform.

or even:

   x = 0;
   loop  {
       if lid == 0 { continue; }   // nonuniform
       x = 1;
       break;
   }
   // here x must be seen to be MayBeNonUniform.

To make the analysis more precise, it seems we can cut the control dependency if the assignment to x occurs in a portion of the body that hasn't potentially seen a Continue, Next, or Return. Schematically:

    loop { s1;   x = ... ; ...} 

In this case if s1 has behaviour Next, then the assignment to x always occurs (as you observed), and then the analysis would work better if we could cut the path from the backedge of the loop. I don't have the rest worked out or generalized, or an estimate of how costly that would be to represent or analyze.

[jimblandy]

@dneto0 Can we treat this issue as editorial? Or would you like the committee to consider it?