uniformity: test uniformity when a loop body only returns

[dneto0]

Consider: loop { s1 continuing { s2 } }

If the statement behaviour of s1 is exactly {Return}, then the statement
behaviour of the entire loop is {Return}. In particular, nonuniformity
effects do not leak out from s2 to the context after or outside of the
whole loop.

See the WGSL spec fix for https://github.com/gpuweb/gpuweb/issues/5100

Test three new statement scenarios, each where s1 starts with `return;`,
but varying where a collective operation is placed:
 - immediately after the return statement
 - in the continuing block of the same loop
 - immediately after the loop, where the continuing block in the loop
   has a break-if. This case requires the analysis to avoid using a
   Next behaviour from the overall loop.
 - immediately after the loop

Also test similar variants where we use `break` instead of `return`.

Also test loop constructs:
 - for with a condition: desugars to loop with initial `if !(cond) { break;}`
 - for without a condition: desugars to loop without that initial
   conditional break.
 - while loop

Fixed: #4476

Issue: #4476

---

Requirements for PR author:

• [x] All missing test coverage is tracked with "TODO" or .unimplemented().
• [x] New helpers are /** documented */ and new helper files are found in helper_index.txt.
• [x] Test behaves as expected in a WebGPU implementation. (If not passing, explain above.)
• [x] Test have be tested with compatibility mode validation enabled and behave as expected. (If not passing, explain above.)

Requirements for reviewer sign-off:

• [ ] Tests are properly located in the test tree.
• [ ] Test descriptions allow a reader to "read only the test plans and evaluate coverage completeness", and accurately reflect the test code.
• [ ] Tests provide complete coverage (including validation control cases). Missing coverage MUST be covered by TODOs.
• [ ] Helpers and types promote readability and maintainability.

When landing this PR, be sure to make any necessary issue status updates.

[dneto0]

Works in dawn.node

[jrprice]

This case fails with Tint locally:

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

Looks like the tests here do not cover the case where the is a break if in the continuing block. Is the corresponding spec PR supposed to cover that or is that a separate issue?

[dneto0]

Looks like the tests here do not cover the case where the is a break if in the continuing block. Is the corresponding spec PR supposed to cover that or is that a separate issue?

Good catch. It should be covered by this rule. I've added your case locally and it fails. So that's good. I'll file a Tint issue to fix it.

[dneto0]

I want to add more cases, as found from testing the Tint fixes.

See the DeadLoopTests in https://dawn-review.googlesource.com/c/dawn/+/267474 that cover loop, for-with-cond, for-without-cond, and while. It's tricky enough that it needs coverage.

[dneto0]

I've reworked the statement uniformity checks to use an explicit sensitive data field.

And I've added the variants from the Tint unit test, as described in the updated commit comment.

[dneto0]

Passes on dawn.node when Dawn is patched with https://dawn-review.googlesource.com/c/dawn/+/267474

[dneto0]

I have a followup which greatly expands coverage of loops. I'll land this now as it is still valid and does fix the #4476 issue