#include <metal_stdlib>
using namespace metal;

struct ComputeInputs {
  uint subgroup_invocation_id;
  uint subgroup_size;
};

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 tint_module_vars_struct {
  device tint_array<uint, 1>* output;
};

void tint_symbol_inner(ComputeInputs inputs, tint_module_vars_struct tint_module_vars) {
  (*tint_module_vars.output)[inputs.subgroup_invocation_id] = inputs.subgroup_size;
}

kernel void tint_symbol(uint ComputeInputs_subgroup_invocation_id [[thread_index_in_simdgroup]], uint ComputeInputs_subgroup_size [[threads_per_simdgroup]], device tint_array<uint, 1>* output [[buffer(0)]]) {
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.output=output};
  tint_symbol_inner(ComputeInputs{.subgroup_invocation_id=ComputeInputs_subgroup_invocation_id, .subgroup_size=ComputeInputs_subgroup_size}, tint_module_vars);
}