Four ways to talk to a GPU: CUDA, ROCm, Vulkan and OpenCL on one pricing kernel

Here is the result that made this article worth writing.

On an AMD integrated GPU, running one Monte-Carlo pricing kernel, Vulkan — a graphics API, from a vendor-neutral standards body, on a driver AMD did not write — is 3.4× faster than either of AMD's own compute stacks. Not 5% faster. Three and a half times, on AMD's own silicon, against the software AMD ships specifically for this purpose.

That is not a result you can generalise into a rule, and most of this article is about why. But it is a very good reason to stop treating "put it on the GPU" as a single decision. It is at least four decisions, and the four options are not four dialects of the same language — they are four different kinds of object, with different owners, different histories, and wildly different answers to the one question a risk engine actually cares about: can I have double precision, and what does it cost me?

So: one payoff, recorded once as a tape of primitive operations, compiled into a fused forward-plus-adjoint kernel, and run on every accelerator to hand. Two matrices — one for fp32, one for fp64 — with 8-thread cpu-jit and 8-thread SIMD in every row so the GPU numbers have something honest to stand next to. Then a tour of what CUDA, ROCm, Vulkan and OpenCL each actually are, where each came from, the history that shaped them, and which one you should reach for.

A GPU is not one thing, and neither are its APIs

Before any numbers, the orientation that most "GPU computing" writing skips. These four are not comparable objects:

	what it actually is	who owns it
CUDA	a whole platform — language, compiler, runtime, libraries, profiler, debugger	NVIDIA
ROCm / HIP	a deliberate near-clone of that platform, aimed at different silicon	AMD
Vulkan	a thin, explicit driver API for the GPU as a device; compute is one thing you can ask it to do	Khronos (a committee)
OpenCL	a compute-only standard of CUDA's own vintage, designed to span every kind of processor	Khronos (a committee)

The asymmetry is the point. Two of them are specifications with many implementations; two of them are implementations that happen to have a spec. That single structural difference explains most of what follows.

Underneath all four, though, sits the same machine, and it helps to know what it is. A GPU is not "a fast CPU with more cores." It is a very wide, very dumb, very patient machine:

• Lanes run in lockstep. Threads are grouped — 32 to a warp on NVIDIA, 32 or 64 to a wavefront on AMD — and the whole group executes one instruction at a time. If half the group takes an if and half takes the else, the hardware runs both halves and masks off the lanes that shouldn't have run. Branchy code wastes the machine.
• Latency is hidden by oversubscription, not by caches. A CPU spends its transistor budget on out-of-order execution and a deep cache hierarchy so one thread never waits. A GPU spends it on registers, so that when one wavefront stalls on memory the scheduler simply runs a different one. The number of wavefronts it can keep resident is called occupancy, and occupancy is limited by how many registers each one needs.
• Registers are the scarce resource. This matters enormously here, and we will come back to it: an adjoint sweep roughly doubles the live values a kernel is carrying, which halves occupancy, which is why the "value + Greeks" column behaves so differently from the "price-only" column on the GPU rows.

A Monte-Carlo path is close to the ideal workload for that machine. Every path is independent, so there is no divergence and no synchronisation. The arithmetic is dense and transcendental-heavy. And — in this engine's design — the entire path lives in registers, so a hundred million paths touch almost no memory at all.

What is actually being measured

The thing being replayed is identical regardless of backend. A European call priced by Monte Carlo, one time step, recorded once against ADouble scalars:

try (MonteCarlo<EquityMarket> mc = MonteCarlo.of(Products.europeanCall())
        .market(EquityMarket.atmOneYear())   // S0 = K = 100, sigma 20%, r 3%, T 1y
        .steps(1)
        .greeks()                            // forward pass + one reverse sweep
        .threads(8)                          // CPU engines: eight physical cores
        .on(engine)                          // "cpu-jit" | "simd" | "vulkan" | "opencl" | "rocm" | "cuda"
        .fp32()                              // or .fp64()
        .build()) {

    var v = mc.run(100_000_000L, /*seed*/ 42L);
    // v.price(), v.greek(EquityMarket::spot), ::vol, ::rate, ::strike, ::maturity
}

MonteCarlo.of(...) runs the payoff once against a recording scalar type and flattens it into a tape — here 26 nodes. Each backend then turns that tape into something its hardware can run, once, at build time:

backend	the tape becomes
cpu	nothing — an interpreter walks the ops[] / argA[] / argB[] arrays node by node, per path
cpu-jit	generated straight-line Java bytecode, which HotSpot then JITs
simd	a JDK Vector API plan over Float/Double lanes
vulkan	GLSL → SPIR-V compute shader
opencl	OpenCL C — a mechanical dialect rewrite of the CUDA-C source
rocm	CUDA C, compiled by HIPRTC at runtime
cuda	the same CUDA C string, compiled by NVRTC at runtime

Every GPU kernel here has the same shape: one scenario per thread, a fully unrolled forward sweep followed by the adjoint sweep, both entirely in registers, with the per-scenario random stream generated in-thread by a counter-based Philox draw. A replay touches no memory beyond a handful of input scalars and one per-workgroup reduction at the end. That structure is why the reverse sweep costs so little on top of the price, and it holds on every backend below.

Vulkan is the one that needs its own generator, because GLSL is a different language with a different set of holes. Two are worth knowing about: core GLSL has no 64-bit integer type, so the Philox counter's 64-bit path arithmetic is synthesised from pairs of uints with an explicit carry; and the per-invocation accumulators are float, carried with a Kahan compensation term, because one invocation strides over thousands of scenarios and a naive fp32 running sum loses the low bits of every later term.

The rules of this benchmark

Enough GPU benchmarks are quietly rigged that the method deserves its own section.

• One fresh JVM per row. Running every configuration in one process pollutes HotSpot's profiles and lets the iGPU heat-soak across rows; an earlier single-JVM sweep understated Vulkan by ~3× and SIMD fp32 by ~2×.
• Warmed hard, then best of fifteen. Six shake-out runs, a path-count ramp, twenty more warm runs at the final size, then fifteen timed runs and the best wall-clock is reported. This is the settled rate, not the cold one. (The cold-versus-settled question gets its own treatment in the 13-billion-paths piece — on an iGPU the first call after an idle period can be 40× slower than the settled rate, and that is worth disclosing rather than warming away.)
• CPU rows pinned to 8 threads. Eight physical Zen 4 cores, SMT excluded. This is deliberately not the fastest CPU configuration available — using all sixteen hardware threads is about 25% better, and that number is published below rather than hidden. Eight is the conservative comparison point: it is what one well-configured box gives you before you reach for an accelerator.
• Same seed, same path count, same tape. Every backend replays the same Philox stream path-for-path, so values can be compared with no statistical allowance at all.
• Repeated, not sampled once. Every configuration that looked at all suspicious was re-run in several fresh JVMs. One of them turned out to be bimodal, which changed what this article says about it — see the row that would not sit still.
• The machine. Every row but cuda is an AMD Ryzen 7 8845HS — 8 cores, 16 threads, AVX-512 — with a Radeon 780M (gfx1103, RDNA3) integrated on the same die. JDK 26, 512-bit vectors (16 float / 8 double lanes), Linux 7.0. The cuda rows are a Tesla T4 (2 vCPU host); the same code generator feeds both, NVRTC on one side and HIPRTC on the other.
• Provenance. The CPU rows — cpu, cpu-jit, simd, at one and eight threads, in both precisions — were measured fresh for this article. The opencl, rocm and vulkan rows are carried over from an earlier run of the same harness a fortnight before; cuda was measured on the T4.

Throughput below is in millions of Monte-Carlo paths per second, reported twice: price-only, and value plus all five first-order Greeks from one reverse sweep.

The fp32 matrix

This is where accelerators earn the name.

backend / config	price-only	value + 5 Greeks	Greek cost
cpu-jit — generated bytecode, 8T	171	97	1.8×
simd — JDK Vector API, 8T	545	150–315¹	1.7–3.6×
opencl — Radeon 780M iGPU	~9 000	3 560	~2.5
rocm / HIP — Radeon 780M iGPU	~9 000	3 590	~2.5×
vulkan — Radeon 780M iGPU	15 640	12 180	1.3×
cuda — Tesla T4	~38 000	8 600	~4.4×

¹ Not a rounding range. The SIMD adjoint row is bimodal across fresh JVMs — it has its own section below.

The T4 cuda row is a discrete card on a different host — compare it to the other GPU rows only. On this one-step kernel it trails the on-die 780M, and that is not a fluke (I re-ran the Vulkan row on the laptop to be sure): the tape is 26 nodes, the kernel is over in microseconds, and a discrete GPU's per-launch cost — PCIe, driver — dominates where an integrated GPU has almost none. Its price-only rate is still enormous (~38 000), but the adjoint sweep costs it ~4.4× a price pass against Vulkan's 1.3× — register pressure cuts T4 occupancy harder. Give it real work and the ranking flips: on a 252-fixing Asian tape the T4 runs ~2× the 780M's fp32 rate.

Four things to read out of that.

Vulkan is 3.4× either of AMD's own compute stacks, on AMD's own GPU. On the adjoint sweep — the number that matters for a risk vector — Vulkan does 12.2 billion paths a second against ROCm's and OpenCL's ~3.6 billion. The vendor-neutral API, running on Mesa's open-source RADV driver, beats the vendor's own HIP and OpenCL runtimes on this kernel by a wide margin. This is an integrated-GPU, consumer-driver result and it would not necessarily hold on an Instinct card with a tuned ROCm build. But on the hardware most people actually have in front of them, it is the pattern — and it is a pattern people outside finance keep rediscovering too. In the local-LLM world, llama.cpp's Vulkan backend regularly trades wins with ROCm on consumer AMD parts, sometimes decisively, depending on the phase of the workload.

Why would the committee standard win? Three reasons that have nothing to do with API design and everything to do with where the engineering went. First, the Vulkan compute path is thin: you hand the driver a SPIR-V module and a dispatch, and there is almost nothing between your kernel and the shader cores — where ROCm and OpenCL both reach the device through the heavier HSA queue machinery in amdkfd. Second, the shader compiler. RADV's back-end, ACO, was written and is maintained largely by Valve, because every game on the Steam Deck depends on it; it is one of the most exercised compilers in the Linux graphics stack. Third, and bluntly: on a gfx1103 APU, RADV's compute path runs a large fraction of the games on Steam, and ROCm runs approximately nothing. Testing volume is a performance feature.

The GPU price-only numbers are noisy; the adjoint numbers are not. The iGPU ramps its clock hard for the lean price-only kernel — best-of-fifteen caught it anywhere from 6 to 9 billion paths a second for OpenCL and ROCm across repeats — while the register-heavier adjoint kernel is occupancy-bound and sits steady near 3.6 billion. That is why the OpenCL and ROCm "Greek cost" column reads ~2.5× rather than Vulkan's 1.3×: it is the price-only number that ran away, not the sweep that got slower. This is the register-pressure story from the primer, showing up as a measurement artefact.

fp32 does nothing at all for cpu-jit. The price-only figure is 171 in single precision and 171 in double — identical to the digit, which is the tell. The generated bytecode kernel widens float to double on the JVM operand stack anyway, so narrowing the declared type only adds conversions. The JVM has no float fast path in the way the hardware does; a float on the JVM is a storage format, not an execution mode.

fp32 buys SIMD about 1.2× on the price pass — 545 against fp64's ~460 — which is roughly what you would expect from sixteen float lanes in a 512-bit vector against eight double lanes, minus the transcendentals that do not narrow as cleanly. The adjoint column is a different story, and it gets its own section.

The fp64 tax, and why finance keeps paying it

Before the second matrix: the single most useful fact about GPUs that quantitative people tend to learn the expensive way.

Double precision on a GPU is a product-segmentation decision, not a physics one. Vendors run fp64 at whatever fraction of the single-precision rate the market segment justifies:

hardware class	fp64 : fp32 rate	in practice
NVIDIA data-centre (A100, H100, and successors)	1 : 2	full-speed double; what HPC buys
AMD Instinct MI250X	1 : 1	vector fp64 at the fp32 rate
AMD Instinct MI300X	1 : 2	full-speed double
AMD consumer Radeon / this 780M	~1 : 16	usable, not fast
NVIDIA consumer GeForce (RTX)	1 : 64	deliberately crippled
Apple M-series GPU	none	Metal has no double type at all
Vulkan, portably	none	shaderFloat64 is an optional feature

So a fp64 Monte Carlo on a gaming card is running on about 6% of the silicon, and on an Apple GPU it is not running at all. This is not an accident of manufacturing — the same die, sold as a data-centre part, would do it at half rate.

The finance angle is that fp64 is usually not negotiable. Not because Monte Carlo needs fifteen significant figures — it emphatically does not; the sampling error at a hundred million paths swamps float rounding by orders of magnitude — but because a number that goes into a regulatory capital calculation has to be reproducible and defensible, and "we ran it in single precision" is a conversation nobody wants to have with a model-validation team. Whether that instinct is right is a real debate. The matrices let you cost it.

The fp64 matrix

Same recording, same seed, .fp64().

backend / config	price-only	value + 5 Greeks	Greek cost
cpu — scalar interpreter, 8T	87	48	1.8×
cpu-jit — generated bytecode, 8T	171	107	1.6×
simd — JDK Vector API, 8T	442–487	190–222	2.2–2.5×
opencl — Radeon 780M iGPU	360	294	1.2×
rocm / HIP — Radeon 780M iGPU	404	323	1.3×
vulkan	fp32 only — no portable fp64 in SPIR-V
cuda — Tesla T4	~1 000	690	~1.5×

The read-out:

On an integrated GPU, fp64 is barely a GPU workload. ROCm and OpenCL land at 300–400 Mpath/s — level with the fastest CPU configuration's price-only rate. Against cpu-jit on the adjoint sweep it is about 3×: real, but not the kind of number anyone buys hardware for. The whole 1:16 tax is visible in that one comparison. On a discrete data-centre card the ratio changes completely, and the fp32 column shows what the same silicon does when you let it use its full width.

cpu-jit is the number to beat if you cannot ship a native dependency. 171 / 107 Mpath/s on a stock LTS JDK: no incubator module, no driver, no .so, no container gymnastics. Every GPU row above needs a working Vulkan or ROCm userspace on the deployment box. This one needs Java.

The Greek sweep costs about 1.6–1.8× a price pass on CPU and 1.2–1.3× on the GPU — never the ~11× a five-Greek central bump would cost (1 + 2×5 full repricings). That is the adjoint-AD payoff from the bump-and-revalue post, and the striking thing is that it survives every backend switch. Changing where the tape runs changes the throughput by two orders of magnitude and leaves the adjoint-versus-bump ratio essentially alone.

The interpreter earns its keep by being slow. cpu walks the tape node by node instead of compiling it, and pays about 2× for the privilege (87 / 48 against cpu-jit's 171 / 107). It is not in the table to be fast. It is in the table because it is the bit-exact oracle every other backend gets checked against — which is the subject of the next section.

CUDA and OpenCL on the Tesla T4

Two notes on the cuda rows above — same workload and unit as every other row, just a different machine.

CUDA and its OpenCL shim are level. 8 600 against 8 550 Mpath/s at fp32, 690 against 660 at fp64. NVIDIA's OpenCL is a thin layer over the same driver, and on a 26-node kernel the layer costs next to nothing. It is the AMD result inverted and shrunk: on NVIDIA the vendor stack is ahead by a hair, never the 3.4× the open API opened on AMD. The gap widens on heavier tapes, where the mature toolchain has more to compile against.

cuda does real fp64, where the fp64 matrix leaves vulkan blank — SPIR-V has no portable double. The T4 runs it at about a twelfth of its fp32 scenario rate, well short of the tax raw FLOP ratios predict, because the kernel is launch- and occupancy-bound rather than double-precision-FLOP-bound.

The row that would not sit still

The simd adjoint figures above are ranges, and honesty about why is worth more than a tidy number.

Running the same 8-thread SIMD configuration six times, each in its own fresh JVM, warmed identically, best-of-fifteen each time, fp32, value + 5 Greeks:

147   314   177   303   304   311      Mpath/s, value + 5 Greeks
548   536   551   545   548   543      Mpath/s, price-only, same six runs

That top row is not noise. Noise is a spread around a mean. This is bimodal: every run lands either near 165 or near 310, nothing in between, and whichever one it picks it holds for the whole life of that JVM. The price-only row underneath, from the same six processes, is boringly stable inside 3%. So it is not thermal, not the scheduler, and not the measurement harness — it is something about how the adjoint kernel specifically gets compiled.

A ~1.9× cliff that is decided once per process and then never moves is the classic signature of a Vector API intrinsification failure: when a hot method grows past HotSpot's inlining budget, FloatVector operations stop being compiled to single AVX-512 instructions and fall back to allocating boxed vector objects on the heap. The adjoint kernel is exactly the method that would be on that boundary — it is roughly twice the size of the price-only one, because it carries the reverse sweep. Whether C2 gets under the budget on a given run appears to depend on the order it happens to compile things in.

That is a hypothesis, not a confirmed diagnosis — nailing it down means reading compilation logs, which is a different article. Two things are certain enough to act on:

1. The published range is real and both ends of it happen. An earlier draft of this comparison reported the fp32 adjoint row as a single stable figure near 300 and asserted that only the fp64 path wobbled. Six repeats showed that was luck, not stability — and, incidentally, that the fp64 adjoint row it had flagged as wildly unstable is the steadier of the two today (190–222 across six runs). A benchmark you have run once is a benchmark you have not run.
2. The fastest CPU engine is also the least predictable one, and that is a deployment fact. simd at fp32 is somewhere between 1.7× and 3.2× cpu-jit on the adjoint sweep, and you do not get to choose which. cpu-jit is slower and lands on 107 every single time. For a nightly batch, take the speed. For a latency budget with a number in a contract, take the one that does not have a cliff in it.

Threads, and why the tables say eight

Every CPU engine splits the path range across .threads(n) workers, and every CPU row above uses eight — the eight physical Zen 4 cores, SMT excluded.

That is a deliberate choice, and it is worth being explicit that it is not the fastest choice. Letting the same kernel take all sixteen hardware threads is measurably better, and the number belongs in the article rather than in a footnote (cpu-jit, fp64, Mpath/s):

threads	price-only	value + 5 Greeks
1	24	15
8 — physical cores	171	107
16 — with SMT	205	134

So simultaneous multithreading buys about 25% on top of the physical cores. This kernel has enough transcendental and dependency latency in it that a second thread per core does find work to do — the intuition that a dense floating-point loop gets nothing from SMT turns out to be wrong here, which is exactly the sort of thing worth measuring rather than assuming.

simd gains from SMT too, and more interestingly: at fp32, sixteen threads gave 654 price-only and 369 on the adjoint sweep — better than even the good half of the bimodal 8-thread distribution.

Eight is in the tables anyway, for two reasons. It is the number that transfers to a sized production box, where you rarely get to assume an idle machine. And it is the conservative side of the comparison to be making: quoting a GPU speedup against a CPU baseline you have not maximised is how benchmarks flatter themselves. If you would rather read the accelerator ratios against a fully loaded CPU, divide them by about 1.25.

Scaling from a single core is close to linear, which is what a workload with no shared state should look like — every path is independent, the only synchronisation is one reduction at the end (fp64, value + 5 Greeks):

engine	1 thread	8 threads	scaling
cpu — interpreter	7	48	6.9×
cpu-jit	15	107	7.1×
simd	33	190–222	5.8–6.7×

Do the numbers agree?

A speedup is worthless if the fast path computes a different number. Because every backend replays the same tape against the same Philox stream path-for-path, at an equal seed and an equal path count the results can be compared with no statistical allowance whatsoever — any difference is arithmetic reordering, not sampling noise. At 20 million paths, seed 42:

backend        precision   price          delta        vega
cpu-jit        fp64        9.4122353      0.5986992    38.660793
simd           fp64        9.4122353      0.5986992    38.660793
opencl         fp64        9.4122353      0.5986992    38.660793
rocm  / HIP    fp64        9.4122353      0.5986992    38.660793
cuda           fp64        9.4122353      0.5986992    38.660793
cpu-jit        fp32        9.4122351      0.5986992    38.660789
simd           fp32        9.4122351      0.5986991    38.660790
vulkan         fp32        9.4122346      0.5986992    38.660787
opencl         fp32        9.4122343      0.5986991    38.660787
rocm  / HIP    fp32        9.4122343      0.5986991    38.660787
cuda           fp32        9.4122330      0.5986993    38.660776

The fp64 backends are bit-identical: generated Java bytecode, a Vector API plan, a HIPRTC-compiled HIP kernel, an OpenCL kernel and an NVRTC-compiled CUDA kernel all land on the same 9.4122353. The fp32 rows differ from fp64 only in the sixth or seventh significant figure — float rounding in the per-path arithmetic — and note that the reduction accumulators stay in double on every backend, so that error does not grow with path count. There is no backend here whose price you would not trust to five figures, and none whose Greeks disagree past the sixth.

That is worth dwelling on for a second, because it is the argument for keeping a deliberately slow reference implementation around forever. The fast kernels are not trusted; they are checked.

CUDA — the one with the ecosystem

CUDA is NVIDIA's proprietary parallel-computing platform: a C/C++ dialect with __global__ kernels, a runtime, a profiler, a debugger, and — the part that matters for a record-and-replay engine — NVRTC, a library that compiles kernel source to GPU code at runtime with no external toolchain on the box. It runs only on NVIDIA hardware, and that is the deal. In exchange for the lock-in you get the most mature compiler, the deepest library stack (cuBLAS, cuDNN, cuRAND, Thrust, and a hundred others), and the better part of two decades of numerical-computing projects that targeted it first and everything else, if ever, afterwards.

In finance specifically, CUDA has a long incumbency, and one press release did much of the work. In August 2011 NVIDIA announced that J.P. Morgan's Equity Derivatives Group had moved more than half its risk computation onto Tesla M2070 cards and was seeing a 40× end-to-end speedup — risk calculations "in minutes, not hours", and an 80% cut in the cost of running them. That story, repeated at every quant conference for a decade, is a large part of why "GPU" and "CUDA" became synonyms on trading floors. (NVIDIA's announcement.)

Pick CUDA when your deployment hardware is NVIDIA and you want the library ecosystem, the profiler, and the path every piece of documentation assumes. If you need fp64 at full rate, this is also — with AMD's Instinct line — one of only two places to get it.

ROCm / HIP — AMD's answer, deliberately CUDA-shaped

ROCm is AMD's open-source GPU-compute stack; HIP is its kernel language, and HIP is a near-clone of CUDA C entirely on purpose. A hipify tool does most of the mechanical translation from CUDA source, and AMD's claim is that 90%+ of a typical CUDA codebase converts automatically. This engine leans on that directly and about as hard as it is possible to lean: the ROCm backend feeds the same generated CUDA-C string the CUDA backend uses straight into HIPRTC, ROCm's runtime compiler, with no modification at all. One code generator, two vendors.

Pick ROCm when you are on AMD data-centre hardware or one of the officially supported Radeon cards, and you want fp64 on the device with a CUDA-like programming model. It is the only route to full-rate double precision that isn't NVIDIA.

Vulkan — the portable one that was a graphics API first

Vulkan is a Khronos open standard, and unlike the other three it is not primarily a compute API at all. It is a low-level graphics-and-compute API that happens to expose compute shaders. You write the kernel in GLSL (or any language with a suitable compiler), compile it to SPIR-V — a portable binary intermediate representation — and hand that to the driver.

The payoff is reach. The same SPIR-V runs on NVIDIA, AMD, Intel, Apple through MoltenVK, Android phones, the Steam Deck, a Raspberry Pi, and a CPU software rasteriser for CI machines with no GPU at all. Anything with a Vulkan driver, which by now is very nearly everything.

Pick Vulkan when you want one GPU path that runs on whatever hardware the user happens to have, fp32 is acceptable, and you do not want a vendor toolchain in your build or your container. On the numbers above it is also, by a wide margin, the fastest thing available.

OpenCL — the original open standard, still everywhere and nowhere

OpenCL is the oldest of the four: a Khronos standard for writing kernels in a C dialect that run on CPUs, GPUs, FPGAs and DSPs from any vendor. The promise was write-once-run-anywhere compute — CUDA's capability without CUDA's collar, arriving barely a year after CUDA itself.

The reality is that runtime quality varies enormously by vendor, fp64 and most interesting features are optional extensions you have to query for at runtime, and the tooling never got a fraction of the investment CUDA did. It is still the most broadly supported option — if you need to span an old Intel iGPU, an AMD card and an NVIDIA card with one binary, OpenCL is often the only thing that works at all — and it remains the only vendor-neutral route to double precision on a GPU.

Pick OpenCL when portability across vendors matters more than peak performance and you need fp64 on a GPU that is neither NVIDIA nor a supported-ROCm AMD part. Here the OpenCL device is AMD's own HSA-based runtime on the gfx1103; it works, it reconciles to seven figures, and on the adjoint sweep it lands at roughly a third of Vulkan — which is still an order of magnitude over the best CPU configuration.

cpu-jit and simd — why the GPU rows have context

Neither is an accelerator API, and both stay in every table on purpose.

cpu-jit is the tape compiled to plain Java bytecode. No native library, no driver, no incubator module, no container changes; it runs on any LTS JDK and does 107 Mpath/s of full risk vector on eight cores in fp64, 97 in fp32. It is the floor — and, across every repeat run in this article, the single most predictable number in any of the tables. For a great many real workloads the floor is enough.

simd is the same tape mapped onto the JDK Vector API — still pure JVM, still no native code, one --add-modules jdk.incubator.vector flag — and at fp32 it does 150–315 Mpath/s with the adjoint sweep, or 1.7–3.2× cpu-jit depending on which side of the cliff that JVM landed on. On the price pass, where it is stable, it is a clean 3.2×.

The AMD GPU compute stacks are roughly 12–24× the SIMD adjoint number at fp32; Vulkan is roughly 40–80×. Whether that gap matters depends entirely on whether another order of magnitude of "already fast enough" is worth a driver dependency in production. For a nightly risk batch it very much is. For a pricing service that answers in 20 ms and spends most of that waiting on the network, it very much is not — and shipping a JAR that runs anywhere has a value that does not appear in any of these tables.

Pros and cons, on one screen

	CUDA	ROCm / HIP	Vulkan	OpenCL	simd	cpu-jit
Vendors	NVIDIA only	AMD only	all GPUs	all vendors	any CPU	any CPU
fp64 on device	yes, full rate on DC parts	yes, full rate on Instinct	no (optional feature)	optional extension	yes	yes
Runtime kernel compile	NVRTC	HIPRTC	GLSL → SPIR-V	driver	n/a (JIT)	n/a (JIT)
Native dependency	CUDA driver	ROCm userspace	Vulkan loader + shaderc	ICD loader	none	none
In a plain container	toolkit needed	~20 GB stack, or a hand-staged subset	two .sos, a few MB	ICD, tiny	nothing	nothing
Library ecosystem	vast	growing, CUDA-shaped	none for compute	thin	—	—
Tooling / profiler	best in class	decent, improving	graphics-oriented	poor	JFR, async-profiler	JFR, async-profiler
Run-to-run stability	—	price-only noisy	good	price-only noisy	bimodal (see above)	excellent
fp32 val+5G (Mpath/s)	8 600	3 590	12 180	3 560	150–315	97
fp64 val+5G (Mpath/s)	690	323	—	294	190–222	107

So which one should you actually use?

• Shipping to the JVM with no native dependency at all: cpu-jit. It needs nothing, and ~107 Mpath/s of full risk vector on a modest CPU is enough for a great many real workloads.
• CPU throughput without a driver, and you can add one JVM flag: simd at fp32. Up to ~3× cpu-jit, still pure Java — but read the bimodality section first if you are sizing a latency budget rather than a batch window.
• One GPU path that runs on whatever hardware is in the box, and fp32 is acceptable: Vulkan. Portable, tiny dependency, no vendor toolchain — and on this evidence frequently the fastest thing available anyway.
• You are on NVIDIA and want the ecosystem, the profiler and full-rate fp64: CUDA. Nothing else gives you the libraries.
• You are on supported AMD hardware and need device fp64: ROCm. On Instinct parts it is the fastest double precision money can buy.
• You must span vendors — some Intel, some AMD, some NVIDIA — with one build, and need device fp64: OpenCL, eyes open about the performance ceiling and the tooling.

The genuinely nice property of a record-and-replay design is that this is a one-word decision, made at deployment time, rather than an architectural commitment made at the start of the project. The tape does not know or care which of these ends up executing it — and, as the reconciliation table shows, neither does the answer.

Try it yourself

git clone https://github.com/nablatensor-dev/nablatensor && cd nablatensor
mvn -o -q compile

# every backend with a driver present, side by side (fp64)
MAVEN_OPTS="--add-modules jdk.incubator.vector" mvn -o -q -pl nablatensor-examples exec:java \
  -Dexec.mainClass=com.nablatensor.examples.AsianGreeksBackends -Dscenarios=2000000 -Dsteps=252

# the European-call showpiece, pinned to one backend at a time
mvn -o -q -pl nablatensor-examples exec:java \
  -Dexec.mainClass=com.nablatensor.examples.BlackScholesBothWays -Dengine=cpu-jit   # or simd / vulkan / opencl

The two matrices in this article come from a smaller harness that runs exactly one configuration per JVM — OneCfg.java, in the engine repo's docs-internal/bs-perf/ — because, as the SIMD section shows, sharing a JVM between configurations is how you publish a number that is not there:

CP="$(find . -path '*/target/classes' -type d | tr '\n' ':')"
java -cp "$CP" --enable-native-access=ALL-UNNAMED --add-modules=jdk.incubator.vector \
  --source 25 docs-internal/bs-perf/OneCfg.java cpu-jit fp64 8     # engine, precision, threads

AsianGreeksBackends loops over availableEngines() — cpu, cpu-jit, and whichever GPU backends have a working userspace. To stay on the CPU, pin one backend with .on("cpu-jit"). GPU setup notes are in docs/install/vulkan.md and docs/install/rocm.md.