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 gpuweb/gpuweb#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

Fixed: #4476

Issue: #4476

---

Requirements for PR author:

• All missing test coverage is tracked with "TODO" or .unimplemented().
• New helpers are /** documented */ and new helper files are found in helper_index.txt.
• Test behaves as expected in a WebGPU implementation. (If not passing, explain above.)
• 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.