TornadoVM — Java on NVIDIA GPUs, Apple Silicon, and more

CUDA · OpenCL · Metal · Tensor Cores · Native Libraries

TornadoVM JIT-compiles JVM bytecode to CUDA, OpenCL and Metal — and now natively integrates cuBLAS, cuFFT, and cuDNN inside your Java code. Write kernels in Java, keep your build on Maven, and let the runtime generate the GPU code. No JNI glue, no second toolchain, extreme performance.

Java · TornadoVM Kernel APIMatMulVec.java

// You write this
public static void matMulVec(KernelContext ctx,
        HalfFloatArray w, FloatArray x,
        FloatArray out, int dim) {
    int row = ctx.globalIdx;
    float sum = 0f;
    for (int j = 0; j < dim; j++) {
        sum += w.get(row * dim + j).getFloat32()
             * x.get(j);
    }
    out.set(row, sum);
}

Generated CUDA CmatMulVec.cu → NVRTC → cubin

// TornadoVM generates this
extern "C" __global__ void matMulVec(
    const __half* w, const float* x,
    float* out, int dim)
{
  int row = blockIdx.x * blockDim.x + threadIdx.x;
  float sum = 0.0f;
  for (int j = 0; j < dim; j++)
    sum = fmaf(__half2float(w[row * dim + j]),
               x[j], sum);
  out[row] = sum;
}

GPULlama3

The first LLM inference engine written in pure Java — running on your GPU.

Llama 3, Mistral, Qwen 3, Phi-3, and DeepSeek models, loaded from GGUF and executed through TornadoVM's CUDA backend. FP16 and Q8 with fused dequantize–compute kernels. Watch it stream:

gpullama3 — CUDA backend 0.0 tok/s

Hybrid API — new

The first Java-native way to tap into the CUDA ecosystem: cuBLAS, cuFFT, cuDNN.

Library tasks put native NVIDIA library calls inside a TornadoVM TaskGraph. They share TornadoVM-managed device buffers with your JIT-compiled kernels and run on the same CUDA stream — a kernel can feed cuBLAS or cuDNN and consume the output with no extra copies and no host synchronization.

Java · TornadoVM TaskGraphHybridSgemv.java

TaskGraph tg = new TaskGraph("cuBLAS")
    .transferToDevice(EVERY_EXECUTION, matrix, vector)
    .task("preprocess", MyClass::preprocess, matrix)      // JIT-compiled Java
    .libraryTask("sgemv", CuBlas::cublasSgemv,           // native cuBLAS
            CUBLAS_OP_T.operation(), m, n, alpha,
            matrix, lda, vector, incx, beta, output, incy)
    .task("postprocess", MyClass::postprocess, output)   // JIT-compiled Java
    .transferToHost(EVERY_EXECUTION, output);

SGEMM 4096×4096 · TFLOP/s · RTX 4090 · CUDA 12.6 — why library tasks matter for compute-bound ops

TornadoVM JIT, Kernel API (tiled)

5.8

TornadoVM Hybrid, cuBLAS FP32

57.0

TornadoVM Hybrid, cuBLAS TF32

81.1

TornadoVM Hybrid, cuBLAS FP16

160.6

Benchmark SGEMM · cross-validated against sequential Java 27.4× over the best hand-written Java kernel — one .libraryTask() away

Performance

52× from pure Java. 74× when cuBLAS joins the graph.

Matrix–vector (y = W·x), a memory-bound example, using the same device buffers three ways: the Loop Parallel API, the Kernel API, and the Hybrid API to call cuBLAS.

1 · Loop Parallel API MatMulVec.java

public static void matMulVecParallel(
        FloatArray w, FloatArray x,
        FloatArray out, int dim) {
    for (@Parallel int row = 0; row < dim; row++) {
        float sum = 0f;
        for (int j = 0; j < dim; j++) {
            sum += w.get(row * dim + j) * x.get(j);
        }
        out.set(row, sum);
    }
}

2 · Kernel API MatMulVec.java

public static void matMulVecKernel(KernelContext ctx,
        FloatArray w, FloatArray x,
        FloatArray out, int dim) {
    int row = ctx.globalIdx;
    float sum = 0f;
    for (int j = 0; j < dim; j++) {
        sum += w.get(row * dim + j) * x.get(j);
    }
    out.set(row, sum);
}

3 · Hybrid API HybridSgemv.java

// no kernel to write — the "task" is a
// direct call into native cuBLAS:
CuBlas::cublasSgemv(
    CUBLAS_OP_T.operation(), dim, dim, alpha,
    w, dim, x, 1, beta, out, 1);

SGEMV 8192×2048 · FP32 · effective bandwidth (GB/s) — higher is better

SGEMV (Sequential Java)baseline, 8.551 ms

7.9

SGEMV (Loop Parallel API)naive, 3.309 ms · 2.6×

20.3

SGEMV (Kernel API)workgroup-per-row, 0.164 ms · 52×

410

SGEMV (Hybrid API calling cuBLAS)0.116 ms · 74×

578

NVIDIA GeForce RTX 4090 · CUDA 12.6 · 120 iterations, end-to-end execute MatrixVectorRowMajorWithCuBlas 8192 2048 32

The NVIDIA path

Native to the NVIDIA stack, dual path.

TornadoVM can dynamically generate CUDA C and also directly bind CUDA libraries inside your code. Mix and match as you feel, TornadoVM will reuse memory, combine operations and give you the maximum performance.

Quickstart

Run your first example on the GPU in three commands.

1Install with SDKMAN!

$ sdk install tornadovm

2Verify devices

$ tornado --devices

3Run an example

$ java @$TORNADOVM_HOME/tornado-argfile -cp $TORNADOVM_HOME/share/java/tornado/tornado-examples-5.0.0-jdk21.jar uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

Runs everywhere

LinuxmacOSWindowsOpenCLPTX / CUDASPIR-VApple Metal