[inureyes]

Summary

Windows --features cuda binaries now load every CUDA, cuDNN and NVRTC DLL from one runtime directory and can run from a relocatable bundle, following the ROCm design of #2606/#2607. No NVIDIA GPU was available, so this PR is verified by building and by GPU-less runtime checks only (see below); runs on an NVIDIA host are still open.

• overlays/cuda/.../runtime_dir.h (new, header-only, host-tested): choose_cuda_runtime_dir() picks MLXCEL_CUDA_RUNTIME_DIR (absolute, must hold the runtime, no fallback), else the executable directory when it holds the CUDA runtime DLL named in the binary's own delay-import table, else the build-time layout. nvcc links the static cudart, so the marker is cublasLt64_13.dll.
• delayload.cpp: in a bundle, the DLLs the libraries open by plain name (nvrtc64_*, nvrtc-builtins*, cublas64_*, cublasLt64_*, nvJitLink_*, seven required cudnn_* sub-libraries; cusparse64_* checked only) are checked before any CUDA DLL loads and preloaded by full path; a failed load raises the delay-load exception naming the expected path instead of returning null. The build-time path keeps upstream's per-library lookup and its fallback; it only adds the [mlx-cuda] CUDA runtime: line, plus a message naming the missing marker path when a bundle-looking exe directory lacks it.
• jit_module.cpp: NVRTC headers, CCCL and CUTLASS come from the bundle's include\ (in or beside the runtime directory), ranked after CUDA_HOME and ahead of CUDA_PATH.
• scripts/package_windows_cuda_runtime.py: stages or zips bin\ (DLLs + OpenBLAS), include\ (JIT headers, CCCL from the build's _deps, CUTLASS), licenses\ (CUDA EULA, cuDNN, CCCL, CUTLASS); fails naming every missing DLL, header or license. Licenses are told apart by content, because the redistributable archives all ship a LICENSE that overwrites the others when unpacked into one directory.
• make verify-cuda-runtime-dir (in make verify), docs in installation.md and environment-variables.md, overlay README record.

Decisions and open points

• The DLL closure is derived statically (import tables and embedded DLL names of the x64 CUDA 13.4.3 / cuDNN 9.27 DLLs), not from a loaded-module list of a generate run; to be confirmed on hardware.
• The JIT header set was found by compiling each of the 17 headers MLX embeds for the JIT with NVRTC 13.4 against only the staged bundle's include\ (needs no GPU). That added cuda_runtime.h, mma.h and crt/mma.h(pp), which CUTLASS includes. Several of these headers (cooperative_groups.h, math_constants.h, cuda_runtime.h, mma.h, ...) are not on the CUDA EULA Attachment A list; the docs say so and advise confirming terms with NVIDIA before publishing a bundle. Prebuilt archives stay out of scope.
• delayload.cpp is 595 lines, over the 500-line guideline; it is an overlay of an upstream file (like jit_module.cpp at 805), so it was not split.

Changes during review

• CUDA_HOME now really outranks the bundle's headers: the bundle include\ was placed ahead of the chosen CUDA header dir on the NVRTC include list, so NVRTC took every CUDA header from the bundle even with CUDA_HOME set. With CUDA_HOME set the bundle include now goes last and only supplies CUTLASS/CuTe (implementation review, HIGH).
• cusparse64_*, which only cuSOLVER imports, is checked before any CUDA DLL loads, so a missing cuSPARSE is reported under its own name instead of as cuSOLVER (implementation review).
• MLXCEL_CUDA_RUNTIME_DIR is normalized (./.., trailing separator) before use; C:\b\bin\ previously made the header lookup miss the bundle's include\ and fall back to CUDA_PATH (security review).
• The baked-mode cuDNN log line no longer converts through the ANSI code page unguarded inside the delay-load hook (security review).
• Documented, not fixed: the CUDA libraries also open nvcuda.dll, nvcudart_hybrid64.dll, zlibwapi.dll, cask_profile_interface.dll, nvtileiras64.dll and the two optional cuDNN sub-libraries by plain name; the bundle neither ships nor checks them, so the standard search order (application directory, System32, then current directory and PATH) applies, as it does in developer builds. Closing that needs a decision and is left to the DLL-planting follow-up (fix(security): Close remaining ROCm DLL and JIT-cache planting paths #2641).
• On Linux the only behavior difference is that an empty CUDA_HOME now counts as unset.

Test plan

Toolkit: NVIDIA redistributables for windows-x86_64, CUDA 13.4.3 (cudart/nvcc/crt/nvrtc/nvjitlink/nvvm 13.4.92, cccl 13.3.4.3.1, cublas 13.8.0.4, cufft 12.4.0.43, cusolver 12.3.5.3, cusparse 12.8.6.72) and cuDNN 9.27.0.42 for CUDA 13, sha256-checked. Host: Windows 11 x64, Radeon 8060S, no NVIDIA GPU.

• x64 cargo build --release --features cuda (MLX_CUDA_ARCHITECTURES=120): cold build 32 min, rebuilt after the review fixes (17 min 37 s, mlxcel.exe 133,756,928 bytes) with the harness re-run, no build-glue change needed; mlxcel.exe 128 MiB, mlxcel-server.exe 118 MiB. llvm-readobj --coff-imports: cuBLASLt, cuFFT, cuSOLVER, cuDNN and NVRTC are DelayImport, no CUDA DLL in the regular import table.
• cargo clippy --release --features cuda -- -D warnings (x64), cargo fmt --all --check, cargo test -p mlxcel-mlx-pin --lib (49 passed).
• make verify-cuda-runtime-dir host test (clang++), checked to fail on mutated selection logic; 25 unit tests in tests/test_package_windows_cuda_runtime.py; env-var docs check; license header check.
• Packaging on the real archives: 1830 files, 1968 MiB; fails naming the file when one is renamed, and refuses cuDNN's LICENSE as the CUDA EULA.
• Hook exercised with no GPU: a harness linking this delayload.cpp unchanged calls cublasLtGetVersion, nvrtcVersion, cufftGetVersion and cudnnGetVersion through delay-loaded DLLs. With the toolkit renamed away and CUDA_PATH unset, every CUDA/cuDNN/NVRTC module loads from the bundle (exe directory and MLXCEL_CUDA_RUNTIME_DIR), nvcuda.dll from System32; relative and wrong-folder overrides, a removed delay-loaded DLL, a removed cudnn_ops and a removed marker each fail with the expected path. With the toolkit present the harness loads from the toolkit as before.
• mlxcel generate and mlxcel-server from the bundle (toolkit hidden) and the developer build both stop with 0xC0000409 at the first device query after Detected 1 GPU(s)., before any delay-loaded CUDA DLL is requested (this host has a stale nvcuda.dll but no NVIDIA GPU). Everything after device initialization is unverified.
• On a Windows NVIDIA host (RTX Spark N1X or x64): generate (including the first NVRTC JIT) and the server from a fresh bundle folder with the toolkit moved away; module list; removed-DLL failure; developer build unchanged.

Closes #2640

[inureyes]

Implementation Review Summary

Intent

Windows CUDA builds load every CUDA, cuDNN and NVRTC DLL from one runtime directory and run from a relocatable bundle built by scripts/package_windows_cuda_runtime.py, mirroring the ROCm design of #2606/#2607, with developer builds and Linux unchanged.

Findings Addressed

• With CUDA_HOME set, NVRTC still took every CUDA header from the bundle, because the bundle's include was added as an include path ahead of the CUDA header directory. The docs and choose_cuda_include_dir() say CUDA_HOME outranks the bundle. With CUDA_HOME set, the bundle's include now goes last and only supplies CUTLASS/CuTe. Include paths on Linux and in builds without a bundle are unchanged (HIGH, ae955bc)
• A bundle without cusparse64_12.dll (imported only by cuSOLVER) failed with a message naming cusolver64_12.dll instead of the missing file. cuSPARSE is now part of the up-front check (checked, not preloaded), and the host test covers it (MEDIUM, ae955bc)

Remaining Items

• The packaging script requires cudnn_ext64_* and cudnn_engines_tensor_ir64_*, which the hook treats as optional, so packaging an earlier cuDNN 9 fails. This is conservative and documented as the cuDNN 9.27 closure (LOW), left as is
• Runs on NVIDIA hardware: generate with the first NVRTC JIT compile, the server, the module list, and an unchanged developer build. Needs an RTX Spark N1X or an x64 NVIDIA host; this is not a failure of the PR

Verification

• All stated requirements implemented (hardware criteria verified only by building and with a harness)
• No placeholder/mock code remaining
• Integrated into project code flow (delay-load hook, NVRTC include path, make verify)
• Project conventions followed
• Existing modules reused where applicable
• No unintended structural changes (Linux CUDA paths unchanged)
• Tests pass: host test, 25 packaging tests, env-var docs check, x64 cargo build --release --features cuda (15m 50s incremental), and a hook harness with the toolkit hidden. Bundle, MLXCEL_CUDA_RUNTIME_DIR, relative and wrong-folder overrides, a missing cuFFT and a missing cuSPARSE all give the expected path or message, and the developer layout still loads from the toolkit

[inureyes]

Security and performance review

No CRITICAL or HIGH findings. Fixed in d73cb01:

• MEDIUM: MLXCEL_CUDA_RUNTIME_DIR with a trailing separator (C:\b\bin\) made parent_path() return bin itself, so the NVRTC JIT missed the bundle's include\ and fell back to CUDA_PATH or the baked toolkit (bundle NVRTC paired with other headers). The override is now normalized (., .., trailing separator); host test added, fails without the fix.
• MEDIUM (docs): the libraries open nvcuda.dll, nvcudart_hybrid64.dll, zlibwapi.dll, cask_profile_interface.dll and nvtileiras64.dll themselves (embedded names), plus the optional cudnn_ext64_9 / cudnn_engines_tensor_ir64_9 when a bundle lacks them. The hook neither ships nor checks these, so the standard search (application directory, System32, then the current directory and PATH) applies. docs/installation.md now says so.
• LOW: the baked-mode cuDNN log line converted a narrow path inside the hook's static initializer; now guarded against throwing.

Checked and fine: the regular-import closure of every bundle DLL stays inside the bundle or the system; no fallback after a bad override or an incomplete bundle; the upstream return-null path is reachable only in the build-time layout; the env override and exe-dir candidates add no planting surface beyond the application directory Windows already searches first; NVRTC include order; the packaging script's destinations come from fixed names, with no traversal, and the zip is written through .partial.

Remaining (follow-up, not fixed here):

• Plain-name loads outside the bundle can reach the current directory. Closing this needs either a packaging manifest that the hook enforces, for the optional cuDNN sub-libraries, or a process-wide SetDllDirectoryW(L""), a narrower relative of the settled no-SetDefaultDllDirectories decision. A planting probe could not be run here, and none of it is verified on an NVIDIA host.
• LOW: --stage-dir leaves a partial tree on failure; rglob over _deps follows directory symlinks on Python 3.12 and older; the cuDNN license marker cuDNN is loose but does not collide with the CUDA EULA at 13.4.

Verification: host test, 25 packaging tests, env-var docs check, x64 cargo build --release --features cuda (17m 37s; same five DelayImports), and the hook harness with the patched delayload.cpp. The exe-dir bundle and the trailing-separator override both load every CUDA DLL from the bundle, and a wrong-folder override fails naming it.