# TornadoVM 8.0: Meet TC, TornadoVM's Own Compiler

> Source: <https://www.tornadovm.org/blogs/tornadovm-8-0-0-tc-tornadovm-compiler>
> Published: 2026-10-30 10:17:36+00:00

**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](https://github.com/beehive-lab/TornadoVM/pull/1228) touches 1,218 files for**+40,166 / −119,699** lines. The relocated Graal module, the vendored`jdk.internal.vm.ci` and the three backends' Graal packages are all deleted.
- **Test parity or better on every backend.** TC passes more of`tornado-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 of`jdk.internal.vm.ci` exports. That made every JDK upgrade awkward and ruled out Leyden AOT caches and GraalVM Native Image.
- **Size.** 644 TornadoVM files imported`tornado.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](https://github.com/beehive-lab/TornadoVM/issues/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](https://github.com/mikepapadim/tornado-compiler-spec).

## 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 loading) 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-aware`min` /`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 rebuilds`if` /`while` /`break` itself instead of emitting a flat control-flow graph full of`goto` 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](https://github.com/mikepapadim/tornado-compiler-spec/blob/main/compiler-replacement/examples.md) 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](https://github.com/beehive-lab/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](fast-start-up-for-gpu-java-tornadovm-leyden-graalvm-native-image.html).

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` builds`tc` 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](https://github.com/mikepapadim/tornadovm-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!](https://sdkman.io/):

```
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](../downloads.html). 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](https://github.com/beehive-lab/TornadoVM/issues) with the `tc emit` output or a `tornado-selftest --json` report attached.

To dig deeper, the [TC specification](https://github.com/mikepapadim/tornado-compiler-spec) 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](https://github.com/beehive-lab/TornadoVM/blob/master/tornado-compiler/README.md) explains how to run TC's own tests.

[Read the 6.0.0 release post →](tornadovm-6-0-0-zero-jni-faster-runtime.html) for how TornadoVM got to one build across JDK 21–27 and dropped JNI.
