a.metal

#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 tint_symbol_7 {
  /* 0x0000 */ tint_array<float, 1> arr;
};

struct tint_symbol_10 {
  /* 0x0000 */ tint_array<float, 1> arr;
};

struct tint_symbol_12 {
  /* 0x0000 */ tint_array<float, 1> arr;
};

struct tint_symbol {
  /* 0x0000 */ tint_array<uint4, 1> buffer_size;
};

float float_from_bits(uint x) {
  return as_type<float>(x);
}

void foo_inner(uint3 local_id, uint3 group_id, device tint_array<float, 1>* const tint_symbol_2, const constant tint_symbol* const tint_symbol_3, device tint_array<float, 1>* const tint_symbol_4, device tint_array<float, 1>* const tint_symbol_5) {
  int const _f8_s0___outermost___outermost_v3___block_id_x = int(group_id[0]);
  int const v___thread_id_x = int(local_id[0]);
  {
    tint_array<float, 1> _sum_0 = {};
    int const _1 = 0;
    float const _2 = float_from_bits(0u);
    _sum_0[min(uint(_1), 0u)] = _2;
    int const _3 = 0;
    int const _4 = 100;
    for(int _sum_s1_r8__x = _3; (_sum_s1_r8__x < as_type<int>((as_type<uint>(_3) + as_type<uint>(_4)))); _sum_s1_r8__x = as_type<int>((as_type<uint>(_sum_s1_r8__x) + as_type<uint>(1)))) {
      int const _5 = 0;
      float const _6 = (*(tint_symbol_2))[min(uint(_sum_s1_r8__x), (((*(tint_symbol_3)).buffer_size[0u][0u] / 4u) - 1u))];
      float const _7 = (*(tint_symbol_4))[min(uint(_sum_s1_r8__x), (((*(tint_symbol_3)).buffer_size[0u][1u] / 4u) - 1u))];
      float const _8 = (_6 - _7);
      float const _9 = float_from_bits(0u);
      float const _10 = (_9 - _8);
      bool const _11 = (_8 > _9);
      float const _12 = select(_10, _8, _11);
      float const _13 = _12;
      float const _14 = _sum_0[min(uint(_5), 0u)];
      float const _15 = (_13 + _14);
      _sum_0[min(uint(_5), 0u)] = _15;
    }
    int const _16 = 0;
    float const _17 = _sum_0[min(uint(_16), 0u)];
    (*(tint_symbol_5))[min(uint(_16), (((*(tint_symbol_3)).buffer_size[0u][2u] / 4u) - 1u))] = _17;
  }
}

kernel void foo(device tint_symbol_7* tint_symbol_6 [[buffer(0)]], const constant tint_symbol* tint_symbol_8 [[buffer(30)]], device tint_symbol_10* tint_symbol_9 [[buffer(1)]], device tint_symbol_12* tint_symbol_11 [[buffer(2)]], uint3 local_id [[thread_position_in_threadgroup]], uint3 group_id [[threadgroup_position_in_grid]]) {
  foo_inner(local_id, group_id, &((*(tint_symbol_6)).arr), tint_symbol_8, &((*(tint_symbol_9)).arr), &((*(tint_symbol_11)).arr));
  return;
}

b.metal

#define INDEX_WITH_CLAMPING 1

#include <metal_stdlib>

using namespace metal;
struct Result {
  /* 0x0000 */ uint value;
};
struct tint_array_wrapper {
  uint arr[10];
};
struct tint_array_wrapper_1 {
  float arr[3];
};
struct S {
  tint_array_wrapper startCanary;
  tint_array_wrapper_1 data;
  tint_array_wrapper endCanary;
};

#if INDEX_WITH_CLAMPING
#define INDEX(i) min(i, 9u)
#else
#define INDEX(i) i
#endif

uint runTest() {
  S s = {};
  for(uint i = 0u; (i < 10u); i = (i + 1u)) {
    s.startCanary.arr[INDEX(i)] = 4294967295u;
    s.endCanary.arr[INDEX(i)] = 4294967295u;
  }
  s.data.arr[0] = float();
  for(uint i = 0u; (i < 10u); i = (i + 1u)) {
    if ((s.startCanary.arr[INDEX(i)] != 4294967295u)) {
      return s.startCanary.arr[INDEX(i)];
    }
    if ((s.endCanary.arr[INDEX(i)] != 4294967295u)) {
      return s.endCanary.arr[INDEX(i)];
    }
  }
  return 0u;
}

kernel void foo(device Result& result [[buffer(0)]]) {
  result.value = runTest();
  return;
}

msl-compile.mm

@import Metal;

int main(int argc, char *argv[]) {
  NSString* path;
  NSString* entry_point;
  if (argc == 2) {
    path = [NSString stringWithUTF8String:argv[1]];
    entry_point = nil;
  }
  else if (argc == 3) {
    path = [NSString stringWithUTF8String:argv[1]];
    entry_point = [NSString stringWithUTF8String:argv[2]];
  }
  else if (argc == 8) {
    path = [NSString stringWithUTF8String:argv[7]];
    entry_point = nil;
  }
  else if (argc == 9) {
    path = [NSString stringWithUTF8String:argv[7]];
    entry_point = [NSString stringWithUTF8String:argv[8]];
  }
  else {
    printf("Usage:\n");
    printf("\tmsl-compile FILE [ENTRY_POINT]\n");
    printf("\tmsl-compile '' '' '' '' '' '' FILE [ENTRY_POINT]\n");
    exit(1);
  }

  NSError* error = nil;

  // Load source.
  NSString* source = [NSString stringWithContentsOfFile:path encoding:NSUTF8StringEncoding error:&error];
  if(!source) {
    NSString *output = [error localizedDescription];
    printf("Failed to load source: %s\n", [output UTF8String]);
    return 1;
  }

  // Build MSL library.
  id<MTLDevice> device = MTLCopyAllDevices()[0];
  printf("Device = %s\n", [device.name UTF8String]);
  id<MTLLibrary> library = [device newLibraryWithSource:source options:nil error:&error];
  if(!library) {
    NSString *output = [error localizedDescription];
    printf("Library creation failed: %s\n", [output UTF8String]);
    return 1;
  } else {
    printf("Library created successfully.\n");
  }

  if (entry_point) {
    // TODO: Assumes compute entry point

    // Create a shader function.
    id<MTLFunction> function = [library newFunctionWithName:entry_point];
    if(!function) {
      printf("Failed to make entry point\n");
      return 1;
    }

    // Build compute pipeline state.
    id<MTLComputePipelineState> pipeline =
      [device newComputePipelineStateWithFunction:function error:&error];
    if (!pipeline) {
      NSString *output = [error localizedDescription];
      printf("Pipeline creation failed: %s\n", [output UTF8String]);
      return 1;
    } else {
      printf("Pipeline created successfully.\n");
    }
  }

  return 0;
}