[RFC] Add an Apple Metal/AIR backend target

LLVM Project
11

@asindarov Hi thank you for that, I attempted opaque pointers+bitcode v2 and implementation seems straightforward but I ran into something that at the moment seems to me like an AGX bug that if true would be a blocker, bit off-topic but repro in case you are interested:

#include <metal_stdlib>
#include <metal_simdgroup_matrix>
using namespace metal;

kernel void tg_mma(device float* out [[buffer(0)]],
                   uint tid [[thread_position_in_threadgroup]]) {
    threadgroup float tile[64];
    tile[tid % 64] = float(tid);
    threadgroup_barrier(mem_flags::mem_threadgroup);
    simdgroup_float8x8 m;
    simdgroup_load(m, tile, 8);
    simdgroup_float8x8 r = m;
    simdgroup_store(r, out, 8);
}

With -opaque-pointers it errors out as XPC_ERROR_CONNECTION_INTERRUPTED (.ips shows EXC_BAD_ACCESS (SIGSEGV) KERN_INVALID_ADDRESS at 0x0).