TornadoVM Annotations: @Parallel and @Reduce Guide

TornadoVM provides two source-level annotations that let you express GPU parallelism by decorating ordinary Java for loops, without writing a single line of OpenCL or CUDA. The @Parallel annotation marks a loop iteration variable as a parallel dimension — the TornadoVM JIT compiler replaces the loop with a GPU thread launch, using the loop bounds to determine the total thread count. The @Reduce annotation marks an accumulator variable as the output of a parallel reduction — the compiler inserts the necessary tree-reduction hardware instructions automatically. Both annotations live in the package uk.ac.manchester.tornado.api.annotations and are retained at runtime so the JIT compiler can inspect them via reflection during code generation.

Annotations and KernelContext are mutually exclusive programming models. Do not mix @Parallel loop annotations with KernelContext thread-ID reads inside the same task method. If you need explicit thread IDs, use KernelContext exclusively and omit all annotations from that method.

@Parallel

@Parallel targets the loop iteration variable (ElementType.LOCAL_VARIABLE) of a simple counted for loop. At JIT-compile time, TornadoVM replaces the loop body with a parallel GPU kernel where each thread handles one iteration independently.

Declaration

package uk.ac.manchester.tornado.api.annotations;

@Target({ ElementType.LOCAL_VARIABLE, ElementType.TYPE, ElementType.TYPE_USE, ElementType.TYPE_PARAMETER })
@Retention(RetentionPolicy.RUNTIME)
public @interface Parallel { }

Annotation Properties

@Target

ElementType[]

LOCAL_VARIABLE, TYPE, TYPE_USE, TYPE_PARAMETER — applied to the loop iteration variable declared in the for initializer.

@Retention

RetentionPolicy

RUNTIME — the annotation is retained in bytecode so the TornadoVM JIT compiler can inspect it via reflection during code generation.

How TornadoVM uses it

When the compiler encounters a for loop whose iteration variable is annotated with @Parallel, it:

Uses the loop’s upper bound as the global work size for that dimension.
Replaces the sequential iteration with a GPU thread index read (get_global_id(n) in OpenCL; blockIdx.x * blockDim.x + threadIdx.x in CUDA).
Strips the loop structure entirely from the generated kernel — each GPU thread executes the body exactly once.

Restrictions

The loop must be a simple counted for loop with a constant or array-length upper bound.
The iteration variable must be of type int or long.
The loop increment must be +1 (i.e., i++ or i += 1).
No break, continue, or non-loop-bound return inside the annotated loop.
Nesting up to three @Parallel loops is supported for 1D, 2D, and 3D parallelism.

TornadoVM does not check for loop-carried data dependencies at the annotation level. It is the programmer’s responsibility to ensure that iterations are independent before applying @Parallel. Annotating a loop with a dependency produces incorrect results silently.

Usage Parameter

loop iteration variable

@Parallel

The annotation is placed immediately before the type in the for loop initializer. The variable type must be int or long. TornadoVM uses the loop’s upper bound expression as the global work size for the corresponding GPU dimension.

for (@Parallel int i = 0; i < N; i++) { ... }

1D Parallel Example

import uk.ac.manchester.tornado.api.annotations.Parallel;

public class VectorAdd {

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

TornadoVM compiles add into a 1D kernel with a.length threads. Each thread computes a single element of c.

2D Parallel Example

Two nested @Parallel annotations map to a 2D GPU thread grid:

public static void matrixAdd(float[] a, float[] b, float[] c, int M, int N) {
    for (@Parallel int i = 0; i < M; i++) {
        for (@Parallel int j = 0; j < N; j++) {
            c[i * N + j] = a[i * N + j] + b[i * N + j];
        }
    }
}

The outer @Parallel loop maps to globalIdy (Y dimension); the inner loop maps to globalIdx (X dimension). The global work size becomes M × N.

3D Parallel Example

public static void stencil3D(float[] in, float[] out, int X, int Y, int Z) {
    for (@Parallel int z = 0; z < Z; z++) {
        for (@Parallel int y = 0; y < Y; y++) {
            for (@Parallel int x = 0; x < X; x++) {
                int idx = z * Y * X + y * X + x;
                out[idx] = in[idx] * 0.5f;
            }
        }
    }
}

@Reduce

@Reduce marks an accumulator variable (parameter, local variable, or field) as the output of a parallel reduction. TornadoVM inspects the loop body to infer the reduction operation (sum, max, or min) from the binary operator applied to the accumulator, then generates a hardware-efficient tree-reduction kernel.

Declaration

package uk.ac.manchester.tornado.api.annotations;

@Target({ ElementType.PARAMETER, ElementType.LOCAL_VARIABLE, ElementType.FIELD })
@Retention(RetentionPolicy.RUNTIME)
public @interface Reduce { }

Annotation Properties

@Target

ElementType[]

PARAMETER, LOCAL_VARIABLE, FIELD — applied to the accumulator variable, which must be a single-element primitive array parameter.

@Retention

RetentionPolicy

RUNTIME — retained in bytecode so the JIT compiler can detect and specialise the reduction pattern.

How TornadoVM uses it

When the compiler detects @Reduce on an accumulator:

It generates a two-phase reduction: a parallel phase where each thread reduces a sub-range of the input, followed by a serial phase on the host that combines the per-thread partial results.
The reduction operation is inferred from the update expression in the loop body (+= → sum, = Math.max(…) → max, = Math.min(…) → min).
The partial result array written by the kernel has a size determined by the number of thread blocks; the host phase is transparent to the user.

Usage Parameter

accumulator parameter

@Reduce

Applied to a single-element primitive array that serves as the reduction output. The array must be passed as a method parameter (not a local variable) so TornadoVM can size the per-thread partial-result buffer correctly.

public static void sum(float[] input, @Reduce float[] result) { ... }

Supported Reduction Operations

Show Inferred reduction operations and identity values

Expression pattern	Inferred operation	Identity value
`result[0] += a[i]`	Sum	`0`
`result[0] = result[0] + a[i]`	Sum	`0`
`result[0] = Math.max(result[0], a[i])`	Maximum	`Float.MIN_VALUE` / `Integer.MIN_VALUE`
`result[0] = Math.min(result[0], a[i])`	Minimum	`Float.MAX_VALUE` / `Integer.MAX_VALUE`

TornadoVM infers the operation from the assignment expression inside the @Parallel loop body. The identity value must be assigned to result[0] by the kernel method itself (see examples below) — the runtime does not inject an initialisation.

Restrictions and Limitations

Only one @Reduce accumulator is allowed per task method in the current TornadoVM release.
The accumulator must be a primitive scalar (int, long, float, double) held in a single-element array (e.g., float[] sum = { 0.0f }) passed as a parameter. This enables the JIT to allocate the per-thread partial-results buffer at the correct size.
The reduction loop must also be annotated with @Parallel; the two annotations work together.
Mixed reduction types (e.g., computing both sum and max in the same loop) are not supported in a single task.

The accumulator array passed to a reduction task must be initialised to the identity element for the operation before each execute() call: 0 for sum, Integer.MIN_VALUE / Float.MIN_VALUE for max, Integer.MAX_VALUE / Float.MAX_VALUE for min. TornadoVM does not reset the accumulator automatically between invocations.

Reduction Examples

import uk.ac.manchester.tornado.api.*;
import uk.ac.manchester.tornado.api.annotations.Parallel;
import uk.ac.manchester.tornado.api.annotations.Reduce;
import uk.ac.manchester.tornado.api.enums.DataTransferMode;

public class ReductionSum {

    // The accumulator is a single-element array so the JIT can
    // allocate the per-block partial-results buffer at the right size.
    public static void sum(float[] input, @Reduce float[] result) {
        result[0] = 0.0f;
        for (@Parallel int i = 0; i < input.length; i++) {
            result[0] += input[i];
        }
    }

    public static void main(String[] args) throws Exception {
        int n = 1 << 20;
        float[] input  = new float[n];
        float[] result = new float[]{ 0.0f };

        // fill input ...
        for (int i = 0; i < n; i++) input[i] = 1.0f;

        TaskGraph tg = new TaskGraph("reduction")
            .transferToDevice(DataTransferMode.EVERY_EXECUTION, input)
            .task("sum", ReductionSum::sum, input, result)
            .transferToHost(DataTransferMode.EVERY_EXECUTION, result);

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

        System.out.println("Sum = " + result[0]); // expected: 1048576.0
    }
}

Annotations vs. KernelContext — When to Use Which

Use @Parallel / @Reduce when...

Your kernel is a simple element-wise map or reduction
You want minimal boilerplate and automatic thread-count derivation
You do not need explicit control over workgroup size or local memory
You are porting existing sequential Java code incrementally

Use KernelContext when...

You need explicit thread IDs (globalIdx, localIdx, etc.)
You want to allocate and use shared/local memory for tiling
You require barriers, atomics, or warp-level SIMD intrinsics
You are targeting Tensor Core MMA operations
You need precise control over workgroup dimensions via GridScheduler

You can mix annotated tasks and KernelContext tasks within the same TaskGraph — just not within the same task method. For example, a graph might have a @Parallel element-wise normalization task followed by a KernelContext-based tiled matrix multiply task.

Task Graph

Kernel Context

Explore the explicit thread-programming alternative to annotations.

Core API

Data Types

Library Tasks

TornadoVM Annotations: @Parallel and @Reduce Guide

@Parallel

Declaration

Annotation Properties

How TornadoVM uses it

Restrictions

Usage Parameter

1D Parallel Example

2D Parallel Example

3D Parallel Example

@Reduce

Declaration

Annotation Properties

How TornadoVM uses it

Usage Parameter

Supported Reduction Operations

Restrictions and Limitations

Reduction Examples

Annotations vs. KernelContext — When to Use Which

Use @Parallel / @Reduce when...

Use KernelContext when...

Task Graph

Kernel Context

Build docs developers (and LLMs) love

Core API

Data Types

Library Tasks

Documentation Index

​@Parallel

​Declaration

​Annotation Properties

​How TornadoVM uses it

​Restrictions

​Usage Parameter

​1D Parallel Example

​2D Parallel Example

​3D Parallel Example

​@Reduce

​Declaration

​Annotation Properties

​How TornadoVM uses it

​Usage Parameter

​Supported Reduction Operations

​Restrictions and Limitations

​Reduction Examples

​Annotations vs. KernelContext — When to Use Which

Use @Parallel / @Reduce when...

Use KernelContext when...

Task Graph

Kernel Context

Build docs developers (and LLMs) love

@Parallel

Declaration

Annotation Properties

How TornadoVM uses it

Restrictions

Usage Parameter

1D Parallel Example

2D Parallel Example

3D Parallel Example

@Reduce

Declaration

Annotation Properties

How TornadoVM uses it

Usage Parameter

Supported Reduction Operations

Restrictions and Limitations

Reduction Examples

Annotations vs. KernelContext — When to Use Which