SKIP: FAILED

#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 {
  float a;
};

thread bool bool_var = false;
thread int i32_var = 0;
thread uint u32_var = 0u;
thread float f32_var = 0.0f;
thread int2 v2i32_var = 0;
thread uint3 v3u32_var = 0u;
thread float4 v4f32_var = 0.0f;
thread float2x3 m2x3_var = float2x3(0.0f);
thread tint_array<float, 4> arr_var = {};
thread S struct_var = {};
threadgroup float wg_var;
kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
  if ((tint_local_index == 0u)) {
    wg_var = 0.0f;
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  bool_var = false;
  i32_var = 0;
  u32_var = 0u;
  f32_var = 0.0f;
  v2i32_var = int2(0);
  v3u32_var = uint3(0u);
  v4f32_var = float4(0.0f);
  m2x3_var = float2x3(float3(0.0f), float3(0.0f));
  arr_var = tint_array<float, 4>{};
  struct_var = S{};
  wg_var = 42.0f;
}
program_source:19:13: error: program scope variable must reside in constant address space
thread bool bool_var = false;
            ^
program_source:20:12: error: program scope variable must reside in constant address space
thread int i32_var = 0;
           ^
program_source:21:13: error: program scope variable must reside in constant address space
thread uint u32_var = 0u;
            ^
program_source:22:14: error: program scope variable must reside in constant address space
thread float f32_var = 0.0f;
             ^
program_source:23:13: error: program scope variable must reside in constant address space
thread int2 v2i32_var = 0;
            ^
program_source:24:14: error: program scope variable must reside in constant address space
thread uint3 v3u32_var = 0u;
             ^
program_source:25:15: error: program scope variable must reside in constant address space
thread float4 v4f32_var = 0.0f;
              ^
program_source:26:17: error: program scope variable must reside in constant address space
thread float2x3 m2x3_var = float2x3(0.0f);
                ^
program_source:27:29: error: program scope variable must reside in constant address space
thread tint_array<float, 4> arr_var = {};
                            ^
program_source:28:10: error: program scope variable must reside in constant address space
thread S struct_var = {};
         ^
program_source:29:19: error: program scope variable must reside in constant address space
threadgroup float wg_var;
                  ^
program_source:31:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
  if ((tint_local_index == 0u)) {
       ~~~~~~~~~~~~~~~~~^~~~~
program_source:31:25: note: remove extraneous parentheses around the comparison to silence this warning
  if ((tint_local_index == 0u)) {
      ~                 ^    ~
program_source:31:25: note: use '=' to turn this equality comparison into an assignment
  if ((tint_local_index == 0u)) {
                        ^~
                        =