CubeCL GPU backend: kernels as Rust, evaluated on CUDA/Metal/RADV - #99
Conversation
The CUDA path is NVIDIA-only. This adds a second backend that runs on any
Vulkan 1.2 device with shaderInt64 -- AMD, Intel, NVIDIA, and llvmpipe --
selected at runtime with --gpu-backend {auto,cuda,vulkan}.
Both cudarc and ash dlopen their driver, so one binary can carry both
backends and require neither at build time. `auto` tries CUDA first, so
behaviour on an NVIDIA machine is unchanged; an explicitly named backend
that fails to initialize is fatal rather than silently falling back to a
much slower path.
Detailed mode only for now; niceonly falls back to the CPU. The detailed
kernel takes no input buffers (each thread derives its own n), so it needs
neither the MSD prefix pipeline nor the stride residue table on-device.
Kernels are generated as WGSL per base and compiled to SPIR-V by naga at
runtime -- the same JIT-per-base design NVRTC gives the CUDA path, which
matters because every divisor in the hot loop must be a compile-time
constant.
The one place the port could not be literal is the scan_digits chunk
split. nvcc strength-reduces `u64 / CHUNK_DIV` to a multiply-high; ACO
does that only for 32-bit divisors (NIR's nir_opt_idiv_const is width
limited), and expands the 64-bit form into a ~220-instruction restoring
shift-subtract loop. So CHUNK_DIV is kept under 2^16 and the split is done
as two 32-bit constant divisions over 16-bit halves, which is exact.
Measured on base 40 (Radeon 860M): 3.88e8 vs 1.36e8 candidates/s, i.e.
2.85x, despite needing ~1.7x more chunk iterations; shader code size
halves and the division expansion leaves the ISA entirely.
Per-base constants move to a new backend-neutral gpu_config module, since
client_process_gpu is cfg(feature = "gpu") and unreachable from a Vulkan
build. The CUDA path keeps every one of its existing tests.
Measured end to end, 1e9 numbers at base 40: 4.33e8/s on the integrated
Radeon 860M against 6.50e7/s for the 16-thread CPU on the same machine.
CPU/GPU parity is exact (distribution and near-miss list) at bases 10, 40,
62 and 80 -- and unlike the CUDA parity tests, it also passes against
lavapipe, so correctness can be checked with no GPU at all.
Assisted-by: Claude Code:claude-opus-5
`--gpu-backend auto` is meant to try CUDA and fall through to Vulkan, but cudarc panics rather than returning an error when libcuda.so is absent (`panic_no_lib_found` returns `!`). That is the ordinary case on a machine with no NVIDIA driver -- precisely the case the fallback exists for -- so `nice_client --gpu` died with cudarc's 20-line message instead of using the Vulkan backend sitting right there in the same binary. Wrap CUDA init in catch_unwind with the panic hook silenced for the duration, so a routine fallback is quiet and the reason is still available at debug level. Upstream catches the same panic for the same reason in nvrtc_compiles_kernels_for_all_supported_bases. An explicit --gpu-backend cuda still fails hard, now with the underlying message and a pointer at --gpu-backend vulkan, rather than a panic. This relies on unwinding; the workspace does not set panic = "abort". Assisted-by: Claude Code:claude-opus-5
Bases with b % 4 == 3 have an empty residue set R_b. That means there are provably no solutions -- but it also means StrideTable's first_valid_at_or_after indexes valid_residues[idx] on an empty vec and panics, so the answer has to be decided before any table is built. process_range_niceonly_vulkan built one unconditionally. The CUDA path has the same guard inside process_range_niceonly_gpu; this puts it ahead of the CPU fallback rather than after it, so it also covers bases the GPU itself cannot take. The server never hands out such bases, so this was not reachable in production -- but --benchmark and any direct call are. Also document --gpu-backend in the README, note that the two backend features are additive and need neither library at build time, and build all three GPU feature combinations in `just test`. Assisted-by: Claude Code:claude-opus-5
Ports niceonly_ranges_kernel to WGSL and lifts the host side of the niceonly path out of client_process_gpu.rs into a new backend-neutral module, so CUDA and Vulkan drive one pipeline instead of two copies. This is the gpu_config.rs split applied to behaviour rather than constants: backends supply a RangeSink (launch + sync), everything above that -- the adaptive MSD floor, the worker/channel/batching loop, the per-field stats -- is written once. The CUDA path was moved onto it; its kernel source is unchanged and it keeps every one of its tests. Correctness is exact against client_process::process_range_niceonly at bases 10, 12, 25, 40, 45 and 62, on RADV and on lavapipe, and base 10 returns exactly [69] through the whole pipeline. That is the real claim: the GPU checks a superset of the CPU's candidates (its MSD floor is coarser), so the two only agree if the on-device candidate reconstruction and the digit scan are both right. One construct did not port, and it is the same one as last time. CUDA computes n mod M with two 64-bit divisions by a constant, which RADV/ACO expands into a ~220-instruction shift-subtract loop rather than strength-reducing. The shader now takes field_start mod M as a push constant and reduces only the 64-bit offset, byte by byte: acc = (acc << 8 | byte) % M, eight 32-bit constant divisions. Exact because acc < M <= 2^24, and real moduli (b-1)*b^2 top out near 2^21. RADV_DEBUG=asm confirms it: zero s_brev_b64, one v_cmp_le_u64 (the loop bound), no scratch, no spills. Two things this turned up that are worth calling out. AdaptiveFloor's metric was wrong on any backend whose launch blocks. It inferred GPU time by subtracting MSD time from the total, which assumes the two phases are disjoint -- true for an async CUDA launch, false for a Vulkan dispatch that waits on a fence inside launch. The same 1e13 base-50 field read "msd 122.6s, gpu 0.15s" before and "msd 5.3s, gpu 133.0s" after: not slightly off, inverted. It would have driven the floor to MSD_FLOOR_MAX every field, which measures >2000s against 129s at the optimum. The pipeline now times its own calls into the sink and subtracts them; CUDA's numbers barely move. Bounding the descriptor channel (needed because a blocking consumer lets the MSD workers queue over a gigabyte) introduced a deadlock on the launch-error path: the consumer breaks, workers parked in send never wake, and thread::scope waits on them forever. Fixed by dropping the receiver inside the scope and breaking on send failure, and pinned by a timeout-guarded test that hangs for 120s without the fix. Not included, deliberately: the modular prefilter (base <= 40 only) and async dispatch. See vulkan-nice-spec.md for both, and for the floor sweep -- end to end this buys ~1.16x over the 16-thread CPU on an iGPU, not detailed mode's 6.7x, because the floor that keeps the GPU fed also multiplies the work it is fed. Assisted-by: Claude:claude-opus-5 [RADV_DEBUG] [clippy]
The niceonly kernel's uniform pre-check — are the lowest p base-b digits of n^2 and n^3 all distinct? — rejects ~99% of candidates in fixed, branch-free work, so whole lanes skip the full check. It was the last piece of niceonly_ranges_kernel not ported. __umul64hi, the difficulty the port expected here, is not needed at all, because the construct it appears in does not survive for a different reason. reduce_pre ends in `lo % PRE_MOD`, and the digit peel is `sq % BASE` / `sq /= BASE` on a u64: three more 64-bit divisions by a compile-time constant, which is exactly the expansion NIR's width-limited nir_opt_idiv_const leaves behind (~220 instructions), in the one place that runs on every candidate. So carry the residue as three chunks of base^chunk_digits < 2^16 — the same chunk the digit scan already uses — instead of in a u64. `n mod b^p` is three rounds of the existing split16 step, the multiply is truncated schoolbook over those chunks (a*b + acc + carry <= chunk_div^2 - 1 < 2^32, so every intermediate stays a u32), and the digits fall out of each chunk with 32-bit divisions. Nothing 64-bit is ever divided. Three chunks also lands on exactly CUDA's digit count at every base that enables either filter. Measured on an AMD Radeon 860M, 1e12 numbers at base 40, MSD floor pinned at 4000: the device phase drops from 31.8-34.7s to 19.0-20.7s (1.68x on the kernel, 1.30e9 vs 7.7e8 candidates/s) and the field from 38.3-41.3s to 25.4-27.2s, against 67.7-71.0s for the same field on the 16-thread CPU. RADV_DEBUG=asm,shaderstats confirms the expansion never appeared: 0 s_brev_b64, scratch 0, 56 VGPRs and 18 subgroups/SIMD with the filter and without, for 362 extra instructions. Correctness needed its own test. The niceonly parity tests cannot see this filter — the only nice number any of them finds is 69, in base 10, where the prefilter is disabled — so a filter that rejected every candidate would pass all of them, which is the bug the CUDA kernel shipped in v3.2.14. A probe build of the shader reports the prefilter's survivors directly and they must match a Rust mirror of the chunk arithmetic candidate for candidate, on RADV and on lavapipe. It agrees at bases 30/34/40, and the 1.12% survival it measures at b40 matches the CUDA crossover study's 1.1% on other hardware. Also adds the massive-low-base benchmark (1e12 @ base 40): every range large enough to measure niceonly was base 50, where the prefilter is compiled out on profitability grounds, so nothing exercised it. NICE_VULKAN_PREFILTER=0 compiles it out, which is how the A/B above was taken. Assisted-by: Claude:claude-opus-5
…patch Phase 3, both changes identified and measured for by earlier phases. The lane tiling was CUDA's warp carried over as a constant 32. Nothing here is a warp — the lanes stride a range's candidates by index and never communicate — but they do all repeat the range's setup, the residue reduction and a ~12-iteration binary search. At MSD floor 250 a base-40 range holds ~39 candidates, so 32 lanes bought 32 copies of that setup to share out 1.2 candidates each. The shader now takes log2(lanes) as a push constant and the host sizes it per dispatch from the batch's mean range length. Device time at floor 250 by pinned width: 30.0s at 32 lanes, 24.6 at 16, 22.8 at 8, 21.6 at 4, 21.1 at 2 and 1 — monotone where ranges are short, flat within noise everywhere else, which is the whole shape of the result. Dispatch is now asynchronous over a ring of command buffers and fences, with per-slot range descriptor buffers so the host fills batch i+1 while the device reads batch i. The reservation is a guard that holds the submitter lock across the caller's fill: without it two threads could be handed the same slot, since nothing marks a slot taken until its submit. Holding it costs no device time — it blocks other hosts, not the queue. Measured on an AMD Radeon 860M, totals at the floor each configuration is best at. Base 50, 1e13: 129.4s -> 86.5s, against 149.7-180.5s for the 16-thread CPU, so this base goes from 1.16x the CPU to 1.7-2.1x. Base 40, 1e12: 25.4s -> 22.5s pinned, ~23.6s adaptive. The prefilter does not apply at base 50, so that base's whole gain is this phase. It also repairs the control loop rather than the constant. AdaptiveFloor aims at msd == gpu, which minimizes a maximum — right only when the phases overlap. They now do: the same base-40 field that used to walk the floor from 24.6s down to 57s settles into a limit cycle at 23.5-25.0s against 22.3s for the best fixed floor, ~5% off where it was 2.3x off. The residue is the hard +/-1.5x step, not the target. Two constants flagged as never swept are now swept and both are flat: HIST_COPIES (1/2/4/8: 2.66/2.58/2.62/2.65s) and MAX_WORKGROUPS (1024..16384: 22.6-22.8s). Left at 4 and 4096. One caveat the async path introduces: device_secs is time the host spent blocked in launch/sync, not the device's work, and an asynchronous backend drives it to nearly zero on a field the device keeps up with. Compare totals. Assisted-by: Claude:claude-opus-5
The bound was quoted as LAUNCH_BATCH_RANGES * PIPELINE_DEPTH descriptors, which names the wrong quantity. A channel item is the whole output of one PROCESSING_CHUNK_SIZE chunk; LAUNCH_BATCH_RANGES is the consumer's flush threshold and never bounds what is queued. The real cap is PIPELINE_DEPTH times the most ranges one chunk can yield. get_valid_ranges_recursive returns a range whole once it is at or below the floor, so that is about PROCESSING_CHUNK_SIZE / floor -- ~4000 at MSD_FLOOR_MIN, fewer at any coarser floor, and the depth-22 subdivision limit never binds first. Roughly 3 MB rather than the 50 MB the old arithmetic implied; the stated bound was conservative, so nothing was unsafe, but this constant is one of the two things that changed for CUDA users (the channel used to be unbounded) and it should say what it does. Comment only; no behaviour change. Assisted-by: Claude:claude-opus-5
…nt LSD bitmap Brings wasabipesto/nice 0ed8652 onto vulkan-backend. Two upstream filter fixes, both of which the Vulkan backend inherits without shader changes: - has_duplicate_msd_prefix loses its cross MSD-LSD collision check, which was unsound (it treated the range start's low square/cube digits as fixed across a range whose n mod b^k varies). Ranges containing nice numbers were being skipped, e.g. [68, 70) in base 10 around 69. The Vulkan niceonly path calls get_valid_ranges_recursive, so it was skipping the same ranges the CPU was. - get_valid_multi_lsd_bitmap now checks all 2k fixed low digits of n^2 and n^3 for pairwise distinctness at fixed width. The GPU uploads this host-built stride table verbatim, so the tighter bitmap reaches the device with no kernel or codegen change. Expect niceonly to slow down where the old filter was wrongly aggressive (upstream measured ~1.6x at b40, ~2.6x at b50); that cost is the fix.
…2 storage Brings wasabipesto/nice up to a066ab3 (origin/main). wasabipesto#87 came in on the previous merge; this one is wasabipesto#88 plus the version/changelog commits. Conflict: wasabipesto#88 edited GPU_LSD_K and the MSD-floor machinery in place in client_process_gpu.rs, but this branch had already moved both into gpu_niceonly.rs so the Vulkan backend could share them. Resolved by keeping the shared module and carrying wasabipesto#88's semantic change to it — GPU_LSD_K 2 -> 3. The duplicated AdaptiveFloor block upstream re-added was dropped; it is identical to the one already in gpu_niceonly.rs. The k=3 table needed one real fix on the Vulkan side. The niceonly shader reduces a range offset with a Horner loop, acc = (acc << c | chunk) % M, which requires M <= 2^(32-c); c was a fixed 8, sound only because M = (b-1)*b^k stayed under 2^24 at k=2. At k=3 every modulus gains a factor of b: base 65 hits 17 576 000 and base 128 hits 266 338 304, so NiceonlyConfig::new would have refused every base >= 65 — base 80, which is in production here, among them. stride_chunk_bits now picks c per base (8 while M <= 2^24, else 4, bound 2^28), and the emitted shader and the host mirror derive it from the same modulus. It costs nothing measurable: this reduction runs once per range descriptor, not per candidate. Also adapted this branch's stride-table test mirrors to the u32 residue and gap types, the same way wasabipesto#88 did for the CUDA ones, and added base 80 to the device parity list — every base below 65 keeps the byte chunk, so nothing else here executes the 4-bit path on hardware. Verified: 120 unit tests plus all 6 device tests pass on the AMD iGPU; niceonly and detailed both match the CPU at bases 10/12/25/40/45/62/80. Assisted-by: Claude:claude-opus-5
…ark sweep, telemetry Upstream replaced the fixed named benchmark fields (common/src/benchmark.rs) with a structured adaptive sweep in client/src/bench.rs (wasabipesto#90), added benchmark uploads and opt-in submission telemetry (wasabipesto#91), a POST /estimate endpoint with GPU model canonicalization (wasabipesto#93, wasabipesto#95), and swapped the MSD prefix checks for an interval digit-domain (Hall) analysis (wasabipesto#89). One conflict and two integration fixes: * common/src/benchmark.rs was deleted upstream and carried our MassiveLowBase mode, which existed only because every other niceonly benchmark was base 50, where the modular prefilter is compiled out (GPU_PREFILTER_MAX_BASE). Took the deletion: the new sweep's niceonly scenarios include two base-40 windows (b40_msd_strong, b40_msd_weak), so the prefilter is exercised again. * DataToServer gained a telemetry field; client_process_vulkan.rs builds two of them and needed the same `telemetry: None` wasabipesto#91 gave the CUDA path. * gpu_name() was CUDA-only, so a Vulkan run reported gpu:true with a null gpu_model, and on a box with both backends a --gpu-backend vulkan run would have reported the CUDA device's name. Both feed the new estimator's matching. It now asks the live GpuHandle, which is the only thing that knows which backend actually initialized; Vulkan already recorded device_name at init. wasabipesto#89 needed no change on our side: the Vulkan pipeline calls get_valid_ranges_recursive rather than reimplementing the filter, so it inherits the stronger prune. Device parity re-verified — all 6 device tests and 131 unit tests pass, including vulkan_matches_cpu_niceonly at base 80. No GPU performance number was re-measured. wasabipesto#89 trades -2 to -12% must-process work for +0 to +8% CPU wall clock, which shifts the host/device balance in the niceonly pipeline and invalidates the pre-merge Vulkan speedups again. Assisted-by: Claude:claude-opus-5
…eardown `DetailedRun::drop` and `NiceonlyRun::drop` called `vkDeviceWaitIdle`, which is specified as a `vkQueueWaitIdle` on *every* queue of the device and so requires external synchronization against all of them. That is a precondition neither `Drop` is in a position to guarantee, and violating it is undefined behaviour whose observed cost in nice-count was not a crash but whole dispatches silently doing nothing -- the host fenced, read zeroed counters, and folded the work away. No such violation is reachable today: one field is processed at a time, and only `run_range_pipeline`'s consumer thread submits -- the MSD workers just produce descriptors down a channel. So this is not a bug fix, it is the removal of an unstated invariant. Nothing in either `Drop` records that the call is safe only because no second thread submits, and a future overlapping field or second context would break it silently rather than loudly. `wait_all` needs no such invariant: it takes the submitter lock and waits the fences of the slots actually in flight, which is a superset of the dropping run's own dispatches and the only work that can still be reading the descriptor pools and buffers freed immediately below. `VulkanContext::drop` keeps `device_wait_idle`. It owns the queue, and the `'a` borrow both run types take makes "every run is already dropped" a compile-time guarantee rather than a convention. Teardown-only, so there is nothing to measure. All 131 unit tests and all 6 device tests pass on the AMD iGPU, including `vulkan_matches_cpu_detailed` and `vulkan_matches_cpu_niceonly` through base 80, both of which exercise each changed `Drop`. Assisted-by: Claude:claude-opus-5
…itten kernels The vulkan/codegen.rs detailed kernel expressed as a #[cube] Rust function: same per-base constants, same split16 chunk scan (every division 32-bit by a comptime-constant divisor), same 4-copy workgroup histogram, same MISS_STRIDE near-miss records. Per-base WGSL string generation becomes #[comptime] parameters; CubeCL JIT-specializes one kernel per base at first launch. Selected with --gpu-backend cubecl (feature "cubecl"; detailed mode only). The benchmark sweep runs through it unchanged, so rates are directly comparable with the cuda and vulkan backends on the same scenarios. Parity: cubecl_matches_cpu_detailed checks exact distribution + near-miss agreement with the CPU at bases 10/40/62/80, passing on lavapipe (same harness as vulkan_matches_cpu_detailed). This is a benchmark-grade evaluation port, not a merge candidate: niceonly is not ported, and device selection takes wgpu's default adapter. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The kernel was already runtime-generic; this makes the host body generic too and adds a CUDA context variant (feature cubecl-cuda), so the NVIDIA comparison can race CubeCL's CUDA codegen against the hand-written kernels directly, without depending on container Vulkan. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…lity cubecl-cuda references CUDA 12.8 tensor-map symbols unconditionally; with the workspace's cuda-12000 floor unioned in, cudarc compiles the 12.0 surface and the build fails. The hand-CUDA path calls none of the new APIs and symbol lookup is lazy under dynamic loading. Evaluation branch only; revisit the floor before any upstreaming. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
CUDA gets the wide flavor (one u64 constant division per limb, 5 digits per chunk at base 40, matching the hand kernel); wgpu keeps split16 for drivers that cannot strength-reduce 64-bit constant division. One kernel source, one comptime bool — this is the per-device choice PR wasabipesto#96 lists as future work for the WGSL generator. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Loaders >= 1.3.207 hide portability (non-conformant) drivers unless the instance enables VK_KHR_portability_enumeration, and the spec requires VK_KHR_portability_subset on devices that advertise it. Both are detected at runtime, so this is a no-op on conformant Linux/Windows stacks. Verified on an Apple M4 mini through MoltenVK 1.4.2: all six device parity tests pass (detailed, niceonly, prefilter mirror), and the detailed benchmark runs at 6.7e8 n/s (b40) / 3.9e8 n/s (b50) -- CubeCL-on-Metal measures 0.86x of this hand-WGSL backend on the same device. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The CUDA backend's feature and symbols lose their generic names: feature gpu -> cuda, client_process_gpu -> client_process_cuda, GpuContext -> CudaContext, GPU_BATCH_SIZE -> CUDA_BATCH_SIZE, process_*_gpu -> process_*_cuda. (cudarc's own CudaContext is imported as DriverContext where the two would collide.) Names that are genuinely backend-neutral — gpu_config, gpu_niceonly, --gpu-backend, the gpu CLI flag — keep them. A new umbrella feature takes the gpu name: it enables every backend (cuda, vulkan, cubecl, cubecl-cuda), so the release recipe is two builds — default features for a lightweight CPU-only binary, --features gpu for one binary that runs on the CPU or any supported GPU. No backend needs its GPU libraries at build time; verified by building the full umbrella on a machine with no CUDA toolkit. --gpu-backend auto now ranks backends by measured speed: CUDA, then CubeCL, then Vulkan (CubeCL measured 1.4x the hand-WGSL backend on RDNA4/RADV and 2.5x on NVIDIA). In niceonly mode auto skips CubeCL, which has no niceonly kernel yet, and falls through to Vulkan; an explicitly named backend that fails to init stays fatal. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Two jobs: lavapipe on Linux (vulkan + cubecl through software Vulkan, the same device the suites were developed against) and macOS arm64 (cubecl through wgpu->Metal on the runner's paravirtualized Apple GPU, and the hand-WGSL backend through MoltenVK, exercising the new portability-enumeration path). The workflow header notes the backends CI cannot cover: cuda and cubecl-cuda need NVIDIA silicon, so their result parity stays a manual step on real hardware — NVRTC compile-tests of every base's kernel still run in the normal suite. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The niceonly kernel joins detailed as a #[cube] Rust function: the same on-device candidate reconstruction from the residue table, offset reduction (chunked u32 Horner, every divisor comptime), lane tiling, and low-digit modular prefilter as the generated WGSL, reviewable side by side with it. The stride modulus and residue count ride along as comptime parameters, so CubeCL specializes one kernel per base exactly as the string generators bake literals. The host end is a RangeSink like the other backends', driving the shared gpu_niceonly pipeline; the tiling helpers (lane_shift_for, stride_chunk_bits, MAX_STRIDE_MODULUS, MAX_LANES_PER_RANGE) and the residue-empty guard move from vulkan/codegen.rs and client_process_vulkan.rs into gpu_niceonly, where the Vulkan backend re-imports them — they were never Vulkan-specific, and the CubeCL backend must build without the vulkan feature. Device tests mirror the Vulkan suite: CPU parity at bases 10/12/25/40/45/62/80 (the GPU checks a superset of the CPU's candidates; the sets must still be identical), 69-in-base-10 through the whole pipeline, a CUDA-runtime variant, and a prefilter probe that reports survivors at every lane width against an independent u128 host mirror — an over-rejecting prefilter passes every parity test (the v3.2.14 bug class), so the filter's own output has to be observable. All pass on lavapipe; survivor counts match the hand-WGSL filter's exactly (396/159058 b30, 2099/350205 b34, 4472/288591 b40). With both modes served, --gpu-backend auto drops its niceonly carve-out: the order is CUDA, CubeCL, Vulkan in every mode. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
=0.10.0: CubeCL is pre-1.0 and its host API drifts between releases; bumps should be deliberate, gated on the device parity suite. cuda-12080 stays, now with the argument written down: it selects which bindings are compiled, not which driver is required — cudarc's dynamic loading resolves symbols lazily, the hand-CUDA path calls nothing newer than 12.0, and cubecl-cuda's 12.8 declarations (tensor-map) are never launched by our kernels. Parity-verified on real silicon. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Ubuntu ships the lavapipe manifest as lvp_icd.json; upstream Mesa builds name it lvp_icd.<arch>.json. The job's glob only matched the latter, so VK_ICD_FILENAMES exported empty and the loader found no drivers. Glob both, and fail the step loudly when no ICD is found instead of exporting an empty loader list. Verified in an ubuntu:24.04 container. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Three backends at very different throughputs can share one device name (the 9070 XT differs 1.4x between vulkan and cubecl), so hardware rows without the backend cannot be compared. hardware.gpu_backend carries the --gpu-backend value of the backend that actually won init — read from the live handle, since auto decides at init among the compiled backends. Additive JSON; absent on CPU runs; no schema bump. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
CubeclNiceonlyRun rebuilt the stride table and re-uploaded the residue buffer on every field. StrideTable::new walks the whole modulus on one CPU thread — 8.6e6 steps at base 52 — and at benchmark-window field sizes that walk dominated the pipeline: on an RTX 4060 the per-field logs showed the device phase at 0.65-1.0x of hand-CUDA while the scenario rates cratered to 0.24-0.57x, the difference being wall time spent between fields on table rebuilds the hand backends never do (both cache their per-base plans in the context; now this backend does too). Explains the base-dependence (modulus grows with base) and the device spread (the slower the box's CPU, the worse the hit). Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Same shape of knob as the hand backend's NICE_VULKAN_LANES: a power of two up to 32 pins every dispatch's lane width; unset keeps the adaptive per-dispatch sizing. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The pass-loop structure copied sq into a zero-padded cu_limbs window, built cu in its own array and copied it into the scratch, and branched on a runtime pass selector — per candidate. At bases with the low-digit prefilter (<= 40) almost nothing reaches the full check and none of that mattered; above it every candidate pays, which is where the remaining device-phase gap against hand-CUDA lived. Now the sq scan reads a sq_limbs window and the cube is built directly in the scratch and destroyed by its scan — no cu array, no copies, no pass branch, matching check_is_nice in nice_kernels.cu shape for shape. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Batches launch asynchronously; the blocking histogram read is the only stream sync, and doing it per 50M batch put a GPU-idle bubble after every ~12ms of kernel work — measured as most of the hand-CUDA backend's remaining detailed-mode lead on an RTX 4060. Each candidate increments exactly one u32 bin, so a drain per 64 batches keeps every bin under 64 * 50e6 < u32::MAX (const-asserted). The hand-CUDA kernel sidesteps the bound with u64 bins; WGSL has no u64 atomics, and this host loop serves both runtimes. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The kernels paired every constant division with a matching %. nvcc lowers each u64 constant division independently — a full multiply-high correction sequence per operator — so every div+mod pair cost double, and the digit peel pays the pair per digit. SASS for the b50 detailed kernel showed +42% instructions and ~100 extra ISETP correction compares vs the hand kernel, whose source computes every remainder as cur - q * DIV for exactly this reason. All sixteen div+mod sites in the detailed scan, niceonly scans, and prefilter now do the same. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
ptxas register-promotes the generated kernels' scratch array, turning every runtime-indexed access in the split loop into a compare/select chain over all words — ~100 extra SASS compares per kernel at b50, where the hand kernels' local-memory array pays one LDL/STL. Shared memory gives the same one-instruction indexed access without the local round trip: one word per thread per limb, padded to an odd stride so a warp's accesses spread across banks. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
|
Tested this on four machines: an AMD RDNA3.5 iGPU (Linux/RADV), an RTX 4060, an Short version: it builds, device parity passes everywhere hardware allowed it, There is one bug I'd call blocking, described first because it is a Blocking:
|
| A1000 (broken) | RTX 4060 (working) | ratio | |
|---|---|---|---|
b40_detailed |
1.16e10/s (206x CPU) | 5.84e9/s | 1.98 |
b50_detailed |
1.03e10/s (280x CPU) | 4.32e9/s | 2.38 |
An A1000 laptop GPU is not 2x an RTX 4060. Against the same box's wgpu backend
on the same card, the CUDA runtime claimed 3.4-6.5x in niceonly.
The hand-written CUDA path gets this right on the identical machine: it
detects the missing NVRTC, logs an error with troubleshooting steps, and exits
- Only the CubeCL CUDA path swallows the per-thread panics.
Reproduction: NVIDIA driver installed, CUDA toolkit not installed,
--gpu --gpu-backend cubecl-cuda --gpu-device 0. With no NVIDIA card at all it
fails cleanly instead (init errors, exit 101), so the driver-without-toolkit
combination is what triggers it — a very common configuration.
Build
All four backends build into one binary with no GPU libraries present at build
time, as advertised. Stripped sizes (lto = true, codegen-units = 1):
| build | stripped | vs cpu-only |
|---|---|---|
| cpu-only | 6.3 MiB | — |
cuda |
6.7 MiB | +0.3 |
vulkan |
8.2 MiB | +1.9 |
cubecl |
13.1 MiB | +6.7 |
cubecl-cuda |
14.8 MiB | +8.4 |
gpu (all four) |
16.4 MiB | +10.1 |
CubeCL roughly doubles a CPU-only binary. Worth knowing for
docker-build-push.yml, which publishes with --features gpu: the image grows
~8 MiB, and on the nvidia/cuda runtime base (no Vulkan ICD) the wgpu half is
inert. Not an objection — a separate CUDA-only tag would avoid it.
Correctness
Device parity passes on every backend each machine could run:
| backend | RTX 4060 | A1000 | Intel | RADV |
|---|---|---|---|---|
| cpu | pass (157) | pass (157) | pass (157) | pass (157) |
| cuda | pass (3) | no toolkit | no device | no device |
| vulkan | pass (5) | pass (5) | pass (5) | pass (5) |
| cubecl | pass (5) | pass (5) | pass (5) | pass (5) |
| cubecl-cuda | pass (2) | no toolkit | no device | no device |
Benchmarks
nice_client --benchmark, 30 s budget, 3 interleaved repeats, medians. Fixed
windows, so every arm grinds identical work. Implicit Vulkan layers disabled.
Ratios are per machine — absolute rates are not comparable across boxes
(different CPUs, OSes, rustc 1.94.1 vs 1.97.1).
CubeCL-over-wgpu vs the hand-written Vulkan kernel
| scenario | RTX 4060 | A1000 | Intel | RADV |
|---|---|---|---|---|
b40_detailed |
2.47 | 2.27 | 2.44 | 1.76 |
b50_detailed |
2.46 | 2.18 | 2.57 | 1.81 |
b40_msd_strong |
0.97 | 1.42 | (noisy) | 0.82 |
b40_msd_weak |
1.11 | 1.50 | (noisy) | 0.61 |
b50_residue_dense |
1.38 | 1.41 | (noisy) | 0.97 |
b50_msd_weak |
1.34 | 1.43 | (noisy) | 0.92 |
b52_msd_weak |
1.42 | 1.57 | (noisy) | 0.96 |
Detailed mode: CubeCL wins on all four vendors, by 1.8-2.6x. That is the
clearest result in the set.
Niceonly splits by vendor: CubeCL wins on both NVIDIA cards (1.11-1.57) and
loses slightly on RADV (0.61-0.97). The Intel box's niceonly figures are
withheld deliberately — its Vulkan arm varied by up to 2.6x between repeats
of identical work, so its medians are not a measurement. Its detailed figures
repeat to 0.5% and are quoted above.
I do not have an explanation for the vendor split, and the obvious one is wrong:
c1ae5e5 picks the chunk width per runtime, but wgpu keeps split16 on both
vendors, so it cannot be the cause (it is the right explanation for the
cubecl-cuda column below). One thing that would help diagnose it: the client
logs CubeCL wgpu device (name) without recording which graphics API
AutoGraphicsApi resolved to, so on Windows this may be DX12 against the hand
kernel's Vulkan. Happy to re-run with WGPU_BACKEND pinned.
CubeCL-over-CUDA vs the hand-written CUDA kernel (RTX 4060)
| mode | scenario | cubecl-cuda / cuda |
|---|---|---|
| detailed | b40_detailed |
1.08 |
| detailed | b50_detailed |
1.10 |
| niceonly | b40_msd_weak |
0.97 |
| niceonly | b40_msd_strong |
0.86 |
| niceonly | b50_residue_dense |
0.86 |
| niceonly | b50_msd_weak |
0.85 |
| niceonly | b52_msd_weak |
0.85 |
A portable Rust kernel beating hand-written CUDA in detailed mode is the
headline. Niceonly still trails by ~15%.
Sanity check on the setup: the vulkan/cuda ratios on the 4060 reproduce
figures measured on that card before this PR (0.51/0.49/0.48 niceonly, 0.35/0.30
detailed, against 0.51/0.50/0.47 and 0.34/0.29 recorded earlier).
What the last six commits bought (RADV)
I benchmarked 16779af first. The two runs are in different sessions and the
second came in 8-11% slower across the board — including the CPU and Vulkan
arms, whose code did not change — so absolute rates across them are not
comparable. Since the commits touch only common/src/cubecl_backend.rs, Vulkan
is a control, and its speedup-over-CPU is flat within noise (niceonly 6.51 ->
6.46, 6.96 -> 6.77, 6.96 -> 6.86, 2.73 -> 2.64; detailed 7.53 -> 7.61, 9.42 ->
9.53). Against that control, CubeCL/Vulkan moved:
| mode | scenario | 16779af | 30994e2 |
|---|---|---|---|
| niceonly | b40_msd_weak |
0.51 | 0.61 |
| niceonly | b50_residue_dense |
0.61 | 0.97 |
| niceonly | b50_msd_weak |
0.69 | 0.92 |
| niceonly | b52_msd_weak |
0.62 | 0.96 |
| detailed | b40_detailed |
1.37 | 1.76 |
| detailed | b50_detailed |
1.33 | 1.81 |
The niceonly work did what it set out to do.
Three smaller items
1. --features nice_client/cubecl-cuda alone does not compile.
error[E0433]: failed to resolve: use of undeclared type `CubeclContext`
--> client/src/main.rs:437:15
error[E0425]: cannot find value `CUBECL_BATCH_SIZE` in this scope
error[E0599]: no variant or associated item named `Cubecl` found for enum `GpuHandle`
client/Cargo.toml has cubecl-cuda = ["nice_common/cubecl-cuda"], which
forwards to nice_common's cubecl feature and so never sets
cfg(feature = "cubecl") on nice_client — the gate on the CubeclContext
import, while the new_cuda call site is gated on cubecl-cuda. The gpu
umbrella enables both, which is why the combined build hides it.
cubecl-cuda = ["cubecl", "nice_common/cubecl-cuda"]Same shape as the umbrella caveat already in the tree: cargo feature cfgs do not
propagate upward.
2. cubecl-cuda is unreachable from --gpu-backend auto. The cuda, cubecl
and vulkan blocks in init_gpu all match Auto | <backend>, but the
cubecl-cuda block is if want == GpuBackend::CubeclCuda. So the fastest
detailed-mode backend on the 4060 (59x/91x against CUDA's 55x/83x) is only
reachable by naming it. Possibly deliberate — it trails CUDA by ~15% in
niceonly, so the best pick is mode-dependent — but Auto's doc comment ("CUDA
if available, then CubeCL, then Vulkan") does not mention the exclusion.
3. --gpu-device is a per-backend namespace, and wgpu ignores it entirely.
On the A1000 laptop the A1000 is Vulkan ordinal 1 (0 is the Intel iGPU) but CUDA
ordinal 0, so no single value is right for all backends — passing 1 gave
CUDA_ERROR_INVALID_DEVICE. And CubeclContext::new_default() takes no ordinal
at all, so the wgpu backend picks its own adapter with no override. On a
multi-GPU box that silently benchmarks two backends on two different GPUs (it
did here, until caught). Worth a --gpu-device equivalent for wgpu, or at least
logging the chosen adapter distinctly.
Machines
| # | GPU | OS / rustc | backends exercised |
|---|---|---|---|
| A | Radeon 860M (RADV GFX1152) | Linux, 1.94.1 | cpu, vulkan, cubecl |
| B | RTX 4060 | Windows 11, 1.97.1 | all five |
| C | RTX A1000 6GB Laptop | Windows 11, 1.97.1 | cpu, vulkan, cubecl (no CUDA toolkit) |
| D | Intel(R) Graphics | Windows 11, 1.97.1 | cpu, vulkan, cubecl |
|
Of course right after posting that I found another bug when trying to run a real field from the server. Dead ~1 second after the claim, every time. The kernel log shows the GPU being What it actually depends onIt is the number of sequential dispatches on one stream, and nothing else I
Ruled out:
Minimal reproduction: let ctx = CubeclContext::new_default()?;
let r = get_base_range_u128(50)?.unwrap();
let start = r.start() + (r.end() - r.start()) / 2;
let range = FieldSize::new(start, start + 1_800_000_000); // 36 batches
process_range_detailed_cubecl(&ctx, &range, 50)?; // device lostI have not identified the mechanism. It looks like per-dispatch resources Why
|
Review found the silent-wrong-results case: with an NVIDIA driver but no CUDA toolkit, CubeCL compiles lazily at first launch, a missing NVRTC panics its server thread, and launches keep reporting success — the client completes and reports numbers measured from zeroed buffers. new_cuda now proves the runtime end-to-end with a one-thread smoke kernel and a verified readback before reporting success (the hand backends' init-time compile checks, ported), and detailed_impl checks histogram conservation — every candidate lands in exactly one bin — so any future silently-dropped work refuses to produce results on any runtime. Also from review: the client's cubecl-cuda feature now enables its own cubecl gate (feature cfgs do not propagate upward, so the standalone feature did not compile); the wgpu adapter log and telemetry name now include the resolved graphics API (AutoGraphicsApi may pick DX12 on Windows, and backend comparisons need to know); Auto's docs state that cubecl-cuda is deliberately explicit-only. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The M4's prefilter probe caught survivors diverging from the host mirror after the mul-sub commit — Metal only; the CUDA and SPIR-V paths agree with the mirror, and the digit scan's mul-sub sites pass full parity on Metal, so the miscompile is specific to these four prefilter sites under naga's MSL backend (or Apple's MSL compiler). The prefilter runs once per candidate, so pairing / with % here costs nothing measurable; the scan keeps its mul-sub form and its measured gains. One more entry for the naga-on-Metal ledger next to the select(ulong, uint, bool) division-guard bug. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
|
Found it, and I owe you a correction: my previous comment's diagnosis was The mechanism
I measured the driver's tolerance directly. The process aborts 2.06 s in: Time is the discriminator, not count. Same dispatch count, different bases:
That last row is the whole proof: 80 dispatches — more than twice what "fails at Three corrections to what I posted1. My threshold table is not reproducible. I reported 32 dispatches as 2. The task cap is not the bound that matters — the field is. This is the
The times are measured; the last column is measured at base 50 and inferred from 3. OS and driver are probably not confounded after all. Windows' default TDR The patchSubmit each batch on its own. --- a/common/src/cubecl_backend.rs
+++ b/common/src/cubecl_backend.rs
@@ -1117,6 +1117,30 @@ fn detailed_impl<R: cubecl::prelude::Runtime>(
);
}
+ // Submit this batch on its own, because a GPU driver kills a
+ // submission that runs too long. CubeCL aggregates dispatches into one
+ // command buffer and submits only at `CUBECL_WGPU_MAX_TASKS` tasks (32
+ // by default) or at the next read — a bound tuned for ML kernels that
+ // run in microseconds. This one is a persistent grid-stride kernel over
+ // a whole `CUBECL_BATCH_SIZE`, ~0.12 s per dispatch at base 50 on an
+ // AMD 860M, so without this flush a single submission holds the *whole
+ // field*: `DETAILED_SEARCH_MAX_FIELD_SIZE` is 20 batches, which is
+ // 2.4 s at base 50 and 6.3 s at base 80. Measured on Linux/RADV,
+ // amdgpu resets a gfx job at ~2 s ("ring gfx_0.0.0 timeout ... device
+ // wedged") and the device is lost for the rest of the process; Windows
+ // has the same 2 s TDR, and only survives because the GPUs it was
+ // tested on run a field inside it. So the task cap is not the bound
+ // that matters — the field is — and raising or lowering it does not
+ // help.
+ //
+ // Flushing per batch is the granularity the hand Vulkan backend has
+ // always used, and it is free: `flush` submits without waiting, and one
+ // batch already saturates the device (measured interleaved, 80 batches
+ // at base 50: 9.76 s median against 10.00 s for the aggregated form).
+ client
+ .flush()
+ .map_err(|e| anyhow::anyhow!("stream flush failed: {e:?}"))?;
+
undrained += 1;
if undrained == DRAIN_INTERVAL {
drain(hist_handle, &mut histogram)?;Cost: none. Interleaved, three pairs, base 50, 80 batches — new It is also robust against the env override: with niceonly was never affected, and the reason is worth knowingI flagged it as untested. It is fine — I ran 5e11-candidate fields both at the The reason is incidental rather than designed: Regression testThe fix is invisible to every existing test; here is one that catches it. --- a/common/src/cubecl_backend.rs
+++ b/common/src/cubecl_backend.rs
@@ -1583,6 +1607,52 @@ mod tests {
}
}
+ /// Many batches in one field must not lose the device. Without the
+ /// per-batch flush in `detailed_impl`, `CubeCL` packs persistent
+ /// grid-stride dispatches into one submission until it reaches
+ /// `CUBECL_WGPU_MAX_TASKS` (32), and the driver's watchdog resets the GPU
+ /// long before that — ~2 s on Linux/RADV, the same 2 s TDR on Windows.
+ ///
+ /// 40 batches is deliberately over that cap so the test also covers the
+ /// aggregated path, but it is not the interesting size: a real field is
+ /// `DETAILED_SEARCH_MAX_FIELD_SIZE` = 20 batches, under the cap and still
+ /// over the watchdog at every base past ~46.
+ ///
+ /// Skipped on software rasterizers: CI runs the ignored tests on lavapipe,
+ /// where 2e9 candidates would take hours and no watchdog is involved.
+ #[test]
+ #[ignore = "requires a wgpu device"]
+ fn cubecl_detailed_survives_more_batches_than_the_task_cap() {
+ let ctx = CubeclContext::new_default().expect("CubeCL init");
+ let device = ctx.device_name();
+ let squashed = device.to_lowercase().replace(' ', "");
+ let software = ["llvmpipe", "lavapipe", "swiftshader", "softwarerasterizer"];
+ if software.iter().any(|s| squashed.contains(s)) {
+ println!("skipping on software rasterizer: {device}");
+ return;
+ }
+
+ let base = 40u32;
+ let batches = 40u128;
+ let count = batches * CUBECL_BATCH_SIZE;
+ let start = crate::base_range::get_base_range_u128(base)
+ .unwrap()
+ .unwrap()
+ .range_start;
+ let range = FieldSize::new(start, start + count);
+ let results = process_range_detailed_cubecl(&ctx, &range, base).expect("cubecl run");
+
+ // Non-vacuous: every candidate lands in exactly one bin, so the
+ // histogram must account for the whole field. A submission the driver
+ // killed would take its batch's counts with it.
+ let counted: u128 = results.distribution.iter().map(|d| d.count).sum();
+ assert_eq!(
+ counted, count,
+ "{batches} batches on {device}: histogram covers {counted} of {count} candidates"
+ );
+ println!("{batches} batches ({count} candidates) survived on {device}");
+ }
+
/// CPU/CubeCL parity through the native CUDA runtime, on real silicon.
#[test]
#[cfg(feature = "cubecl-cuda")]Mutation-checked: with the The software-rasterizer skip is aimed at your CI: Rebased onto
|
…eview) Applied from Janzert's diagnosis on wasabipesto#99, with the regression test: CubeCL aggregates dispatches into one command buffer and submits at CUBECL_WGPU_MAX_TASKS (32) or the next read — so one submission held a whole field, and GPU watchdogs (amdgpu ~2s, Windows TDR 2s) reset the device on any GPU slow enough that a field outruns them: every detailed field past base ~46 wedged an AMD 860M on its first claim. flush() submits without waiting, so the drain pipelining is unchanged; measured free on the reporting machine. Niceonly is safe incidentally (its per-dispatch buffer writes flush the stream) — now documented at that site so pooling those buffers doesn't quietly inherit the bug. Also from the same review pass, two defects in b825b2e's conservation check: a stray duplicate of the ensure inside the drain branch compared a partially-drained histogram against the whole field (fired at 65+ batches; a replace-all editing accident), and the failure hint said 'install the CUDA toolkit' regardless of runtime — it now names the watchdog on wgpu. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Co-Authored-By: Janzert <janzert@janzert.com>
Detailed: cubecl-cuda -> cubecl -> cuda -> vulkan. The CubeCL family wins detailed on every vendor tested (1.03-1.07x over hand-CUDA on NVIDIA, 1.2-2.6x over hand-Vulkan elsewhere), and since b825b2e a toolkit-less cubecl-cuda init fails cleanly, so NVIDIA-without-toolkit falls through to wgpu speed instead of the hand-WGSL fallback. Niceonly: cuda -> cubecl -> vulkan. Hand-CUDA keeps a slim edge at the b50+ bases where long-run wall time concentrates (0.91-0.97 for cubecl-cuda); everywhere without a toolkit CubeCL is the best available: it wins RADV b50+ (1.28-1.38), wins NVIDIA-over-wgpu outright, and runs out of the box on Apple. The known concession is Apple niceonly via MoltenVK (0.55-0.79), which needs optional brew packages and stays reachable with --gpu-backend vulkan. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…gain The evaluation's wasm build fixes lived only in the throwaway probe crate; at head, cargo check -p nice_common --features cubecl --target wasm32-unknown-unknown failed in getrandom (no entropy backend named). Target-gated dependencies activate getrandom's wasm_js backend and cubecl-common's serde feature on wasm builds, tied to the cubecl feature; native builds are unaffected. First step of readying the core for the browser/WebGPU client. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
|
@Janzert thanks for the comments, feel free to test again and benchmark this version. I may remove the vulkan backend before merge based on performance and auditablity, unless you have hardware where it performs better. |
…ming Three changes so the coming wasm/WebGPU PR does not have to touch this backend again (its remaining work is the u32-only kernel variant, which can live in its own module, plus wasm-client plumbing): - CubeclContext::new_default_async over init_setup_async — the only form a browser permits (the sync init_setup panics on wasm by design). Both constructors share one process-wide client. - process_range_detailed_cubecl_async: the detailed path is now an async body using read_async (a wasm future cannot block); the sync entry is that future under block_on, so native device tests exercise the same code. Niceonly stays sync — its MSD thread pipeline is not browser scope. - Instant now comes from web-time, which re-exports std on native and works on wasm32, where std::time::Instant panics at runtime. Also --gpu-wgpu-device (env NICE_GPU_WGPU_DEVICE), from review item 3: --gpu-device indexes per-backend namespaces (CUDA ordinals != Vulkan ordinals != wgpu adapters), so on a multi-GPU box no single number is right for every backend. The flag takes CubeCL's own device spelling (DiscreteGpu(0), IntegratedGpu(1), Cpu, ...) and is translated to the CUBECL_WGPU_DEFAULT_DEVICE env var before init, keeping one parser. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
One place covering the GPU path: the four backends and what each needs at runtime per platform, auto's per-mode selection order and its explicit-is-fatal rule, multi-GPU selection (--gpu-device's per-backend namespaces and the typed --gpu-wgpu-device), and the troubleshooting answers for the failure modes we have actually seen — missing NVRTC, watchdog resets, MoltenVK setup, wrong laptop adapter, CPU fallback bases. The embedded --help output was three backends out of date; regenerated from the umbrella build. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
|
Looking good here, below are the latest numbers. My Claude is rather fixated on niceonly processing and really doesn't want to give it up to drop the vulkan backend. I'm not really sure it's worth keeping around for the relatively small AMD and Apple gains (and possibly Intel). All four boxes at
Detailed: CubeCL wins everywhere, 1.74–2.57. Niceonly: a split — both NVIDIA boxes have CubeCL ahead by 1.34–1.55, AMD has it behind, Intel can't resolve it. The Vulkan backend still performs better somewhere — AMD niceonly. And the specific claim in P.S. Part of Claude's response after reviewing the above:
Like I said it really likes niceonly. It also disputes my characterization that the Apple (and potentially the Intel) gains are relatively small. |
|
That's pretty funny. I'm still leaning towards removing the Vulkan backend, mainly because most of the compute I'm buying right now is from cuda-only markets (Vast and my local HPC). These advances are super cool and will allow some people with non-nvidia cards to contribute but I'm skeptical that we'll see more than two or three users per month using the non-cuda backends (just based on historic patterns). If you knew of a market where I could rent a dozen or so cheap preemptible AMD cards or mac minis it may be worth it to maintain a moderately-faster backend, but since most of the wall-time is going to hit cuda anyways I'm not so concerned about these cases. |
The gpu umbrella no longer includes vulkan: every platform it serves is covered by cubecl, which beats it in detailed mode on all vendors measured, so the standard GPU binary and the -gpu docker image drop the hand-WGSL backend and its ash/naga dependencies. The feature stays in the tree, marked experimental, for --features gpu,vulkan builds — to be promoted or removed once its remaining niceonly niches (Metal via MoltenVK, RADV at prefilter-heavy bases) stop mattering or get closed. No user-facing change otherwise: auto's order already ends with vulkan only in builds that include it, and CI's device-parity jobs name their features explicitly and still cover it. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
3 participants
Note
Stacked on #96 — the first 11 commits here are #96's (
vulkan-backend); only the last 6 are new. Review fromAdd a CubeCL detailed-mode backend for evaluation against the hand-written kernelsonward.What this is
An evaluation of CubeCL as the GPU kernel layer, motivated by auditability: the detailed-mode kernel is one
#[cube]Rust function replacing ~500 lines of WGSL string generation (and, potentially, the CUDA C codegen too). Same per-base constants fromgpu_config, same 4-copy workgroup histogram, same near-miss records, sameFieldResults. Per-base specialization is#[comptime]— CubeCL JITs one kernel per base at first launch, exactly like the NVRTC/naga paths.Two runtime targets behind one kernel source:
--gpu-backend cubecl— wgpu (Vulkan on Linux, Metal on macOS)--gpu-backend cubecl-cuda— native CUDA runtimeThe port also implements the per-device chunk-scan width #96 lists as future work: a comptime bool selects wide u64 chunks on CUDA vs split16 on wgpu (
c1ae5e5).Correctness
CPU parity tests (distribution + near-miss, bases 10/40/62/80) pass exactly on every device and runtime tested:
The hand-written backends' own device suites (detailed, niceonly, prefilter mirror) also pass on all of the above, so both arms of each benchmark were parity-checked on the same silicon.
Benchmarks
nice_client detailed --benchmark --benchmark-secs 60 --gpu, 3 interleaved rounds per arm (A,B,A,B,A,B), median-of-3. Spreads ≤0.8% everywhere.Context:
Arrays landing in local memory where the hand kernel forces scalar registers (checkable viaCUBECL_DEBUG_LOG=1).Findings along the way
988a902): Add Vulkan GPU backend #96's ash instance creation didn't opt intoVK_KHR_portability_enumeration, so the hand-Vulkan backend found no device on macOS at all. Fixed with runtime detection (no-op on Linux/Windows); all 6 device tests then pass through MoltenVK 1.4.2 on an M4, includingshaderInt64.SHADER_INT64— naga (wgpu 29) emits an ambiguousmetal::select(ulong, uint, bool)in its u64 checked-division guard. Split16 stays the wgpu flavor for a concrete reason; recheck on future naga releases.OnceLockguard around non-idempotent init, host API drifted from docs in a few places, and the wasm build needs three dependency-feature unifications. All solvable in minutes, none blocking, but worth knowing.Update 2026-08-17 — feature restructure, auto order, CI
a845244): per-backend flagscuda/vulkan/cubecl/cubecl-cudaunder a newgpuumbrella that enables all of them. Release recipe is two builds: default features = lightweight CPU-only binary,--features gpu= one binary that runs on CPU or any supported GPU. Verified the full umbrella builds on a machine with no CUDA toolkit. CUDA-specific_gpunames renamed to_cuda(client_process_cuda,CudaContext,process_*_cuda, …); backend-neutral names (gpu_config,gpu_niceonly,--gpu-backend) keep theirs.autonow ranks by measured speed: CUDA → CubeCL → Vulkan. In niceonly modeautoskips CubeCL (no niceonly kernel yet) and falls through to Vulkan; explicitly named backends still fail fatally rather than falling back.01d4763): newgpu-parity.ymlruns the device suites on every backend a hosted runner can drive — lavapipe on Linux (vulkan + cubecl) and macOS arm64 (cubecl via Metal, hand-WGSL via MoltenVK using the new portability path). The workflow header notes inline thatcuda/cubecl-cudacannot run in CI (NVIDIA silicon required); their NVRTC compile-tests still run in the normal suite.Update — niceonly ported
0c9665aports niceonly: the kernel does the same on-device candidate reconstruction, offset reduction, lane tiling, and low-digit modular prefilter as the generated WGSL, with the stride modulus/residue count as comptime parameters. The backend-neutral tiling helpers move fromvulkan/codegen.rsintogpu_niceonly(the Vulkan backend re-imports them). Device tests mirror the Vulkan suite — CPU parity at bases 10/12/25/40/45/62/80, 69-in-base-10 through the whole pipeline, a CUDA-runtime variant, and a prefilter probe against an independent u128 host mirror at every lane width (the v3.2.14 over-rejection class is untestable any other way). All pass on lavapipe; prefilter survivor counts match the hand-WGSL filter exactly.autodrops its niceonly carve-out.Update — round 2 benchmarks: detailed + niceonly on all three devices
Same interleaved 3-round methodology at head
16779af, parity suites first (all pass everywhere they can run;hardware.gpu_backendconfirmed in reports).Detailed reproduces round 1 on every device: 1.43/1.38 (9070 XT, vs hand-WGSL), 0.86 (M4, vs MoltenVK), 0.75 (4060 cubecl-cuda vs hand-CUDA).
Niceonly: the fresh port loses to the hand backends everywhere (cubecl/hand medians):
The 4060 gap is strongly base-dependent while hand-CUDA holds its rate; ranked suspects: per-dispatch
client.create()of the descriptor buffers in the CubeCLRangeSink(both hand backends reuse persistent buffers), the flag-guarded early exit vs the hand kernels' true earlyreturn, and the lane tiling vs hand-CUDA's fixed warp. Correctness unaffected — parity exact everywhere. This folds into the perf-chase item below, which is now clearly load-bearing for the niceonly half of the backend-lineup decision.Update — perf chase closed the gap; review blocker fixed
Four A/B-verified root causes on an RTX 4060, parity-checked at every step (
fee54a8…b825b2e): the per-field stride-table rebuild (now a per-base plan cache), the pass-loop copies in the full check, div+mod pairs (nvcc lowers each u64 constant division independently — remainders are now mul-sub like the hand kernels), and ptxas register-promoting the scan scratch into compare/select chains (now shared memory). Record run, 3 interleaved rounds:Review response (
b825b2e): the driver-without-toolkit blocker is fixed —new_cudaproves the runtime with a smoke kernel + verified readback before reporting success (verified both ways by stashinglibnvrtcon the test instance), and detailed mode checks histogram conservation so silently-dropped device work can never report results on any runtime. Also: standalone--features cubecl-cudacompiles, and the wgpu adapter log/telemetry includes the resolved graphics API (AutoGraphicsApimay pick DX12 on Windows).Note for the wgpu columns: the mul-sub and shared-memory changes affect WGSL codegen too and postdate the review's RADV/Intel runs — those numbers need a re-measure.
Remaining before this is mergeable
0c9665a) —RangeSinkpipeline, prefilter included, probe-tested;autocarve-out removed64c14a7,e493c90):autopicks the measured-fastest order per mode — detailedcubecl-cuda → cubecl → cuda, niceonlycuda → cubecl— and hand-Vulkan is now an experimental opt-in feature outside thegpuumbrella (every platform it serves is covered bycubecl; build with--features gpu,vulkanwhile its promote-or-remove fate is decided). Hand-CUDA stays as the niceonly-b50+ leader and independent oracle1f6260c):cuda-12080selects the compiled binding surface, not a driver floor — dynamic loading resolves symbols lazily, the hand-CUDA path calls nothing newer than 12.0, and cubecl-cuda's 12.8 declarations are never launched by our kernels (parity-verified on silicon). Rationale now lives in the Cargo.toml comments1f6260c):=0.10.0with the bump policy (deliberate, parity-suite-gated) in the manifest comment01d4763)--gpu-backenddocumentation (a845244)06ddd5e): detailed cubecl/hand-vulkan 1.90 (RX 9070 XT) and 1.17–1.20 (M4); niceonly wins RADV b50+ (1.28–1.38), trails RADV b40 (0.85–0.91) and Metal (0.55–0.79). The M4 prefilter probe caught a second Metal-only naga-MSL miscompile (mul-sub in the prefilter sites; reverted there,06ddd5e)ffe33a6), plus the two defects it found inb825b2e's conservation check (duplicated ensure in the drain branch; runtime-appropriate failure hint)--gpu-devicefor wgpu (70d753e): new--gpu-wgpu-device/NICE_GPU_WGPU_DEVICEtakes CubeCL's own device spelling (DiscreteGpu(0),IntegratedGpu(1),Cpu, …) — a separate selector because--gpu-deviceindexes per-backend namespaces and no flat ordinal is right for every backend on a multi-GPU box; the chosen adapter + graphics API are always logged8a132db,70d753e): wasm32 dependency unification (the graph had regressed),new_default_async(browsers cannot init synchronously), async detailed path viaread_async(sync entry wraps the same body, so native device tests cover it),web-timeInstant. The coming wasm/WebGPU PR should only need its u32-only kernel variant (own module) andwasm-clientplumbing-gpuimage with all backends, no CUDA-only variant (simplicity beats 8 MiB)SHADER_INT64in WebGPU; the evaluation probe proved parity for that shape) +wasm-clientintegration — core prep above means it should not touch the backend filesBenchmark raw data (per-round stdout + JSON reports for all three devices) is archived outside the repo; happy to attach it here if useful.