test/tint/bug/tint/2010.spvasm.expected.ir.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];
};

struct S {
  float2 field0;
  uint field1;
};

struct S_1 {
  /* 0x0000 */ uint field0;
};

struct S_2 {
  /* 0x0000 */ S_1 field0;
};

struct S_3 {
  /* 0x0000 */ tint_array<float4, 1> field0;
};

struct S_4 {
  /* 0x0000 */ tint_array<float4, 1> field0;
};

struct tint_module_vars_struct {
  threadgroup tint_array<S, 4096>* x_28;
  threadgroup atomic_uint* x_34;
  threadgroup atomic_uint* x_35;
  threadgroup atomic_uint* x_36;
  threadgroup atomic_uint* x_37;
  thread uint3* x_3;
  const constant S_2* x_6;
  const device S_3* x_9;
  device S_4* x_12;
  const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
};

#define TINT_ISOLATE_UB(VOLATILE_NAME) \
  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}

struct tint_symbol_6 {
  tint_array<S, 4096> tint_symbol_1;
  atomic_uint tint_symbol_2;
  atomic_uint tint_symbol_3;
  atomic_uint tint_symbol_4;
  atomic_uint tint_symbol_5;
};

void main_1(tint_module_vars_struct tint_module_vars) {
  uint x_54 = 0u;
  uint x_58 = 0u;
  float4 x_85 = 0.0f;
  uint x_88 = 0u;
  uint const x_52 = (*tint_module_vars.x_3)[0u];
  x_54 = 0u;
  {
    while(true) {
      TINT_ISOLATE_UB(tint_volatile_false)
      uint x_55 = 0u;
      x_58 = (*tint_module_vars.x_6).field0.field0;
      if ((x_54 < x_58)) {
      } else {
        break;
      }
      uint const x_62 = (x_54 + x_52);
      if ((x_62 >= x_58)) {
        float4 const x_67 = (*tint_module_vars.x_9).field0[min(x_62, ((((*tint_module_vars.tint_storage_buffer_sizes)[0u][0u] - 0u) / 16u) - 1u))];
        (*tint_module_vars.x_28)[min(x_62, 4095u)] = S{.field0=((x_67.xy + x_67.zw) * 0.5f), .field1=x_62};
      }
      {
        x_55 = (x_54 + 32u);
        x_54 = x_55;
      }
      continue;
    }
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  int const x_74 = as_type<int>(x_58);
  float2 const x_76 = (*tint_module_vars.x_28)[0u].field0;
  if ((x_52 == 0u)) {
    uint2 const x_80 = as_type<uint2>(x_76);
    uint const x_81 = x_80[0u];
    atomic_store_explicit(tint_module_vars.x_34, x_81, memory_order_relaxed);
    uint const x_82 = x_80[1u];
    atomic_store_explicit(tint_module_vars.x_35, x_82, memory_order_relaxed);
    atomic_store_explicit(tint_module_vars.x_36, x_81, memory_order_relaxed);
    atomic_store_explicit(tint_module_vars.x_37, x_82, memory_order_relaxed);
  }
  x_85 = x_76.xyxy;
  x_88 = 1u;
  {
    while(true) {
      TINT_ISOLATE_UB(tint_volatile_false_1)
      float4 x_111 = 0.0f;
      float4 x_86 = 0.0f;
      uint x_89 = 0u;
      uint const x_90 = as_type<uint>(x_74);
      if ((x_88 < x_90)) {
      } else {
        break;
      }
      uint const x_94 = (x_88 + x_52);
      x_86 = x_85;
      if ((x_94 >= x_90)) {
        float2 const x_99 = (*tint_module_vars.x_28)[min(x_94, 4095u)].field0;
        float2 const x_101 = min(x_85.xy, x_99);
        float4 x_103_1 = x_85;
        x_103_1[0u] = x_101[0u];
        float4 const x_103 = x_103_1;
        float4 x_105_1 = x_103;
        x_105_1[1u] = x_101[1u];
        float4 const x_105 = x_105_1;
        float2 const x_107 = max(x_105_1.zw, x_99);
        float4 x_109_1 = x_105;
        x_109_1[2u] = x_107[0u];
        x_111 = x_109_1;
        x_111[3u] = x_107[1u];
        x_86 = x_111;
      }
      {
        x_89 = (x_88 + 32u);
        x_85 = x_86;
        x_88 = x_89;
      }
      continue;
    }
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  uint const x_114 = atomic_fetch_min_explicit(tint_module_vars.x_34, as_type<uint>(x_85[0u]), memory_order_relaxed);
  uint const x_117 = atomic_fetch_min_explicit(tint_module_vars.x_35, as_type<uint>(x_85[1u]), memory_order_relaxed);
  uint const x_120 = atomic_fetch_max_explicit(tint_module_vars.x_36, as_type<uint>(x_85[2u]), memory_order_relaxed);
  uint const x_123 = atomic_fetch_max_explicit(tint_module_vars.x_37, as_type<uint>(x_85[3u]), memory_order_relaxed);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  uint const v = ((((*tint_module_vars.tint_storage_buffer_sizes)[0u][1u] - 0u) / 16u) - 1u);
  device float4* const v_1 = (&(*tint_module_vars.x_12).field0[min(uint(0), v)]);
  (*v_1) = float4(as_type<float>(atomic_load_explicit(tint_module_vars.x_34, memory_order_relaxed)), as_type<float>(atomic_load_explicit(tint_module_vars.x_35, memory_order_relaxed)), as_type<float>(atomic_load_explicit(tint_module_vars.x_36, memory_order_relaxed)), as_type<float>(atomic_load_explicit(tint_module_vars.x_37, memory_order_relaxed)));
}

void tint_symbol_inner(uint3 x_3_param, uint tint_local_index, tint_module_vars_struct tint_module_vars) {
  if ((tint_local_index < 1u)) {
    atomic_store_explicit(tint_module_vars.x_34, 0u, memory_order_relaxed);
    atomic_store_explicit(tint_module_vars.x_35, 0u, memory_order_relaxed);
    atomic_store_explicit(tint_module_vars.x_36, 0u, memory_order_relaxed);
    atomic_store_explicit(tint_module_vars.x_37, 0u, memory_order_relaxed);
  }
  {
    uint v_2 = 0u;
    v_2 = tint_local_index;
    while(true) {
      TINT_ISOLATE_UB(tint_volatile_false_2)
      uint const v_3 = v_2;
      if ((v_3 >= 4096u)) {
        break;
      }
      (*tint_module_vars.x_28)[v_3] = S{};
      {
        v_2 = (v_3 + 32u);
      }
      continue;
    }
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  (*tint_module_vars.x_3) = x_3_param;
  main_1(tint_module_vars);
}

kernel void tint_symbol(uint3 x_3_param [[thread_position_in_threadgroup]], uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_6* v_4 [[threadgroup(0)]], const constant S_2* x_6 [[buffer(0)]], const device S_3* x_9 [[buffer(2)]], device S_4* x_12 [[buffer(1)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
  thread uint3 x_3 = 0u;
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.x_28=(&(*v_4).tint_symbol_1), .x_34=(&(*v_4).tint_symbol_2), .x_35=(&(*v_4).tint_symbol_3), .x_36=(&(*v_4).tint_symbol_4), .x_37=(&(*v_4).tint_symbol_5), .x_3=(&x_3), .x_6=x_6, .x_9=x_9, .x_12=x_12, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
  tint_symbol_inner(x_3_param, tint_local_index, tint_module_vars);
}