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

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.

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;

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

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

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.

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 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 cudnnSoftmax(input, output, rows, cols), Operation per-row numerically-stable softmax

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().

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).