test/vk-gl-cts/non_robust_buffer_access/unexecuted_oob_underflow/0-opt.wgsl.expected.msl

#include <metal_stdlib>

using namespace metal;
struct tint_array_wrapper {
  /* 0x0000 */ int arr[8];
};
struct In2 {
  /* 0x0000 */ tint_array_wrapper data_in2;
};
struct tint_array_wrapper_1 {
  /* 0x0000 */ int arr[1024];
};
struct Out0 {
  /* 0x0000 */ tint_array_wrapper_1 data_out0;
};
struct tint_array_wrapper_2 {
  /* 0x0000 */ int arr[512];
};
struct In0 {
  /* 0x0000 */ tint_array_wrapper_2 data_in0;
};
struct In1 {
  /* 0x0000 */ tint_array_wrapper_2 data_in1;
};

void main_1(const device In2& x_13, const device In0& x_17, device Out0& x_15, const device In1& x_19, thread uint3* const tint_symbol_1) {
  uint base_index_in = 0u;
  uint base_index_out = 0u;
  int index_in0 = 0;
  int index_in1 = 0;
  int index_out0 = 0;
  int index_out1 = 0;
  int condition_index = 0;
  int i = 0;
  int temp0 = 0;
  int temp1 = 0;
  uint const x_56 = (*(tint_symbol_1)).x;
  base_index_in = (128u * x_56);
  uint const x_59 = (*(tint_symbol_1)).x;
  base_index_out = (256u * x_59);
  index_in0 = 0;
  index_in1 = -128;
  index_out0 = 0;
  index_out1 = -128;
  condition_index = 0;
  i = 0;
  while (true) {
    int const x_65 = i;
    if ((x_65 < 256)) {
    } else {
      break;
    }
    int const x_68 = condition_index;
    int const x_70 = x_13.data_in2.arr[x_68];
    if ((x_70 == 0)) {
      uint const x_75 = base_index_out;
      int const x_76 = index_out0;
      uint const x_79 = base_index_in;
      int const x_80 = index_in0;
      int const x_84 = x_17.data_in0.arr[(x_79 + as_type<uint>(x_80))];
      x_15.data_out0.arr[(x_75 + as_type<uint>(x_76))] = x_84;
      int const x_86 = index_out0;
      index_out0 = as_type<int>((as_type<uint>(x_86) + as_type<uint>(1)));
      int const x_88 = index_in1;
      index_in1 = as_type<int>((as_type<uint>(x_88) + as_type<uint>(1)));
    } else {
      uint const x_90 = base_index_out;
      int const x_91 = index_out1;
      uint const x_94 = base_index_in;
      int const x_95 = index_in1;
      int const x_99 = x_19.data_in1.arr[(x_94 + as_type<uint>(x_95))];
      x_15.data_out0.arr[(x_90 + as_type<uint>(x_91))] = x_99;
      int const x_101 = index_out1;
      index_out1 = as_type<int>((as_type<uint>(x_101) + as_type<uint>(1)));
      int const x_103 = index_in1;
      index_in1 = as_type<int>((as_type<uint>(x_103) + as_type<uint>(1)));
    }
    int const x_105 = condition_index;
    int const x_108 = x_13.data_in2.arr[as_type<int>((as_type<uint>(x_105) + as_type<uint>(1)))];
    int const x_109 = condition_index;
    condition_index = as_type<int>((as_type<uint>(x_109) + as_type<uint>(x_108)));
    int const x_111 = index_in0;
    temp0 = x_111;
    int const x_112 = index_in1;
    index_in0 = x_112;
    int const x_113 = temp0;
    index_in1 = x_113;
    int const x_114 = index_out0;
    temp1 = x_114;
    int const x_115 = index_out1;
    index_out0 = x_115;
    int const x_116 = temp1;
    index_out1 = x_116;
    {
      int const x_117 = i;
      i = as_type<int>((as_type<uint>(x_117) + as_type<uint>(1)));
    }
  }
  return;
}

void tint_symbol_inner(const device In2& x_13, const device In0& x_17, device Out0& x_15, const device In1& x_19, uint3 gl_WorkGroupID_param, thread uint3* const tint_symbol_2) {
  *(tint_symbol_2) = gl_WorkGroupID_param;
  main_1(x_13, x_17, x_15, x_19, tint_symbol_2);
}

kernel void tint_symbol(uint3 gl_WorkGroupID_param [[threadgroup_position_in_grid]], const device In2& x_13 [[buffer(2)]], const device In0& x_17 [[buffer(0)]], device Out0& x_15 [[buffer(3)]], const device In1& x_19 [[buffer(1)]]) {
  thread uint3 tint_symbol_3 = 0u;
  tint_symbol_inner(x_13, x_17, x_15, x_19, gl_WorkGroupID_param, &(tint_symbol_3));
  return;
}