test/benchmark/metaball-isosurface.wgsl.expected.msl

#include <metal_stdlib>

using namespace metal;

template<typename T, int N, int M>
inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
  return lhs * vec<T, N>(rhs);
}

template<typename T, int N, int M>
inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
  return vec<T, M>(lhs) * rhs;
}

struct tint_array_wrapper {
  /* 0x0000 */ uint arr[256];
};
struct tint_array_wrapper_1 {
  /* 0x0000 */ int arr[4096];
};
struct Tables {
  /* 0x0000 */ tint_array_wrapper edges;
  /* 0x0400 */ tint_array_wrapper_1 tris;
};
struct IsosurfaceVolume {
  /* 0x0000 */ packed_float3 min;
  /* 0x000c */ int8_t tint_pad[4];
  /* 0x0010 */ packed_float3 max;
  /* 0x001c */ int8_t tint_pad_1[4];
  /* 0x0020 */ packed_float3 step;
  /* 0x002c */ int8_t tint_pad_2[4];
  /* 0x0030 */ packed_uint3 size;
  /* 0x003c */ float threshold;
  /* 0x0040 */ float values[1];
  /* 0x0044 */ int8_t tint_pad_3[12];
};
struct PositionBuffer {
  /* 0x0000 */ float values[1];
};
struct NormalBuffer {
  /* 0x0000 */ float values[1];
};
struct IndexBuffer {
  /* 0x0000 */ uint tris[1];
};
struct DrawIndirectArgs {
  /* 0x0000 */ uint vc;
  /* 0x0004 */ atomic_uint vertexCount;
  /* 0x0008 */ uint firstVertex;
  /* 0x000c */ uint firstInstance;
  /* 0x0010 */ atomic_uint indexCount;
  /* 0x0014 */ uint indexedInstanceCount;
  /* 0x0018 */ uint indexedFirstIndex;
  /* 0x001c */ uint indexedBaseVertex;
  /* 0x0020 */ uint indexedFirstInstance;
};
struct tint_array_wrapper_2 {
  float3 arr[12];
};
struct tint_array_wrapper_3 {
  uint arr[12];
};

float valueAt(uint3 index, device IsosurfaceVolume* const tint_symbol) {
  if (any((index >= (*(tint_symbol)).size))) {
    return 0.0f;
  }
  uint const valueIndex = ((index[0] + (index[1] * (*(tint_symbol)).size[0])) + ((index[2] * (*(tint_symbol)).size[0]) * (*(tint_symbol)).size[1]));
  return (*(tint_symbol)).values[valueIndex];
}

float3 positionAt(uint3 index, device IsosurfaceVolume* const tint_symbol_1) {
  return ((*(tint_symbol_1)).min + ((*(tint_symbol_1)).step * float3(uint3(index).xyz)));
}

float3 normalAt(uint3 index, device IsosurfaceVolume* const tint_symbol_2) {
  return float3((valueAt((index - uint3(1u, 0u, 0u)), tint_symbol_2) - valueAt((index + uint3(1u, 0u, 0u)), tint_symbol_2)), (valueAt((index - uint3(0u, 1u, 0u)), tint_symbol_2) - valueAt((index + uint3(0u, 1u, 0u)), tint_symbol_2)), (valueAt((index - uint3(0u, 0u, 1u)), tint_symbol_2) - valueAt((index + uint3(0u, 0u, 1u)), tint_symbol_2)));
}

void interpX(uint index, uint3 i, float va, float vb, device IsosurfaceVolume* const tint_symbol_3, thread tint_array_wrapper_2* const tint_symbol_4, thread uint* const tint_symbol_5, thread tint_array_wrapper_2* const tint_symbol_6, thread tint_array_wrapper_3* const tint_symbol_7) {
  float const mu = (((*(tint_symbol_3)).threshold - va) / (vb - va));
  (*(tint_symbol_4)).arr[*(tint_symbol_5)] = (positionAt(i, tint_symbol_3) + float3(((*(tint_symbol_3)).step[0] * mu), 0.0f, 0.0f));
  float3 const na = normalAt(i, tint_symbol_3);
  float3 const nb = normalAt((i + uint3(1u, 0u, 0u)), tint_symbol_3);
  (*(tint_symbol_6)).arr[*(tint_symbol_5)] = mix(na, nb, float3(mu, mu, mu));
  (*(tint_symbol_7)).arr[index] = *(tint_symbol_5);
  *(tint_symbol_5) = (*(tint_symbol_5) + 1u);
}

void interpY(uint index, uint3 i, float va, float vb, device IsosurfaceVolume* const tint_symbol_8, thread tint_array_wrapper_2* const tint_symbol_9, thread uint* const tint_symbol_10, thread tint_array_wrapper_2* const tint_symbol_11, thread tint_array_wrapper_3* const tint_symbol_12) {
  float const mu = (((*(tint_symbol_8)).threshold - va) / (vb - va));
  (*(tint_symbol_9)).arr[*(tint_symbol_10)] = (positionAt(i, tint_symbol_8) + float3(0.0f, ((*(tint_symbol_8)).step[1] * mu), 0.0f));
  float3 const na = normalAt(i, tint_symbol_8);
  float3 const nb = normalAt((i + uint3(0u, 1u, 0u)), tint_symbol_8);
  (*(tint_symbol_11)).arr[*(tint_symbol_10)] = mix(na, nb, float3(mu, mu, mu));
  (*(tint_symbol_12)).arr[index] = *(tint_symbol_10);
  *(tint_symbol_10) = (*(tint_symbol_10) + 1u);
}

void interpZ(uint index, uint3 i, float va, float vb, device IsosurfaceVolume* const tint_symbol_13, thread tint_array_wrapper_2* const tint_symbol_14, thread uint* const tint_symbol_15, thread tint_array_wrapper_2* const tint_symbol_16, thread tint_array_wrapper_3* const tint_symbol_17) {
  float const mu = (((*(tint_symbol_13)).threshold - va) / (vb - va));
  (*(tint_symbol_14)).arr[*(tint_symbol_15)] = (positionAt(i, tint_symbol_13) + float3(0.0f, 0.0f, ((*(tint_symbol_13)).step[2] * mu)));
  float3 const na = normalAt(i, tint_symbol_13);
  float3 const nb = normalAt((i + uint3(0u, 0u, 1u)), tint_symbol_13);
  (*(tint_symbol_16)).arr[*(tint_symbol_15)] = mix(na, nb, float3(mu, mu, mu));
  (*(tint_symbol_17)).arr[index] = *(tint_symbol_15);
  *(tint_symbol_15) = (*(tint_symbol_15) + 1u);
}

void computeMain_inner(uint3 global_id, device IsosurfaceVolume* const tint_symbol_18, const device Tables* const tint_symbol_19, thread tint_array_wrapper_2* const tint_symbol_20, thread uint* const tint_symbol_21, thread tint_array_wrapper_2* const tint_symbol_22, thread tint_array_wrapper_3* const tint_symbol_23, device DrawIndirectArgs* const tint_symbol_24, device PositionBuffer* const tint_symbol_25, device NormalBuffer* const tint_symbol_26, device IndexBuffer* const tint_symbol_27) {
  uint3 const i0 = global_id;
  uint3 const i1 = (global_id + uint3(1u, 0u, 0u));
  uint3 const i2 = (global_id + uint3(1u, 1u, 0u));
  uint3 const i3 = (global_id + uint3(0u, 1u, 0u));
  uint3 const i4 = (global_id + uint3(0u, 0u, 1u));
  uint3 const i5 = (global_id + uint3(1u, 0u, 1u));
  uint3 const i6 = (global_id + uint3(1u, 1u, 1u));
  uint3 const i7 = (global_id + uint3(0u, 1u, 1u));
  float const v0 = valueAt(i0, tint_symbol_18);
  float const v1 = valueAt(i1, tint_symbol_18);
  float const v2 = valueAt(i2, tint_symbol_18);
  float const v3 = valueAt(i3, tint_symbol_18);
  float const v4 = valueAt(i4, tint_symbol_18);
  float const v5 = valueAt(i5, tint_symbol_18);
  float const v6 = valueAt(i6, tint_symbol_18);
  float const v7 = valueAt(i7, tint_symbol_18);
  uint cubeIndex = 0u;
  if ((v0 < (*(tint_symbol_18)).threshold)) {
    cubeIndex = (cubeIndex | 1u);
  }
  if ((v1 < (*(tint_symbol_18)).threshold)) {
    cubeIndex = (cubeIndex | 2u);
  }
  if ((v2 < (*(tint_symbol_18)).threshold)) {
    cubeIndex = (cubeIndex | 4u);
  }
  if ((v3 < (*(tint_symbol_18)).threshold)) {
    cubeIndex = (cubeIndex | 8u);
  }
  if ((v4 < (*(tint_symbol_18)).threshold)) {
    cubeIndex = (cubeIndex | 16u);
  }
  if ((v5 < (*(tint_symbol_18)).threshold)) {
    cubeIndex = (cubeIndex | 32u);
  }
  if ((v6 < (*(tint_symbol_18)).threshold)) {
    cubeIndex = (cubeIndex | 64u);
  }
  if ((v7 < (*(tint_symbol_18)).threshold)) {
    cubeIndex = (cubeIndex | 128u);
  }
  uint const edges = (*(tint_symbol_19)).edges.arr[cubeIndex];
  if (((edges & 1u) != 0u)) {
    interpX(0u, i0, v0, v1, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  if (((edges & 2u) != 0u)) {
    interpY(1u, i1, v1, v2, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  if (((edges & 4u) != 0u)) {
    interpX(2u, i3, v3, v2, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  if (((edges & 8u) != 0u)) {
    interpY(3u, i0, v0, v3, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  if (((edges & 16u) != 0u)) {
    interpX(4u, i4, v4, v5, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  if (((edges & 32u) != 0u)) {
    interpY(5u, i5, v5, v6, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  if (((edges & 64u) != 0u)) {
    interpX(6u, i7, v7, v6, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  if (((edges & 128u) != 0u)) {
    interpY(7u, i4, v4, v7, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  if (((edges & 256u) != 0u)) {
    interpZ(8u, i0, v0, v4, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  if (((edges & 512u) != 0u)) {
    interpZ(9u, i1, v1, v5, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  if (((edges & 1024u) != 0u)) {
    interpZ(10u, i2, v2, v6, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  if (((edges & 2048u) != 0u)) {
    interpZ(11u, i3, v3, v7, tint_symbol_18, tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23);
  }
  uint const triTableOffset = ((cubeIndex << 4u) + 1u);
  uint const indexCount = uint((*(tint_symbol_19)).tris.arr[(triTableOffset - 1u)]);
  uint firstVertex = atomic_fetch_add_explicit(&((*(tint_symbol_24)).vertexCount), *(tint_symbol_21), memory_order_relaxed);
  uint const bufferOffset = ((global_id[0] + (global_id[1] * (*(tint_symbol_18)).size[0])) + ((global_id[2] * (*(tint_symbol_18)).size[0]) * (*(tint_symbol_18)).size[1]));
  uint const firstIndex = (bufferOffset * 15u);
  for(uint i = 0u; (i < *(tint_symbol_21)); i = (i + 1u)) {
    (*(tint_symbol_25)).values[((firstVertex * 3u) + (i * 3u))] = (*(tint_symbol_20)).arr[i][0];
    (*(tint_symbol_25)).values[(((firstVertex * 3u) + (i * 3u)) + 1u)] = (*(tint_symbol_20)).arr[i][1];
    (*(tint_symbol_25)).values[(((firstVertex * 3u) + (i * 3u)) + 2u)] = (*(tint_symbol_20)).arr[i][2];
    (*(tint_symbol_26)).values[((firstVertex * 3u) + (i * 3u))] = (*(tint_symbol_22)).arr[i][0];
    (*(tint_symbol_26)).values[(((firstVertex * 3u) + (i * 3u)) + 1u)] = (*(tint_symbol_22)).arr[i][1];
    (*(tint_symbol_26)).values[(((firstVertex * 3u) + (i * 3u)) + 2u)] = (*(tint_symbol_22)).arr[i][2];
  }
  for(uint i = 0u; (i < indexCount); i = (i + 1u)) {
    int const index = (*(tint_symbol_19)).tris.arr[(triTableOffset + i)];
    (*(tint_symbol_27)).tris[(firstIndex + i)] = (firstVertex + (*(tint_symbol_23)).arr[index]);
  }
  for(uint i = indexCount; (i < 15u); i = (i + 1u)) {
    (*(tint_symbol_27)).tris[(firstIndex + i)] = firstVertex;
  }
}

kernel void computeMain(device IsosurfaceVolume* tint_symbol_28 [[buffer(0)]], const device Tables* tint_symbol_29 [[buffer(5)]], device DrawIndirectArgs* tint_symbol_34 [[buffer(1)]], device PositionBuffer* tint_symbol_35 [[buffer(2)]], device NormalBuffer* tint_symbol_36 [[buffer(3)]], device IndexBuffer* tint_symbol_37 [[buffer(4)]], uint3 global_id [[thread_position_in_grid]]) {
  thread tint_array_wrapper_2 tint_symbol_30 = {};
  thread uint tint_symbol_31 = 0u;
  thread tint_array_wrapper_2 tint_symbol_32 = {};
  thread tint_array_wrapper_3 tint_symbol_33 = {};
  computeMain_inner(global_id, tint_symbol_28, tint_symbol_29, &(tint_symbol_30), &(tint_symbol_31), &(tint_symbol_32), &(tint_symbol_33), tint_symbol_34, tint_symbol_35, tint_symbol_36, tint_symbol_37);
  return;
}