mirror of
https://github.com/gfx-rs/wgpu.git
synced 2024-11-29 02:03:35 +00:00
3fda684eb9
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.
226 lines
4.3 KiB
Plaintext
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;
|
|
}
|
|
}
|
|
}
|