Quickstart: Your First GPU-Accelerated Java Program

This guide walks you through verifying your TornadoVM installation, running a built-in example, and then writing your own GPU-accelerated program from scratch — first using the simple @Parallel annotation style, then using the lower-level KernelContext API. By the end you will have a complete, runnable program and understand exactly how to launch it with both the tornado command and the java @tornado-argfile approach.

Step 0: Prerequisites Check

Before diving in, confirm that TornadoVM and a suitable JDK are ready on your machine.

Check your JDK version

java -version

You need JDK 21 (or JDK 25). If JAVA_HOME is not set, TornadoVM will not start.

Check that TORNADOVM_HOME is set

echo $TORNADOVM_HOME   # Linux / macOS
echo %TORNADOVM_HOME%  # Windows

This variable is set automatically by SDKMAN! or by sourcing setvars.sh after a source build.

List available devices

tornado --devices

You must see at least one device. If the list is empty, revisit the Installation guide to install the correct GPU driver and backend.

If tornado is not on your PATH, add $TORNADOVM_HOME/bin (Linux/macOS) or %TORNADOVM_HOME%\bin (Windows) to your PATH environment variable and open a new terminal.

Step 1: Run the Built-In Example

TornadoVM ships with a suite of ready-to-run examples. Running one now confirms that the runtime, native libraries, and GPU driver are all wired up correctly — before you write any code.

java @$TORNADOVM_HOME/tornado-argfile \
  -cp $TORNADOVM_HOME/share/java/tornado/tornado-examples-5.2.0.jar \
  uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

You should see timing output and a correctness check confirming that the GPU result matches the CPU reference. If you see an error, run with --debug for more information:

tornado --debug \
  -cp $TORNADOVM_HOME/share/java/tornado/tornado-examples-5.2.0.jar \
  uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

Step 2: Write Your First GPU Program — @Parallel Style

The @Parallel annotation is the fastest way to accelerate an existing Java loop. TornadoVM inspects the annotated index variable and automatically maps each loop iteration to a separate GPU thread — you never write a thread ID calculation. The following program performs an element-wise addition of two FloatArray buffers. FloatArray is TornadoVM’s off-heap array type, allocated outside the Java heap so it can be transferred to the GPU without an extra copy step.

Create the kernel method

Write a plain static method with @Parallel on the loop variable. This method is valid Java and can run sequentially on the JVM with no changes — TornadoVM only accelerates it when it appears inside a TaskGraph.

import uk.ac.manchester.tornado.api.annotations.Parallel;
import uk.ac.manchester.tornado.api.types.arrays.FloatArray;

public class VectorAdd {

    /**
     * Each iteration of this loop runs as an independent GPU thread
     * when TornadoVM JIT-compiles the method.
     */
    public static void add(FloatArray a, FloatArray b, FloatArray c) {
        for (@Parallel int i = 0; i < c.getSize(); i++) {
            c.set(i, a.get(i) + b.get(i));
        }
    }
}

Allocate off-heap arrays and fill with data

int size = 1_000_000;

FloatArray a = new FloatArray(size);
FloatArray b = new FloatArray(size);
FloatArray c = new FloatArray(size);

// FloatArray.init() fills every element with the given value
a.init(1.0f);
b.init(2.0f);

Build a TaskGraph

A TaskGraph declares what to run. transferToDevice moves buffers to the GPU; task registers the kernel method; transferToHost copies results back.

import uk.ac.manchester.tornado.api.TaskGraph;
import uk.ac.manchester.tornado.api.enums.DataTransferMode;

TaskGraph tg = new TaskGraph("s0")
    // Transfer a and b to the device only on the first execution
    // (they don't change between iterations)
    .transferToDevice(DataTransferMode.FIRST_EXECUTION, a, b)
    // Register the kernel — method reference + arguments
    .task("t0", VectorAdd::add, a, b, c)
    // Transfer c back to the host after every execution
    .transferToHost(DataTransferMode.EVERY_EXECUTION, c);

Snapshot the graph and create an execution plan

snapshot() produces an ImmutableTaskGraph — a frozen, thread-safe view of the task graph that cannot be modified. The TornadoExecutionPlan wraps it and provides all runtime controls (device selection, profiling, batching, etc.).

import uk.ac.manchester.tornado.api.ImmutableTaskGraph;
import uk.ac.manchester.tornado.api.TornadoExecutionPlan;

ImmutableTaskGraph itg = tg.snapshot();

// TornadoExecutionPlan implements AutoCloseable — use try-with-resources
try (TornadoExecutionPlan plan = new TornadoExecutionPlan(itg)) {
    plan.execute();  // First call JIT-compiles the kernel; subsequent calls reuse the binary
}

Verify the output

boolean correct = true;
for (int i = 0; i < size; i++) {
    if (Math.abs(c.get(i) - 3.0f) > 1e-5f) {
        correct = false;
        break;
    }
}
System.out.println(correct ? "Result is CORRECT" : "Result is WRONG");

Complete @Parallel Program

import uk.ac.manchester.tornado.api.*;
import uk.ac.manchester.tornado.api.annotations.Parallel;
import uk.ac.manchester.tornado.api.enums.DataTransferMode;
import uk.ac.manchester.tornado.api.types.arrays.FloatArray;

public class VectorAdd {

    public static void add(FloatArray a, FloatArray b, FloatArray c) {
        for (@Parallel int i = 0; i < c.getSize(); i++) {
            c.set(i, a.get(i) + b.get(i));
        }
    }

    public static void main(String[] args) {
        int size = 1_000_000;

        FloatArray a = new FloatArray(size);
        FloatArray b = new FloatArray(size);
        FloatArray c = new FloatArray(size);

        a.init(1.0f);
        b.init(2.0f);

        TaskGraph tg = new TaskGraph("s0")
            .transferToDevice(DataTransferMode.FIRST_EXECUTION, a, b)
            .task("t0", VectorAdd::add, a, b, c)
            .transferToHost(DataTransferMode.EVERY_EXECUTION, c);

        ImmutableTaskGraph itg = tg.snapshot();

        try (TornadoExecutionPlan plan = new TornadoExecutionPlan(itg)) {
            plan.execute();
        }

        boolean correct = true;
        for (int i = 0; i < size; i++) {
            if (Math.abs(c.get(i) - 3.0f) > 1e-5f) {
                correct = false;
                break;
            }
        }
        System.out.println(correct ? "Result is CORRECT" : "Result is WRONG");
    }
}

Step 3: The KernelContext Variant

KernelContext gives you explicit control over thread IDs, local (shared) memory, and synchronisation barriers — the same programming model as CUDA’s threadIdx/blockIdx or OpenCL’s get_local_id(). Use it when you need maximum performance, want to use local memory tiling, or need to call Tensor Core intrinsics. The same vector addition rewritten with KernelContext:

Write the kernel with explicit thread indexing

ctx.globalIdx is TornadoVM’s equivalent of blockIdx.x * blockDim.x + threadIdx.x in CUDA. The bounds check (if (i < c.getSize())) is essential because the GPU may launch more threads than there are elements.

import uk.ac.manchester.tornado.api.KernelContext;
import uk.ac.manchester.tornado.api.types.arrays.FloatArray;

public class VectorAddKernel {

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

Create a WorkerGrid and GridScheduler

With KernelContext, you must specify the thread grid explicitly. WorkerGrid1D sets the total number of global threads; setLocalWork sets the work-group (block) size.

import uk.ac.manchester.tornado.api.GridScheduler;
import uk.ac.manchester.tornado.api.WorkerGrid;
import uk.ac.manchester.tornado.api.WorkerGrid1D;

KernelContext ctx = new KernelContext();
WorkerGrid worker = new WorkerGrid1D(size);
worker.setLocalWork(256, 1, 1);  // 256 threads per work-group

// "s0.t0" must match the TaskGraph name "s0" and task name "t0"
GridScheduler grid = new GridScheduler("s0.t0", worker);

Build the TaskGraph (same structure, extra ctx argument)

The only difference from the @Parallel version is that ctx is passed as the first argument to task(), and the execution plan receives the grid scheduler.

import uk.ac.manchester.tornado.api.*;
import uk.ac.manchester.tornado.api.enums.DataTransferMode;
import uk.ac.manchester.tornado.api.types.arrays.FloatArray;

FloatArray a = new FloatArray(size);
FloatArray b = new FloatArray(size);
FloatArray c = new FloatArray(size);
a.init(1.0f);
b.init(2.0f);

TaskGraph tg = new TaskGraph("s0")
    .transferToDevice(DataTransferMode.FIRST_EXECUTION, a, b)
    .task("t0", VectorAddKernel::add, ctx, a, b, c)
    .transferToHost(DataTransferMode.EVERY_EXECUTION, c);

ImmutableTaskGraph itg = tg.snapshot();

try (TornadoExecutionPlan plan = new TornadoExecutionPlan(itg)) {
    plan.withGridScheduler(grid).execute();
}

When using KernelContext, TornadoVM will not infer a thread grid automatically. If you omit withGridScheduler(grid), execution will throw an error. Always pair a KernelContext kernel with a GridScheduler.

Step 4: Running Your Program

TornadoVM requires a set of JVM flags (--module-path, --add-exports, JVMCI settings, etc.) that would be tedious to type by hand. Two approaches handle this automatically.

tornado command
java @tornado-argfile

The tornado launcher is a wrapper script that sets all required flags and then invokes java. It is the simplest option for running from a terminal.

# Compile your program first
javac -cp $TORNADOVM_HOME/share/java/tornado/tornado-api-5.2.0.jar \
      VectorAdd.java

# Run with the tornado wrapper
tornado -cp . VectorAdd

Useful tornado flags:

Flag	Effect
`--devices`	List discovered devices and exit
`--printKernel` / `-pk`	Print the generated OpenCL/CUDA/MSL kernel source
`--enableProfiler console`	Print per-task profiling data to stdout
`--debug`	Enable verbose debug logging
`--jvm="<opts>"`	Pass extra JVM options (e.g., `-Xmx8g`, device selectors)

The argfile approach uses Java’s native @file syntax to inject all TornadoVM flags from a pre-generated file. It integrates cleanly with Maven, Gradle, and IDEs.

# Compile
javac -cp $TORNADOVM_HOME/share/java/tornado/tornado-api-5.2.0.jar \
      VectorAdd.java

# Run — @$TORNADOVM_HOME/tornado-argfile injects all required JVM flags
java @$TORNADOVM_HOME/tornado-argfile -cp . VectorAdd

On Windows:

java @%TORNADOVM_HOME%\tornado-argfile -cp . VectorAdd

You can append any standard JVM option after the @argfile:

# Select a specific device for task graph "s0", task "t0"
java @$TORNADOVM_HOME/tornado-argfile \
     -Ds0.t0.device=1:1 \
     -cp . VectorAdd

For IntelliJ IDEA or Eclipse, import the project as a Maven project and set @$TORNADOVM_HOME/tornado-argfile as a VM option in your run configuration. This gives you full IDE debugging support while TornadoVM manages the GPU execution.

Step 5: What to Explore Next

More Built-In Examples

The tornado-examples JAR ships with NBody, DFT, KMeans, matrix multiplications, reductions, and more. Explore them in $TORNADOVM_HOME/share/java/tornado/.

Run on a Specific Device

Use -D<graphName>.<taskName>.device=<driver>:<device> to target a specific GPU or CPU. Run tornado --devices first to find the right IDs.

Enable the Profiler

Add --enableProfiler console (tornado) or -Dtornado.profiler=true (argfile) to print per-kernel dispatch, data transfer, and JIT compilation times.

Print the Generated Kernel

Pass --printKernel to tornado or -Dtornado.printKernel=true to the argfile invocation to see the CUDA PTX, OpenCL C, or Metal MSL that TornadoVM generated from your Java method.

Get Started

Programming Guide

Backends

NVIDIA Integration

Guides

Contributing

Quickstart: Your First GPU-Accelerated Java Program

Step 0: Prerequisites Check

Step 1: Run the Built-In Example

Step 2: Write Your First GPU Program — @Parallel Style

Complete @Parallel Program

Step 3: The KernelContext Variant

Step 4: Running Your Program

Step 5: What to Explore Next

More Built-In Examples

Run on a Specific Device

Enable the Profiler

Print the Generated Kernel

Build docs developers (and LLMs) love

Get Started

Programming Guide

Backends

NVIDIA Integration

Guides

Contributing

Documentation Index

​Step 0: Prerequisites Check

​Step 1: Run the Built-In Example

​Step 2: Write Your First GPU Program — @Parallel Style

​Complete @Parallel Program

​Step 3: The KernelContext Variant

​Step 4: Running Your Program

​Step 5: What to Explore Next

More Built-In Examples

Run on a Specific Device

Enable the Profiler

Print the Generated Kernel

Build docs developers (and LLMs) love

Step 0: Prerequisites Check

Step 1: Run the Built-In Example

Step 2: Write Your First GPU Program — @Parallel Style

Complete @Parallel Program

Step 3: The KernelContext Variant

Step 4: Running Your Program

Step 5: What to Explore Next