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