Skip to main content

Documentation Index

Fetch the complete documentation index at: https://mintlify.com/Deepak-Sangle/TornadoVM/llms.txt

Use this file to discover all available pages before exploring further.

The TornadoVM Metal backend brings native Apple Silicon GPU acceleration to Java by generating Metal Shading Language (MSL) source from your TaskGraph kernels and compiling it through Apple’s Metal framework into an MTLLibrary pipeline at runtime. Introduced in TornadoVM 5.2.0, it targets the unified-memory GPU cores in M1, M2, M3, and M4 chips — the same hardware used by tools like Core ML and Accelerate — without any Objective-C or Swift. The identical Java kernel code that runs on NVIDIA or AMD hardware via the CUDA or OpenCL backends executes unchanged on Apple Silicon through Metal, while Metal-specific additions such as simdgroup_multiply_accumulate matrix operations are available to kernels that want them.

​
Prerequisites

Apple Silicon Mac

macOS 12 (Monterey) or later running on an M1, M2, M3, or M4 chip. Intel Macs are not supported by the Metal backend.

Xcode / Metal Framework

The Metal framework ships with macOS. Install Xcode Command Line Tools to get the Metal compiler toolchain: xcode-select --install.

JDK 21 or JDK 25

JAVA_HOME must point to a JDK 21 or JDK 25 build compiled for aarch64. GraalVM for JDK 21 on ARM64 is recommended.

macOS Deployment Target

Set MACOSX_DEPLOYMENT_TARGET at build time to control the minimum macOS version the resulting SDK can run on (e.g., 11.0 for maximum compatibility).
Native libraries compiled on a newer macOS version may not run on older releases. Set MACOSX_DEPLOYMENT_TARGET=11.0 before building if you need the SDK to run on macOS 11 Big Sur or later.

​
Installation

The prebuilt Metal SDK requires only an Apple Silicon Mac with the Metal framework present (ships with macOS):
sdk install tornadovm 5.2.0-metal
Activate and verify:
sdk use tornadovm 5.2.0-metal
tornado --devices

​
Verifying the Metal Backend

Run tornado --devices after installation to confirm the Metal GPU is visible: 
tornado --devices
Expected output on an M2 MacBook Pro: 
Number of Tornado drivers: 1
Driver: Metal
  0: Metal -- Apple M2 Pro (available)
       Global Memory Size: 16384 MB
       Local Memory Size: 32 KB
       Max Work-Group Size: 1024
       Max Compute Units: 19
To print the generated MSL source for any kernel: 
tornado --printKernel \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

​
JIT Compilation Pipeline

1

Java Bytecode → Graal IR

The task method is parsed by GraalVM’s bytecode parser via MetalHotSpotBackendFactory. @Parallel loop annotations and KernelContext accesses are lifted to Metal-specific parallel IR nodes, and SIMDGROUP intrinsic calls are recognised as MetalSimdgroupMatrix* graph nodes.
2

Graal IR → Metal Shading Language (MSL)

MetalBackend and MetalAssembler serialise the LIR to MSL source. Thread-indexing maps ctx.globalIdx to thread_position_in_grid.x, local memory to threadgroup storage, and barriers to threadgroup_barrier(mem_flags::mem_threadgroup). SIMDGROUP matrix nodes emit simdgroup_load, simdgroup_multiply_accumulate, and simdgroup_store intrinsics.
3

MSL → MTLLibrary via Metal Compiler

The MSL source string is compiled through the Metal framework’s runtime compiler (MTLDevice.makeLibrary) into an MTLLibrary and then an MTLComputePipelineState. This compilation step is performed at runtime against the active GPU.
4

Pipeline Cached & Dispatched

MetalCodeCache caches the compiled MTLComputePipelineState. On subsequent executions the cache is checked first; recompilation is skipped unless the kernel changes.

​
SIMDGROUP Matrix Multiply-Accumulate

Apple Metal exposes hardware-accelerated matrix multiply-accumulate through simdgroup_multiply_accumulate. TornadoVM’s Metal backend maps this to KernelContext.simdgroupMatrixMultiplyAccumulate, generating MetalSimdgroupMatrixMmaNode IR nodes that lower to native MSL intrinsics: 
void matmulKernel(KernelContext ctx,
                  FloatArray a, FloatArray b, FloatArray c,
                  int m, int n, int k) {
    // Tile origins for this SIMD group (one 8×8 output tile per group)
    int tilesPerRow = n / 8;
    int tileIdx   = ctx.groupIdx;
    int tileRow   = tileIdx / tilesPerRow;
    int tileCol   = tileIdx % tilesPerRow;

    int aBase = tileRow * 8 * k;      // A[M][K] row-major, tile origin
    int bBase = tileCol * 8;          // B[K][N] row-major, tile origin
    int cBase = tileRow * 8 * n + tileCol * 8;

    // Accumulate over the K dimension in 8-column steps
    Matrix8x8Float acc = ctx.simdgroupMatrixZero();
    for (int p = 0; p < k; p += 8) {
        // Load 8×8 fragments — lowers to simdgroup_load in MSL
        Matrix8x8Float fragA = ctx.simdgroupMatrixLoad(a, aBase + p, k);
        Matrix8x8Float fragB = ctx.simdgroupMatrixLoad(b, bBase + p * n, n);
        // Multiply-accumulate: acc = fragA * fragB + acc
        // Lowers to simdgroup_multiply_accumulate in MSL
        acc = ctx.simdgroupMatrixMultiplyAccumulate(fragA, fragB, acc);
    }
    // Store the result fragment — lowers to simdgroup_store in MSL
    ctx.simdgroupMatrixStore(acc, c, cBase, n);
}
SIMDGROUP matrix operations (simdgroup_multiply_accumulate) require Apple Silicon hardware and macOS 13 or later. Each operation computes a * b + c for three Matrix8x8Float fragments and returns the new accumulator fragment. The surrounding tiling loop is ordinary Java compiled by TornadoVM’s Metal backend; only the simdgroupMatrixLoad, simdgroupMatrixMultiplyAccumulate, and simdgroupMatrixStore calls are replaced by hardware-level Metal intrinsics (MetalSimdgroupMatrixLoadNode, MetalSimdgroupMatrixMmaNode, MetalSimdgroupMatrixStoreNode IR nodes). The local work size must be 32 (one full Apple Silicon SIMD group) per output tile.

​
macOS Deployment Target and SDK Compatibility

SDKs built on a newer macOS version may not load on older releases because the Metal JNI library (libTornadoMetal.dylib) is linked against the system’s libstdc++ and Metal framework version at build time. 
# Build for maximum compatibility (supports macOS 11+)
export MACOSX_DEPLOYMENT_TARGET=11.0
make clean && make BACKEND=metal
Built OnMACOSX_DEPLOYMENT_TARGETRuns On
macOS 14 (Sonoma)Not setmacOS 14+ only
macOS 14 (Sonoma)13.0macOS 13+
macOS 14 (Sonoma)11.0macOS 11+ (all Apple Silicon)

​
Limitations Compared to the CUDA Backend

The Metal backend provides full JIT kernel compilation and SIMDGROUP matrix operations. The following CUDA-backend-only features are not available on Metal:
FeatureCUDAMetal
JIT kernel compilation✅✅
KernelContext (thread IDs, local mem, barriers)✅✅
SIMDGROUP matrix multiply-accumulate❌✅
cuBLAS / cuBLASLt library tasks✅❌
cuFFT library tasks✅❌
cuDNN library tasks✅❌
cuSPARSE / CUTLASS library tasks✅❌
Tensor Core mma.sync intrinsics✅❌
CUDA Graphs (withCUDAGraph())✅❌
For deep-learning workloads on Apple Silicon that need library-level acceleration, consider combining TornadoVM’s custom kernel execution (via the Metal backend) with Apple’s Accelerate framework for BLAS operations outside the TornadoVM TaskGraph.

​
Running Examples

# List all Metal devices
tornado --devices

# Run matrix-vector multiply on the Metal backend
tornado \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

# Print the generated MSL kernel source
tornado --printKernel \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

# Run the full unit test suite
tornado-test --ea --verbose

Build docs developers (and LLMs) love

Get started for freeTalk to us