OpenElm 270M 4-bit crashes on 2018 iPad Pro (A12X Bionic)

[DePasqualeOrg]

When I try to generate output with OpenElm 270M on a 2018 iPad Pro, the LLMEval demo app crashes with the following error:

Compiler failed to build request
libc++abi: terminating due to uncaught exception of type std::runtime_error: [metal::Device] Unable to load kernel rmsfloat16
Encountered unlowered function call to air.simd_sum.f32

It works on my iPhone 13 and on my M3 MacBook Pro, so perhaps there's an issue with the A12X Bionic chip?

[davidkoski]

Possibly. This issue has some discussion on the subject:

• https://github.com/ml-explore/mlx-swift-examples/issues/21

I have used it on an iPhone 12 which has an A14 (Apple7 metal support). A12 series is Apple5 so there are certainly a bunch of features missing.

My guess is this one is "SIMD-scoped reduction operations" which is first supported in Apple7.

[awni]

Yea it looks like it. Since we won't support old devices without SIMD ops, maybe the best thing to do is disable compilation for older GPU families? I think it makes sense to draw the line at Apple 7, since it's been tested (and it corresponds to M1).

[davidkoski]

I don't know if we can limit compilation. I think we can target different gpu families in the App Store :-)

Maybe a runtime check to at least print a warning?

[awni]

We might be able to static assert on __METAL__ which is only set when using Metal 3.0:

From the metal spec:

__METAL__ // Set if compiled with the unified Metal language
// Set with -std=metal3.0 or above.

From the feature set tables:

Presumably if you build for A12 then __METAL__ will not be set by the Metal compiler so the build should break.

[davidkoski]

I will have to play around with that -- the same binary is used on all platforms. The only check I know of is the minimum OS version.

[awni]

I was thinking this would in MLX core rather than in the Swift front-end. E.g. just a static assert in one of the metal header files (like utils.h). I think you are right that there is no environment variable for the GPU family. A runtime warning is workable, but compile is nicer if we can finagle it.

I added this to kernels/utils.h. If you have an older device, maybe you can check that it correctly breaks the build for Apple GPUs <= A13. I don't have on easily accessible to test on.

#ifndef __METAL__
static_assert(
    false,
    "Building the Metal back-end requires an Apple GPU supporting Metal >= 3.0");
#endif

[DePasqualeOrg]

I can test it on my iPad, if you let me know what package/commit I need to import.