Skip to content

beehive-lab/TornadoVM

Folders and files

NameName
Last commit message
Last commit date

Latest commit

Β 

History

10,608 Commits
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 
Β 

TornadoVM logo


Write Java. Run on GPUs. Fast.

TornadoVM is a GPU programming framework for Java that works with JDK 21+ (currently JDK 21 and JDK 25). It JIT-compiles Java bytecode into NVIDIA CUDA PTX, OpenCL C, SPIR-V, and Apple Metal (MSL) at runtime, so your existing Java code runs on NVIDIA GPUs (via CUDA/PTX), AMD, Intel, and Apple Silicon GPUs, integrated GPUs, FPGAs, and multi-core CPUs. On NVIDIA hardware it goes further: beyond generating PTX, TornadoVM now calls straight into the NVIDIA library ecosystem β€” cuBLAS, cuFFT, cuDNN β€” and exposes Tensor Core mma.sync instructions from pure Java. No CUDA C. No JNI bindings to maintain. No native toolchain in your application.

Build & Test JDK 21 Build & Test JDK 25 Tornado API Install with SDKMAN! Docs Slack

Latest release: TornadoVM 4.0.1 (JDK 21 / JDK 25) β€” native NVIDIA library integration (cuBLAS / cuFFT / cuDNN) and Tensor Core intrinsics, plus a native Apple Metal backend for Apple Silicon. Changelog Β· Website Β· Documentation


This is the whole programming model

Write the kernel in Java with the same thread-indexing model you'd use in CUDA β€” then build a task graph and execute. TornadoVM JIT-compiles the bytecode to a GPU kernel at runtime and manages all host↔device data transfers for you. On NVIDIA GPUs that kernel is emitted as CUDA PTX and compiled through NVRTC to a native cubin.

Java + TornadoVM (Kernel API) The same kernel in CUDA C
void mxv(KernelContext ctx,
         FloatArray m,
         FloatArray v,
         FloatArray out,
         int rows, int cols) {
  int i = ctx.globalIdx;
  if (i < rows) {
    float sum = 0f;
    for (int j=0; j<cols; j++)
      sum += m.get(i*cols+j)*v.get(j);
    out.set(i, sum);
  }
}

// TornadoVM auto-manages device
// memory + kernel dispatch via a
// TaskGraph β€” no host plumbing,
// runs unchanged on all 4 backends.
__global__ void mxv(
    const float *m,
    const float *v,
    float *out,
    int rows, int cols) {
  int i = blockIdx.x*blockDim.x
        + threadIdx.x;
  if (i < rows) {
    float sum = 0.f;
    for (int j=0; j<cols; j++)
      sum += m[i*cols+j]*v[j];
    out[i] = sum;
  }
}

// ...and on the host you STILL write:
//   cudaMalloc/cudaMemcpy per buffer
//   grid/block dims, launch, sync
//   cudaMemcpy back, then cudaFree
//   nvcc build + per-GPU binaries
//   + a rewrite for non-NVIDIA GPUs

…and the host side β€” wrap data, map a thread grid, execute (click to expand)

// TornadoVM off-heap arrays (flat, row-major)
FloatArray m   = new FloatArray(rows * cols);
FloatArray v   = new FloatArray(cols);
FloatArray out = new FloatArray(rows);
// ...fill m and v...

KernelContext ctx  = new KernelContext();
WorkerGrid worker  = new WorkerGrid1D(rows);
GridScheduler grid = new GridScheduler("compute.mxv", worker);

TaskGraph tg = new TaskGraph("compute")
    .transferToDevice(DataTransferMode.FIRST_EXECUTION, m, v)
    .task("mxv", Kernels::mxv, ctx, m, v, out, rows, cols)
    .transferToHost(DataTransferMode.EVERY_EXECUTION, out);

try (TornadoExecutionPlan plan = new TornadoExecutionPlan(tg.snapshot())) {
    plan.withGridScheduler(grid).execute();   // JIT-compiled to your GPU
}

KernelContext gives you the full GPU programming model β€” global/local thread IDs, local memory, and barriers, the same semantics as CUDA/OpenCL/SYCL β€” while TornadoVM handles memory management and runs the identical code across all four backends. Don't need that control? Drop KernelContext and just annotate the loop with @Parallel β€” TornadoVM infers the thread mapping for you. Both styles combine in the same TaskGraph. Programming guide β†’


🟩 The NVIDIA ecosystem, native to Java

On NVIDIA hardware, TornadoVM is more than a PTX code generator β€” it's an open-source on-ramp to the whole CUDA software stack, callable from the same TaskGraph you already use. Generated kernels and native library calls share TornadoVM-managed device buffers on one CUDA stream, so a JIT-compiled kernel can feed a cuBLAS call and consume its output with no extra copies and no host synchronization.

Capability What it gives Java developers
CUDA PTX backend Java bytecode β†’ Graal IR β†’ CUDA PTX β†’ NVRTC β†’ cubin, JIT-compiled and specialized to your data sizes and GPU at runtime.
cuBLAS / cuBLASLt library tasks SGEMV, SGEMM, strided-batched, TF32 and FP16 GemmEx on Tensor Cores, and cuBLASLt with plan caching and fused BIAS / GELU_BIAS epilogues β€” one library task replaces a GEMM plus a separate activation kernel.
cuFFT library tasks C2C, R2C / C2R, and Z2Z transforms (1D and 2D) with capture-safe plan caching; FFT-filter pipelines mix with JIT kernels in one graph.
cuDNN library tasks Deep-learning primitives through the cuDNN graph API, including fused scaled-dot-product (flash) attention via cudnn-frontend.
Tensor Core MMA intrinsics mma.sync exposed through KernelContext (mmaLoadA/B, mma, mmaStore) β€” FP16 (m16n8k16 β†’ FP32) and INT8 (m16n8k32 β†’ INT32), with swizzled shared-memory staging. Not a binding: real PTX/CUDA generated from Java.
CUDA Graphs executionPlan.withCUDAGraph() records kernels, library calls, and transfers into a captured graph and replays them with a single cuGraphLaunch.

Mixing your own kernels with NVIDIA's tuned libraries looks like this:

TaskGraph tg = new TaskGraph("hybrid")
    .transferToDevice(DataTransferMode.EVERY_EXECUTION, matrix, vector)
    .task("preprocess", MyKernels::preprocess, matrix)              // JIT-compiled Java kernel
    .libraryTask("gemv", CuBlas::cublasSgemv,                       // native cuBLAS call
            CUBLAS_OP_T, m, n, alpha, matrix, lda, vector, incx, beta, output, incy)
    .task("postprocess", MyKernels::activate, output)              // JIT-compiled Java kernel
    .transferToHost(DataTransferMode.EVERY_EXECUTION, output);

try (TornadoExecutionPlan plan = new TornadoExecutionPlan(tg.snapshot())) {
    plan.withCUDAGraph().execute();   // captured once, replayed each iteration
}

Library bindings are discovered via Java ServiceLoader β€” implement TornadoLibraryProvider, add a JNI module, and any native library joins the graph with no core runtime changes. TornadoVM is a member of the NVIDIA Inception Program and has presented this work at NVIDIA GTC. Hybrid API guide β†’


What people build with it

Project What it shows
πŸ¦™ GPULlama3.java LLM inference (Llama 3, Qwen 3, Mistral, Phi-3, Granite, DeepSeek distills) in pure Java β€” 117 tok/s on an RTX 5090, official GPU engine for LangChain4j and Quarkus
πŸ”† TornadoVM-Ray-Tracer Real-time ray tracing in Java, interactive frame rates on consumer GPUs
πŸ“· kfusion-tornadovm KinectFusion 3D reconstruction β€” a full computer-vision pipeline on integrated and discrete GPUs
πŸ“· Gaia Mission (ESA) High-performance algorithms for space exploration at the European Space Agency use TornadoVM

TornadoVM is used to accelerate machine learning and deep learning, computer vision, physics simulations, financial applications, computational photography, and signal processing. Building something with TornadoVM? Tell us β€” we feature community projects.


⚑ Quick start

Prerequisites

  • JDK 21 (or GraalVM based on JDK 21) β€” JAVA_HOME must point to it
  • GCC/G++ β‰₯ 13, plus the driver for your target (OpenCL runtime, CUDA Toolkit, Level Zero, or macOS for Metal)
  • For the NVIDIA library tasks (cuBLAS / cuFFT / cuDNN): the CUDA Toolkit with the corresponding libraries; on systems with multiple toolkits, /usr/local/cuda (or $CUDA_PATH) is preferred

Install via SDKMAN!

sdk install tornadovm

Pick a backend-specific build if you prefer a smaller install:

Backend SDKMAN! version Targets
OpenCL (default) 5.0.0-opencl NVIDIA / AMD / Intel GPUs, multi-core CPUs, FPGAs
CUDA 5.0.0-cuda NVIDIA GPUs (CUDA) β€” PTX codegen, Tensor Cores, cuBLAS/cuFFT/cuDNN library tasks
PTX 5.0.0-ptx NVIDIA GPUs (CUDA) β€” PTX codegen, Tensor Cores, cuBLAS/cuFFT/cuDNN library tasks
SPIR-V 5.0.0-spirv Intel GPUs (Level Zero / oneAPI)
Metal πŸ†• 5.0.0-metal Apple Silicon GPUs (M1–M4), natively via MSL
All backends 5.0.0-full Everything above

Binaries are also on the official website. For Docker and AWS (CPUs/GPUs/FPGAs) see the linked guides.

Verify your devices

tornado --devices

Run your first program

# Unix (Linux/macOS)
java @$TORNADOVM_HOME/tornado-argfile \
  -cp $TORNADOVM_HOME/share/java/tornado/tornado-examples-5.0.0.jar \
  uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

# Windows 10+
java @%TORNADOVM_HOME%\tornado-argfile ^
  -cp %TORNADOVM_HOME%\share\java\tornado\tornado-examples-5.0.0.jar ^
  uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

More examples β€” NBody, DFT, KMeans, matrix kernels, reductions: tornado-examples


πŸ“¦ Use TornadoVM in your project

<dependencies>
  <dependency>
    <groupId>io.github.beehive-lab</groupId>
    <artifactId>tornado-api</artifactId>
    <version>5.0.0-jdk21</version>
  </dependency>
  <dependency>
    <groupId>io.github.beehive-lab</groupId>
    <artifactId>tornado-runtime</artifactId>
    <version>5.0.0-jdk21</version>
  </dependency>
</dependencies>

❓ FAQ

How does TornadoVM relate to OpenJDK's Project Babylon / HAT?

Project Babylon is OpenJDK's exploratory work on code reflection, with HAT (Heterogeneous Accelerator Toolkit) as a research vehicle for GPU programming. We think it validates the direction TornadoVM has pursued since 2018 β€” and the projects are complementary rather than competing. The practical difference today: TornadoVM is usable now, with four production backends (OpenCL, PTX, SPIR-V, Metal), FPGA support, a profiler, dynamic reconfiguration, Maven Central artifacts, and years of hardening across vendor hardware β€” and it runs on standard JDK 21/25 releases. We follow Babylon closely and expect the ecosystems to converge over time.

What can TornadoVM do on NVIDIA GPUs specifically?

On NVIDIA hardware TornadoVM JIT-compiles your Java kernels to CUDA PTX and compiles them through NVRTC to native cubins. On top of that, it integrates the NVIDIA software ecosystem directly into the TaskGraph: cuBLAS / cuBLASLt (including TF32 and FP16 GemmEx on Tensor Cores and fused epilogues), cuFFT, and cuDNN (including fused flash attention) are available as library tasks that share device buffers and a CUDA stream with your generated kernels. You can also target Tensor Cores directly from Java via mma.sync intrinsics (FP16 and INT8), and capture the whole pipeline into a CUDA Graph for single-launch replay. See The NVIDIA ecosystem, native to Java.

What's the licensing situation for commercial use?

The parts your application links against β€” Tornado-API and all the modules you code with β€” are Apache 2.0. The runtime and drivers are GPLv2 with Classpath Exception, the same license as OpenJDK itself: it does not impose copyleft obligations on your application, exactly as running on OpenJDK doesn't. Full per-module breakdown below.

Does it replace my JVM?

No. TornadoVM is a plug-in to your existing OpenJDK or GraalVM installation. It complements the JVM with the ability to offload selected tasks to accelerators, handling device memory management and kernel orchestration. Everything else in your application runs on the JVM as normal.

Which hardware is supported?

Multi-core CPUs; dedicated GPUs from NVIDIA, AMD, and Intel; integrated GPUs (Apple Silicon M1–M4, Intel HD Graphics, ARM Mali); and FPGAs from Intel and Xilinx. Backends can be installed individually or together, and tasks can migrate between devices at runtime. NVIDIA GPUs can be driven either through the PTX/CUDA backend (with the library and Tensor Core features above) or through the OpenCL backend.

How fast is it?

It depends on the workload β€” data-parallel kernels with enough arithmetic intensity see order-of-magnitude speedups over sequential Java; see the benchmarking guide and the GPULlama3.java performance tables for end-to-end numbers on real applications (e.g. 117 tok/s for Llama-3 inference on an RTX 5090). For compute-bound linear algebra, the cuBLAS library tasks reach the tuned vendor throughput β€” including full FP16 Tensor Core rates β€” from Java.


🀝 Contributing & community


πŸ“š Resources & publications

Videos, presentations, and articles: resources. Selected academic publications: publications.

If you use TornadoVM β‰₯ 0.2 in research, please cite:

@inproceedings{Fumero:DARHH:VEE:2019,
 author    = {Fumero, Juan and Papadimitriou, Michail and Zakkak, Foivos S. and
              Xekalaki, Maria and Clarkson, James and Kotselidis, Christos},
 title     = {{Dynamic Application Reconfiguration on Heterogeneous Hardware}},
 booktitle = {Proceedings of the 15th ACM SIGPLAN/SIGOPS International
              Conference on Virtual Execution Environments},
 series    = {VEE '19},
 year      = {2019},
 doi       = {10.1145/3313808.3313819},
 publisher = {Association for Computing Machinery}
}

For Tornado 0.1 (initial release), cite Clarkson et al., ManLang '18.

Licenses per module

Link the Tornado-API (Apache 2.0) into your application.

Modules License
Tornado-API, Tornado-Assembly, Tornado-scripts, Tornado-Annotation, Tornado-Unittests, Tornado-Benchmarks, Tornado-Examples, Tornado-Matrices, Tornado-Drivers-OpenCL-Headers Apache 2.0
Tornado-Runtime, Tornado-Drivers GPLv2 with Classpath Exception

Acknowledgments

Partially funded by Intel Corporation and by EU & UKRI grants (most recent first): AERO (101092850), P2CODE (101093069), ENCRYPT (101070670), TANGO (101070052), ELEGANT (957286), E2Data (780245), ACTiCLOUD (732366); and EPSRC grants PAMELA (EP/K008730/1) and AnyScale Apps (EP/L000725/1).

Meet the team.