Supported CUDA Features: a guide to the CUDA-C backend, KernelContext and the hybrid library API

[mikepapadim]

This is a map of everything the CUDA-C backend exposes today, written against develop @ba28f1526. It covers the three ways you can reach NVIDIA hardware from Java in TornadoVM — the JIT compiler, KernelContext, and the hybrid library API — plus the CUDA-specific execution-plan switches and how to profile any of it.

Every snippet below is taken from code in the repository, not written for the occasion; each one links to its source so you can run it.

Everything under "Available today" is on develop with unit tests behind it. Items still in review are listed separately at the end, so nothing here should surprise you when you try it.

---

1. Getting started

make BACKEND=cuda
tornado --devices          # confirm the CUDA driver enumerates your GPU

Three levels of control, from least to most explicit:

level	what you write	when to use it
Loop parallelism	@Parallel / @Reduce on plain Java loops	portable kernels, no GPU knowledge needed
KernelContext	thread indices, shared memory, barriers, warp and tensor-core intrinsics	you want the CUDA programming model without leaving Java
Hybrid library tasks	libraryTask(...) calling cuBLAS, cuDNN, cuFFT, cuSPARSE, CUTLASS	a vendor kernel already solves your problem better than any JIT can

All three compose inside one TaskGraph, sharing the same device buffers — no round trip to the host between them.

The shape every example below plugs into:

KernelContext context = new KernelContext();

WorkerGrid worker = new WorkerGrid1D(size);
worker.setLocalWork(256, 1, 1);                       // block size == CUDA blockDim
GridScheduler grid = new GridScheduler("s0.t0", worker);

TaskGraph taskGraph = new TaskGraph("s0")
        .transferToDevice(DataTransferMode.EVERY_EXECUTION, input)
        .task("t0", MyKernels::myKernel, context, input, output)
        .transferToHost(DataTransferMode.EVERY_EXECUTION, output);

try (TornadoExecutionPlan plan = new TornadoExecutionPlan(taskGraph.snapshot())) {
    plan.withGridScheduler(grid).execute();
}

---

2. KernelContext — the CUDA programming model in Java

2.1 Thread indexing and synchronisation

globalIdx / globalIdy / globalIdz, localIdx / localIdy / localIdz, groupIdx / groupIdy / groupIdz, the grid dimensions globalGroupSizeX/Y/Z and localGroupSizeX/Y/Z, plus localBarrier() (__syncthreads()) and globalBarrier().

2.2 Shared (local) memory

float[] tile = context.allocateFloatLocalArray(256);

Available for int, byte, long, float, double, HalfFloat and Half2.

2.3 Atomics

A histogram is the whole story in two lines — TestHistogram.java:

private static void histogramKernel(KernelContext ctx, IntArray input, IntArray globalHist) {
    int bin = input.get(ctx.globalIdx) % NUM_BINS;
    ctx.atomicAdd(globalHist, bin, 1);
}

API	lowering
atomicAdd(IntArray / LongArray / FloatArray / DoubleArray / int[], index, value)	native atomicAdd (double needs compute ≥ 6.0)
atomicCAS(int[], index, expected, value)	atomicCAS
atomicExchange / atomicMin / atomicMax on int[]	atomicExch / atomicMin / atomicMax

2.4 Warp-level primitives

simdSum(float), simdShuffleDown(float, delta) and simdBroadcastFirst(float) give you warp shuffles without shared memory. A block-wide suffix scan, warp-level fold plus one shared-memory carry pass — TestWarpShuffleScan.java:

private static void suffixScan(KernelContext ctx, FloatArray data) {
    int tid    = ctx.localIdx;
    int lane   = tid % WARP_SIZE;
    int warpId = tid / WARP_SIZE;

    float val = data.get(ctx.globalIdx);
    for (int stride = 1; stride < WARP_SIZE; stride <<= 1) {
        float neighbour = ctx.simdShuffleDown(val, stride);
        if (lane + stride < WARP_SIZE) {
            val += neighbour;
        }
    }
    // val is now the suffix sum from this lane to the end of its warp.

    float[] warpTotals = ctx.allocateFloatLocalArray(NUM_WARPS);
    if (lane == 0) {
        warpTotals[warpId] = val;                 // lane 0 holds the whole-warp sum
    }
    ctx.localBarrier();

    float carry = 0.0f;
    for (int w = warpId + 1; w < NUM_WARPS; w++) {
        carry += warpTotals[w];
    }
    data.set(ctx.globalIdx, val + carry);
}

Note simdShuffleDown only pulls from a higher lane (shfl.sync.down) and has no int overload — which is why that kernel computes a suffix scan on float rather than a conventional prefix scan.

2.5 Tensor cores (mma.sync)

MMA fragments are first-class, in shapes M16N8K16 and M16N8K32:

precision	entry points
FP16 → FP32	mmaFragment, mmaLoadA, mmaLoadB, mmaLoadBSwizzled, mma, mmaStore
BF16	mmaBF16
INT8	mmaFragmentInt, mmaLoadAInt8, mmaLoadBInt8, mmaInt8, mmaStoreInt
FP8 (E4M3 / E5M2)	mmaLoadAFP8, mmaLoadBFP8, mmaFP8E4M3, mmaFP8E5M2

One warp computing a 16×16 output tile, two M16N8K16 accumulators side by side — condensed from TestMatrixMultiplicationMMA.java:

public static void gemmMMA(KernelContext ctx,
                           HalfFloatArray a, HalfFloatArray b, FloatArray c,
                           int dimM, int dimN, int dimK) {

    int warpId = ctx.groupIdx;
    int lane   = ctx.localIdx;

    int numTilesN = dimN / WMMA_N;
    int tileRow   = (warpId / numTilesN) * WMMA_M;
    int tileCol   = (warpId % numTilesN) * WMMA_N;

    // Shared tiles hold two FP16 values packed per int (a b16 pair).
    int[] aTile  = ctx.allocateIntLocalArray(WMMA_M * WMMA_K / 2);
    int[] bTile0 = ctx.allocateIntLocalArray(WMMA_K * WMMA_N / 2);
    int[] bTile1 = ctx.allocateIntLocalArray(WMMA_K * WMMA_N / 2);

    float[] fragC0 = ctx.mmaFragment(0.0f);
    float[] fragC1 = ctx.mmaFragment(0.0f);

    for (int kBase = 0; kBase < dimK; kBase += WMMA_K) {

        // Cooperative packed load of A: each lane packs two adjacent FP16 into one int.
        for (int idx = lane; idx < (WMMA_M * WMMA_K) / 2; idx += WARP_SIZE) {
            int elemBase   = idx * 2;
            int r          = elemBase / WMMA_K;
            int kk         = elemBase % WMMA_K;
            int globalBase = (tileRow + r) * dimK + kBase + kk;
            int lo = a.get(globalBase).getHalfFloatValue()     & 0xFFFF;
            int hi = a.get(globalBase + 1).getHalfFloatValue() & 0xFFFF;
            aTile[r * (WMMA_K / 2) + kk / 2] = lo | (hi << 16);
        }
        // ... same idea for bTile0 / bTile1 (two 8-wide B panels) ...
        ctx.localBarrier();

        HalfFloat[] fragA  = ctx.mmaLoadA(aTile,  WMMA_K);
        HalfFloat[] fragB0 = ctx.mmaLoadB(bTile0, WMMA_K);
        fragC0 = ctx.mma(fragA, fragB0, fragC0, MMAShape.M16N8K16);

        HalfFloat[] fragB1 = ctx.mmaLoadB(bTile1, WMMA_K);
        fragC1 = ctx.mma(fragA, fragB1, fragC1, MMAShape.M16N8K16);

        ctx.localBarrier();
    }

    ctx.mmaStore(fragC0, c, tileRow, tileCol,     dimN);
    ctx.mmaStore(fragC1, c, tileRow, tileCol + 8, dimN);
}

Supporting cast: bank-conflict-free shared access via swizzleLoadFp16Stride16/32, swizzleStoreFp16Stride16/32, swizzleLoadInt8, swizzleStoreInt8; and Matrix8x8Float helpers simdgroupMatrixZero/Load/MultiplyAccumulate/Store for a portable 8×8 tile API.

More end-to-end examples: MatrixMultiplicationMMA.java, LowPrecisionGemmBenchmark.java, FP8GemmStage.java.

2.6 Asynchronous global → shared copies (cp.async)

The simple form — stage a tile, commit, wait, publish:

public static void stageBytesKernel(KernelContext ctx, ByteArray in, IntArray out) {
    int[] tile = ctx.allocateIntLocalArray(SLOTS);
    int lane = ctx.localIdx;

    ctx.asyncCopyToLocal(tile, lane, in, lane * 4);
    ctx.asyncCopyCommit();
    ctx.asyncCopyWaitGroup(0);
    ctx.localBarrier();

    out.set(ctx.globalIdx, tile[lane]);
}

And the reason cp.async exists — a double-buffered pipeline that stages chunk k+1 while consuming chunk k, from TestAsyncCopyToLocal.java:

public static void pipelinedStageKernel(KernelContext ctx, ByteArray in, IntArray out) {
    int[] tile = ctx.allocateIntLocalArray(2 * SLOTS);   // two buffers, alternating
    int lane = ctx.localIdx;

    // Prologue: get chunk 0 in flight before the loop.
    ctx.asyncCopyToLocal(tile, lane, in, lane * 4);
    ctx.asyncCopyCommit();

    int acc = 0;
    for (int k = 0; k < CHUNKS - 1; k++) {
        int next = k + 1;
        ctx.asyncCopyToLocal(tile, (next & 1) * SLOTS + lane, in, next * BYTES + lane * 4);
        ctx.asyncCopyCommit();
        ctx.asyncCopyWaitGroup(1);   // chunk k has landed; chunk k+1 stays in flight
        ctx.localBarrier();

        acc += tile[(k & 1) * SLOTS + lane];
        ctx.localBarrier();          // all reads done before this buffer is restaged
    }

    // Epilogue: the last chunk has no successor to overlap with.
    ctx.asyncCopyWaitGroup(0);
    ctx.localBarrier();
    acc += tile[((CHUNKS - 1) & 1) * SLOTS + lane];

    out.set(ctx.globalIdx, acc);
}

Source overloads: ByteArray, HalfFloatArray, FP8Array.

2.7 Low-precision conversion, dot products, printf

float16ToFloat, bf16FromFloat / bf16ToFloat, e4m3ToFloat / e5m2ToFloat, dp4a / dp4a_packed (4-way INT8 dot-product-accumulate), and printf(...) lowered to device-side printf.

---

3. Data types

Off-heap, MemorySegment-backed native arrays — no copies, no boxing:

ByteArray, CharArray, ShortArray, IntArray, LongArray, FloatArray, DoubleArray, HalfFloatArray (FP16), BFloat16Array (BF16), FP8Array (E4M3 / E5M2), Int8Array.

Plus vector types (Float2/3/4/8/16, Int*, Double*, Byte3/4, Half2) and matrix types (Matrix2DFloat, Matrix3DFloat, Matrix4x4Float, Matrix8x8Float, and int/double/short variants).

---

4. Hybrid library tasks — NVIDIA libraries inside a TaskGraph

A libraryTask runs a native vendor kernel on the buffers already resident on the device, in the middle of a normal task graph — from TestCuBlas.java:

TaskGraph taskGraph = new TaskGraph("s0")
        .transferToDevice(DataTransferMode.EVERY_EXECUTION, matrix, vector)
        .libraryTask("sgemv", CuBlas::cublasSgemv,
                     CuBlasOperation.CUBLAS_OP_T.operation(), SIZE, SIZE,
                     1.0f, matrix, SIZE, vector, 1, 0.0f, output, 1)
        .task("activation", MyKernels::relu, context, output)   // JIT kernel, same buffer
        .transferToHost(DataTransferMode.EVERY_EXECUTION, output);

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

Available providers on develop:

library	operations
cuBLAS / cuBLASLt (12)	cublasSgemv, cublasSgemm, cublasSgemmTF32, cublasSgemmStridedBatched, cublasGemmExFP16, cublasGemmExFP16FP32, cublasGemmExBF16; cuBLASLt: ltMatmulFP32, ltMatmulFP16, ltMatmulFP8, ltMatmulBiasFP16, ltMatmulGeluBiasFP16
cuDNN (8)	cudnnSoftmax, cudnnRelu, cudnnSigmoid, cudnnTanh, cudnnMaxPool2d, cudnnConv2d, sdpaForward (fused scaled-dot-product attention)
cuFFT (8)	cufftForwardC2C / cufftInverseC2C, cufftForwardR2C / cufftInverseC2R, cufftForwardZ2Z / cufftInverseZ2Z (double), cufftForward2dC2C / cufftInverse2dC2C
cuSPARSE (2)	cusparseSpMV, cusparseSpMM (CSR)
CUTLASS (11)	cutlassSgemm, cutlassHgemm, cutlassBgemm, cutlassHgemmBatched, and fused epilogues cutlassGemmBiasRelu / Gelu / Silu / Sigmoid / Tanh / HardSwish

Library tasks are automatically wrapped in an NVTX range named <library>/<function>, so they show up as labelled spans on an Nsight timeline alongside the JIT kernels.

Tests: unittests.cublas.TestCuBlas, TestCuBlasLt, unittests.cudnn.TestCuDnn, unittests.cufft.TestCuFft, unittests.cusparse.TestCusparse, unittests.cutlass.TestCutlass.

---

5. CUDA-specific execution-plan options

try (TornadoExecutionPlan plan = new TornadoExecutionPlan(taskGraph.snapshot())) {
    plan.withGridScheduler(grid)
        .withCUDAGraph()              // capture once, replay afterwards
        .withIntraPlanConcurrency()   // independent ops on separate streams
        .withWarmUpIterations(10)     // move JIT cost out of the measured region
        .execute();
}

option	what it does
withCUDAGraph()	captures the task graph into a CUDA Graph on first execution and replays it, collapsing per-launch dispatch
withIntraPlanConcurrency()	issues independent operations of one plan on several streams instead of one
withStagedTransfers()	pinned-memory staging ring for large host↔device copies (tornado.staged.* tunes chunk size, ring depth, threshold, fill threads)
withConcurrentDevices()	runs task graphs of a plan on several devices at once
withBatch("512MB")	splits a workload that does not fit in device memory
withMemoryLimit("2GB")	caps device memory for the plan
withPreCompilation(), withWarmUpIterations(n), withWarmUpTime(ms)	move JIT cost out of the measured region
withProfiler(ProfilerMode)	per-task kernel / copy / compile timings
withPrintKernel(), withThreadInfo()	dump the generated CUDA C and the launch geometry

Useful properties: -Dtornado.cuda.priority, -Dtornado.max.events, -Dtornado.reuse.device.buffers, -Dtornado.eventpool.size.

Compiled kernels are cached as cubins on disk (#1008), so repeat runs of the same application skip NVRTC entirely.

---

6. Profiling

6.1 See what was generated

tornado --printKernel     ...   # the generated CUDA C
tornado --printBytecodes  ...   # the TornadoVM bytecode driving the device
tornado --threadInfo      ...   # grid / block geometry per task

6.2 TornadoVM's own profiler

tornado --enableProfiler console ...
tornado --enableProfiler silent --dumpProfiler out.json ...

Reports per task: TASK_COMPILE_GRAAL_TIME, TASK_COMPILE_DRIVER_TIME (NVRTC), TASK_CODE_GENERATION_TIME, TASK_KERNEL_TIME (device events), COPY_IN_TIME / COPY_OUT_TIME, TOTAL_DISPATCH_KERNEL_TIME, allocation and transfer byte counts, plus power via NVML.

Programmatically: plan.withProfiler(ProfilerMode.SILENT) then executionResult.getProfilerResult().

Caveat worth knowing. The profiler reads each event's timestamp inline, which serialises host and device. On short kernels that both slows the run down several times over and inflates its own transfer numbers — measured at 3.3× on a 512-element saxpy, reporting copyInAvg = 12.9 µs for a copy Nsight measures at 0.65 µs. #1024 addresses this. For microsecond-scale work, trust Nsight over the built-in numbers.

6.3 Nsight Systems — the timeline

nsys profile -t cuda,nvtx --sample=none --cpuctxsw=none -o report \
  tornado -m tornado.benchmarks/uk.ac.manchester.tornado.benchmarks.BenchmarkRunner \
  --params="sgemm 100 1024"

nsys stats --report cuda_api_sum,cuda_gpu_kern_sum,cuda_gpu_mem_time_sum,nvtx_sum report.nsys-rep

The reports that answer the usual questions:

report	question it answers
cuda_api_sum	how many driver calls per iteration, and where host time goes
cuda_gpu_kern_sum	actual kernel duration on the device
cuda_gpu_mem_time_sum	H2D / D2H time and volume
nvtx_sum	per-operation host cost, using the ranges TornadoVM emits

TornadoVM labels its own NVTX ranges, so the timeline is readable without guesswork: transfers appear as H2D 2.0 KB / D2H 24.0 MB, kernels under their generated name, and hybrid library calls as <library>/<function> (e.g. cublas/cublasSgemm).

A worked example of reading these — GPU work was 3.31 µs of a 17.46 µs iteration, with the rest in host-side dispatch, which is what #1022–#1024 came from.

6.4 Nsight Compute — inside a kernel

ncu --set full --kernel-name regex:myKernel -o profile \
  tornado -m ... --params="..."

Kernel names in the report are the generated CUDA-C function names, so --printKernel first if you are unsure what to match. -lineinfo (see #1010) makes Nsight attribute samples back to the generated source.

6.5 Profiling the Java side

Standard JFR works, and is the right tool when you suspect the runtime rather than the kernel:

tornado --jvm="-XX:StartFlightRecording=filename=app.jfr,dumponexit=true \
               -XX:FlightRecorderOptions=stackdepth=192 \
               -XX:+UnlockDiagnosticVMOptions -XX:+DebugNonSafepoints" ...
jfr print --events jdk.ExecutionSample,jdk.NativeMethodSample,jdk.ObjectAllocationSample app.jfr

stackdepth=192 matters — the default of 64 truncates TornadoVM's execute() stack into unattributable roots. The ratio of jdk.ExecutionSample (in Java) to jdk.NativeMethodSample (inside a JNI/CUDA call) tells you immediately whether you are runtime-bound or device-bound.

---

7. Where to look in the tree

Every feature above has runnable code in the repository. Paths are relative to the repo root.

Runnable examples

topic	file
FP16 tensor-core GEMM, end to end	tornado-examples/.../examples/compute/MatrixMultiplicationMMA.java
FP16 / BF16 / INT8 / FP8 GEMM benchmark	tornado-examples/.../examples/kernelcontext/matrices/LowPrecisionGemmBenchmark.java
FP8 GEMM stage with cp.async	tornado-examples/.../examples/kernelcontext/matrices/FP8GemmStage.java
KernelContext basics, reductions, matrices	tornado-examples/.../examples/kernelcontext/
@Parallel / @Reduce style kernels	tornado-examples/.../examples/compute/, .../examples/reductions/

Tests, which double as the specification

topic	file
MMA FP16 (plus swizzled and offset variants)	tornado-unittests/.../kernelcontext/matrices/TestMatrixMultiplicationMMA.java
MMA BF16 / INT8 / FP8	TestMatrixMultiplicationMMABF16.java, TestMatrixMultiplicationMMAInt8.java, TestMatrixMultiplicationMMAFP8.java (same directory)
MMA fed by cp.async	TestMatrixMultiplicationMMACpAsync.java (same directory)
cp.async semantics, including the double-buffered pipeline	tornado-unittests/.../kernelcontext/api/TestAsyncCopyToLocal.java
atomicAdd on global memory	tornado-unittests/.../kernelcontext/reductions/TestHistogram.java
atomicCAS / atomicExchange / atomicMin / atomicMax	tornado-unittests/.../kernelcontext/api/TestAtomicRmw.java, TestGlobalAtomics.java
warp shuffles composed into a scan	tornado-unittests/.../kernelcontext/reductions/TestWarpShuffleScan.java
what KernelContext does not support per backend	tornado-unittests/.../kernelcontext/api/TestUnsupportedKernelContextOps.java
hybrid library tasks	tornado-unittests/.../{cublas,cudnn,cufft,cusparse,cutlass}/

Run any of them with:

tornado-test -V --printKernel uk.ac.manchester.tornado.unittests.kernelcontext.api.TestAsyncCopyToLocal

Backend internals, if you want to add an intrinsic

what	file
the public API surface and its javadoc	tornado-api/.../api/KernelContext.java
which intrinsics are registered, and how they lower	tornado-drivers/cuda/.../graal/compiler/plugins/CUDAGraphBuilderPlugins.java
the emitted CUDA-C statements	tornado-drivers/cuda/.../graal/lir/CUDALIRStmt.java
driver calls: launch, transfers, events, graphs	tornado-drivers/cuda-jni/src/main/cpp/source/CUDACommandQueue.cpp
NVRTC compilation and the cubin cache	tornado-drivers/cuda-jni/src/main/cpp/source/CUDAProgram.cpp

---

Corrections welcome

If something here is wrong, out of date, or you expected a CUDA feature that is not listed, reply on this thread and I will fold it in.