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.

TornadoVM exposes two parallel systems for controlling runtime behaviour: CLI flags passed directly to the tornado Python wrapper (e.g. --printKernel), and JVM system properties passed with the -D prefix via --jvm (e.g. -Dtornado.printKernel=true). CLI flags are convenience shortcuts — each one maps internally to one or more -D properties. The properties themselves are all defined in TornadoOptions.java inside the tornado-runtime module. Understanding which flag to use and when is key to diagnosing compilation failures, profiling kernel execution time, and tuning device scheduling.
In flag examples below, s0 refers to a task graph name and t0 to a specific task within that graph. These match the names you use when constructing TaskGraph objects in your application code.

Basic Usage

# Using a CLI flag
tornado --printKernel -m tornado.examples/uk.ac.manchester.tornado.examples.compute.Montecarlo 1024

# Using a JVM property via --jvm
tornado --jvm "-Dtornado.printKernel=true" \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.Montecarlo 1024

# Combining multiple flags
tornado --debug --printKernel --threadInfo \
  --jvm "-Dtornado.fullDebug=true" \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.Montecarlo 1024

Debugging and Logging

FlagDescription
--fullDebugEnables full debug mode. Maps to -Dtornado.fullDebug=true.
--debugEnables basic debug output: compilation status, device info, execution events.
--printKernelPrints the generated OpenCL/CUDA/Metal kernel source to stdout.
--threadInfoDisplays the number of threads (global and local work sizes) used for each kernel launch.
--devicesLists all available hardware devices across all enabled backends.

Debugging Workflow: Printing Generated Kernels

When a kernel behaves unexpectedly, the first step is to inspect the generated source:
tornado --printKernel \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixMultiplication2D 128
To save the kernel for further analysis:
tornado --jvm "-Dtornado.print.kernel.dir=/tmp/matmul.cl" \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixMultiplication2D 128

Profiling

FlagDescription
--enableProfiler consoleEnables profiling and prints metrics as JSON to stdout after execution.
--enableProfiler silentCollects profiling metrics internally; access them via the TornadoVM Profiler API.
--dumpProfiler FILENAMESaves profiling output (JSON) to the specified file.

Profiling Workflow: Measuring Kernel Execution Time

# Print profiling JSON to console
tornado --enableProfiler console \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.NBody 1024 1

# Save profiling results to file
tornado --enableProfiler console --dumpProfiler /tmp/nbody-profile.json \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.NBody 1024 1
The JSON output includes fields for task-graph total time, kernel execution time, data-transfer time (host-to-device and device-to-host), and compilation time.

Device Selection

FlagDescription
--devicesPrints all available backends and devices with their index pairs (B:D).
# List all available devices first
tornado --devices

# Example output:
# Number of Tornado drivers: 2
# Driver: OpenCL
#   0: 0 -- GeForce GTX 1050
#   0: 1 -- Intel(R) Core(TM) i7-7700HQ CPU @ 2.80GHz
# Driver: CUDA
#   1: 0 -- GeForce GTX 1050

# Force task t0 in graph s0 onto CUDA device 0
tornado --jvm "-Ds0.t0.device=1:0" \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixMultiplication2D 512

Performance and Scheduling

PropertyDescription
-Ds0.t0.global.workgroup.size=X,Y,ZSets a custom global work-group size for task t0 in graph s0.
-Ds0.t0.local.workgroup.size=X,Y,ZSets a custom local (thread-block) work-group size.
-Dtornado.scheduler.block=truePartitions the iteration space into blocks, one per visible CPU core (default: false). Useful for CPU OpenCL devices.
-Dtornado.concurrent.devices=trueEnables concurrent execution of tasks across multiple devices (default: false).
PropertyDescription
-Dtornado.reuse.device.buffers=falseDisables reusing device buffers across executions of the same task graph (default: true).
-Dtornado.deallocate.buffers=falseDisables freeing device resources when the execution plan closes (default: true).
-Dtornado.ns.time=trueUses nanoseconds for all timing measurements (default: true).

Optimizations

PropertyDefaultDescription
-Dtornado.enable.fma=truetrueEnables fused multiply-add instructions. May cause precision differences on some platforms.
-Dtornado.enable.mathOptimizations=truetrueEnables math simplifications, e.g. 1/sqrt(x)rsqrt.
-Dtornado.enable.fastMathOptimizations=truetrueEnables aggressive fast-math optimisations.
-Dtornado.enable.nativeFunctions=truetrueEnables native math intrinsics in the generated kernel.
-Dtornado.experimental.partial.unroll=truefalseEnables partial loop unrolling. Combine with -Dtornado.unroll.factor=N (default: 4).

CUDA C Backend

These flags apply only when using the CUDA C backend (make BACKEND=cuda).
PropertyDescription
-Dtornado.cuda.compile.profile=PROFILENamed NVRTC compilation profile. Options: default (no extra flags), fast (--use_fast_math --extra-device-vectorization), debug (-lineinfo for Nsight profiling), repro (--fmad=false for reproducibility).
-Dtornado.cuda.compiler.flags=FLAGSAppends additional flags to the NVRTC compiler invocation. Applied after the profile flags.
-Dtornado.cuda.host.pinning=falseDisables host memory pinning for host↔device transfers (default: true). Pinned memory enables faster DMA transfers.
The CUDA backend’s code cache properties still carry the opencl prefix — inherited from the OpenCL backend’s implementation and not yet renamed. These apply to both backends.
PropertyDefaultDescription
-Dtornado.opencl.codecache.enable=truetrueEnables the on-disk kernel code cache.
-Dtornado.opencl.codecache.dump=truefalseDumps compiled binaries to the cache directory.
-Dtornado.opencl.source.dump=truefalseDumps generated source files to the source directory.
-Dtornado.opencl.codecache.dir=PATH/var/opencl-codecacheDirectory for compiled kernel cache.
-Dtornado.opencl.source.dir=PATH/var/opencl-compilerDirectory for generated kernel sources.
-Dtornado.opencl.log.dir=PATH/var/opencl-logsDirectory for OpenCL/CUDA compiler logs.

Metal Backend (macOS / Apple Silicon)

These flags apply only when using the Metal backend (make BACKEND=metal) on macOS.
PropertyDefaultDescription
-Dtornado.metal.fastmath=truefalseCompiles Metal kernels with fast/relaxed math. Trades some FP precision for speed.
-Dtornado.metal.threadgroupHint=truefalseEmits a max_total_threads_per_threadgroup attribute when the local work-group size is statically known, helping the Metal compiler tune occupancy.
-Dtornado.metal.profiling.enable=falsetrueDisables Metal GPU profiling.
-Dtornado.metal.compiler.flags=FLAGS(none)Passes additional flags to the Metal compiler.

Common Debugging Workflows

1

Inspect the generated kernel

tornado --printKernel \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixMultiplication2D 256
Look for correct loop bounds, data types (float vs double), and that @Parallel loops map to the expected get_global_id() calls.
2

Check thread and work-group configuration

tornado --threadInfo \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixMultiplication2D 256
Verify that the global work size matches your array dimensions and that the local work size is a sensible divisor.
3

Profile execution time breakdown

tornado --enableProfiler console \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixMultiplication2D 256
The JSON output separates kernel time from data-transfer time, helping you identify whether your bottleneck is compute or PCIe bandwidth.
4

Enable full debug for compiler diagnostics

tornado --fullDebug \
  --jvm "-Dtornado.print.bytecodes=true" \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixMultiplication2D 256
This prints internal TornadoVM bytecodes before kernel compilation — useful when reporting compiler bugs or understanding how Java code maps to GPU operations.
5

Force execution on a specific device

# List devices first
tornado --devices

# Pin task t0 of graph s0 to device 0:1
tornado --jvm "-Ds0.t0.device=0:1" \
  -m tornado.examples/uk.ac.manchester.tornado.examples.compute.MatrixMultiplication2D 256

Quick-Reference: CLI Flags Summary

CLI FlagJVM EquivalentCategory
--fullDebug-Dtornado.fullDebug=trueDebug
--debug-Dtornado.debug=trueDebug
--printKernel-Dtornado.printKernel=trueDebug
--threadInfo-Dtornado.threadInfo=trueDebug
--devices(prints device list and exits)Device info
--enableProfiler console-Dtornado.profiler=trueProfiling
--enableProfiler silent-Dtornado.log.profiler=trueProfiling
--dumpProfiler FILE-Dtornado.profiler.dump.dir=FILEProfiling
--jvm "FLAGS"(passes FLAGS directly to the JVM)JVM pass-through
-cp CLASSPATH(classpath)Launch
-m MODULE/CLASS(module-qualified main class)Launch

Build docs developers (and LLMs) love