test/tint/bug/tint/221.wgsl.expected.msl

#include <metal_stdlib>

using namespace metal;

template<typename T, size_t N>
struct tint_array {
    const constant T& operator[](size_t i) const constant { return elements[i]; }
    device T& operator[](size_t i) device { return elements[i]; }
    const device T& operator[](size_t i) const device { return elements[i]; }
    thread T& operator[](size_t i) thread { return elements[i]; }
    const thread T& operator[](size_t i) const thread { return elements[i]; }
    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
    T elements[N];
};

#define TINT_ISOLATE_UB(VOLATILE_NAME) \
  volatile bool VOLATILE_NAME = true; \
  if (VOLATILE_NAME)

struct Buf {
  /* 0x0000 */ uint count;
  /* 0x0004 */ tint_array<uint, 50> data;
};

uint tint_mod(uint lhs, uint rhs) {
  return (lhs % select(rhs, 1u, (rhs == 0u)));
}

kernel void tint_symbol(device Buf* tint_symbol_1 [[buffer(0)]]) {
  uint i = 0u;
  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
    if ((i >= (*(tint_symbol_1)).count)) {
      break;
    }
    uint const p_save = i;
    if ((tint_mod(i, 2u) == 0u)) {
      {
        (*(tint_symbol_1)).data[p_save] = ((*(tint_symbol_1)).data[p_save] * 2u);
        i = (i + 1u);
      }
      continue;
    }
    (*(tint_symbol_1)).data[p_save] = 0u;
    {
      (*(tint_symbol_1)).data[p_save] = ((*(tint_symbol_1)).data[p_save] * 2u);
      i = (i + 1u);
    }
  }
  return;
}