Four ways to talk to a GPU

Module 4.1 named vulkan, rocm, and cuda as "compiles to a shader or a device kernel" and stopped there. They're not three flavors of the same thing. Two of them are the same generated string, compiled by two different vendors' compilers; one of them mechanically rewrites that string into a different C dialect; and one of them isn't C at all.

The whole story

Did you know?

The opencl backend is about forty lines of String.replace. NVRTC (CUDA) and HIPRTC (ROCm) both accept the exact same CUDA-C source, so cuda and rocm already share one code generator, verbatim. OpenCL C differs from CUDA C in a small, closed set of ways — __kernel for __global__, get_global_id(0) for blockIdx.x * blockDim.x + threadIdx.x, __local for __shared__, barrier() for __syncthreads() — so the opencl backend takes the same generated CUDA-C string and rewrites exactly those tokens before handing it to the driver's own compiler. Three backends, one generator, nothing to keep in sync by hand.

Same string, three compilers, and one that isn't

Nothing about your code changes except the argument to .on(...):

try (MonteCarlo<EquityMarket> mc = MonteCarlo.of(Products.europeanCall())
        .market(EquityMarket.atmOneYear())
        .steps(1)
        .greeks()
        .on(engine)          // "cuda" | "rocm" | "opencl" | "vulkan"
        .fp32()
        .build()) {
    var v = mc.run(100_000_000L, 42L);
}

Underneath, cuda, rocm, and opencl all start from the identical generated CUDA-C kernel — one forward sweep, one reverse sweep, entirely in registers, one scenario per thread. cuda hands that string to NVRTC; rocm hands the same string to HIPRTC; opencl hands the token-rewritten version to whichever driver's ICD compiler is installed. vulkan never sees that string at all — it compiles hand-written GLSL to SPIR-V, a portable binary instruction format, and that's what the driver actually runs. It's the only one of the four that isn't, underneath, a dialect of C.

The number that makes this page worth reading

On one AMD Radeon 780M integrated GPU, replaying the same 26-node tape 1.3 built (Products.europeanCall(), one step), value plus all five Greeks, fp32:

backend	Mpath/s	vs. price-only
`opencl`	3,560	~2.5×
`rocm` / HIP	3,590	~2.5×
`vulkan`	12,180	1.3×

Vulkan — a graphics API, from a standards committee, running on a driver AMD did not write — beats both of AMD's own compute stacks by 3.4×, on AMD's own silicon. Not a rounding difference: three and a half times.

Did you know?

Why would the committee standard win? Nothing to do with API design, and everything to do with where the engineering went. Vulkan's compute path is thin — a SPIR-V module and a dispatch, almost nothing between your kernel and the shader cores — where ROCm and OpenCL both route through the heavier HSA queue machinery in amdkfd. And the shader compiler behind Vulkan on this GPU, RADV's ACO backend, is written and tuned largely by Valve, because every game on the Steam Deck depends on it. On a gfx1103 APU, that compute path runs a large fraction of the games on Steam; ROCm runs approximately none of them. Testing volume is a performance feature.

`cuda`, and the double-precision catch

cuda isn't in the table above — the numbers there are all one machine's integrated GPU, and NVIDIA's card in this comparison is a Tesla T4 on a different box entirely, so it doesn't belong in the same ranking. On its own hardware it does 8,600 Mpath/s on this same tape, ahead of the 780M's rocm/opencl rows and behind Vulkan's — a genuinely different machine, not an apples-to-apples fourth row.

Precision is where vulkan's story has a real catch. cuda and rocm both do full fp64 on the hardware they're built for; opencl gets it as an optional extension most drivers implement. vulkan cannot — SPIR-V's shaderFloat64 is an optional feature, and there's no double-precision type in core GLSL at all. If a regulator needs the bank's number in double precision, vulkan is off the table before performance enters the conversation.

Try it yourself

Take any GPU example and change one string. If you have an NVIDIA card, switch .on("cuda") to .on("vulkan") and rerun — same seed, same path count, and the price and every Greek should reconcile to the same digits; only throughput moves. On an AMD box like this one, the equivalent flip is .on("rocm") to .on("vulkan") — one word, same tape, same result, different silicon path underneath.

▶️ Run it

mvn -o -q install
mvn -o -q -pl nablatensor-examples exec:java \
  -Dexec.mainClass=com.nablatensor.examples.BlackScholesBothWays \
  -Dengine=cpu-jit   # or simd / vulkan / rocm / opencl / cuda

Defaults to cpu-jit, which needs no driver at all. Swap -Dengine=... for whichever backend has a working userspace on your machine — nothing runs on a GPU unless you name one.

⚠️ What this doesn't do

This page compares what CUDA, ROCm, Vulkan, and OpenCL are and how they compile a kernel; it doesn't re-derive the register-pressure story behind why a GPU's Greek pass can cost more than its price pass — that's 4.2's checkpointing page, and it applies to all three CUDA-C backends here the same way. It also doesn't cover the fp64 throughput matrix in full (the 1:16 tax on a consumer AMD part, the 1:64 one on consumer NVIDIA), or the history of how CUDA, ROCm, and Vulkan each got their names — the deeper link below has both, plus the licensing dispute over running compiled CUDA on non-NVIDIA hardware.

What's next

→ Deeper: Four ways to talk to a GPU: CUDA, ROCm, Vulkan and OpenCL on one pricing kernel has the full fp32/fp64 matrices, the bit-exact reconciliation table across all four backends, the SIMD bimodality digression, and each API's history — Brook, the Boltzmann Initiative, Mantle, and OpenCL's own un-shipped 2.0. → Next: Basket options — correlated underlyings, once you've picked whichever backend actually runs them.