---
## Write Java. Run on GPUs. Fast.
TornadoVM is a GPU programming framework for Java that works with JDK 21+ (currently JDK 21-27). It JIT-compiles Java bytecode into **NVIDIA CUDA, OpenCL C, and Apple Metal (MSL)** at runtime, so your existing Java code runs on **NVIDIA GPUs (via CUDA)**, AMD, Intel, and Apple Silicon GPUs, integrated GPUs, and multi-core CPUs.
On NVIDIA hardware it goes further: beyond generating CUDA, TornadoVM now calls straight into the **NVIDIA library ecosystem (cuBLAS, cuFFT, cuDNN, cuDF) and exposes Tensor Core `mma.sync` instructions from pure Java**. No CUDA C. No JNI bindings to maintain. No native toolchain in your application. In addition, it has support for Tile Programming [cuTile](https://developer.nvidia.com/cuda/tile) via its **TileContext API**.
[](https://github.com/beehive-lab/TornadoVM/actions/workflows/gpu-ci.yml)
[](https://central.sonatype.com/artifact/io.github.beehive-lab/tornado-api)
[](https://sdkman.io/sdks/tornadovm/)
[](https://tornadovm.readthedocs.io/en/latest/)
[](https://join.slack.com/t/tornadovmcommunity/shared_invite/zt-3ai2wyqva-bKz~cQRFlaJ~ZnPrbkwIEw)
**Latest release:** TornadoVM 8.0.0 (JDK 21 - JDK 27) β native **NVIDIA library integration** (cuBLAS / cuFFT / cuDNN) and Tensor Core intrinsics, plus a native **Apple Metal backend** for Apple Silicon. [Changelog](https://tornadovm.readthedocs.io/en/latest/CHANGELOG.html) Β· [Website](https://www.tornadovm.org) Β· [Documentation](https://tornadovm.readthedocs.io/en/latest/)
---
## This is the whole programming model
TornadoVM gives you three Java APIs for writing GPU kernels, from the least to the most explicit about how work maps onto the hardware. All three build the same `TaskGraph`, run through the same `TornadoExecutionPlan`, and can be mixed in a single graph. Here's one computation, dense matrix multiplication, `C = A Γ B` written all three ways:
Loop Parallel API (@Parallel) |
KernelContext API |
|---|---|
```java
void matmul(Matrix2DFloat A,
Matrix2DFloat B,
Matrix2DFloat C, int size) {
for (@Parallel int i=0; i|
```java
void matmul(KernelContext ctx,
FloatArray A, FloatArray B,
FloatArray C, int size) {
int row = ctx.globalIdx;
int col = ctx.globalIdy;
float sum = 0f;
for (int k=0; k | |
| TileContext API (CuTile) π β CUDA backend only | |
| ```java void matmul(TileContext tc, HalfFloatArray A, HalfFloatArray B, FloatArray C, int m, int n, int k) { PartitionView aView = tc.partition(tc.view(A, m, k), TILE, TILE); PartitionView bView = tc.partition(tc.view(B, k, n), TILE, TILE); PartitionView cView = tc.partition(tc.view(C, m, n), TILE, TILE); Tile acc = tc.zeros(DType.F32, TILE, TILE); for (int step = 0; step < k / TILE; step++) { acc = tc.mma(aView.load(tc.bidX(), step), bView.load(step, tc.bidY()), acc); } cView.store(acc, tc.bidX(), tc.bidY()); } // No thread index, no shared memory, no barrier, no instruction shape β // the tile compiler picks them. ``` | |