Loop Parallel API: @Parallel and @Reduce Annotations

TornadoVM’s Loop Parallel API lets you turn ordinary sequential Java methods into GPU kernels without touching any GPU-specific code. Instead of rewriting your algorithms in OpenCL or CUDA, you annotate loop-index variables with @Parallel to declare which loops can execute concurrently across hardware threads, and you annotate accumulator parameters with @Reduce to indicate that TornadoVM should generate a parallel reduction pattern. The compiler and runtime handle all the thread-dispatch, synchronisation, and backend-specific code generation automatically. This API is the recommended starting point for developers who are new to GPU programming.

The `@Parallel` Annotation

@Parallel is a runtime-retained annotation that targets local variables (specifically loop-index declarations). When TornadoVM’s JIT compiler encounters a loop whose counter is marked @Parallel, it maps each iteration to a separate GPU thread, deriving the ND-Range work size from the loop’s bounds.

1D Parallelism
2D Parallelism
3D Parallelism

Annotate a single loop counter to dispatch a 1D grid of threads — one thread per iteration.

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

public static void scaleArray(FloatArray data, float factor) {
    for (@Parallel int i = 0; i < data.getSize(); i++) {
        data.set(i, data.get(i) * factor);
    }
}

TornadoVM launches data.getSize() threads, each executing one iteration in parallel.

Annotate two consecutive loops to create a 2D thread grid, useful for matrix and image operations.

import uk.ac.manchester.tornado.api.annotations.Parallel;
import uk.ac.manchester.tornado.api.types.matrix.Matrix2DFloat;

public static void matMul(Matrix2DFloat A, Matrix2DFloat B,
                          Matrix2DFloat C, int size) {
    for (@Parallel int i = 0; i < size; i++) {
        for (@Parallel int j = 0; j < size; j++) {
            float sum = 0.0f;
            for (int k = 0; k < size; k++) {
                sum += A.get(i, k) * B.get(k, j);
            }
            C.set(i, j, sum);
        }
    }
}

The two @Parallel loops map to get_global_id(0) and get_global_id(1) in OpenCL, and to the XY block-thread indices in CUDA PTX.

Three consecutive @Parallel loops create a 3D ND-Range grid, useful for volumetric computations.

public static void volumeKernel(FloatArray vol,
                                int nx, int ny, int nz) {
    for (@Parallel int x = 0; x < nx; x++) {
        for (@Parallel int y = 0; y < ny; y++) {
            for (@Parallel int z = 0; z < nz; z++) {
                int idx = x * ny * nz + y * nz + z;
                vol.set(idx, vol.get(idx) + 1.0f);
            }
        }
    }
}

The loop bounds must be deterministic at JIT-compile time — they can be method parameters or constants, but not values computed from non-constant expressions that the TornadoVM compiler cannot resolve.

Sequential vs. Annotated: Side-by-Side

The annotation approach keeps the business logic entirely in Java. You can run the exact same method on CPU (by not routing it through a TaskGraph) or on GPU by wrapping it in a task graph — there is no code duplication.

public static void vectorAdd(float[] a, float[] b, float[] c) {
    for (int i = 0; i < a.length; i++) {
        c[i] = a[i] + b[i];
    }
}

The only structural changes are:

Primitive arrays (float[]) replaced with TornadoVM off-heap types (FloatArray).
Direct array indexing (a[i]) replaced with typed accessors (a.get(i)).
The @Parallel annotation on the loop counter.

The `@Reduce` Annotation

@Reduce targets method parameters (or local variables acting as accumulators) to signal that the annotated variable participates in a parallel reduction. TornadoVM automatically generates work-group-level reduction code for GPUs and scalar-fold code for CPUs — the correct pattern is selected at JIT-compile time based on the target device. Supported reduction operators:

Addition (+)
Multiplication (*)
Maximum (Math.max)
Minimum (Math.min)

For int, long, float, and double element types.

Reduction Sum Example

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

public static void sumArray(FloatArray input, @Reduce FloatArray result) {
    for (@Parallel int i = 0; i < input.getSize(); i++) {
        result.set(0, result.get(0) + input.get(i));
    }
}

Reduction Max Example

public static void maxArray(FloatArray input, @Reduce FloatArray result) {
    for (@Parallel int i = 0; i < input.getSize(); i++) {
        result.set(0, TornadoMath.max(result.get(0), input.get(i)));
    }
}

The @Reduce array must be pre-allocated before the TaskGraph is built. TornadoVM may internally resize it based on the number of work-groups and threads chosen by the runtime scheduler. Always read result.get(0) for the final accumulated value, not a higher index.

Complete Reduction Task Graph

import uk.ac.manchester.tornado.api.*;
import uk.ac.manchester.tornado.api.annotations.*;
import uk.ac.manchester.tornado.api.enums.DataTransferMode;
import uk.ac.manchester.tornado.api.types.arrays.FloatArray;
import java.util.Random;

public class ReductionExample {

    public static void reductionSum(FloatArray input,
                                    @Reduce FloatArray result) {
        for (@Parallel int i = 0; i < input.getSize(); i++) {
            result.set(0, result.get(0) + input.get(i));
        }
    }

    public static void main(String[] args) throws Exception {
        final int SIZE = 65_536;

        FloatArray input  = new FloatArray(SIZE);
        FloatArray result = FloatArray.fromElements(0.0f);

        // Populate input with random values
        Random rnd = new Random();
        for (int i = 0; i < SIZE; i++) {
            input.set(i, rnd.nextFloat());
        }

        TaskGraph tg = new TaskGraph("reduce")
            .transferToDevice(DataTransferMode.EVERY_EXECUTION, input)
            .task("sum", ReductionExample::reductionSum, input, result)
            .transferToHost(DataTransferMode.EVERY_EXECUTION, result);

        ImmutableTaskGraph itg = tg.snapshot();

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

        System.out.printf("Sum = %.4f%n", result.get(0));
    }
}

Map/Reduce Pipeline

Multi-task graphs let you chain a map step and a reduce step inside a single execution plan, keeping the intermediate data on the device.

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

public static void reduceSum(IntArray c, @Reduce IntArray result) {
    result.set(0, 0);
    for (@Parallel int i = 0; i < c.getSize(); i++) {
        result.set(0, result.get(0) + c.get(i));
    }
}

TaskGraph tg = new TaskGraph("mapreduce")
    .transferToDevice(DataTransferMode.EVERY_EXECUTION, a, b)
    .task("t0", MyClass::mapAdd,    a, b, c)      // map phase
    .task("t1", MyClass::reduceSum, c, result)    // reduce phase
    .transferToHost(DataTransferMode.EVERY_EXECUTION, result);

Reductions with Data Dependencies

@Reduce works even when each reduction term is computed on-the-fly from other inputs, as in the π estimation below:

public static void computePi(FloatArray input,
                              @Reduce FloatArray result) {
    for (@Parallel int i = 1; i < input.getSize(); i++) {
        float term = (float) (Math.pow(-1, i + 1) / (2 * i - 1));
        result.set(0, result.get(0) + term + input.get(i));
    }
}

Limitations and Best Practices

What TornadoVM Can Parallelise

Loops with simple integer bounds (i < size, i < array.getSize())
Element-wise operations with no cross-iteration data dependencies
Reductions using +, *, max, min on supported element types
Nested @Parallel loops (up to 3 dimensions)

Current Limitations

Loops with non-constant or data-dependent bounds may not parallelize
Pointer aliasing between input and output arrays is not supported
@Parallel on non-outermost loops only when outer loops are also @Parallel
Only static methods are compiled to GPU kernels

Why only static methods?

TornadoVM’s JIT compiler analyses the bytecode of the target method to generate GPU code. Instance methods introduce implicit this references and potential heap accesses that the compiler cannot safely eliminate. Using static methods (with all data passed as explicit parameters) ensures the compiler can produce a pure, side-effect-free GPU kernel.

Can I mix @Parallel and non-@Parallel loops?

Yes. Non-annotated inner loops (like the k loop in matrix multiplication) execute sequentially within each GPU thread. Only the loops marked @Parallel are distributed across the thread grid. This is the standard tiling pattern: outer parallel loops select the tile, inner sequential loops compute within it.

For maximum performance on GPUs, ensure that @Parallel loops iterate over the outermost dimensions. This maximises memory access coalescing and keeps global memory access patterns GPU-friendly.

Get Started

Programming Guide

Backends

NVIDIA Integration

Guides

Contributing

Loop Parallel API: @Parallel and @Reduce Annotations

The `@Parallel` Annotation

Sequential vs. Annotated: Side-by-Side

The `@Reduce` Annotation

Reduction Sum Example

Reduction Max Example

Complete Reduction Task Graph

Map/Reduce Pipeline

Reductions with Data Dependencies

Limitations and Best Practices

What TornadoVM Can Parallelise

Current Limitations

Build docs developers (and LLMs) love

Get Started

Programming Guide

Backends

NVIDIA Integration

Guides

Contributing

Documentation Index

​The @Parallel Annotation

​Sequential vs. Annotated: Side-by-Side

​The @Reduce Annotation

​Reduction Sum Example

​Reduction Max Example

​Complete Reduction Task Graph

​Map/Reduce Pipeline

​Reductions with Data Dependencies

​Limitations and Best Practices

What TornadoVM Can Parallelise

Current Limitations

Build docs developers (and LLMs) love

The `@Parallel` Annotation

Sequential vs. Annotated: Side-by-Side

The `@Reduce` Annotation

Reduction Sum Example

Reduction Max Example

Complete Reduction Task Graph

Map/Reduce Pipeline

Reductions with Data Dependencies

Limitations and Best Practices