Windows owned-CUDA worker: delay-loaded cuBLASLt + sidecar packaging gate

.github/workflows/tests.yml

@@ -812,49 +812,19 @@ jobs:
           "@ | Set-Content build-gates/windows-x86_64-msvc-owned-cuda.txt
           if ($status -ne 0) { exit $status }
 
-      # Hollow-green guard for the compiled backend, with no PE-inspection
-      # dependency. `--features cuda` is the only thing that makes the build
-      # link cudart/cublas, and on Windows those are LOAD-TIME imports: an
-      # exe that imports them cannot start when they are absent from PATH
-      # (STATUS_DLL_NOT_FOUND, 0xC0000135) even though its --version path
-      # never calls into them. So the same probe answers both halves:
-      #   off-PATH must FAIL with 0xC0000135 -> the CUDA backend is baked in
-      #   on-PATH  must PASS                  -> the resolved DLL set is right
-      # A worker that silently compiled CPU-only prints --version both times,
-      # which this step refuses. It also records, in CI, the sidecar fact the
-      # release/installer PR must solve: the worker is not self-contained.
-      # Caveat: the discriminator depends on cudart/cublas staying load-time
-      # imports; if a future build moves them to delay-load, the off-PATH run
-      # exits 0 and this step reports "cuda feature is not compiled" — a false
-      # failure that points at the real cause rather than hiding it.
-      - name: Assert the packaged worker really carries the CUDA backend
+      # Derive sidecars from the same verified runtime components used to build.
+      # No independently maintained DLL manifest; nvcc and driver files are not
+      # redistribution inputs. The packager carries component licenses/hashes.
+      - name: Package and verify Windows CUDA sidecars
         working-directory: synapse
         shell: pwsh
         run: |
-          $ErrorActionPreference = 'Continue'
-          $exe = (Resolve-Path 'target/release/ck-synapse-worker-cuda.exe').Path
-          $isolated = Join-Path $env:RUNNER_TEMP 'cuda-isolated'
-          New-Item -ItemType Directory -Force $isolated | Out-Null
-          Copy-Item $exe $isolated
-          # PATH narrowed to the OS directories: no CUDA toolkit, no sidecars.
-          $env:PATH = "C:\Windows\system32;C:\Windows"
-          & (Join-Path $isolated 'ck-synapse-worker-cuda.exe') --version 2>&1 | Out-Null
-          $isolatedCode = $LASTEXITCODE
-          if ($isolatedCode -eq 0) {
-            throw "owned-CUDA worker ran without CUDA DLLs on PATH (exit 0): the cuda feature is not compiled into this binary — refusing to record a hollow-green gate"
-          }
-          if ($isolatedCode -ne -1073741515) {
-            throw "owned-CUDA worker failed off-PATH with exit $isolatedCode, expected -1073741515 (0xC0000135 STATUS_DLL_NOT_FOUND)"
-          }
-          # CUDA 13's redist layout keeps the runtime DLLs under bin\x64 (the
-          # nvcc drivers sit in bin); both directories are needed, as verified
-          # locally. With them restored the worker must start.
-          $env:PATH = "$env:CUDA_PATH\bin\x64;$env:CUDA_PATH\bin;C:\Windows\system32;C:\Windows"
-          & $exe --version
-          if ($LASTEXITCODE -ne 0) {
-            throw "owned-CUDA worker --version failed with CUDA on PATH (exit $LASTEXITCODE)"
-          }
-          "cuda_import_probe=ok off_path=$isolatedCode on_path=0" >> $env:GITHUB_STEP_SUMMARY
+          ./scripts/package-owned-cuda.ps1 `
+            -Worker target/release/ck-synapse-worker-cuda.exe `
+            -RuntimeComponents @("$env:RUNNER_TEMP/x-cuda_cudart", "$env:RUNNER_TEMP/x-libcublas") `
+            -Output build-gates/owned-cuda-windows-x64.zip
+          ./scripts/test-owned-cuda-package.ps1 -Archive build-gates/owned-cuda-windows-x64.zip
+          'cuda_package_probe=passed; GPU execution requires a GPU runner' >> $env:GITHUB_STEP_SUMMARY
 
       - name: Retain Windows owned-CUDA gate evidence
         if: always()

README.md

@@ -69,3 +69,58 @@ Example user-tier `~/.config/cortexkit/synapse.jsonc` (project configs must omit
 
 Tests can point at a file with `SYNAPSE_CONFIG_PATH`. Only one synapse module
 per machine (singleton lease); a second instance refuses to start.
+
+### Owned-CUDA hardware floor
+
+`ck-synapse-worker-cuda` implements `--probe-floor` (hidden, like the
+`--test-abort*` surfaces). It prints one JSON object and exits 0:
+
+```json
+{"driver_api": 13030, "compute_capability": {"major": 8, "minor": 9}}
+```
+
+The module probes the configured `worker_bin`, the engine's worker-binary
+environment override, or the sibling `ck-synapse-worker-cuda`, in that order.
+It caches one result per process unless both environment readings parse
+successfully. The child wait is bounded to 10 seconds; stdout is capped at
+4096 bytes and each pipe completion wait is bounded to another 100 ms.
+A missing binary, non-zero exit, timeout, or invalid output produces
+`HardwareUnavailable`. Refusal and model evidence carry diagnostic context,
+including the last 4096 bytes of stderr when available, under `observed`.
+Failed probes do not fabricate numeric hardware readings.
+
+The environment overrides the probe only as a complete, parseable pair.
+Otherwise both readings come from the probe; partial overrides are not merged:
+
+- `SYNAPSE_CUDA_DRIVER_API` (alias `CUDA_DRIVER_API`) — the raw CUDA **driver
+  API** integer from `cuDriverGetVersion()`, not the marketing driver version.
+  For example, a measured driver API value is `13030`. `610.88` is not a valid
+  API integer; without a parseable alias, it causes fallback to the probe.
+- `SYNAPSE_CUDA_COMPUTE_CAPABILITY` (alias `CUDA_COMPUTE_CAPABILITY`) — device
+  0's compute capability as `major.minor`, for example `8.9`.
+- `SYNAPSE_CUDA_PACKAGING_DRIVER` — optional; the driver string a packaging
+  build was tested against, carried into the refusal for diagnostics.
+
+### Windows owned-CUDA package
+
+The manual Windows CUDA gate packages the worker with runtime DLLs derived
+from the same pinned `cuda_cudart` and `libcublas` redistribution archives
+used for compilation. `scripts/package-owned-cuda.ps1` places the executable
+and DLLs at the ZIP root, includes component licenses, and records source
+components and SHA-256 hashes in `manifest.json`. The NVIDIA driver is not
+bundled and must already be installed.
+
+Windows worker builds require CUDA 13. The build script rejects other toolkit
+major versions before linking, because this package resolves CUDA 13 DLL names.
+
+`scripts/test-owned-cuda-package.ps1 -Archive <zip> -RequireGpu` extracts a
+fresh copy, verifies hashes, and checks no-sidecar `--version`, actionable
+missing-library refusal, and a real hardware-floor probe with adjacent DLLs
+and CUDA removed from PATH. Without `-RequireGpu`, a runner without NVIDIA
+hardware may report an explicit driver/device refusal; this is not a GPU
+execution pass. Neither mode loads model weights or certifies embeddings.
+
+The worker delays its cuBLASLt import and checks runtime library loading
+before CUDA calls. No global PATH changes or extra DLL search directories
+are needed. Release-matrix publication remains separate from this manual
+gate artifact.

crates/synapse-engine-cuda/src/cuda.rs

@@ -60,8 +60,46 @@ mod enabled {
         context: NonNull<c_void>,
     }
 
+    /// Keep successfully loaded libraries resident for all subsequent FFI calls.
+    /// Resolve explicitly before touching a delay import so failure is a Rust
+    /// error, not an unhandled Windows loader exception.
+    fn ensure_libraries_loaded() -> Result<()> {
+        #[cfg(target_env = "msvc")]
+        {
+            use std::sync::LazyLock;
+            static LOADED: LazyLock<Result<(), String>> = LazyLock::new(|| {
+                #[link(name = "kernel32")]
+                unsafe extern "system" {
+                    fn LoadLibraryW(name: *const u16) -> *mut c_void;
+                }
+                for name in [
+                    "cublasLt64_13.dll",
+                    "cublas64_13.dll",
+                    "cudart64_13.dll",
+                    "nvcuda.dll",
+                ] {
+                    let wide: Vec<u16> = name.encode_utf16().chain(Some(0)).collect();
+                    // Windows searches the executable directory and installed
+                    // system locations without changing PATH or global policy.
+                    if unsafe { LoadLibraryW(wide.as_ptr()) }.is_null() {
+                        return Err(format!(
+                            "cannot load CUDA library {name}: {}",
+                            std::io::Error::last_os_error()
+                        ));
+                    }
+                }
+                Ok(())
+            });
+            if let Err(message) = &*LOADED {
+                anyhow::bail!("{message}");
+            }
+        }
+        Ok(())
+    }
+
     impl DeviceBinding {
         fn capture() -> Result<Self> {
+            ensure_libraries_loaded()?;
             cuda_driver_check(unsafe { cuInit(0) }, "cuInit")?;
             let mut runtime_device = 0;
             cuda_runtime_check(
@@ -110,12 +148,59 @@ mod enabled {
     }
 
     pub fn ensure_available() -> Result<()> {
+        ensure_libraries_loaded()?;
         cuda_driver_check(unsafe { cuInit(0) }, "cuInit")?;
         let version = unsafe { synapse_cuda_cublaslt_version() };
         ensure!(version > 0, "cuBLASLt did not report a version");
         Ok(())
     }
 
+    /// Read the driver API version and device 0's compute capability.
+    ///
+    /// This runs before an owned-CUDA load is admitted, so it deliberately
+    /// touches nothing else: no context is retained, no model is loaded, and
+    /// no weights are mapped. The reading is what the floor predicate is
+    /// applied to, which is why it reports the raw numbers rather than a
+    /// verdict.
+    pub fn probe_hardware_floor() -> Result<crate::HardwareFloorProbe> {
+        ensure_libraries_loaded()?;
+        cuda_driver_check(unsafe { cuInit(0) }, "cuInit")?;
+        let mut driver_api = 0;
+        cuda_driver_check(
+            unsafe { cuDriverGetVersion(&mut driver_api) },
+            "cuDriverGetVersion",
+        )?;
+        let mut device = 0;
+        cuda_driver_check(unsafe { cuDeviceGet(&mut device, 0) }, "cuDeviceGet")?;
+        let mut major = 0;
+        cuda_driver_check(
+            unsafe {
+                cuDeviceGetAttribute(
+                    &mut major,
+                    CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,
+                    device,
+                )
+            },
+            "cuDeviceGetAttribute(COMPUTE_CAPABILITY_MAJOR)",
+        )?;
+        let mut minor = 0;
+        cuda_driver_check(
+            unsafe {
+                cuDeviceGetAttribute(
+                    &mut minor,
+                    CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR,
+                    device,
+                )
+            },
+            "cuDeviceGetAttribute(COMPUTE_CAPABILITY_MINOR)",
+        )?;
+        Ok(crate::HardwareFloorProbe {
+            driver_api: driver_api as u32,
+            compute_major: major as u32,
+            compute_minor: minor as u32,
+        })
+    }
+
     pub struct MiniLmContext {
         binding: DeviceBinding,
         raw: NonNull<c_void>,
@@ -465,8 +550,16 @@ mod enabled {
         }
     }
 
+    /// `CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR` from `cuda.h`.
+    const CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR: i32 = 75;
+    /// `CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR` from `cuda.h`.
+    const CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR: i32 = 76;
+
     unsafe extern "C" {
         fn cuInit(flags: u32) -> i32;
+        fn cuDriverGetVersion(version: *mut i32) -> i32;
+        fn cuDeviceGet(device: *mut i32, ordinal: i32) -> i32;
+        fn cuDeviceGetAttribute(value: *mut i32, attrib: i32, device: i32) -> i32;
         fn cuCtxGetDevice(device: *mut i32) -> i32;
         fn cuCtxSetCurrent(context: *mut c_void) -> i32;
         fn cuDevicePrimaryCtxRetain(context: *mut *mut c_void, device: i32) -> i32;
@@ -551,6 +644,10 @@ mod enabled {
         bail!("owned CUDA requires a non-macOS build with cargo feature `cuda`")
     }
 
+    pub fn probe_hardware_floor() -> Result<crate::HardwareFloorProbe> {
+        bail!("owned CUDA requires a non-macOS build with cargo feature `cuda`")
+    }
+
     pub struct MiniLmContext;
     impl MiniLmContext {
         pub fn new(_graphs: bool) -> Result<Self> {
@@ -629,4 +726,6 @@ mod enabled {
     }
 }
 
-pub use enabled::{ensure_available, MiniLmContext, ModernBertContext, Qwen3Context};
+pub use enabled::{
+    ensure_available, probe_hardware_floor, MiniLmContext, ModernBertContext, Qwen3Context,
+};

crates/synapse-engine-cuda/src/lib.rs

@@ -18,6 +18,8 @@ use synapse_core::{
 mod cuda;
 mod model;
 
+pub use cuda::probe_hardware_floor;
+
 pub const ENGINE_VERSION: &str = "owned-cuda-v1";
 /// The source revision from which the CUDA kernels were ported.
 pub const KERNEL_REVISION: &str = "4d0ded67c30286fe2be37cc7413359ad745dd751";
@@ -183,6 +185,17 @@ pub fn build_identity(family: ModelFamily, dtype: StorageDType) -> CudaBuildIden
     }
 }
 
+/// A hardware-floor reading taken before any owned-CUDA worker is spawned.
+///
+/// Carried separately from [`device_meets_floor`] so the caller can log or
+/// refuse on the observed values rather than on a bare boolean.
+#[derive(Clone, Copy, Debug, Eq, PartialEq)]
+pub struct HardwareFloorProbe {
+    pub driver_api: u32,
+    pub compute_major: u32,
+    pub compute_minor: u32,
+}
+
 /// Hardware-floor predicate used by capability probes before worker creation.
 #[must_use]
 pub fn device_meets_floor(driver_api: u32, compute_major: u32, compute_minor: u32) -> bool {