beehive-lab/TornadoVM — documented rules

Concern Kernel-entry resolution, Where TaskUtils.resolveViaSerializedLambda, Notes lambda → SerializedLambda → Method

Concern SPI (no HotSpot type), Where TornadoBackendProvider.createBackend(OptionValues, TornadoVMConfigAccess), Notes backends construct reflection providers directly

OCLHotSpotBackendFactory.createJITCompiler → new TornadoMetaAccessProvider(), new TornadoConstantReflectionProvider(snippetReflection) (no HotSpotJVMCIRuntime param).

OCLHotSpotBackendFactory.createJITCompiler → new TornadoMetaAccessProvider(), new TornadoConstantReflectionProvider(snippetReflection) (no HotSpotJVMCIRuntime param).

OCLHotSpotBackendFactory.createJITCompiler → new TornadoMetaAccessProvider(), new TornadoConstantReflectionProvider(snippetReflection) (no HotSpotJVMCIRuntime param).

The *GraphBuilderPlugins.registerInvocationPlugins signature must take jdk.vm.ci.meta.MetaAccessProvider, not HotSpotMetaAccessProvider.

OCLFieldBuffer resolves the object type via TornadoCoreRuntime.getTornadoRuntime().getMetaAccess() and the reflection ResolvedJavaType (no HotSpotResolvedJavaType cast).

graal/phases/TornadoOpenCLIntrinsicsReplacements (registered in OCLHighTier right after HighTierLoweringPhase).

Also OCLLoweringProvider#isLocalArrayInput matches surviving contains("LocalArray") invokes at lowering time.

OCLBackend.emitMethodParameters continues past a KernelContext/AtomicInteger parameter so it is not emitted as a __global buffer arg.

OCLBackend.emitMethodParameters continues past a KernelContext/AtomicInteger parameter so it is not emitted as a __global buffer arg.

OCLNodeLIRBuilder.emitPrologue — a KernelContext/AtomicInteger parameter whose only remaining usages are FrameStates (deopt) must not load a buffer;

OCLLoweringProvider uses (int) TornadoOptions.PANAMA_OBJECT_HEADER_SIZE for panama array element addresses;

OCLAssembler.formatConstant escapes a TornadoObjectConstant wrapping a String as a C string literal (a HotSpotObjectConstant no longer appears on the reflection path).

OCLAssembler.formatConstant escapes a TornadoObjectConstant wrapping a String as a C string literal (a HotSpotObjectConstant no longer appears on the reflection path).

OCLGraphBuilderPlugins.registerNativeArrayAccessPlugins → registerNativeArrayGetSet(Int/Float/ Double/Long/Short/Byte/Int8/Char) + registerHalfFloatArrayGetSet + the arch-neutral arrayElementAddress helper (&segment[(baseIndex+index)*elementBytes]).

`cuDNN also exposes cudnnSoftmax, cudnnRelu/cudnnSigmoid/cudnnTanh, cudnnMaxPool2d, and fused FP16 scaled-dot-product (flash) attention via sdpaForward` — see the full provider catalog linked below.

`cuDNN also exposes cudnnSoftmax, cudnnRelu/cudnnSigmoid/cudnnTanh, cudnnMaxPool2d, and fused FP16 scaled-dot-product (flash) attention via sdpaForward` — see the full provider catalog linked below.

`cuDNN also exposes cudnnSoftmax, cudnnRelu/cudnnSigmoid/cudnnTanh, cudnnMaxPool2d, and fused FP16 scaled-dot-product (flash) attention via sdpaForward` — see the full provider catalog linked below.

`cuDNN also exposes cudnnSoftmax, cudnnRelu/cudnnSigmoid/cudnnTanh, cudnnMaxPool2d, and fused FP16 scaled-dot-product (flash) attention via sdpaForward` — see the full provider catalog linked below.

`cuDNN also exposes cudnnSoftmax, cudnnRelu/cudnnSigmoid/cudnnTanh, cudnnMaxPool2d, and fused FP16 scaled-dot-product (flash) attention via sdpaForward` — see the full provider catalog linked below.

`cuDNN also exposes cudnnSoftmax, cudnnRelu/cudnnSigmoid/cudnnTanh, cudnnMaxPool2d, and fused FP16 scaled-dot-product (flash) attention via sdpaForward` — see the full provider catalog linked below.

Note that the TornadoVM profiler works only if enabled in the execution plan (via the `withProfiler` method).

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

**NOTE:** The methods `init() and clear()` are essential because, contrary to their counterpart primitive arrays which are initialized by default with 0, the new types contain garbage values when first created.

`local only means the operation is supported when the array was allocated with KernelContext.allocateIntLocalArray(int)`;

The Task-Graph API also defines a method, named `transferToDevice` to set which arrays need to be copied to the target accelerator.

Similar to `transferToDevice, the TaskGraph API also offers a call to sync the data back from the device to the host. The API call is transferToHost` with the following parameters:

TornadoVM supports batch processing through the following API call `withBatch of the TornadoExecutionPlan` API.

`Tile exposes getDType, getRank, getDimension and getElement` (JVM path only, for tests and host debugging);

`Tile exposes getDType, getRank, getDimension and getElement` (JVM path only, for tests and host debugging);

`Tile exposes getDType, getRank, getDimension and getElement` (JVM path only, for tests and host debugging);

`Tile exposes getDType, getRank, getDimension and getElement` (JVM path only, for tests and host debugging);

`TensorView exposes getDType, getRank, getExtent, and getElementOffset/getRowStride` for a strided view;

`TensorView exposes getDType, getRank, getExtent, and getElementOffset/getRowStride` for a strided view;

`TensorView exposes getDType, getRank, getExtent, and getElementOffset/getRowStride` for a strided view;

`TensorView exposes getDType, getRank, getExtent, and getElementOffset/getRowStride` for a strided view;

`TensorView exposes getDType, getRank, getExtent, and getElementOffset/getRowStride` for a strided view;

`PartitionView adds getTileDimension and getBlockCount`.

`PartitionView adds getTileDimension and getBlockCount`.

`DType names F16, BF16, F32, F64, TF32, FP8_E4M3, FP8_E5M2, S8, S32 and the comparison-only PRED`.

`DType names F16, BF16, F32, F64, TF32, FP8_E4M3, FP8_E5M2, S8, S32 and the comparison-only PRED`.

`DType names F16, BF16, F32, F64, TF32, FP8_E4M3, FP8_E5M2, S8, S32 and the comparison-only PRED`.

`DType names F16, BF16, F32, F64, TF32, FP8_E4M3, FP8_E5M2, S8, S32 and the comparison-only PRED`.

`DType names F16, BF16, F32, F64, TF32, FP8_E4M3, FP8_E5M2, S8, S32 and the comparison-only PRED`.

`DType names F16, BF16, F32, F64, TF32, FP8_E4M3, FP8_E5M2, S8, S32 and the comparison-only PRED`.

`DType names F16, BF16, F32, F64, TF32, FP8_E4M3, FP8_E5M2, S8, S32 and the comparison-only PRED`.

`DType names F16, BF16, F32, F64, TF32, FP8_E4M3, FP8_E5M2, S8, S32 and the comparison-only PRED`.

`DType names F16, BF16, F32, F64, TF32, FP8_E4M3, FP8_E5M2, S8, S32 and the comparison-only PRED`.

`DType names F16, BF16, F32, F64, TF32, FP8_E4M3, FP8_E5M2, S8, S32 and the comparison-only PRED`.

Factories (uk.ac.manchester.tornado.cublas.CuBlas):

Factory cublasSgemv(op, m, n, alpha, A, lda, x, incx, beta, y, incy), Operation y = alpha·op(A)·x + beta·y

Factory cublasSgemm(opA, opB, m, n, k, alpha, A, lda, B, ldb, beta, C, ldc), Operation C = alpha·op(A)·op(B) + beta·C

Factory cublasSgemmTF32(...), Operation SGEMM using TF32 tensor cores

Factory cublasGemmExFP16(...), Operation FP16 inputs and output, tensor-core GEMM

Factory cublasGemmExFP16FP32(...), Operation FP16 inputs, FP32 output — keeps accumulation precision at the boundary

Factory cublasGemmExBF16(...), Operation BF16 inputs and output ( BFloat16Array ), tensor-core GEMM

Factory cublasSgemmStridedBatched(...), Operation batched SGEMM

Factory ltMatmulFP32 / ltMatmulFP16, Epilogue none

Factory ltMatmulFP32 / ltMatmulFP16, Epilogue none

Factory ltMatmulFP8, Epilogue none (E4M3 operands, FP16 output, TN form, ld multiple of 16 B)

Factory ltMatmulBiasFP16, Epilogue BIAS

Factory ltMatmulGeluBiasFP16, Epilogue GELU_BIAS (tanh approximation)

Factory cufftForwardC2C / cufftInverseC2C(input, output, n, batch), Transform 1D complex→complex

Factory cufftForwardC2C / cufftInverseC2C(input, output, n, batch), Transform 1D complex→complex

Factory cufftForwardR2C / cufftInverseC2R(input, output, n, batch), Transform real↔complex (Hermitian, n ↔ n/2+1 )

Factory cufftForwardR2C / cufftInverseC2R(input, output, n, batch), Transform real↔complex (Hermitian, n ↔ n/2+1 )

Factory cufftForwardZ2Z / cufftInverseZ2Z(input, output, n, batch), Transform FP64 complex→complex, 1D

Factory cufftForwardZ2Z / cufftInverseZ2Z(input, output, n, batch), Transform FP64 complex→complex, 1D

Factory cufftForward2dC2C / cufftInverse2dC2C(input, output, nx, ny), Transform 2D complex→complex

Factory cufftForward2dC2C / cufftInverse2dC2C(input, output, nx, ny), Transform 2D complex→complex

Factory cudnnSoftmax(input, output, rows, cols), Operation per-row numerically-stable softmax

Factory cudnnRelu / cudnnSigmoid / cudnnTanh(input, output, size), Operation activations

Factory cudnnRelu / cudnnSigmoid / cudnnTanh(input, output, size), Operation activations

Factory cudnnRelu / cudnnSigmoid / cudnnTanh(input, output, size), Operation activations

Factory cudnnMaxPool2d(input, output, n, c, h, w, window, stride), Operation 2D max pooling

Factory cudnnConv2d(input, filter, output, n, c, h, w, k, r, s, pad, stride), Operation 2D convolution

Factory sdpaForward(q, k, v, o, b, h, sQ, sKv, d, scale, causal), Operation fused scaled-dot-product attention (FP16, BHSD)

Factory cutlassSgemm(m, n, k, alpha, A, B, beta, C), Operation FP32 SIMT GEMM

Factory cutlassHgemm(m, n, k, alpha, A, B, beta, D), Operation FP16 tensor-core GEMM, FP32 accumulate

Factory cutlassHgemmBatched(m, n, k, alpha, A, B, beta, C, batchCount), Operation batched FP16 tensor-core GEMM

Factory cutlassBgemm(m, n, k, alpha, A, B, beta, C), Operation BF16 tensor-core GEMM ( BFloat16Array )

Factory cutlassGemmBiasRelu(m, n, k, A, B, bias, D), Operation fused relu(A·B + bias)

Factory cutlassGemmBiasGelu(m, n, k, A, B, bias, D), Operation fused gelu(A·B + bias)

Factory cutlassGemmBiasSilu(m, n, k, A, B, bias, D), Operation fused silu(A·B + bias)

Factory cutlassGemmBiasSigmoid(m, n, k, A, B, bias, D), Operation fused sigmoid(A·B + bias)

Factory cutlassGemmBiasTanh(m, n, k, A, B, bias, D), Operation fused tanh(A·B + bias)

Factory cutlassGemmBiasHardSwish(m, n, k, A, B, bias, D), Operation fused hardswish(A·B + bias)

Factory cusparseSpMV(rows, cols, nnz, csrRowOffsets, csrColInd, csrValues, x, y), Operation y = A·x

Factory cusparseSpMM(rows, k, n, nnz, csrRowOffsets, csrColInd, csrValues, b, c), Operation C = A·B , dense B / C row-major

Option withPreCompilation(), What it does JIT every task graph up front For hybrid plans this also runs each provider's prepare() , so native contexts, plans and workspaces exist before the first execution — the same reason it is a prerequisite for graph capture, Backends all

Option withIntraPlanConcurrency(), What it does Routes DAG-independent work (H2D, kernels, D2H) to separate role streams so it can overlap, with ordering preserved through device events derived from the bytecode dependency DAG Off by default, Backends CUDA no-op on OpenCL/Metal

Option withStagedTransfers(), What it does Stages non-batched read-only transfers through pinned host memory instead of the direct path, Backends CUDA no-op on OpenCL/Metal

Option withMemoryLimit("8GB"), What it does Caps device memory for the plan Worth setting explicitly when library workspaces sit alongside large operands, Backends all

Option withWarmUpIterations(n), What it does Runs the whole plan n times first — transfers, compilation and execution — before the measured run, Backends all

TornadoExecutionPlan.transferToDevice(Object...) uploads the current host contents of specific objects **without running the plan**.

…and add the same class name to src/main/resources/META-INF/services/uk.ac.manchester.tornado.runtime.library.spi.TornadoLibraryProvider.

**LibraryContext** — marker;

**LibraryInvocation** — per-call payload: getArg(i), getDevicePointer(i), isReference(i), getContext(), getTuning().

getArg(i), getDevicePointer(i), isReference(i), getContext(), getTuning().

getArg(i), getDevicePointer(i), isReference(i), getContext(), getTuning().

getArg(i), getDevicePointer(i), isReference(i), getContext(), getTuning().

getArg(i), getDevicePointer(i), isReference(i), getContext(), getTuning().

getArg(i), getDevicePointer(i), isReference(i), getContext(), getTuning().

**TornadoNativeStreamSupport** — getNativeStream(planId) / getNativeContext(planId), implemented by the CUDA backend device.

**TornadoNativeStreamSupport** — getNativeStream(planId) / getNativeContext(planId), implemented by the CUDA backend device.

Factory cublasSgemv, cuBLAS function cublasSgemv, Types FloatArray, Semantics y = α·op(A)·x + β·y

Factory cublasSgemm, cuBLAS function cublasSgemm, Types FloatArray, Semantics C = α·op(A)·op(B) + β·C

Factory cublasSgemmTF32, cuBLAS function cublasSgemm + TF32 math mode, Types FloatArray, Semantics Same as cublasSgemm , executed on TF32 Tensor Cores (~1e-4 rel error, up to 1.6x faster)

Factory cublasSgemmStridedBatched, cuBLAS function cublasSgemmStridedBatched, Types FloatArray, Semantics C[i] = α·op(A[i])·op(B[i]) + β·C[i] for a whole batch in one call operand i lives at base + i stride in one flat array

Factory cublasGemmExFP16, cuBLAS function cublasGemmEx, Types HalfFloatArray in/out, Semantics FP16 GEMM with FP32 Tensor Core accumulation

Factory cublasGemmExFP16FP32, cuBLAS function cublasGemmEx, Types HalfFloatArray in, FloatArray out, Semantics FP16 inputs, FP32 output (standard inference config)

Factory CuBlasLt.ltMatmulFP32, cuBLAS function cublasLtMatmul, Types FloatArray, Semantics FP32 matmul (heuristic algorithm, plan cached per shape)

Factory CuBlasLt.ltMatmulFP16, cuBLAS function cublasLtMatmul, Types HalfFloatArray, Semantics FP16 matmul, FP32 Tensor Core accumulation

Factory CuBlasLt.ltMatmulBiasFP16, cuBLAS function cublasLtMatmul + BIAS, Types HalfFloatArray, Semantics C = op(A)·op(B) + bias, fused

Factory CuBlasLt.ltMatmulGeluBiasFP16, cuBLAS function cublasLtMatmul + GELU_BIAS, Types HalfFloatArray, Semantics C = GELU(op(A)·op(B) + bias), fully fused transformer MLP block

**CUDA Graphs:** supported — plans (whose creation allocates a device work area) are built in the provider's prepare() hook, which the runtime invokes before capture starts, so the FFT pipeline is captured and replayed like any other library task (see TestCuFft#testRoundTripWithCudaGraph).

Factory Cusparse.cusparseSpMV(rows, cols, nnz, csrRowOffsets, csrColInd, csrValues, x, y), Operation y = A · x

Factory Cusparse.cusparseSpMM(rows, k, n, nnz, csrRowOffsets, csrColInd, csrValues, B, C), Operation C = A · B (dense B k×n, C rows×n)

those have no C ABI to call, so here they are Java records in registries, and the long handles the backend passes around for programs, kernels and events are registry keys rather than raw pointers (real Metal objects -- devices, queues, buffers, libraries, pipelines, command buffers -- stay raw pointers).

The backend classes (Metal, MetalPlatform, MetalDevice, MetalContext, MetalCommandQueue, MetalProgram, MetalKernel, MetalEvent, NativeCommandQueue) delegate to it in place of their old native methods.

The backend classes (Metal, MetalPlatform, MetalDevice, MetalContext, MetalCommandQueue, MetalProgram, MetalKernel, MetalEvent, NativeCommandQueue) delegate to it in place of their old native methods.

The backend classes (Metal, MetalPlatform, MetalDevice, MetalContext, MetalCommandQueue, MetalProgram, MetalKernel, MetalEvent, NativeCommandQueue) delegate to it in place of their old native methods.

The backend classes (Metal, MetalPlatform, MetalDevice, MetalContext, MetalCommandQueue, MetalProgram, MetalKernel, MetalEvent, NativeCommandQueue) delegate to it in place of their old native methods.

The backend classes (Metal, MetalPlatform, MetalDevice, MetalContext, MetalCommandQueue, MetalProgram, MetalKernel, MetalEvent, NativeCommandQueue) delegate to it in place of their old native methods.

The backend classes (Metal, MetalPlatform, MetalDevice, MetalContext, MetalCommandQueue, MetalProgram, MetalKernel, MetalEvent, NativeCommandQueue) delegate to it in place of their old native methods.

The backend classes (Metal, MetalPlatform, MetalDevice, MetalContext, MetalCommandQueue, MetalProgram, MetalKernel, MetalEvent, NativeCommandQueue) delegate to it in place of their old native methods.