Apple (MacOS) MPS back-end?

[jeffhammond]

Given that you like to support ALL THE COMPUTERS, I am making a feature request for an Apple MPS backend, since Apple refuses to do the right thing with OpenCL/SPIR-V and there is no other programming model for GPUs on their systems, whether they have Intel or Apple iGPUs in them.

I have low expectations here but you've always gone above and beyond, so maybe the gods will shine upon me. I am sure that there are dozens if not hundreds of Apple laptop owners out there who will greatly appreciate this.

I am making this request in a personal capacity, unrelated to any professional affiliations.

[illuhad]

Hi Jeff,

thanks for the suggestion. I agree that supporting notebooks is important for developers, but I don't personally use Apple machines, nor do I have access to one. So my personal interest in Apple support is limited at the moment, but I will keep this in mind.
Currently the plan is to focus my "new backends"-efforts on backends that can expand hardware coverage on Linux.

[vrpascuzzi]

Given that you like to support ALL THE COMPUTERS, I am making a feature request for an Apple MPS backend, since Apple refuses to do the right thing with OpenCL/SPIR-V and there is no other programming model for GPUs on their systems, whether they have Intel or Apple iGPUs in them.

Works for me with Ubuntu 20.04 installed. :)

[jeffhammond]

@vrpascuzzi Ubuntu supports OpenCL and SPIR-V for the Apple M1 iGPU?

[vrpascuzzi]

@vrpascuzzi Ubuntu supports OpenCL and SPIR-V for the Apple M1 iGPU?

Ah, sorry. Missed the bit about M1. Back in the MBP 2016 days, Ubuntu served me well. Can't say anything about the newer hardware. I moved away from Apple when getting a new machine I wanted to "just work."

[jeffhammond]

I updated the title to say MacOS to indicate that this request is not about Linux on Apple hardware.

[Michoumichmich]

@jeffhammond It would be really nice to support Apple Metal. I was able to run SPIR-V on Apple Metal using C++ So it should be achievable as hipSYCL already provides a spirv backend. However MPS do not support double precision so It would impose these restrictions in hipSYCL/Apple Metal :/ I guess we could emulate these instructions, but performance wise...

[jeffhammond]

I use this check when relevant.

        bool has_fp64(const sycl::queue & q) {
            auto device = q.get_device();
            return device.has_extension(sycl::string_class("cl_khr_fp64"));
        }

I think it's fine to not support FP64 and not emulate it. It is better to have partial support of Apple GPUs than none at all.

[illuhad]

@Michoumichmich Are we talking SPIR-V in OpenCL kernel model or shader model? One of the big issues of shader model is that it only has opaque pointers of undefined size, which makes it very difficult (or impossible in general) to support data structures that rely on USM pointers in any way.

[Michoumichmich]

@Michoumichmich Are we talking SPIR-V in OpenCL kernel model or shader model? One of the big issues of shader model is that it only has opaque pointers of undefined size, which makes it very difficult (or impossible in general) to support data structures that rely on USM pointers in any way.

Oh yes you're right. I was JIT-ing the SPV shaders to MSL using SPIRV-Cross. It can parse OpenCL SPIR-V, but It's disabled as it unfortunately cannot compile that to MSL (not fully implemented). It just reminded me of hipSYCL/macOS thus the previous. It would be nicer to avoid the round trip through SPIR-V though. MSL is "mostly" C++14 (minus lambdas, goto and unions). Maybe if we were able to extract the C++ SYCL kernels/functors with our mangled names and add their TON of attributes. And then compile that to produce the integration header. Even their metallib is mostly just LLVM bytecode (to claim IP) with metadata, but going down that route seems hacky/unreliable. That's so frustrating

[illuhad]

hmm...I'm currently working towards a generic toolchain-independent LLVM backend, so maybe we could eventually leverage this to generate MSL.. Although this sounds like reinventing the LLVM C backend. Is there some IR for Metal that we can target or is there only this MSL C++ language?

[Michoumichmich]

hmm...I'm currently working towards a generic toolchain-independent LLVM backend, so maybe we could eventually leverage this to generate MSL.. Although this sounds like reinventing the LLVM C backend. Is there some IR for Metal that we can target or is there only this MSL C++ language?

Without reinventing the LLVM C, maybe "just" source-to-source translation where we extract the SYCL kernels from the AST and generate the C++ MSL shader. AFAIK for Apple Metal there is no documented IR. When you compile the MSL with metal you don't even specify the GPU architecture, only the language standard + platform. This generates a .air file (Apple IR?) that are disguised LLVM bytecode (and now with a different "magic" so llvm tools don't work straight forward on it, but there are articles where they disassemble it and it's LLVM IR). Everything happens at runtime when you use makeLibrary to compile the shaders (and kernels) for a specific device using the .air. I guess it's because they have to run on AMD, Intel and Apple Silicon. So the only reliable way I see is to produce MSL C++.

(I guess that the SPV path has the advantage of optimising away the unsupported features. But it's messy and SPV-Cross is really focused on shaders, so I'm not sure if there is interest (on their side) to maintain&support the SPV/OCL -> MSL translation :/ And it would be at the end a LLVM C++ backend.)

[woachk]

hmm...I'm currently working towards a generic toolchain-independent LLVM backend, so maybe we could eventually leverage this to generate MSL.. Although this sounds like reinventing the LLVM C backend. Is there some IR for Metal that we can target or is there only this MSL C++ language?

Metal uses the AIR intermediate representation (an LLVM IR dialect) and does support offline compilation, so AIR is stable from release to release.

AIR is a set of annotations and constraints on LLVM IR. As such, it can be parsed just fine by off the shelf LLVM, it doesn't use a custom bytecode format.

However, Apple explicitly chose to not document AIR, making MSL (C++ 14 dialect, documented at https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf) the only officially supported input language. However, they do ship metal-as that explicitly takes AIR in, a disassembler for AIR (metal-objdump) and a linker (metallib).

For online/runtime compilation, they do not take raw AIR files however, they have to be linked with metallib (aka air-lld) beforehand. (or give MSL)

An example of a simple addition kernel compiled to AIR:

; Function Attrs: norecurse nounwind
define void @add_arrays(float addrspace(1)* noalias nocapture readonly, float addrspace(1)* noalias nocapture readonly, float addrspace(1)* noalias nocapture, i32) local_unnamed_addr #0 {
  %5 = zext i32 %3 to i64
  %6 = getelementptr inbounds float, float addrspace(1)* %0, i64 %5
  %7 = load float, float addrspace(1)* %6, align 4, !tbaa !22
  %8 = getelementptr inbounds float, float addrspace(1)* %1, i64 %5
  %9 = load float, float addrspace(1)* %8, align 4, !tbaa !22
  %10 = fadd fast float %9, %7
  %11 = getelementptr inbounds float, float addrspace(1)* %2, i64 %5
  store float %10, float addrspace(1)* %11, align 4, !tbaa !22
  ret void
}

!air.kernel = !{!8}
!8 = !{void (float addrspace(1)*, float addrspace(1)*, float addrspace(1)*, i32)* @add_arrays, !9, !10}
!9 = !{}
!10 = !{!11, !12, !13, !14}
!11 = !{i32 0, !"air.buffer", !"air.location_index", i32 0, i32 1, !"air.read", !"air.arg_type_size", i32 4, !"air.arg_type_align_size", i32 4, !"air.arg_type_name", !"float", !"air.arg_name", !"inA"}
!12 = !{i32 1, !"air.buffer", !"air.location_index", i32 1, i32 1, !"air.read", !"air.arg_type_size", i32 4, !"air.arg_type_align_size", i32 4, !"air.arg_type_name", !"float", !"air.arg_name", !"inB"}
!13 = !{i32 2, !"air.buffer", !"air.location_index", i32 2, i32 1, !"air.read_write", !"air.arg_type_size", i32 4, !"air.arg_type_align_size", i32 4, !"air.arg_type_name", !"float", !"air.arg_name", !"result"}
!14 = !{i32 3, !"air.thread_position_in_grid", !"air.arg_type_name", !"uint", !"air.arg_name", !"index"}

Note that if you go the AIR route, you will also need a downlevel converter to older LLVM IR versions... (or maybe target SPIR-V at the LLVM stage like Intel does - and then convert back to LLVM IR/AIR later).

[woachk]

@Michoumichmich Are we talking SPIR-V in OpenCL kernel model or shader model? One of the big issues of shader model is that it only has opaque pointers of undefined size, which makes it very difficult (or impossible in general) to support data structures that rely on USM pointers in any way.

As a note, this has been fixed with VK_KHR_buffer_device_address later on, which is mandatory in Vulkan 1.3 (and was implemented quite widely even before that).

Otherwise (if you want to support Vulkan implementations without it), you can do heroic hacks like what https://github.com/google/clspv does, but I can't recommend going that way... it's too brittle. For reference, MoltenVK for Vulkan on Metal (for Apple platforms) does not implement that extension.