test/tint/bug/chromium/1434271.wgsl.expected.ir.msl

#include <metal_stdlib>
using namespace metal;

struct VertexOutput {
  float4 position;
  float4 color;
  float2 quad_pos;
};

struct VertexInput {
  float3 position;
  float4 color;
  float2 quad_pos;
};

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 RenderParams_packed_vec3 {
  /* 0x0000 */ float4x4 modelViewProjectionMatrix;
  /* 0x0040 */ packed_float3 right;
  /* 0x004c */ tint_array<int8_t, 4> tint_pad;
  /* 0x0050 */ packed_float3 up;
  /* 0x005c */ tint_array<int8_t, 4> tint_pad_1;
};

struct SimulationParams {
  /* 0x0000 */ float deltaTime;
  /* 0x0004 */ tint_array<int8_t, 12> tint_pad_2;
  /* 0x0010 */ float4 seed;
};

struct Particle_packed_vec3 {
  /* 0x0000 */ packed_float3 position;
  /* 0x000c */ float lifetime;
  /* 0x0010 */ float4 color;
  /* 0x0020 */ float2 velocity;
  /* 0x0028 */ tint_array<int8_t, 8> tint_pad_3;
};

struct Particles_packed_vec3 {
  /* 0x0000 */ tint_array<Particle_packed_vec3, 1> particles;
};

struct UBO {
  /* 0x0000 */ uint width;
};

struct Buffer {
  /* 0x0000 */ tint_array<float, 1> weights;
};

struct tint_module_vars_struct {
  thread float2* rand_seed;
  const constant RenderParams_packed_vec3* render_params;
  const constant SimulationParams* sim_params;
  device Particles_packed_vec3* data;
  texture1d<float, access::sample> tint_symbol;
  const constant UBO* ubo;
  const device Buffer* buf_in;
  device Buffer* buf_out;
  texture2d<float, access::sample> tex_in;
  texture2d<float, access::write> tex_out;
  const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
};

struct Particle {
  float3 position;
  float lifetime;
  float4 color;
  float2 velocity;
};

struct vertex_main_outputs {
  float4 tint_symbol_1 [[position]];
};

struct vs_main_outputs {
  float4 VertexOutput_position [[position]];
  float4 VertexOutput_color [[user(locn0)]];
  float2 VertexOutput_quad_pos [[user(locn1)]];
};

struct vs_main_inputs {
  float3 VertexInput_position [[attribute(0)]];
  float4 VertexInput_color [[attribute(1)]];
  float2 VertexInput_quad_pos [[attribute(2)]];
};

void asinh_468a48() {
  half arg_0 = 0.0h;
  half res = asinh(arg_0);
}

float4 vertex_main_inner() {
  asinh_468a48();
  return float4(0.0f);
}

fragment void fragment_main() {
  asinh_468a48();
}

kernel void rgba32uintin() {
  asinh_468a48();
}

VertexOutput vs_main_inner(VertexInput in, tint_module_vars_struct tint_module_vars) {
  float3 const v = float3((*tint_module_vars.render_params).right);
  float3 quad_pos = (float2x3(v, float3((*tint_module_vars.render_params).up)) * in.quad_pos);
  float3 position = (in.position - (quad_pos + 0.00999999977648258209f));
  VertexOutput out = {};
  float4x4 const v_1 = (*tint_module_vars.render_params).modelViewProjectionMatrix;
  out.position = (v_1 * float4(position, 1.0f));
  out.color = in.color;
  out.quad_pos = in.quad_pos;
  return out;
}

void tint_store_and_preserve_padding(device Particle_packed_vec3* const target, Particle value_param) {
  (*target).position = packed_float3(value_param.position);
  (*target).lifetime = value_param.lifetime;
  (*target).color = value_param.color;
  (*target).velocity = value_param.velocity;
}

Particle tint_load_struct_packed_vec3(device Particle_packed_vec3* const from) {
  float3 const v_2 = float3((*from).position);
  return Particle{.position=v_2, .lifetime=(*from).lifetime, .color=(*from).color, .velocity=(*from).velocity};
}

void simulate_inner(uint3 GlobalInvocationID, tint_module_vars_struct tint_module_vars) {
  float2 const v_3 = (*tint_module_vars.sim_params).seed.xy;
  float2 const v_4 = (v_3 * float2(GlobalInvocationID.xy));
  (*tint_module_vars.rand_seed) = (v_4 * (*tint_module_vars.sim_params).seed.zw);
  uint const idx = GlobalInvocationID.x;
  Particle particle = tint_load_struct_packed_vec3((&(*tint_module_vars.data).particles[min(idx, ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].x - 0u) / 48u) - 1u))]));
  tint_store_and_preserve_padding((&(*tint_module_vars.data).particles[min(idx, ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].x - 0u) / 48u) - 1u))]), particle);
}

void export_level_inner(uint3 coord, tint_module_vars_struct tint_module_vars) {
  if (all((coord.xy < uint2(uint2(tint_module_vars.tex_out.get_width(0u), tint_module_vars.tex_out.get_height(0u)))))) {
    uint const dst_offset = (coord.x << ((coord.y * (*tint_module_vars.ubo).width) & 31u));
    uint const src_offset = ((coord.x - 2u) + ((coord.y >> (2u & 31u)) * (*tint_module_vars.ubo).width));
    float const a = (*tint_module_vars.buf_in).weights[min((src_offset << (0u & 31u)), ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].y - 0u) / 4u) - 1u))];
    float const b = (*tint_module_vars.buf_in).weights[min((src_offset + 1u), ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].y - 0u) / 4u) - 1u))];
    float const c = (*tint_module_vars.buf_in).weights[min(((src_offset + 1u) + (*tint_module_vars.ubo).width), ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].y - 0u) / 4u) - 1u))];
    float const d = (*tint_module_vars.buf_in).weights[min(((src_offset + 1u) + (*tint_module_vars.ubo).width), ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].y - 0u) / 4u) - 1u))];
    float const sum = dot(float4(a, b, c, d), float4(1.0f));
    (*tint_module_vars.buf_out).weights[min(dst_offset, ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].z - 0u) / 4u) - 1u))] = fmod(sum, 4.0f);
    float4 const probabilities = (float4(a, (a * b), ((a / b) + c), sum) + max(sum, 0.0f));
    tint_module_vars.tex_out.write(probabilities, uint2(int2(coord.xy)));
  }
}

vertex vertex_main_outputs vertex_main() {
  vertex_main_outputs tint_wrapper_result = {};
  tint_wrapper_result.tint_symbol_1 = vertex_main_inner();
  return tint_wrapper_result;
}

vertex vs_main_outputs vs_main(vs_main_inputs inputs [[stage_in]], const constant RenderParams_packed_vec3* render_params [[buffer(0)]]) {
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.render_params=render_params};
  VertexOutput const v_5 = vs_main_inner(VertexInput{.position=inputs.VertexInput_position, .color=inputs.VertexInput_color, .quad_pos=inputs.VertexInput_quad_pos}, tint_module_vars);
  vs_main_outputs tint_wrapper_result = {};
  tint_wrapper_result.VertexOutput_position = v_5.position;
  tint_wrapper_result.VertexOutput_color = v_5.color;
  tint_wrapper_result.VertexOutput_quad_pos = v_5.quad_pos;
  return tint_wrapper_result;
}

kernel void simulate(uint3 GlobalInvocationID [[thread_position_in_grid]], const constant SimulationParams* sim_params [[buffer(1)]], device Particles_packed_vec3* data [[buffer(2)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
  thread float2 rand_seed = 0.0f;
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.rand_seed=(&rand_seed), .sim_params=sim_params, .data=data, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
  simulate_inner(GlobalInvocationID, tint_module_vars);
}

kernel void export_level(uint3 coord [[thread_position_in_grid]], const constant UBO* ubo [[buffer(3)]], const device Buffer* buf_in [[buffer(4)]], device Buffer* buf_out [[buffer(0)]], texture2d<float, access::write> tex_out [[texture(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.ubo=ubo, .buf_in=buf_in, .buf_out=buf_out, .tex_out=tex_out, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
  export_level_inner(coord, tint_module_vars);
}