Library Tasks and the Hybrid API for Native GPU Libraries

TornadoVM’s Hybrid API lets a single TaskGraph mix JIT-compiled Java kernels with calls into vendor-optimized native GPU libraries—cuBLAS, cuBLASLt, cuFFT, cuDNN, cuSPARSE, and CUTLASS—without any manual memory management or host synchronization between them. A library task is a SchedulableTask node that never passes through the JIT compiler. Instead, at execution time the TornadoVM interpreter resolves every reference argument to the raw device pointer of the associated TornadoVM buffer (past the array header), then dispatches the call through a registered TornadoLibraryProvider. Scalars are passed as boxed Java objects. Because library tasks share the same TornadoVM-managed buffers as surrounding JIT tasks, data produced on the GPU by a Java kernel is consumed by cuBLAS with no copies, and vice versa.

Library tasks require the CUDA backend (make BACKEND=cuda) and an NVIDIA GPU. On OpenCL, SPIR-V, and Metal backends the library task is silently reported as UNSUPPORTED.

`TaskGraph.libraryTask`

The entry point for library tasks is the libraryTask method on TaskGraph, a sibling of the standard .task(...) method. It takes a string task identifier, a factory method reference that returns a LibraryTaskDescriptor, and the arguments to forward to that factory.

TaskGraph libraryTask(String id, LibraryTask factory, A1 arg1, ... AN argN)

There are 20 typed overloads covering 1 through 20 arguments. All are otherwise identical—pick the overload that matches your argument count.

import uk.ac.manchester.tornado.api.TaskGraph;
import uk.ac.manchester.tornado.api.enums.DataTransferMode;
import uk.ac.manchester.tornado.api.types.arrays.FloatArray;
import uk.ac.manchester.tornado.cublas.CuBlas;
import uk.ac.manchester.tornado.cublas.enums.CuBlasOperation;

TaskGraph graph = new TaskGraph("gemm")
    .transferToDevice(DataTransferMode.EVERY_EXECUTION, matA, matB)
    .libraryTask("sgemm", CuBlas::cublasSgemm,
            CuBlasOperation.CUBLAS_OP_N.operation(),
            CuBlasOperation.CUBLAS_OP_N.operation(),
            size, size, size,
            1.0f, matB, size, matA, size,
            0.0f, output, size)
    .transferToHost(DataTransferMode.EVERY_EXECUTION, output);

cuBLAS is column-major. For row-major FloatArray inputs, either pass the transpose operation (CUBLAS_OP_T) or swap the A/B operands as shown above so that C_cm = B_cm · A_cm computes the row-major C = A · B.

`LibraryTaskDescriptor`

LibraryTaskDescriptor is a simple builder that factory methods (e.g., CuBlas::cublasSgemm) construct and return to the TornadoVM runtime. You interact with it directly only when writing a custom provider. Package: uk.ac.manchester.tornado.api.common

withLibrary(String libraryName)

LibraryTaskDescriptor

Sets the provider ID string (e.g. "nvidia/cublas"). The runtime matches this string against the registered TornadoLibraryProvider.libraryName() values.

withFunction(String functionName)

LibraryTaskDescriptor

Names the specific entry point within the library (e.g. "cublasSgemm"). Passed verbatim to TornadoLibraryProvider.dispatch(String, LibraryInvocation).

withParameters(Object[] parameters)

LibraryTaskDescriptor

Sets the argument array. Reference arguments (off-heap arrays) are resolved to device pointers by the interpreter; scalars are passed as boxed values.

withAccess(Access[] access)

LibraryTaskDescriptor

Declares the data-flow access mode for each argument using Access.READ_ONLY, Access.WRITE_ONLY, or Access.READ_WRITE. The runtime uses these to schedule transfers and determine dependencies.

withTuning(Object tuning)

LibraryTaskDescriptor

Attaches library-specific tuning options (algorithm selection, workspace size, math mode, etc.) as an opaque object. Ignored by the runtime; interpreted by the matching provider at dispatch time.

Getters

getLibraryName()

String

Returns the provider ID.

getFunctionName()

String

Returns the function entry-point name.

getParameters()

Object[]

Returns the argument array.

getAccess()

Access[]

Returns the per-argument access declarations.

getTuning()

Object

Returns the tuning object, or null.

`Access` enum

uk.ac.manchester.tornado.api.common.Access controls how the runtime transfers data for each buffer argument.

Value	Bit mask	Meaning
`NONE`	`0b00`	No data movement (internal use).
`READ_ONLY`	`0b01`	Buffer is only read by the task; transferred to device but never back.
`WRITE_ONLY`	`0b10`	Buffer is only written; allocated on device, transferred to host after the task.
`READ_WRITE`	`0b11`	Buffer is both read and written; transferred both directions.

When beta != 0 in cuBLAS/CUTLASS GEMM calls, the output matrix is read before being updated. In those cases the runtime automatically marks the output as READ_WRITE and requires it to be included in transferToDevice.

Provider catalog

TornadoVM ships six built-in library providers for NVIDIA GPUs.

Provider ID: nvidia/cublas · Module: tornado-cublasDense linear algebra in FP32 and FP16, using NVIDIA BLAS. All factories are in uk.ac.manchester.tornado.cublas.CuBlas.

Factory	Operation
`cublasSgemv(op, m, n, α, A, lda, x, incx, β, y, incy)`	`y = α·op(A)·x + β·y`
`cublasSgemm(opA, opB, m, n, k, α, A, lda, B, ldb, β, C, ldc)`	`C = α·op(A)·op(B) + β·C`
`cublasSgemmTF32(...)`	SGEMM using TF32 tensor cores
`cublasGemmExFP16(...)`	FP16 inputs, tensor-core GEMM
`cublasSgemmStridedBatched(...)`	Batched SGEMM

// SGEMM: row-major C = A·B by swapping operands for column-major cuBLAS
taskGraph.libraryTask("sgemm", CuBlas::cublasSgemm,
    CuBlasOperation.CUBLAS_OP_N.operation(),
    CuBlasOperation.CUBLAS_OP_N.operation(),
    size, size, size,
    1.0f, matrixB, size, matrixA, size,
    0.0f, output, size);

Provider ID: nvidia/cublaslt · Module: tornado-cublasFP16/FP32/FP8 GEMM with fused bias and activation epilogues. Factories are in uk.ac.manchester.tornado.cublas.CuBlasLt.

Factory	Epilogue
`ltMatmulFP32` / `ltMatmulFP16`	None
`ltMatmulFP8`	None (E4M3 inputs → FP16 output, `ld` must be multiple of 16 B)
`ltMatmulBiasFP16`	`BIAS`
`ltMatmulGeluBiasFP16`	`GELU_BIAS` (tanh approximation)

taskGraph.libraryTask("mlp", CuBlasLt::ltMatmulGeluBiasFP16,
    transa, transb, m, n, k,
    alpha, aFP16, lda, bFP16, ldb,
    beta, dFP16, ldd, bias);

Fused epilogues deliver a 1.1–1.7× speedup over separate GEMM + bias + activation passes, with the largest gains in the launch-bound regime.

Provider ID: nvidia/cufft · Module: tornado-cufftComplex and real FFT transforms. Plans are cached per (n, batch). Factories are in uk.ac.manchester.tornado.cufft.CuFft.

Factory	Transform
`cufftForwardC2C` / `cufftInverseC2C(input, output, n, batch)`	1D complex→complex
`cufftForwardR2C` / `cufftInverseC2R(input, output, n, batch)`	Real ↔ complex (Hermitian, `n ↔ n/2+1`)
`cufftForwardZ2Z` / `cufftInverseZ2Z(input, output, n, batch)`	FP64 complex→complex
`cufftForward2dC2C` / `cufftInverse2dC2C(input, output, nx, ny)`	2D complex→complex

// FFT → JIT low-pass filter → IFFT → JIT normalize, all on-device
new TaskGraph("filter")
    .transferToDevice(DataTransferMode.EVERY_EXECUTION, signal)
    .libraryTask("fwd",  CuFft::cufftForwardC2C, signal, freq, n, 1)
    .task("lowpass",     Filters::lowPass, freq, cutoff)
    .libraryTask("inv",  CuFft::cufftInverseC2C, freq, out, n, 1)
    .task("normalize",   Filters::scale, out, 1.0f / n)
    .transferToHost(DataTransferMode.EVERY_EXECUTION, out);

Provider ID: nvidia/cudnn · Module: tornado-cudnnDeep-learning primitives and fused scaled-dot-product attention. Factories are in uk.ac.manchester.tornado.cudnn.CuDnn.

Factory	Operation
`cudnnSoftmax(input, output, rows, cols)`	Per-row numerically stable softmax
`cudnnRelu` / `cudnnSigmoid` / `cudnnTanh(input, output, size)`	Activation functions
`cudnnMaxPool2d(input, output, n, c, h, w, window, stride)`	2D max pooling
`cudnnConv2d(input, filter, output, n, c, h, w, k, r, s, pad, stride)`	2D convolution
`sdpaForward(q, k, v, o, b, h, sQ, sKv, d, scale, causal)`	Fused FP16 scaled-dot-product attention (BHSD layout)

// Conv → JIT bias-add → ReLU → MaxPool, one captured graph
new TaskGraph("cnn")
    .transferToDevice(DataTransferMode.EVERY_EXECUTION, input, filter)
    .libraryTask("conv", CuDnn::cudnnConv2d,
            input, filter, convOut, n, c, h, w, k, r, s, pad, stride)
    .task("bias",        Layers::addBias, convOut, bias)
    .libraryTask("relu", CuDnn::cudnnRelu, convOut, reluOut, n * k * outH * outW)
    .libraryTask("pool", CuDnn::cudnnMaxPool2d,
            reluOut, pooled, n, k, outH, outW, 2, 2)
    .transferToHost(DataTransferMode.EVERY_EXECUTION, pooled);

Provider ID: nvidia/cutlass · Module: tornado-cutlassOpen-template GEMM in FP32 (SIMT) and FP16 (tensor-core) with fused epilogues. Row-major — no operand-swap needed. Factories are in uk.ac.manchester.tornado.cutlass.Cutlass.

Factory	Operation
`cutlassSgemm(m, n, k, α, A, B, β, C)`	FP32 SIMT GEMM
`cutlassHgemm(m, n, k, α, A, B, β, D)`	FP16 tensor-core GEMM, FP32 accumulate
`cutlassGemmBiasRelu(m, n, k, A, B, bias, D)`	Fused `relu(A·B + bias)`
`cutlassGemmBiasGelu(m, n, k, A, B, bias, D)`	Fused `gelu(A·B + bias)`

// Two-layer FFN: relu(x·W1+b1) then gelu(h·W2+b2)
new TaskGraph("ffn")
    .transferToDevice(DataTransferMode.EVERY_EXECUTION, x, w1, b1, w2, b2)
    .libraryTask("l1", Cutlass::cutlassGemmBiasRelu, m, n1, k1, x, w1, b1, h)
    .libraryTask("l2", Cutlass::cutlassGemmBiasGelu, m, n2, n1, h, w2, b2, y)
    .transferToHost(DataTransferMode.EVERY_EXECUTION, y);

TornadoVM device pointers are 8-byte aligned. FP16 CUTLASS kernels require k and n to be multiples of 4. FP32 SIMT has no alignment constraint.

Provider ID: nvidia/cusparse · Module: tornado-cusparseSparse matrix–vector and sparse matrix–dense matrix products over CSR format (32-bit, zero-based indices), alpha = 1.0, beta = 0.0. Factories are in uk.ac.manchester.tornado.cusparse.Cusparse.

Factory	Operation
`cusparseSpMV(rows, cols, nnz, csrRowOffsets, csrColInd, csrValues, x, y)`	`y = A·x`
`cusparseSpMM(rows, k, n, nnz, csrRowOffsets, csrColInd, csrValues, B, C)`	`C = A·B` (dense B/C, row-major)

new TaskGraph("spmv")
    .transferToDevice(DataTransferMode.EVERY_EXECUTION,
            rowOffsets, colIndices, values, x)
    .libraryTask("spmv", Cusparse::cusparseSpMV,
            rows, cols, nnz, rowOffsets, colIndices, values, x, y)
    .transferToHost(DataTransferMode.EVERY_EXECUTION, y);

The provider pre-allocates an 8 MiB workspace in prepare(). Problems requiring more workspace cannot grow inside a CUDA graph capture region—reduce the problem size or run without CUDA graphs first.

`TornadoLibraryProvider` SPI

The TornadoLibraryProvider interface in uk.ac.manchester.tornado.runtime.library.spi is the extension point for adding new native library bindings. Providers are discovered via java.util.ServiceLoader.

libraryName()

String

Returns the unique provider identifier (e.g. "nvidia/cublas"). Matched against LibraryTaskDescriptor.getLibraryName().

canHandle(TornadoXPUDevice device)

boolean

Returns true when the provider can execute on the given device. CUDA-only providers test instanceof TornadoNativeStreamSupport.

createContext(TornadoXPUDevice device, long executionPlanId)

LibraryContext

Creates a native execution context for the given device and execution plan (e.g. creates a cuBLAS handle and binds it to the device’s CUDA stream). Called once per (library, device, plan) tuple.

prepare(LibraryTaskDescriptor descriptor, LibraryContext context)

default void

Optional hook invoked before CUDA graph capture starts. Providers that allocate per-shape plans or workspaces (cuFFT, cuDNN, CUTLASS, cuSPARSE) create them here so dispatch is capture-safe. Must be idempotent—typically a plan-cache lookup.

dispatch(String functionName, LibraryInvocation invocation)

void

Executes the named function. Arguments are accessed via invocation.getArg(i) (scalars) and invocation.getDevicePointer(i) (device pointer for reference args).

destroyContext(LibraryContext context)

void

Releases all native resources (plans, workspace, library handle) held by the context.

`LibraryInvocation`

A single dispatch call, with all arguments already resolved by the TornadoVM interpreter.

getNumArgs()

int

Total number of arguments.

getArg(int index)

Object

The Java argument at index—a boxed scalar or the host-side array object for reference arguments.

getDevicePointer(int index)

long

Raw device pointer for a reference argument at index, pointing past the TornadoVM array header to the first data element.

isReference(int index)

boolean

Returns true if argument index is a buffer reference (has a device pointer) rather than a scalar.

getContext()

LibraryContext

The per-(device, plan) context created by createContext.

getTuning()

Object

Library-specific tuning object from LibraryTaskDescriptor.withTuning(...), or null.

isCapturing()

boolean

Returns true when the call is being recorded into a CUDA graph. Device allocations and host synchronization are not capture-safe; reject them here and do the sizing work in prepare() instead.

Complete cuBLAS SGEMM example

The following self-contained example adds a cuBLAS SGEMM library task to a TaskGraph between two JIT preprocessing/postprocessing steps, then executes it with a TornadoExecutionPlan.

import uk.ac.manchester.tornado.api.*;
import uk.ac.manchester.tornado.api.enums.DataTransferMode;
import uk.ac.manchester.tornado.api.annotations.Parallel;
import uk.ac.manchester.tornado.api.types.arrays.FloatArray;
import uk.ac.manchester.tornado.cublas.CuBlas;
import uk.ac.manchester.tornado.cublas.enums.CuBlasOperation;

public class HybridGemmExample {

    // JIT kernel: scale matrix elements before GEMM
    public static void scaleMatrix(FloatArray m, float factor) {
        for (@Parallel int i = 0; i < m.getSize(); i++) {
            m.set(i, m.get(i) * factor);
        }
    }

    // JIT kernel: apply ReLU after GEMM
    public static void relu(FloatArray out) {
        for (@Parallel int i = 0; i < out.getSize(); i++) {
            float v = out.get(i);
            out.set(i, v < 0 ? 0 : v);
        }
    }

    public static void main(String[] args) throws Exception {
        final int M = 512, N = 512, K = 512;

        FloatArray matA   = new FloatArray(M * K);
        FloatArray matB   = new FloatArray(K * N);
        FloatArray output = new FloatArray(M * N);

        // Host initialization
        for (int i = 0; i < M * K; i++) matA.set(i, (float) Math.random());
        for (int i = 0; i < K * N; i++) matB.set(i, (float) Math.random());

        TaskGraph graph = new TaskGraph("hybrid_gemm")
            // Transfer inputs to device once
            .transferToDevice(DataTransferMode.FIRST_EXECUTION, matA, matB)
            // JIT: scale A by 0.5 on the GPU
            .task("preprocess", HybridGemmExample::scaleMatrix, matA, 0.5f)
            // Native cuBLAS SGEMM: output = 1.0 * matB * matA + 0.0 * output
            // (operands swapped: row-major A·B == column-major B·A)
            .libraryTask("sgemm", CuBlas::cublasSgemm,
                    CuBlasOperation.CUBLAS_OP_N.operation(),
                    CuBlasOperation.CUBLAS_OP_N.operation(),
                    N, M, K,
                    1.0f, matB, N, matA, K,
                    0.0f, output, N)
            // JIT: apply ReLU on the GPU
            .task("relu", HybridGemmExample::relu, output)
            // Transfer result back
            .transferToHost(DataTransferMode.EVERY_EXECUTION, output);

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

        System.out.printf("output[0] = %.6f%n", output.get(0));
    }
}

Error handling

When a library task runs on a backend that has no matching provider (OpenCL, SPIR-V, Metal), the TornadoVM runtime reports the task as UNSUPPORTED in profiling output and silently skips the dispatch. Production code should either:

Confirm the CUDA backend is active via TornadoRuntimeProvider.getTornadoRuntime().getBackendType(0), or
Guard the libraryTask call path with TornadoVMBackendType.CUDA checks and provide a CPU fallback using the standard .task(...) API.

All six providers are captured into CUDA Graphs automatically on the second execution of a plan. CUDA Graphs eliminate per-launch kernel submission overhead and are particularly beneficial for chains of small library tasks (e.g. a multi-layer transformer block).

Core API

Data Types

Library Tasks

Library Tasks and the Hybrid API for Native GPU Libraries

`TaskGraph.libraryTask`

`LibraryTaskDescriptor`

Getters

`Access` enum

Provider catalog

`TornadoLibraryProvider` SPI

`LibraryInvocation`

Complete cuBLAS SGEMM example

Error handling

Build docs developers (and LLMs) love

Core API

Data Types

Library Tasks

Documentation Index

​TaskGraph.libraryTask

​LibraryTaskDescriptor

​Getters

​Access enum

​Provider catalog

​TornadoLibraryProvider SPI

​LibraryInvocation

​Complete cuBLAS SGEMM example

​Error handling

Build docs developers (and LLMs) love

`TaskGraph.libraryTask`

`LibraryTaskDescriptor`

Getters

`Access` enum

Provider catalog

`TornadoLibraryProvider` SPI

`LibraryInvocation`

Complete cuBLAS SGEMM example

Error handling