[perf] Host-side runtime overhead: where the time actually goes, the fixes, and what is left #1028

Context

For small and medium task graphs, and for anything that runs many short executions per iteration (LLM decode, filters, iterative solvers), TornadoVM's wall-clock is dominated by host-side per-operation cost, not by GPU work. This issue collects the profiling campaign that quantified that cost, the PRs that remove it, and what is left.

A result worth stating up front, because it redirects where optimisation effort belongs:

The Java bytecode interpreter and its data structures account for ~0.9% of wall time on a dispatch-bound graph (118 jdk.ExecutionSample vs 13,545 jdk.NativeMethodSample, and most of the Java samples are the benchmark harness's own Arrays.sort). The cost is in the per-operation CUDA dispatch sequence and runtime bookkeeping, not in interpreting bytecodes.

Method

• JFR — 1 ms sampling, stackdepth=192, -XX:+DebugNonSafepoints; samples phase-split by filtering stacks (decode vs start-up).
• Nsight Systems — nsys profile -t cuda --sample=none --cpuctxsw=none, then nsys stats --report cuda_api_sum for per-iteration API call counts, used as ground truth for what the driver actually did.
• Benchmarks — BenchmarkRunner saxpy 100000 512 (dispatch-bound: one kernel, one H2D, one blocking D2H, 2 KB moved, 1.1 µs kernel) for per-execute() medians; saxpy 2000 16777216 and nbody 500 16384 as GPU-bound controls; GPULlama3.java on Qwen1.5-MoE-A2.7B-Chat Q8_0 as the real end-to-end workload.
• Hardware — NVIDIA RTX 4090, Ubuntu 24.04, JDK 21.0.2, CUDA backend.

What the profiles found

On develop, a saxpy over 512 floats spent 17.46 µs per execute() while the GPU did 3.31 µs of work. Per iteration the runtime issued 4 stream operations, 12 event calls and 3 cuStreamSynchronize to move 2 KB and run a 1.1 µs kernel. Four separate causes, each with its own PR:

A fifth PR came out of the same campaign as a correctness prerequisite for trusting any of these numbers:

Combined measurements

Three builds, same machine, same session:

• A = develop @4d90e2e9f
• B = A + [cuda] Skip the per-launch kernel stack-frame upload when the device copy is unchanged (~2x on short task graphs) #1022 + [cuda] Skip stream synchronisation when nothing has been enqueued since the last drain (3 syncs per execution to 1) #1023 + [profiler] Harvest device timestamps after the stream drains instead of blocking on every operation (profiler-on cost -34%) #1024
• C = B + [runtime] Clear only the written prefix of each wait-list row between executions (1.57x on dependency-tracked decode) #1027

MoE decode = Qwen1.5-MoE-A2.7B-Chat Q8_0, 128 generated tokens, greedy-ish (--temperature 0.1 --seed fixed), 2 runs per cell.

Driver traffic, nsys, 20,000 iterations of saxpy 512:

Across the small-size benchmark sweep the win tracks how dispatch-dominated the graph is, and nothing regresses:

GPU-bound controls — unaffected, as intended:

PRs

All five are independent and separately reviewable. #1023 is measured both standalone and stacked on #1022 because its value only shows once the dominant per-launch cost is gone.

Testing

Every PR was checked with tornado-test --quickPass on the CUDA backend against a same-machine develop baseline, with identical per-class results (all failures pre-existing/whitelisted: Metal-only simdgroup tests, multi-backend virtualization tests, half-float vector types). #1027 was additionally run with -Dtornado.vm.deps=True across the whole suite, again with identical results. CUDA-graph capture paths (TestCuBlas, TestCuFft, TestExecutor) were exercised specifically for #1022 and #1023, since both touch state that must survive stream capture. ./mvnw checkstyle:check clean throughout.

What is left

JFR on build C (saxpy 200000 512, now 7.7 µs/execute()) shows the remaining host time is no longer bookkeeping:

• 65% CUDACommandQueue.readArrayFromDeviceOffHeap — the blocking device-to-host read-back, i.e. genuine device wait
• 14% clEnqueueNDRangeKernel
• 8% writeArrayToDevice
• Java-side samples: 2, against 63 native samples

Candidate follow-ups, roughly in value order:

1. Non-blocking read-back / deferred output materialisation. The terminal blocking D2H is now the single largest host cost on dispatch-bound graphs. Anything that lets the host enqueue the next execution before the previous read-back completes attacks it directly.
2. Remaining profiler overhead (~2.2x versus profiler-off). It is now the timing events themselves, not the deferral: cuEventRecord costs 615 ns with CU_EVENT_DEFAULT against 137 ns with CU_EVENT_DISABLE_TIMING, six per iteration, and timing events stretch the terminal cuStreamSynchronize from 3.0 µs to 14.0 µs. Reusing operation N-1's end event as operation N's start would halve it but is incompatible with deferred harvesting; an opt-in timing granularity (kernels only) is the viable route.
3. Wait-list row sizing. [runtime] Clear only the written prefix of each wait-list row between executions (1.57x on dependency-tracked decode) #1027 removes the clear cost, but a row is still allocated at tornado.max.events (32768 ints = 128 KB) on first write. Growing rows on demand would cut the allocation too.
4. tornado.eventpool.maxwaitevents vs tornado.max.events. Two separate knobs with confusingly similar names and very different cost profiles; only the former is exposed by front-ends such as llama-tornado. Worth consolidating or at least documenting.
5. Start-up, not steady state. On the MoE run, start-up dominates whole-run wall (18 s total, decode only ~3.9 s): 47% NVRTC clBuildProgram, 38% cuMemHostRegister. Separate chapter, but the larger absolute number for short runs.

Reproduction

# dispatch-bound per-execute median
tornado --jvm="-Dtornado.benchmarks.skipserial=True" \
  -m tornado.benchmarks/uk.ac.manchester.tornado.benchmarks.BenchmarkRunner \
  --params="saxpy 100000 512"

# same, profiler on
tornado --jvm="-Dtornado.benchmarks.skipserial=True -Dtornado.profiler=True" \
  -m tornado.benchmarks/uk.ac.manchester.tornado.benchmarks.BenchmarkRunner \
  --params="saxpy 100000 512"

# driver call counts
nsys profile -t cuda --sample=none --cpuctxsw=none -o saxpy512 \
  tornado --jvm="-Dtornado.benchmarks.skipserial=True" \
  -m tornado.benchmarks/uk.ac.manchester.tornado.benchmarks.BenchmarkRunner \
  --params="saxpy 20000 512"
nsys stats --report cuda_api_sum saxpy512.nsys-rep