CUDA backend fixes: CUBIN emission (toolkit newer than driver) + launch-bounds honoring for NULL localSize

[veljjanoski]

Context

We ported the Aevum CUDA backend (make CUDA=1) to Windows (MSYS2/MinGW + CUDA 13.3 on an
RTX 3090, driver 591.86 = CUDA 13.1 runtime) and hit two genuine bugs, fixed in the attached
patch. With both fixes the CUDA backend runs a full PRP correctly — interim res64 values are
bit-identical to the OpenCL backend at every logged iteration.

(Benchmark note, honest: on Ampere the CUDA path measured ~4% SLOWER than OpenCL —
749 vs 717 us/it at p=146,313,533 — both stacks share NVIDIA's PTX->SASS compiler, so the
port is about portability/debuggability, not speed.)

Bug 1 — PTX version mismatch when NVRTC is newer than the driver

NvrtcProgram::compile always emits PTX. With CUDA toolkit 13.3 against a 13.1 driver,
cuModuleLoadData(Ex) fails: CUDA_ERROR_UNSUPPORTED_PTX_VERSION (222),
"Unsupported .version 9.3; current version is '9.1'". Since the code already compiles for a
real arch (--gpu-architecture=sm_NN), the fix is to prefer nvrtcGetCUBIN (native SASS,
no driver JIT, no PTX version coupling) and fall back to PTX only if no CUBIN is available.

Bug 2 — NULL localSize launches ignore launch_bounds

clEnqueueNDRangeKernel defaults to 256 threads/block when localSize == NULL. Kernels
translated from KERNEL(64) carry __launch_bounds__(64); launching them at 256 fails with
CUDA_ERROR_INVALID_VALUE (observed on transposeIn: maxThreads=64, launched at 256).
OpenCL semantics require the implementation-chosen size to satisfy reqd_work_group_size, so
the fix uses the kernel's known workgroup size for NULL launches. A companion fix derives
reqWorkGroupSize via cuFuncGetAttribute(CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK) when the
module is a CUBIN (the existing .maxntid text parse only works on PTX modules).

The patch also upgrades the launch-failure log line to print blocks/threads and the kernel's
compiled maxThreads/regs/shmem — this is what made Bug 2 diagnosable.

Build note for Windows (may be worth a README line)

make CUDA=1 ... COMMON_FLAGS=... overrides must preserve $(CUDAFLAGS), otherwise
-DCUDA_BACKEND is silently dropped from KernelCompiler.cpp and the OpenCL-vs-CUDA
header-array convention (skip-first-file) desynchronizes, producing ~100 confusing NVRTC
errors ("uint undefined", duplicate GF61 definitions) because base.cl gets registered under
the opencl_compat.cuh header slot.

[cherubrock-seb]

Thanks for the detailed report, the correctness comparison against OpenCL, and the honest benchmark results.

Both reported bugs look genuine, and I agree with the general direction of the patch:

• preferring a native CUBIN when NVRTC is compiling for a real sm_XX architecture is the right way to avoid a newer-PTX/older-driver mismatch;
• the current hard-coded 256-thread fallback for localSize == NULL is incorrect for kernels originating from KERNEL(64).

Before merging, could you please make two small adjustments?

1. Please do not use CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK as the exact required workgroup size.

That attribute is the maximum launchable block size for the compiled function and device, not necessarily the original OpenCL reqd_work_group_size.

The wrapper already carries preprocessedSource into the linked program. Could the exact KERNEL(N) / __launch_bounds__(N) value be parsed from that source and stored in reqWorkGroupSize? CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK could then be used to validate that the compiled kernel supports at least that value.

2. Please preserve the existing --maxrregcount behavior.

The current CUDA wrapper injects .maxnreg into the generated PTX. Once the result is a CUBIN, that PTX rewrite no longer works. The safest solution may be to keep using PTX whenever maxregcount != 0, unless the CUBIN path is separately verified to honor the requested register limit.

A compile-time fallback for NVRTC versions without nvrtcGetCUBIN would also be useful if practical.

With those points addressed, I would be happy to merge this. The Windows CUDA correctness results and the improved launch diagnostics are especially valuable. Thanks again for the contribution.

[veljjanoski]

Thanks for the review, both points addressed in the new commit:

1. reqWorkGroupSize now parses the exact launch_bounds(N) from the carried preprocessedSource (with the PTX .maxntid parse as a fallback). CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK is only used to validate/clamp, except for source-less programs (aux kernels loaded via the binary-cache path never populate preprocessedSource), where it remains the last-resort value.

2. The CUBIN fast path is now skipped whenever the compile options request --maxrregcount, preserving the existing .maxnreg PTX-rewrite enforcement (confirmed live that some kernels are register-capped in normal runs, your concern was real, the original patch would have silently dropped the cap). nvrtcGetCUBIN is additionally compile-time guarded for CUDA < 11.2.

Possible future refinement: --maxrregcount only takes effect at the PTX→SASS step (PTX has virtual registers), which is why the .maxnreg rewrite is needed on the PTX path. Since nvrtcGetCUBIN runs ptxas internally, passing --maxrregcount as an NVRTC option should apply the cap directly to the CUBIN, if that checks out on a matched toolkit/driver, the CUBIN path could be re-enabled for capped kernels with a post-load NUM_REGS validation. Left out here pending that verification to keep the change minimal.