wgpu/naga/tests/out/msl/control-flow.msl
Jim Blandy 3fda684eb9 [naga msl-out] Defeat the MSL compiler's infinite loop analysis.
See the comments in the code for details.

This patch emits the definition of the macro only when the first loop
is encountered. This does make that first loop's code look a bit odd:
it would be more natural to define the macro at the top of the
file. (See the modified files in `naga/tests/out/msl`.)

Rejected alternatives:

- We could emit the macro definition unconditionally at the top of the
  file. But this changes every MSL snapshot output file, whereas only
  eight of them actually contain loops.

- We could have the validator flag modules that contain loops. But the
  changes end up being not small, and spread across the validator, so
  this seems disproportionate. If we had other consumers of this
  information, it might make sense.

- We could change the MSL backend to allow text to be generated out of
  order, so that we can decide whether to define the macro after we've
  generated all the function bodies. But at the moment this seems like
  unnecessary complexity, although it might be worth doing in the
  future if we had additional uses for it - say, to conditionally emit
  helper function definitions.

Fixes #4972.
2024-09-18 11:01:51 -04:00

226 lines
4.3 KiB
Plaintext

// language: metal1.0
#include <metal_stdlib>
#include <simd/simd.h>
using metal::uint;
void switch_default_break(
int i
) {
switch(i) {
default: {
break;
}
}
}
void switch_case_break(
) {
switch(0) {
case 0: {
break;
}
default: {
break;
}
}
return;
}
void loop_switch_continue(
int x
) {
#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
LOOP_IS_REACHABLE while(true) {
switch(x) {
case 1: {
continue;
}
default: {
break;
}
}
}
return;
}
void loop_switch_continue_nesting(
int x_1,
int y,
int z
) {
LOOP_IS_REACHABLE while(true) {
switch(x_1) {
case 1: {
continue;
}
case 2: {
switch(y) {
case 1: {
continue;
}
default: {
LOOP_IS_REACHABLE while(true) {
switch(z) {
case 1: {
continue;
}
default: {
break;
}
}
}
break;
}
}
break;
}
default: {
break;
}
}
switch(y) {
default: {
continue;
}
}
}
LOOP_IS_REACHABLE while(true) {
switch(y) {
case 1:
default: {
switch(z) {
default: {
continue;
}
}
break;
}
}
}
return;
}
void loop_switch_omit_continue_variable_checks(
int x_2,
int y_1,
int z_1,
int w
) {
int pos_1 = 0;
LOOP_IS_REACHABLE while(true) {
switch(x_2) {
case 1: {
pos_1 = 1;
break;
}
default: {
break;
}
}
}
LOOP_IS_REACHABLE while(true) {
switch(x_2) {
case 1: {
break;
}
case 2: {
switch(y_1) {
case 1: {
continue;
}
default: {
switch(z_1) {
case 1: {
pos_1 = 2;
break;
}
default: {
break;
}
}
break;
}
}
break;
}
default: {
break;
}
}
}
return;
}
struct main_Input {
};
kernel void main_(
metal::uint3 global_id [[thread_position_in_grid]]
) {
int pos = {};
metal::threadgroup_barrier(metal::mem_flags::mem_device);
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
switch(1) {
default: {
pos = 1;
break;
}
}
int _e4 = pos;
switch(_e4) {
case 1: {
pos = 0;
break;
}
case 2: {
pos = 1;
break;
}
case 3:
case 4: {
pos = 2;
break;
}
case 5: {
pos = 3;
break;
}
default:
case 6: {
pos = 4;
break;
}
}
switch(0u) {
case 0u: {
break;
}
default: {
break;
}
}
int _e11 = pos;
switch(_e11) {
case 1: {
pos = 0;
break;
}
case 2: {
pos = 1;
return;
}
case 3: {
pos = 2;
return;
}
case 4: {
return;
}
default: {
pos = 3;
return;
}
}
}