TornadoVM 8.0: Meet TC, TornadoVM's Own Compiler

wpnews.pro

TornadoVM v8.0.0 compiles your kernels with TC, TornadoVM's own compiler. Since the beginning, TornadoVM has turned Java bytecode into GPU code through a vendored, relocated copy of the Graal compiler running on JVMCI. As of this release that compiler is gone. TC reads the class files itself, optimises them in its own SSA IR, and prints CUDA C++, OpenCL C or Metal Shading Language for the device driver to build.

Your code doesn't change. TaskGraph, TornadoExecutionPlan, KernelContext, @Parallel and @Reduce work as before, and the kernel ABI, the drivers and the profiler keys stay the same. What changes is everything underneath.

At a glance:

No Graal, no JVMCI.PR #1228 touches 1,218 files for**+40,166 / −119,699** lines. The relocated Graal module, the vendoredjdk.internal.vm.ci and the three backends' Graal packages are all deleted.
Test parity or better on every backend. TC passes more oftornado-test than Graal did on CUDA and OpenCL, and every test that passed before still passes.
Same or faster kernels. jitLLM produces token-identical output and runs 1.06–1.14× Graal's throughput on CUDA.
Fast start-up. With JVMCI gone, TornadoVM applications can use a Project Leyden AOT cache or be built into a GraalVM Native Image.
Readable kernels and real diagnostics. Generated code keeps your Java variable names and line numbers. A kernel TC can't compile gets an error code that points at the Java line.

1. Why replace Graal #

Graal served TornadoVM well for years, but it was built as a CPU JIT for HotSpot, and keeping a frozen copy of it inside a GPU compiler had become the main thing holding the project back:

A bytecode ceiling. The vendored Graal class reader can't read newer class files, so TornadoVM builds were pinned to--release 22 . TornadoVM couldn't follow the JDK.
JVMCI. Graal needs JVMCI and a long list ofjdk.internal.vm.ci exports. That made every JDK upgrade awkward and ruled out Leyden AOT caches and GraalVM Native Image.
Size. 644 TornadoVM files importedtornado.graal.* , and each of the three backends carried around 160–220 Graal-based classes (nodes, LIR and phases). Graal's sea-of-nodes pipeline was designed to emit machine code on a CPU, not source kernels for a GPU driver.
Debuggability. There was no textual IR you could read, diff or write a test against.

The replacement was proposed in issue #1227 with ten requirements. The main ones were functional parity on every backend, JDK 21 through 27, Leyden and Native Image compatibility, no JVMCI or new third-party dependencies, and kernel performance at least equal to Graal's. The full design, the migration record and all measurements are public in the TC compiler specification.

2. What TC is #

TC is pure Java (JDK 21+) and lives in tornado-compiler/. It doesn't use JVMCI, Graal, or bindings to LLVM, MLIR or NVVM, and it needs neither a GPU nor the TornadoVM runtime to compile. A kernel goes through six stages:

class files ─► frontend ─► TC-SSA ─► passes ─► Tornado lowering ─► structuring ─► C core + dialect
  (bytecode)    (Braun SSA)  (block      (canonicalise,  (@Parallel, KernelContext,  (if / while /   ─► CUDA C++ (NVRTC)
                             arguments)   GVN, LICM,      @Reduce, atomics,           break)          ─► OpenCL C 1.2
                                          DCE, unroll)    kernel ABI)                                 ─► Metal (MSL)

Frontend. TC parses class files itself (no reflection or class ) and builds SSA directly from the operand stack and locals. Every value keeps its Java name, source line and bytecode index.
IR. TC-SSA is a typed SSA form that uses block arguments instead of phis, MLIR-style. It spells out Java semantics explicitly: masked shifts, saturating float-to-int conversion, NaN-awaremin /max . A verifier checks the IR after every pass in tests, and a printer and parser make it round-trippable text.
Passes. Canonicalisation and constant folding, CFG simplification, dead-code elimination, small-loop unrolling, GVN, LICM and math peepholes.
Tornado lowering. Inlining, and specialisation on what the runtime knows about the actual arguments (array lengths, scalar values, object layouts).@Parallel loops become grid-stride loops, and@Reduce becomes work-group reductions.
Structuring. Because the output is source code for a driver compiler, TC rebuildsif /while /break itself instead of emitting a flat control-flow graph full ofgoto s.
Emission. One C-family printer and thin per-target dialects. Adding a target means adding a dialect (types, address spaces, intrinsic spellings) plus a description of what the device can do. The passes and lowering are shared.

3. Kernels you can read #

Here is a grid-stride vector addition with KernelContext:

public static void add(KernelContext ctx, FloatArray a, FloatArray b, FloatArray c) {
    for (int i = ctx.globalIdx; i < a.getSize(); i += ctx.globalGroupSizeX) {
        c.set(i, a.get(i) + b.get(i));
    }
}

TC turns it into this CUDA:

extern "C" __global__ void C3Kernels_add(long long *_kernel_context, unsigned char *_constant_region, unsigned char *_local_region, int *_atomics, unsigned char *a, unsigned char *b, unsigned char *c)
{
  int a_len, gsize_x, i;
  float *a_data, *b_data, *c_data;
  i = (int) (blockIdx.x * blockDim.x + threadIdx.x);  // C3Kernels.java:34
  a_len = (int) *((long long *) a);  // C3Kernels.java:34
  gsize_x = (int) (gridDim.x * blockDim.x);  // C3Kernels.java:34
  a_data = (float *) (a + 16LL);  // C3Kernels.java:35
  b_data = (float *) (b + 16LL);  // C3Kernels.java:35
  c_data = (float *) (c + 16LL);  // C3Kernels.java:35
  while (i < a_len) {
    c_data[i] = a_data[i] + b_data[i];  // C3Kernels.java:35
    i = i + gsize_x;  // C3Kernels.java:34
  }
}

With TC, the loop is a while, the variables are named a, i and a_len, and every line points back to the Java line it came from. When a profiler such as Nsight Compute points at a hot line of generated CUDA, the comment on that line leads straight back to your Java. The examples in the specification follow this kernel and others through every stage: bytecode, TC-SSA before and after ABI lowering, the structured IR, and the final CUDA beside Graal's.

When TC can't compile a kernel, it doesn't fall over inside a compiler phase. You get a stable code (TC-E1xxx frontend, E2xxx IR, E3xxx Tornado lowering, E4xxx target, E9xxx internal), the Java line and bytecode index, and a hint. The runtime maps the code's category to the matching TornadoVM exception.

4. Parity and performance #

The bar for removing Graal was simple: TC must pass every tornado-test that Graal passes, with no fallback to Graal. It does, and on CUDA and OpenCL it passes more. The table shows the number of tornado-test unit tests that pass on each device:

Device	Tests passed, Graal	Tests passed, TC (v8.0.0)
CUDA, RTX 4090	1516	1539
OpenCL, RTX 4090	1027	1057
OpenCL, Intel UHD 770 (no fp64)	932	934
Metal, Apple M3 Pro	942	941 <sup>*</sup>

<sup>*</sup> Measured with TC before the Graal code was deleted; Metal hasn't been rerun since the removal.

The extra passes are kernels Graal refused and TC compiles: deep inlining, a kernel named after an OpenCL keyword, and atomics on local or global arrays. In the final runs, no test that passed before fails, and no launch configuration changes.

On real applications, jitLLM (Llama-3.2-1B Q8 and Qwen3-0.6B F16, greedy decoding) produces token-identical output under TC. Throughput is 1.06–1.14× Graal's on CUDA and 1.04–1.09× on OpenCL. On Metal the two are level, within noise.

5. Fast start-up: Leyden and Native Image #

JVMCI was the blocker for both Project Leyden's AOT cache and GraalVM Native Image. With it gone, a whole TornadoVM application can use either:

tornado --aotCache app.aot --classpath app.jar com.example.Main   # Leyden (JDK 25+): the first run writes the cache
tornado-native-image -cp app.jar -o app com.example.Main          # GraalVM Native Image: run ./app

On an RTX 4090, a saxpy program that compiles and runs one CUDA kernel takes about 0.6 s from launch to exit on the plain JVM, and about 0.3 s from an AOT cache or as a Native Image. TC's own compile time for the kernel drops from about 30 ms to 11 ms with the cache, and to under 1 ms in the Native Image. Native Image covers the OpenCL and CUDA backends; the JNI-based libraries (cuDNN, cuDF, CUTLASS) aren't supported there yet.

The full measurements, and how to use each option, are in our companion post: Fast Start-up for GPU Java with Project Leyden and GraalVM Native Image.

The build gets simpler as well. The argfile that every TornadoVM launch uses dropped from about 75 lines of JVMCI exports, --patch-module and --upgrade-module-path flags to about a dozen. There's now one argfile for every JDK, and the same jdk22plus SDK runs unchanged on JDK 22 through 27.

6. New tools #

Because TC runs without a GPU, the SDK ships tools that work on any machine:

tc emit compiles a single kernel to source with no device attached, which is useful for inspecting code or attaching it to a bug report.build-tc-native buildstc itself as a Native Image, which runs in 0.02 s instead of 0.19 s and produces byte-identical output.

tc emit --target cuda --cp app.jar com.example.Kernels saxpy

tornado-selftest runs about 20 small kernels on every device TornadoVM finds and checks each result against the same method run on the JVM. It takes a few seconds and prints what TC targets on each device: compute capability, fp64/fp16 support, warp size, tensor cores and more.

tornado-selftest                        # every device
tornado-selftest --json report.json     # also write the report as JSON

tcirviz (tornado --tcirviz ) is an interactive visualiser that follows a value from your Java source, through the bytecode and every TC pass, to the target code.

Get it #

TornadoVM v8.0.0 is available now. Install it with SDKMAN!:

sdk install tornadovm 8.0.0-jdk21-cuda      # or -opencl, -metal, -full
sdk install tornadovm 8.0.0-jdk22plus-cuda  # one SDK for JDK 22 through 27
tornado --devices
tornado-selftest

Prebuilt SDKs for every platform are also on the downloads page. To build from source:

git clone https://github.com/beehive-lab/TornadoVM.git && cd TornadoVM
git checkout v8.0.0
make BACKEND=opencl,cuda          # or: make jdk22plus BACKEND=opencl,cuda
source setvars.sh
tornado-test --quickPass

If you're on 7.x today, upgrading needs no change to your application. Kernels you couldn't compile before may now work. If something that used to compile doesn't, the TC error code and Java line will tell you where to look; please open an issue with the tc emit output or a tornado-selftest --json report attached.

To dig deeper, the TC specification covers the IR, the passes, the per-target kernel ABI, the parity matrix and the full record of the move away from Graal. The compiler README explains how to run TC's own tests.

Read the 6.0.0 release post → for how TornadoVM got to one build across JDK 21–27 and dropped JNI.

source & further reading

tornadovm.org — original article