test/tint/samples/compute_boids.wgsl.expected.ir.msl

#include <metal_stdlib>
using namespace metal;

struct SimParams {
  /* 0x0000 */ float deltaT;
  /* 0x0004 */ float rule1Distance;
  /* 0x0008 */ float rule2Distance;
  /* 0x000c */ float rule3Distance;
  /* 0x0010 */ float rule1Scale;
  /* 0x0014 */ float rule2Scale;
  /* 0x0018 */ float rule3Scale;
};

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 Particle {
  /* 0x0000 */ float2 pos;
  /* 0x0008 */ float2 vel;
};

struct Particles {
  /* 0x0000 */ tint_array<Particle, 5> particles;
};

struct tint_module_vars_struct {
  const constant SimParams* params;
  device Particles* particlesA;
  device Particles* particlesB;
};

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

struct vert_main_outputs {
  float4 tint_symbol [[position]];
};

struct vert_main_inputs {
  float2 a_particlePos [[attribute(0)]];
  float2 a_particleVel [[attribute(1)]];
  float2 a_pos [[attribute(2)]];
};

struct frag_main_outputs {
  float4 tint_symbol_1 [[color(0)]];
};

float4 vert_main_inner(float2 a_particlePos, float2 a_particleVel, float2 a_pos) {
  float angle = -(atan2(a_particleVel[0u], a_particleVel[1u]));
  float const v = (a_pos[0u] * cos(angle));
  float const v_1 = (v - (a_pos[1u] * sin(angle)));
  float const v_2 = (a_pos[0u] * sin(angle));
  float2 pos = float2(v_1, (v_2 + (a_pos[1u] * cos(angle))));
  return float4((pos + a_particlePos), 0.0f, 1.0f);
}

float4 frag_main_inner() {
  return float4(1.0f);
}

void comp_main_inner(uint3 gl_GlobalInvocationID, tint_module_vars_struct tint_module_vars) {
  uint index = gl_GlobalInvocationID[0u];
  if ((index >= 5u)) {
    return;
  }
  float2 vPos = (*tint_module_vars.particlesA).particles[index].pos;
  float2 vVel = (*tint_module_vars.particlesA).particles[index].vel;
  float2 cMass = float2(0.0f);
  float2 cVel = float2(0.0f);
  float2 colVel = float2(0.0f);
  int cMassCount = 0;
  int cVelCount = 0;
  float2 pos = 0.0f;
  float2 vel = 0.0f;
  {
    uint i = 0u;
    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
      if ((i < 5u)) {
      } else {
        break;
      }
      if ((i == index)) {
        {
          i = (i + 1u);
        }
        continue;
      }
      pos = (*tint_module_vars.particlesA).particles[i].pos.xy;
      vel = (*tint_module_vars.particlesA).particles[i].vel.xy;
      float const v_3 = distance(pos, vPos);
      if ((v_3 < (*tint_module_vars.params).rule1Distance)) {
        cMass = (cMass + pos);
        cMassCount = (cMassCount + 1);
      }
      float const v_4 = distance(pos, vPos);
      if ((v_4 < (*tint_module_vars.params).rule2Distance)) {
        colVel = (colVel - (pos - vPos));
      }
      float const v_5 = distance(pos, vPos);
      if ((v_5 < (*tint_module_vars.params).rule3Distance)) {
        cVel = (cVel + vel);
        cVelCount = (cVelCount + 1);
      }
      {
        i = (i + 1u);
      }
      continue;
    }
  }
  if ((cMassCount > 0)) {
    float2 const v_6 = cMass;
    float const v_7 = float(cMassCount);
    float2 const v_8 = (v_6 / float2(v_7, float(cMassCount)));
    cMass = (v_8 - vPos);
  }
  if ((cVelCount > 0)) {
    float2 const v_9 = cVel;
    float const v_10 = float(cVelCount);
    cVel = (v_9 / float2(v_10, float(cVelCount)));
  }
  vVel = (((vVel + (cMass * (*tint_module_vars.params).rule1Scale)) + (colVel * (*tint_module_vars.params).rule2Scale)) + (cVel * (*tint_module_vars.params).rule3Scale));
  float2 const v_11 = normalize(vVel);
  vVel = (v_11 * clamp(length(vVel), 0.0f, 0.10000000149011611938f));
  vPos = (vPos + (vVel * (*tint_module_vars.params).deltaT));
  if ((vPos[0u] < -1.0f)) {
    vPos[0u] = 1.0f;
  }
  if ((vPos[0u] > 1.0f)) {
    vPos[0u] = -1.0f;
  }
  if ((vPos[1u] < -1.0f)) {
    vPos[1u] = 1.0f;
  }
  if ((vPos[1u] > 1.0f)) {
    vPos[1u] = -1.0f;
  }
  (*tint_module_vars.particlesB).particles[index].pos = vPos;
  (*tint_module_vars.particlesB).particles[index].vel = vVel;
}

vertex vert_main_outputs vert_main(vert_main_inputs inputs [[stage_in]]) {
  vert_main_outputs tint_wrapper_result = {};
  tint_wrapper_result.tint_symbol = vert_main_inner(inputs.a_particlePos, inputs.a_particleVel, inputs.a_pos);
  return tint_wrapper_result;
}

fragment frag_main_outputs frag_main() {
  frag_main_outputs tint_wrapper_result = {};
  tint_wrapper_result.tint_symbol_1 = frag_main_inner();
  return tint_wrapper_result;
}

kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], const constant SimParams* params [[buffer(0)]], device Particles* particlesA [[buffer(1)]], device Particles* particlesB [[buffer(2)]]) {
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.params=params, .particlesA=particlesA, .particlesB=particlesB};
  comp_main_inner(gl_GlobalInvocationID, tint_module_vars);
}