TornadoVM CUDA Backend: NVIDIA GPU Acceleration Guide

The TornadoVM CUDA backend is the deepest integration with the NVIDIA hardware and software ecosystem available from the JVM. It compiles Java bytecode through the Graal IR to CUDA PTX, hands the PTX to NVRTC for JIT compilation to a device-specific cubin, and then dispatches the resulting native kernel on a CUDA stream — all automatically, with no CUDA C to write. Beyond basic JIT code generation, the CUDA backend also exposes the full NVIDIA library ecosystem — cuBLAS, cuBLASLt, cuFFT, cuDNN, cuSPARSE, and CUTLASS — as first-class TaskGraph library tasks that share device buffers and a CUDA stream with your generated kernels, plus mma.sync Tensor Core intrinsics accessible directly from KernelContext.

Prerequisites

Before installing the CUDA backend, ensure the following are available on your system:

NVIDIA GPU

Any CUDA-capable NVIDIA GPU. Tensor Core mma.sync intrinsics (FP16, BF16, INT8, FP8) and FP16 GemmEx via cuBLASLt require Ampere (sm_80) or newer. FP8 MMA requires Ada/Hopper (sm_89+).

CUDA Toolkit

Install the CUDA Toolkit from developer.nvidia.com/cuda-downloads. The toolkit must include nvrtc (NVIDIA Runtime Compilation). /usr/local/cuda or $CUDA_PATH must be set.

JDK 21 or JDK 25

JAVA_HOME must point to a JDK 21 or JDK 25 installation (OpenJDK or GraalVM). Graal-based JDK 21 enables additional optimisations.

GCC/G++ ≥ 13

Required to build the native JNI bridges (cuda-jni, cublas-jni, cufft-jni, cudnn-jni, cusparse-jni, cutlass-jni).

For systems with multiple CUDA toolkits installed, TornadoVM resolves the toolkit from /usr/local/cuda first, then falls back to the $CUDA_PATH environment variable. Set CUDA_PATH explicitly if your toolkit lives elsewhere.

Installation

SDKMAN!
Installer Script
Make

The fastest way to get started — downloads a prebuilt SDK that includes the CUDA backend, NVRTC JNI bridges, and all library-task JNI modules:

sdk install tornadovm 5.2.0-cuda

After installation, activate the SDK and verify your device:

sdk use tornadovm 5.2.0-cuda
tornado --devices

Builds TornadoVM from source and configures only the CUDA backend:

./bin/tornadovm-installer --jdk jdk21 --backend cuda

To combine CUDA with OpenCL on the same machine:

./bin/tornadovm-installer --jdk jdk21 --backend opencl,cuda

Direct Make target for contributors building from source:

make BACKEND=cuda

Or use the dedicated shortcut:

make cuda

Both targets call bin/compile --jdk jdk21 --backend cuda internally.

Build the CUDA backend in isolation (make BACKEND=cuda) when running CUDA-specific unit tests. With both CUDA and OpenCL backends installed, the test harness may silently dispatch to the OpenCL device and produce false-positive results for CUDA-only features like library tasks and Tensor Core intrinsics.

Verifying the Backend

After installation, confirm that the CUDA driver is visible and your GPU is enumerated:

tornado --devices

Expected output for a system with a single NVIDIA GPU:

Number of Tornado drivers: 1
Driver: CUDA
  0: CUDA -- NVIDIA GeForce RTX 4090 (available)
       Global Memory Size: 24576 MB
       Local Memory Size: 48 KB
       Max Work-Group Size: 1024
       Max Compute Units: 128

To print the generated PTX for any kernel, add --printKernel:

tornado --printKernel -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

JIT Compilation Pipeline

The CUDA backend follows a four-stage compilation chain at runtime:

Java Bytecode → Graal IR

The method annotated as a TornadoVM task is parsed by GraalVM’s bytecode parser. Loops with @Parallel annotations and KernelContext accesses are identified and lifted into GPU-parallel IR nodes.

Graal IR → CUDA PTX

The CUDA-specific lowering pass (CUDABackend, CUDALIRStmt) serialises the IR to CUDA PTX assembly — a portable virtual-ISA for NVIDIA GPUs. MMA nodes (CUDAMMAComputeNode, CUDAMMALoadANode, CUDAMMAStoreNode) are lowered to mma.sync PTX instructions.

PTX → cubin via NVRTC

CUDAJIT hands the PTX string to the NVRTC runtime (nvrtcCompileProgram), which performs JIT compilation to a device-specific cubin binary specialised for the active GPU’s compute capability.

cubin Loaded & Dispatched

CUDACodeCache caches the compiled cubin. On subsequent executions the cache is checked first; compilation is skipped unless the kernel’s inputs or compilation flags change.

CUDA-Specific Features

Library Tasks: cuBLAS, cuFFT, cuDNN, cuSPARSE, CUTLASS

Library tasks let you call tuned NVIDIA libraries from the same TaskGraph as your JIT-compiled Java kernels. All library calls share TornadoVM-managed device buffers on a single CUDA stream — no extra copies, no host synchronisation required between a kernel and a library call.

TaskGraph tg = new TaskGraph("ml_pipeline")
    .transferToDevice(DataTransferMode.EVERY_EXECUTION, input, weights)
    // JIT-compiled Java preprocessing kernel
    .task("preprocess", Kernels::normalize, input)
    // Native cuBLAS SGEMM on the same device buffer
    .libraryTask("gemm", CuBlas::cublasSgemm,
            CuBlasOperation.CUBLAS_OP_N.operation(),
            CuBlasOperation.CUBLAS_OP_N.operation(),
            m, n, k,
            alpha, weights, lda, input, ldb, beta, output, ldc)
    // JIT-compiled Java activation kernel
    .task("activate", Kernels::relu, output)
    .transferToHost(DataTransferMode.EVERY_EXECUTION, output);

Available library task providers

Provider	Operations
cuBLAS	SGEMV, SGEMM (single-precision)
cuBLASLt	TF32 and FP16 GemmEx on Tensor Cores; fused `BIAS` / `GELU_BIAS` epilogues; plan caching
cuFFT	C2C, R2C/C2R, Z2Z transforms (1D and 2D); FFT-filter pipelines with JIT kernels
cuDNN	Deep-learning primitives via the cuDNN graph API; fused scaled-dot-product (flash) attention via `cudnn-frontend`
cuSPARSE	Sparse-matrix operations
CUTLASS	Templated GEMM kernels with Tensor Core support

New providers implement TornadoLibraryProvider and are discovered via Java ServiceLoader — no core runtime changes required.

Tensor Core MMA Intrinsics

KernelContext exposes mma.sync Tensor Core instructions directly from Java. FP16 (m16n8k16 → FP32) and INT8 (m16n8k32 → INT32) shapes are supported. Matrix tiles are staged in shared memory as int-packed fp16 pairs, and the warp-collective MMA call produces a per-lane accumulator fragment:

void tensorCoreKernel(KernelContext ctx,
                      HalfFloatArray a, HalfFloatArray b, FloatArray c,
                      int m, int n, int k) {
    // Allocate shared-memory tiles (int-packed: two fp16 per int)
    int[] aTile = ctx.allocateIntLocalArray(m * k / 2);
    int[] bTile = ctx.allocateIntLocalArray(k * n / 2);

    // Pack two consecutive fp16 values from A into each int slot
    int tid = ctx.localIdx;
    int lo = a.get(tid * 2).getHalfFloatValue() & 0xFFFF;
    int hi = a.get(tid * 2 + 1).getHalfFloatValue() & 0xFFFF;
    aTile[tid] = lo | (hi << 16);
    ctx.localBarrier();

    // Declare a zeroed fp32 accumulator fragment (4 registers per lane)
    float[] fragC = ctx.mmaFragment(0.0f);

    // Load per-lane A and B fragments from the int-packed shared-memory tiles
    HalfFloat[] fragA = ctx.mmaLoadA(aTile, /* wmmaK */ k);
    HalfFloat[] fragB = ctx.mmaLoadB(bTile, /* wmmaK */ k);

    // Warp-collective MMA: fragD = fragA * fragB + fragC  (m16n8k16, fp16→fp32)
    float[] fragD = ctx.mma(fragA, fragB, fragC, MMAShape.M16N8K16);

    // Store the result fragment to global output
    ctx.mmaStore(fragD, c, /* tileRow */ ctx.groupIdx, /* tileCol */ ctx.groupIdy, n);
}

Tensor Core MMA intrinsics require an Ampere (sm_80) or newer GPU. The Graal lowering phase emits CUDAMMAComputeNode, CUDAMMALoadANode, and CUDAMMAStoreNode IR nodes, which are serialised to mma.sync.aligned PTX instructions by CUDATensorCoreSupportPhase. Matrix tiles must be staged in int-packed shared-memory arrays before calling mmaLoadA/mmaLoadB.

CUDA Graphs

CUDA Graphs allow the entire pipeline — JIT kernels, library tasks, and memory transfers — to be recorded once and replayed with a single cuGraphLaunch, eliminating per-iteration kernel launch overhead:

try (TornadoExecutionPlan plan = new TornadoExecutionPlan(tg.snapshot())) {
    // Record the graph on first call; replay on all subsequent calls
    plan.withCUDAGraph().execute();
}

CUDA Graphs provide the largest speedup for pipelines with many short kernels or where host-side Python-style loop overhead is the bottleneck. For long-running single kernels, the benefit is smaller but the API is identical.

Environment Variables and Build Flags

Variable	Default	Purpose
`CUDA_PATH`	`/usr/local/cuda`	Path to the CUDA Toolkit installation
`TORNADO_DEVICE_MEMORY`	System default	Override device memory limit, e.g. `-Dtornado.device.memory=8GB`
`MACOSX_DEPLOYMENT_TARGET`	Host macOS version	Not used for CUDA (Linux/Windows only)

Running Examples

# Matrix-vector multiply on the CUDA backend
tornado --printKernel \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

# NBody simulation
tornado \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.NBody --params="16384 10"

# Run unit tests against the CUDA backend only
tornado-test --ea --verbose

Get Started

Programming Guide

Backends

NVIDIA Integration

Guides

Contributing

TornadoVM CUDA Backend: NVIDIA GPU Acceleration Guide

Prerequisites

NVIDIA GPU

CUDA Toolkit

JDK 21 or JDK 25

GCC/G++ ≥ 13

Installation

Verifying the Backend

JIT Compilation Pipeline

CUDA-Specific Features

Library Tasks: cuBLAS, cuFFT, cuDNN, cuSPARSE, CUTLASS

Tensor Core MMA Intrinsics

CUDA Graphs

Environment Variables and Build Flags

Running Examples

Build docs developers (and LLMs) love

Get Started

Programming Guide

Backends

NVIDIA Integration

Guides

Contributing

Documentation Index

​Prerequisites

NVIDIA GPU

CUDA Toolkit

JDK 21 or JDK 25

GCC/G++ ≥ 13

​Installation

​Verifying the Backend

​JIT Compilation Pipeline

​CUDA-Specific Features

​Library Tasks: cuBLAS, cuFFT, cuDNN, cuSPARSE, CUTLASS

​Tensor Core MMA Intrinsics

​CUDA Graphs

​Environment Variables and Build Flags

​Running Examples

Build docs developers (and LLMs) love

Prerequisites

Installation

Verifying the Backend

JIT Compilation Pipeline

CUDA-Specific Features

Library Tasks: cuBLAS, cuFFT, cuDNN, cuSPARSE, CUTLASS

Tensor Core MMA Intrinsics

CUDA Graphs

Environment Variables and Build Flags

Running Examples