Plans for Metal support on MacOS X (Ventura onward)?

[khyron4eva]

I'm writing as the owner of a MacBook Pro M2 Max running MacOS Ventura (13.3). It appears that my laptop has OpenCL 1.2 still. It also appears that Apple has moved development of the GPU programming environment to Metal (see below).

A bit about my system:

misanthrope:bin $ ./tornado --version
version=0.15.1
branch=master
commit=d3062ac

Backends installed:

• opencl

misanthrope:bin $ uname -a
Darwin misanthrope. 22.4.0 Darwin Kernel Version 22.4.0: Mon Mar 6 20:59:58 PST 2023; root:xnu-8796.101.5~3/RELEASE_ARM64_T6020 arm64

Running Tornado 15 on this system mostly works however I get a multitude of the following errors when running the benchmarks:

.
.
[JNI] uk.ac.manchester.tornado.drivers.opencl> notify error:
[JNI] uk.ac.manchester.tornado.drivers.opencl> [CL_INVALID_KERNEL] : OpenCL Error : clEnqueueNDRangeKernel failed: invalid kernel 0x0
[TornadoVM-OCL-JNI] ERROR : clEnqueueNDRangeKernel -> Returned: -48
.
.

Additionally, "make tests" yields the following output:

misanthrope:TornadoVM $ make tests
tornado --devices
WARNING: Using incubator modules: jdk.incubator.foreign, jdk.incubator.vector

Number of Tornado drivers: 1
Driver: OpenCL
Total number of OpenCL devices : 1
Tornado device=0:0 (DEFAULT)
OPENCL -- [Apple] -- Apple M2 Max
Global Memory Size: 48.0 GB
Local Memory Size: 32.0 KB
Workgroup Dimensions: 3
Total Number of Block Threads: [256]
Max WorkGroup Configuration: [256, 256, 256]
Device OpenCL C version: OpenCL C 1.2

tornado-test.py --ea --verbose
tornado -ea --jvm "-Xmx6g -Dtornado.recover.bailout=False -Dtornado.unittests.verbose=True " -m tornado.unittests/uk.ac.manchester.tornado.unittests.tools.TornadoTestRunner --params "uk.ac.manchester.tornado.unittests.foundation.TestIntegers"
WARNING: Using incubator modules: jdk.incubator.vector, jdk.incubator.foreign
.
.
.

          Unit tests report 

==================================================

{'[PASS]': 408, '[FAILED]': 96, '[UNSUPPORTED]': 23}
Coverage [PASS/(PASS+FAIL)]: 80.95%
Coverage [PASS/(PASS+FAIL+UNSUPPORTED)]: 77.42%

==================================================

Total Time(s): 79.69343376159668

So my ultimate question is whether it is planned to make Tornado work with the latest Macs running Apple Silicon & Ventura (or later operating systems)? What, if anything, could I do to help with this effort? I'm not really a developer but I really would like to see this work however I don't have a personal need as yet.

Thank you Tornado team for the effort on this project! I can't wait to use Tornado in production one day!

References:

Videos:

Explore the new system architecture of Apple silicon Macs -- https://developer.apple.com/videos/play/wwdc2020/10686/
Working with Metal: Overview -- https://developer.apple.com/videos/play/wwdc2014/603/
Working with Metal: Fundamentals -- https://developer.apple.com/videos/play/wwdc2014/604/
Working with Metal: Advanced -- https://developer.apple.com/videos/play/wwdc2014/605/
Metal Compute on MacBook Pro -- https://developer.apple.com/videos/play/tech-talks/10580/
Build GPU binaries with Metal -- https://developer.apple.com/videos/play/wwdc2020/10615/
Metal for Pro Apps -- https://developer.apple.com/videos/play/wwdc2019/608/
Metal Performance Optimization Techniques -- https://developer.apple.com/videos/play/wwdc2015/610/
Optimize Metal Performance for Apple silicon Macs -- https://developer.apple.com/videos/play/wwdc2020/10632/
Harness Apple GPUs with Metal -- https://developer.apple.com/videos/play/wwdc2020/10602/
Discover Metal 3 -- https://developer.apple.com/videos/play/wwdc2022/10066/
Target and optimize GPU binaries with Metal 3 -- https://developer.apple.com/videos/play/wwdc2022/10102/
Load resources faster with Metal 3 -- https://developer.apple.com/videos/play/wwdc2022/10104/
Scale compute workloads across Apple GPUs -- https://developer.apple.com/videos/play/wwdc2022/10159/
Discover Metal Performance HUD -- https://developer.apple.com/videos/play/tech-talks/110339/

Documentation:

Metal Framework Docs: https://developer.apple.com/documentation/metal/
Metal Programming Guide: https://developer.apple.com/library/archive/documentation/Miscellaneous/Conceptual/MetalProgrammingGuide/Introduction/Introduction.html#//apple_ref/doc/uid/TP40014221
Best Practices Guide: https://developer.apple.com/library/archive/documentation/3DDrawing/Conceptual/MTLBestPracticesGuide/index.html#//apple_ref/doc/uid/TP40016642

[jjfumero]

Hi @khyron4eva , thank you very much for your feedback and kind words.

Concerning the errors, we can also reproduce them. These errors are due to two things:

1. Apple M1/M2 GPUs do not support fp64 (doubles). Thus we get an error when we build the code generated with OpenCL. We are currently working to improve the error messages (patch coming this week).
2. Apple M1/M2 GPUs support OpenCL 1.2 while TornadoVM requires at least 2.1. Unfortunately, Apple stopped supporting OpenCL. Besides, there are some errors in Apple regarding the OpenCL C math functions: Several benchmarks fail when running on MacOS due to ambiguous functions #66 . We will also improve the error messages.

Regarding the Metal support, this is something we have been discussing internally with the TornadoVM team. Currently, there are no plans for supporting it, as we do not have a use case yet. TornadoVM already supports three backends (OpenCL C, SPIR-V + Level Zero, and NVIDIA PTX). I do agree with you that a naive Metal backend would be great, but unfortunately, we do not have the resources to do so.
I am not an expert on Metal but I saw that a way to get there using the current TornadoVM SPIR-V JIT compiler. Khronos seems to have a tool for this: https://github.com/KhronosGroup/SPIRV-Cross, but we have never used it, and we are not sure how complex will be to integrate this. The way I see is that to support Metal, TornadoVM should have a new entire backend (compiler and runtime system).

For contributions to the TornadoVM project, we recommend holding on a bit. We need an Open Source contribution agreement from The University of Manchester and this is taking more time than expected. Once we have this contribution agreement, yes, we welcome contributions and we are happy to assist. Apologies for this inconvenience, but unfortunately it is out of our (TornadoVM team) hands.

In the meantime, we can still fire new issues, comments and feedback. We really appreciate it.