Skip to main content

Documentation Index

Fetch the complete documentation index at: https://mintlify.com/Deepak-Sangle/TornadoVM/llms.txt

Use this file to discover all available pages before exploring further.

The OpenCL backend is TornadoVM’s default backend and the broadest in hardware coverage. By generating OpenCL C kernel source and delegating JIT compilation to the vendor’s OpenCL runtime, it can target discrete GPUs from NVIDIA, AMD, and Intel; integrated GPUs including Apple Silicon (M-series), Intel HD Graphics, and ARM Mali; multi-core CPUs through the CPU OpenCL runtime; and even FPGAs via vendor-specific OpenCL flows. If you install TornadoVM without specifying a backend, you get OpenCL. The same Java TaskGraph code runs on all of these device classes without modification — TornadoVM enumerates all available OpenCL platforms and devices at startup and exposes them for selection.

What OpenCL Targets

NVIDIA GPUs

Supported via the NVIDIA OpenCL runtime bundled with the CUDA Toolkit. An alternative to the CUDA backend when library tasks and Tensor Core intrinsics are not needed.

AMD GPUs

Discrete and integrated AMD GPUs via AMD ROCm’s OpenCL runtime (amdgpu-pro or ROCm OpenCL). Covers RDNA and CDNA architectures.

Intel GPUs & CPUs

Intel discrete GPUs (Arc), integrated GPUs (HD/Iris/Xe), and CPUs via Intel OpenCL Runtime or OpenCL NEO. CPU fallback is supported for all Intel platforms.

Apple Silicon & ARM

Apple M1–M4 integrated GPUs via the Apple OpenCL runtime, and ARM Mali GPUs via ARM’s compute runtime. For Apple Silicon, the Metal backend is preferred for best performance.

Multi-Core CPUs

OpenCL CPU devices expose all CPU cores as OpenCL compute units. Useful for debugging kernels or for workloads that don’t have enough parallelism to fill a GPU.

FPGAs

Xilinx and Intel/Altera FPGAs expose OpenCL runtime APIs. TornadoVM’s OpenCL C output can be used as input to FPGA-oriented offline toolflows.

Prerequisites

The OpenCL backend requires an OpenCL runtime installed for your target hardware. OpenCL runtimes are vendor-supplied and separate from TornadoVM itself.
NVIDIA The NVIDIA OpenCL runtime ships with the CUDA Toolkit. Install the CUDA Toolkit or the standalone NVIDIA GPU driver (which includes libOpenCL.so).AMD Install AMD ROCm (recommended for discrete GPUs):
# Ubuntu/Debian — follow the ROCm installer guide at rocm.docs.amd.com
sudo apt install rocm-opencl-runtime
Intel GPU / CPU Install Intel’s OpenCL runtime or the open-source NEO driver:
# Ubuntu
sudo apt install intel-opencl-icd
macOS (Apple Silicon) The Apple OpenCL runtime is included with macOS. No additional install needed, though the Metal backend is recommended for best performance on M-series chips.Generic ICD loader On Linux, an ICD loader (ocl-icd-libopencl1) lets multiple vendor runtimes coexist:
sudo apt install ocl-icd-libopencl1 clinfo
clinfo   # verify installed platforms
You also need:
  • JDK 21 or JDK 25 (JAVA_HOME set accordingly)
  • GCC/G++ ≥ 13 to build the TornadoVM OpenCL JNI bridge (opencl-jni)

Installation

OpenCL is the default backend — the base sdk install tornadovm command installs it:
# Default install (OpenCL)
sdk install tornadovm

# Explicit version with OpenCL backend
sdk install tornadovm 5.2.0-opencl

Enumerating Platforms and Devices

TornadoVM queries all OpenCL platforms (one per installed vendor runtime) and all devices within each platform at startup. Use tornado --devices to inspect what is visible:
tornado --devices
Example output on a system with NVIDIA and Intel OpenCL runtimes:
Number of Tornado drivers: 1
Driver: OpenCL
  0: Intel(R) OpenCL -- Intel(R) Core(TM) i9-13900K  [type=CPU]
  1: Intel(R) OpenCL Graphics -- Intel(R) Iris(R) Xe Graphics [type=GPU]
  2: NVIDIA CUDA   -- NVIDIA GeForce RTX 4090   [type=GPU]
Each line corresponds to a device index you can target from a TornadoExecutionPlan:
// Select device index 2 (NVIDIA GPU via OpenCL)
TornadoDevice device = TornadoExecutionPlan.getDevice(0, 2);
try (TornadoExecutionPlan plan = new TornadoExecutionPlan(tg.snapshot())) {
    plan.withDevice(device).execute();
}

JIT Compilation Pipeline

1

Java Bytecode → Graal IR

The task method is parsed by GraalVM’s bytecode parser. @Parallel loop annotations and KernelContext accesses are lifted into parallel GPU IR nodes via OpenCL-specific Graal plugins (OCLBackendImpl, OCLHotSpotBackendFactory).
2

Graal IR → OpenCL C

The OpenCL-specific backend (OCLBackend, OCLLIRStmt) serialises the LIR to OpenCL C source. Thread-indexing maps ctx.globalIdx to get_global_id(0), local memory to __local, and barriers to barrier(CLK_LOCAL_MEM_FENCE).
3

OpenCL C → Platform Binary

The generated OpenCL C source string is passed to clBuildProgram on the selected device. The vendor runtime performs the final compilation to a device-specific binary.
4

Binary Cached & Dispatched

OCLCodeCache caches the compiled binary. On subsequent calls the cache is checked first; recompilation is triggered only when the kernel or device changes.

Running on CPU as a Fallback

Any OpenCL CPU device can be used as a fallback if no GPU is available or for debugging. The CPU device appears in --devices output with type=CPU:
// Select the first available CPU device
TornadoDevice cpuDevice = TornadoExecutionPlan.getDevice(0, 0); // adjust index
try (TornadoExecutionPlan plan = new TornadoExecutionPlan(tg.snapshot())) {
    plan.withDevice(cpuDevice).execute();
}
Running on a CPU device is a good first debugging step. If a kernel produces incorrect results on GPU but correct results on CPU, the issue is often an unintialised local memory access or a race condition in the thread mapping.

OpenCL-Specific Notes

The OpenCL backend does not support CUDA library tasks (cuBLAS, cuFFT, cuDNN, cuSPARSE, CUTLASS), Tensor Core mma.sync intrinsics, or CUDA Graphs — these are CUDA-backend-only features. Full JIT kernel compilation, KernelContext (thread IDs, local memory, barriers), and @Parallel annotations are fully supported across all OpenCL devices.

Running Examples

# List all OpenCL devices
tornado --devices

# Run matrix-vector multiply (auto-selects first available device)
tornado \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

# Print the generated OpenCL C kernel source
tornado --printKernel \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

# Run all unit tests
tornado-test --ea --verbose

Build docs developers (and LLMs) love