From 749ff87476ea11af176c844d99bed6f3cb76a13a Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Tue, 21 Jul 2026 12:42:05 -0700
Subject: [PATCH 01/14] Harden program-cache perms, binary-path resolution, and
 coredump helper lifetime

- cuda.core: create the on-disk program cache tree owner-only (0o700) and
  re-assert restrictive perms on POSIX, so cached device code cannot be read
  or planted by other local users regardless of the inherited umask.
- cuda.pathfinder: absolutize the resolved binary-utility path to honor the
  documented absolute-path contract; surface AddDllDirectory failures on
  Windows with GetLastError instead of silently swallowing them.
- cuda.bindings: retain the caller's bytes in _HelperCUcoredumpSettings so the
  borrowed pointer cannot outlive its backing buffer, and free the getter's
  1 KiB buffer in __dealloc__; clarify get_buffer_pointer's lifetime contract.

Adds regression tests for the cache permissions, path absolutization, and
coredump helper lifetime.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 .../cuda/bindings/_internal/utils.pyx         | 13 ++++++-
 cuda_bindings/cuda/bindings/_lib/utils.pxd    |  3 ++
 cuda_bindings/cuda/bindings/_lib/utils.pxi    | 12 +++++-
 cuda_bindings/tests/test_cuda.py              | 35 +++++++++++++++++
 .../core/utils/_program_cache/_file_stream.py | 18 +++++++--
 cuda_core/tests/test_program_cache.py         | 39 +++++++++++++++++++
 .../_binaries/find_nvidia_binary_utility.py   |  6 ++-
 .../_dynamic_libs/load_dl_windows.py          | 16 +++++++-
 .../tests/test_find_nvidia_binaries.py        | 25 ++++++++++++
 9 files changed, 160 insertions(+), 7 deletions(-)

diff --git a/cuda_bindings/cuda/bindings/_internal/utils.pyx b/cuda_bindings/cuda/bindings/_internal/utils.pyx
index a5931f1f080..73c4e223a20 100644
--- a/cuda_bindings/cuda/bindings/_internal/utils.pyx
+++ b/cuda_bindings/cuda/bindings/_internal/utils.pyx
@@ -42,7 +42,18 @@ cdef bint is_nested_sequence(data):
 
 
 cdef void* get_buffer_pointer(buf, Py_ssize_t size, readonly=True) except*:
-    """The caller must ensure ``buf`` is alive when the returned pointer is in use."""
+    """Return a raw pointer into ``buf``.
+
+    The caller must ensure ``buf`` is alive AND not reallocated/resized while
+    the returned pointer is in use. For buffer-protocol objects the export lock
+    (``PyBuffer_Release``) is dropped before this function returns, so the
+    stack-local ``Py_buffer`` view cannot keep the export alive past return.
+    Concurrently resizing a mutable exporter (e.g. ``bytearray``, ``array.array``,
+    a growable NumPy view) during a ``nogil`` native call that dereferences the
+    pointer is therefore a use-after-free (#366). Prefer immutable buffers
+    (``bytes``), or hold the buffer export for the pointer's full lifetime at the
+    call site.
+    """
     cdef void* bufPtr
     cdef int flags = cpython.PyBUF_ANY_CONTIGUOUS
     if not readonly:
diff --git a/cuda_bindings/cuda/bindings/_lib/utils.pxd b/cuda_bindings/cuda/bindings/_lib/utils.pxd
index 24b0ae8de93..7a846d21fb9 100644
--- a/cuda_bindings/cuda/bindings/_lib/utils.pxd
+++ b/cuda_bindings/cuda/bindings/_lib/utils.pxd
@@ -162,6 +162,9 @@ cdef class _HelperCUcoredumpSettings:
     cdef cydriver.CUcoredumpSettings_enum _attrib
     cdef bint _is_getter
     cdef size_t _size
+    # Retain the caller's bytes on the setter path so the borrowed _charstar /
+    # _cptr can never outlive their backing buffer (#379).
+    cdef object _references
 
     # Return values
     cdef bint _bool
diff --git a/cuda_bindings/cuda/bindings/_lib/utils.pxi b/cuda_bindings/cuda/bindings/_lib/utils.pxi
index 2dd4d5c1a27..0daf828ae65 100644
--- a/cuda_bindings/cuda/bindings/_lib/utils.pxi
+++ b/cuda_bindings/cuda/bindings/_lib/utils.pxi
@@ -664,6 +664,10 @@ cdef class _HelperCUcoredumpSettings:
                 self._cptr = <void*><void_ptr>self._charstar
                 self._size = 1024
             else:
+                # Assigning bytes to a char* borrows the object's internal
+                # buffer without incref/copy; retain init_value so the borrowed
+                # _charstar/_cptr can never outlive their backing buffer (#379).
+                self._references = init_value
                 self._charstar = init_value
                 self._cptr = <void*><void_ptr>self._charstar
                 self._size = len(init_value)
@@ -680,7 +684,13 @@ cdef class _HelperCUcoredumpSettings:
             raise TypeError('Unsupported attribute: {}'.format(attr.name))
 
     def __dealloc__(self):
-        pass
+        # Only the getter path owns heap: _charstar is a _callocWrapper(1024)
+        # buffer that pyObj() copies out, orphaning the original (#381). The
+        # setter path borrows the caller's bytes (retained via _references) and
+        # the bool path points _cptr at &self._bool, so freeing in those cases
+        # would corrupt memory -- guard the free to the getter branch.
+        if self._is_getter:
+            free(self._charstar)
 
     @property
     def cptr(self):
diff --git a/cuda_bindings/tests/test_cuda.py b/cuda_bindings/tests/test_cuda.py
index 192ad0f72fe..d58bb7d3b6f 100644
--- a/cuda_bindings/tests/test_cuda.py
+++ b/cuda_bindings/tests/test_cuda.py
@@ -554,6 +554,41 @@ def test_cuda_coredump_attr():
     assert attr_list[3] is True
 
 
+@pytest.mark.skipif(
+    driverVersionLessThan(12010)
+    or not supportsCudaAPI("cuCoredumpSetAttributeGlobal")
+    or not supportsCudaAPI("cuCoredumpGetAttributeGlobal"),
+    reason="Coredump API not present",
+)
+@pytest.mark.agent_authored(model="claude-opus-4-8")
+def test_cuda_coredump_attr_buffer_lifetime():
+    """Regression for the _HelperCUcoredumpSettings lifetime fixes.
+
+    #379: the setter aliases the caller's bytes into a ``char*``; the helper
+    must retain the object so the borrowed pointer cannot outlive it. Build the
+    value inline and force a GC before/after the call so a missing retain would
+    surface as a corrupted/empty path.
+
+    #381: the getter allocates a 1024-byte heap buffer that ``__dealloc__`` must
+    free. Loop many getter calls and force GC so a broken free (leak or
+    double-free/crash) is exercised.
+    """
+    import gc
+
+    path = ("/tmp/" + "cuda_python_coredump_" + str(0xC0FFEE)).encode("ascii")
+    (err,) = cuda.cuCoredumpSetAttributeGlobal(cuda.CUcoredumpSettings.CU_COREDUMP_FILE, path)
+    assert err == cuda.CUresult.CUDA_SUCCESS
+    del path
+    gc.collect()
+
+    for _ in range(64):
+        err, value = cuda.cuCoredumpGetAttributeGlobal(cuda.CUcoredumpSettings.CU_COREDUMP_FILE)
+        assert err == cuda.CUresult.CUDA_SUCCESS
+        assert value == b"/tmp/cuda_python_coredump_" + str(0xC0FFEE).encode("ascii")
+        del value
+        gc.collect()
+
+
 def test_get_error_name_and_string():
     err, device = cuda.cuDeviceGet(0)
     _, s = cuda.cuGetErrorString(err)
diff --git a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py b/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
index 1e22a739ae0..fdf0d726ee4 100644
--- a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
+++ b/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
@@ -397,9 +397,21 @@ def __init__(
         self._entries = self._root / _ENTRIES_SUBDIR
         self._tmp = self._root / _TMP_SUBDIR
         self._max_size_bytes = max_size_bytes
-        self._root.mkdir(parents=True, exist_ok=True)
-        self._entries.mkdir(exist_ok=True)
-        self._tmp.mkdir(exist_ok=True)
+        # The cache stores raw executable device code that is later handed to
+        # ``cuLibraryLoadData``. Create the tree owner-only (0o700) so another
+        # local principal can neither read another user's compiled kernels
+        # (#375) nor plant a binary this process would later load (#359).
+        # ``mkdir``'s ``mode`` is masked by the process umask, and
+        # ``exist_ok=True`` silently accepts a pre-existing (possibly
+        # world-writable) directory, so on POSIX also re-assert 0o700 on the
+        # tree rather than trusting the ambient umask or an attacker-supplied
+        # directory's permissions.
+        self._root.mkdir(parents=True, exist_ok=True, mode=0o700)
+        self._entries.mkdir(exist_ok=True, mode=0o700)
+        self._tmp.mkdir(exist_ok=True, mode=0o700)
+        if os.name != "nt":
+            for d in (self._root, self._entries, self._tmp):
+                os.chmod(d, 0o700)
         # Opportunistic startup sweep of orphaned temp files left by any
         # crashed writers. Age-based so concurrent in-flight writes from
         # other processes are preserved.
diff --git a/cuda_core/tests/test_program_cache.py b/cuda_core/tests/test_program_cache.py
index c305562dcba..928da70f0ff 100644
--- a/cuda_core/tests/test_program_cache.py
+++ b/cuda_core/tests/test_program_cache.py
@@ -3,6 +3,7 @@
 # SPDX-License-Identifier: Apache-2.0
 
 import abc
+import os
 import time
 
 import pytest
@@ -2535,3 +2536,41 @@ def reader(tid: int) -> None:
     # Internal accounting must agree with the cap and with __len__.
     assert cache._total_bytes <= 4096
     assert len(cache) == len(cache._entries)  # no orphan entries
+
+
+@pytest.mark.agent_authored(model="claude-opus-4-8")
+@pytest.mark.skipif(os.name == "nt", reason="POSIX permission bits only")
+def test_program_cache_dir_created_owner_only(tmp_path):
+    """#359/#375: the on-disk program cache stores raw executable device code,
+    so its directory tree must be created owner-only (0o700), independent of the
+    inherited umask, so other local principals can neither read cached kernels
+    nor plant a binary this process would later load."""
+    import stat
+
+    from cuda.core.utils._program_cache._file_stream import FileStreamProgramCache
+
+    root = tmp_path / "pc"
+    FileStreamProgramCache(path=root)
+
+    for d in (root, root / "entries", root / "tmp"):
+        mode = stat.S_IMODE(os.stat(d).st_mode)
+        assert mode == 0o700, f"{d} has mode {oct(mode)}"
+
+
+@pytest.mark.agent_authored(model="claude-opus-4-8")
+@pytest.mark.skipif(os.name == "nt", reason="POSIX permission bits only")
+def test_program_cache_dir_tightens_preexisting_world_writable(tmp_path):
+    """#359: a pre-created world-writable cache root (the poisoning precondition)
+    must be re-tightened to 0o700; mkdir(exist_ok=True) alone would keep the
+    attacker-supplied permissions."""
+    import stat
+
+    from cuda.core.utils._program_cache._file_stream import FileStreamProgramCache
+
+    root = tmp_path / "pc"
+    root.mkdir()
+    os.chmod(root, 0o777)
+
+    FileStreamProgramCache(path=root)
+
+    assert stat.S_IMODE(os.stat(root).st_mode) == 0o700
diff --git a/cuda_pathfinder/cuda/pathfinder/_binaries/find_nvidia_binary_utility.py b/cuda_pathfinder/cuda/pathfinder/_binaries/find_nvidia_binary_utility.py
index 4f857b9ac29..6a8c95919fb 100644
--- a/cuda_pathfinder/cuda/pathfinder/_binaries/find_nvidia_binary_utility.py
+++ b/cuda_pathfinder/cuda/pathfinder/_binaries/find_nvidia_binary_utility.py
@@ -63,7 +63,11 @@ def _resolve_in_trusted_dirs(normalized_name: str, dirs: list[str]) -> str | Non
         seen.add(directory)
         candidate = os.path.join(directory, normalized_name)
         if _is_executable_candidate(candidate):
-            return candidate
+            # Honor the docstring's absolute-path contract even when a search
+            # dir was supplied as a relative path (e.g. a relative CUDA_HOME);
+            # a relative result would otherwise re-resolve against a possibly
+            # different CWD at execution time (#374).
+            return os.path.abspath(candidate)
     return None
 
 
diff --git a/cuda_pathfinder/cuda/pathfinder/_dynamic_libs/load_dl_windows.py b/cuda_pathfinder/cuda/pathfinder/_dynamic_libs/load_dl_windows.py
index 5069a624790..c9057c7e4cd 100644
--- a/cuda_pathfinder/cuda/pathfinder/_dynamic_libs/load_dl_windows.py
+++ b/cuda_pathfinder/cuda/pathfinder/_dynamic_libs/load_dl_windows.py
@@ -7,6 +7,7 @@
 import ctypes.wintypes
 import os
 import struct
+import warnings
 from typing import TYPE_CHECKING
 
 from cuda.pathfinder._dynamic_libs.load_dl_common import LoadedDL
@@ -72,7 +73,20 @@ def add_dll_directory(dll_abs_path: str) -> None:
     # Add the DLL directory to the native search path. AddDllDirectory only
     # affects the LOAD_LIBRARY_SEARCH_USER_DIRS search; PATH is updated
     # unconditionally below to also cover legacy dependent-DLL resolution.
-    kernel32.AddDllDirectory(dirpath)
+    cookie = kernel32.AddDllDirectory(dirpath)
+    if not cookie:
+        # Surface the failure instead of masking it with a bare pass: the PATH
+        # widening below is a weaker, broader fallback that py3.8+ loaders
+        # (SetDefaultDllDirectories) may ignore, so a swallowed failure here
+        # would otherwise surface only as an unexplained dependent-DLL
+        # not-found later, with no breadcrumb (#377).
+        error_code = ctypes.GetLastError()  # type: ignore[attr-defined]
+        warnings.warn(
+            f"AddDllDirectory({dirpath!r}) failed (error code: {error_code}); "
+            "falling back to process-global PATH mutation for dependent-DLL resolution.",
+            RuntimeWarning,
+            stacklevel=2,
+        )
 
     # Update PATH as a fallback for dependent DLL resolution
     curr_path = os.environ.get("PATH")
diff --git a/cuda_pathfinder/tests/test_find_nvidia_binaries.py b/cuda_pathfinder/tests/test_find_nvidia_binaries.py
index ae439546859..b4cb907dde8 100644
--- a/cuda_pathfinder/tests/test_find_nvidia_binaries.py
+++ b/cuda_pathfinder/tests/test_find_nvidia_binaries.py
@@ -370,3 +370,28 @@ def test_caching_per_utility():
     # them is None)
     if nvdisasm1 is not None and nvcc1 is not None:
         assert nvdisasm1 != nvcc1
+
+
+@pytest.mark.agent_authored(model="claude-opus-4-8")
+def test_resolve_in_trusted_dirs_returns_absolute_path(tmp_path, monkeypatch, mocker):
+    """#374: a match found under a relative search dir must be absolutized.
+
+    ``find_nvidia_binary_utility`` documents an absolute, separator-resolved
+    result. A relative search dir (e.g. a relative ``CUDA_HOME``) previously
+    leaked a relative path that would re-resolve against a possibly different
+    CWD at execution time.
+    """
+    rel_dir = os.path.join("some", "relative", "bin")
+    candidate = os.path.join(rel_dir, "nvcc")
+    mocker.patch.object(
+        binary_finder_module,
+        "_is_executable_candidate",
+        side_effect=lambda path: path == candidate,
+    )
+
+    # Anchor CWD so os.path.abspath is deterministic for the assertion.
+    monkeypatch.chdir(tmp_path)
+    result = binary_finder_module._resolve_in_trusted_dirs("nvcc", [rel_dir])
+
+    assert os.path.isabs(result)
+    assert result == os.path.abspath(os.path.join(str(tmp_path), candidate))

From a243030b51b5f2d2f36b1293547fa51af00e4924 Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Tue, 21 Jul 2026 13:47:39 -0700
Subject: [PATCH 02/14] Simplify code comments on the hardening fixes

Shorten the explanatory comments to a line or two each and drop the internal
issue-number references (which don't resolve in this repo). No code changes.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 cuda_bindings/cuda/bindings/_internal/utils.pyx     | 13 ++++---------
 cuda_bindings/cuda/bindings/_lib/utils.pxd          |  4 +---
 cuda_bindings/cuda/bindings/_lib/utils.pxi          | 12 ++++--------
 .../cuda/core/utils/_program_cache/_file_stream.py  | 13 ++++---------
 .../_binaries/find_nvidia_binary_utility.py         |  6 ++----
 .../pathfinder/_dynamic_libs/load_dl_windows.py     |  7 ++-----
 6 files changed, 17 insertions(+), 38 deletions(-)

diff --git a/cuda_bindings/cuda/bindings/_internal/utils.pyx b/cuda_bindings/cuda/bindings/_internal/utils.pyx
index 73c4e223a20..ec415f29dcc 100644
--- a/cuda_bindings/cuda/bindings/_internal/utils.pyx
+++ b/cuda_bindings/cuda/bindings/_internal/utils.pyx
@@ -44,15 +44,10 @@ cdef bint is_nested_sequence(data):
 cdef void* get_buffer_pointer(buf, Py_ssize_t size, readonly=True) except*:
     """Return a raw pointer into ``buf``.
 
-    The caller must ensure ``buf`` is alive AND not reallocated/resized while
-    the returned pointer is in use. For buffer-protocol objects the export lock
-    (``PyBuffer_Release``) is dropped before this function returns, so the
-    stack-local ``Py_buffer`` view cannot keep the export alive past return.
-    Concurrently resizing a mutable exporter (e.g. ``bytearray``, ``array.array``,
-    a growable NumPy view) during a ``nogil`` native call that dereferences the
-    pointer is therefore a use-after-free (#366). Prefer immutable buffers
-    (``bytes``), or hold the buffer export for the pointer's full lifetime at the
-    call site.
+    The caller must keep ``buf`` alive AND not resize it while the pointer is in
+    use. The buffer export is released before this returns, so resizing a mutable
+    buffer (e.g. ``bytearray``) during a ``nogil`` call is a use-after-free.
+    Prefer immutable ``bytes``.
     """
     cdef void* bufPtr
     cdef int flags = cpython.PyBUF_ANY_CONTIGUOUS
diff --git a/cuda_bindings/cuda/bindings/_lib/utils.pxd b/cuda_bindings/cuda/bindings/_lib/utils.pxd
index 7a846d21fb9..0d8af74b4ff 100644
--- a/cuda_bindings/cuda/bindings/_lib/utils.pxd
+++ b/cuda_bindings/cuda/bindings/_lib/utils.pxd
@@ -162,9 +162,7 @@ cdef class _HelperCUcoredumpSettings:
     cdef cydriver.CUcoredumpSettings_enum _attrib
     cdef bint _is_getter
     cdef size_t _size
-    # Retain the caller's bytes on the setter path so the borrowed _charstar /
-    # _cptr can never outlive their backing buffer (#379).
-    cdef object _references
+    cdef object _references  # keeps caller bytes alive so _charstar stays valid
 
     # Return values
     cdef bint _bool
diff --git a/cuda_bindings/cuda/bindings/_lib/utils.pxi b/cuda_bindings/cuda/bindings/_lib/utils.pxi
index 0daf828ae65..2796f910798 100644
--- a/cuda_bindings/cuda/bindings/_lib/utils.pxi
+++ b/cuda_bindings/cuda/bindings/_lib/utils.pxi
@@ -664,9 +664,7 @@ cdef class _HelperCUcoredumpSettings:
                 self._cptr = <void*><void_ptr>self._charstar
                 self._size = 1024
             else:
-                # Assigning bytes to a char* borrows the object's internal
-                # buffer without incref/copy; retain init_value so the borrowed
-                # _charstar/_cptr can never outlive their backing buffer (#379).
+                # Keep a reference so the borrowed _charstar buffer stays alive.
                 self._references = init_value
                 self._charstar = init_value
                 self._cptr = <void*><void_ptr>self._charstar
@@ -684,11 +682,9 @@ cdef class _HelperCUcoredumpSettings:
             raise TypeError('Unsupported attribute: {}'.format(attr.name))
 
     def __dealloc__(self):
-        # Only the getter path owns heap: _charstar is a _callocWrapper(1024)
-        # buffer that pyObj() copies out, orphaning the original (#381). The
-        # setter path borrows the caller's bytes (retained via _references) and
-        # the bool path points _cptr at &self._bool, so freeing in those cases
-        # would corrupt memory -- guard the free to the getter branch.
+        # Only the getter path owns heap (the calloc'd 1024-byte buffer). The
+        # setter borrows caller bytes and the bool path points at &self._bool,
+        # so only free for the getter.
         if self._is_getter:
             free(self._charstar)
 
diff --git a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py b/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
index fdf0d726ee4..cc3d877fe4c 100644
--- a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
+++ b/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
@@ -397,15 +397,10 @@ def __init__(
         self._entries = self._root / _ENTRIES_SUBDIR
         self._tmp = self._root / _TMP_SUBDIR
         self._max_size_bytes = max_size_bytes
-        # The cache stores raw executable device code that is later handed to
-        # ``cuLibraryLoadData``. Create the tree owner-only (0o700) so another
-        # local principal can neither read another user's compiled kernels
-        # (#375) nor plant a binary this process would later load (#359).
-        # ``mkdir``'s ``mode`` is masked by the process umask, and
-        # ``exist_ok=True`` silently accepts a pre-existing (possibly
-        # world-writable) directory, so on POSIX also re-assert 0o700 on the
-        # tree rather than trusting the ambient umask or an attacker-supplied
-        # directory's permissions.
+        # Cache holds compiled device code loaded via cuLibraryLoadData, so keep
+        # it owner-only (0o700) so other local users can't read or replace it.
+        # mkdir's mode is umask-masked and exist_ok keeps an existing dir's
+        # perms, so re-chmod on POSIX.
         self._root.mkdir(parents=True, exist_ok=True, mode=0o700)
         self._entries.mkdir(exist_ok=True, mode=0o700)
         self._tmp.mkdir(exist_ok=True, mode=0o700)
diff --git a/cuda_pathfinder/cuda/pathfinder/_binaries/find_nvidia_binary_utility.py b/cuda_pathfinder/cuda/pathfinder/_binaries/find_nvidia_binary_utility.py
index 6a8c95919fb..834db8fe8fa 100644
--- a/cuda_pathfinder/cuda/pathfinder/_binaries/find_nvidia_binary_utility.py
+++ b/cuda_pathfinder/cuda/pathfinder/_binaries/find_nvidia_binary_utility.py
@@ -63,10 +63,8 @@ def _resolve_in_trusted_dirs(normalized_name: str, dirs: list[str]) -> str | Non
         seen.add(directory)
         candidate = os.path.join(directory, normalized_name)
         if _is_executable_candidate(candidate):
-            # Honor the docstring's absolute-path contract even when a search
-            # dir was supplied as a relative path (e.g. a relative CUDA_HOME);
-            # a relative result would otherwise re-resolve against a possibly
-            # different CWD at execution time (#374).
+            # Return an absolute path, as the docstring promises (a relative
+            # search dir would otherwise leak a relative result).
             return os.path.abspath(candidate)
     return None
 
diff --git a/cuda_pathfinder/cuda/pathfinder/_dynamic_libs/load_dl_windows.py b/cuda_pathfinder/cuda/pathfinder/_dynamic_libs/load_dl_windows.py
index c9057c7e4cd..31ad35197b0 100644
--- a/cuda_pathfinder/cuda/pathfinder/_dynamic_libs/load_dl_windows.py
+++ b/cuda_pathfinder/cuda/pathfinder/_dynamic_libs/load_dl_windows.py
@@ -75,11 +75,8 @@ def add_dll_directory(dll_abs_path: str) -> None:
     # unconditionally below to also cover legacy dependent-DLL resolution.
     cookie = kernel32.AddDllDirectory(dirpath)
     if not cookie:
-        # Surface the failure instead of masking it with a bare pass: the PATH
-        # widening below is a weaker, broader fallback that py3.8+ loaders
-        # (SetDefaultDllDirectories) may ignore, so a swallowed failure here
-        # would otherwise surface only as an unexplained dependent-DLL
-        # not-found later, with no breadcrumb (#377).
+        # Warn instead of failing silently; the PATH update below is a weaker
+        # fallback that newer loaders may ignore.
         error_code = ctypes.GetLastError()  # type: ignore[attr-defined]
         warnings.warn(
             f"AddDllDirectory({dirpath!r}) failed (error code: {error_code}); "

From c9e5d5d3f9e38f1d852e5d642f889750b321e715 Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Tue, 21 Jul 2026 13:57:50 -0700
Subject: [PATCH 03/14] Remove agent_authored markers from the new tests

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 cuda_bindings/tests/test_cuda.py                   | 1 -
 cuda_core/tests/test_program_cache.py              | 2 --
 cuda_pathfinder/tests/test_find_nvidia_binaries.py | 1 -
 3 files changed, 4 deletions(-)

diff --git a/cuda_bindings/tests/test_cuda.py b/cuda_bindings/tests/test_cuda.py
index d58bb7d3b6f..e0ae2d4d6d9 100644
--- a/cuda_bindings/tests/test_cuda.py
+++ b/cuda_bindings/tests/test_cuda.py
@@ -560,7 +560,6 @@ def test_cuda_coredump_attr():
     or not supportsCudaAPI("cuCoredumpGetAttributeGlobal"),
     reason="Coredump API not present",
 )
-@pytest.mark.agent_authored(model="claude-opus-4-8")
 def test_cuda_coredump_attr_buffer_lifetime():
     """Regression for the _HelperCUcoredumpSettings lifetime fixes.
 
diff --git a/cuda_core/tests/test_program_cache.py b/cuda_core/tests/test_program_cache.py
index 928da70f0ff..722b55d24d5 100644
--- a/cuda_core/tests/test_program_cache.py
+++ b/cuda_core/tests/test_program_cache.py
@@ -2538,7 +2538,6 @@ def reader(tid: int) -> None:
     assert len(cache) == len(cache._entries)  # no orphan entries
 
 
-@pytest.mark.agent_authored(model="claude-opus-4-8")
 @pytest.mark.skipif(os.name == "nt", reason="POSIX permission bits only")
 def test_program_cache_dir_created_owner_only(tmp_path):
     """#359/#375: the on-disk program cache stores raw executable device code,
@@ -2557,7 +2556,6 @@ def test_program_cache_dir_created_owner_only(tmp_path):
         assert mode == 0o700, f"{d} has mode {oct(mode)}"
 
 
-@pytest.mark.agent_authored(model="claude-opus-4-8")
 @pytest.mark.skipif(os.name == "nt", reason="POSIX permission bits only")
 def test_program_cache_dir_tightens_preexisting_world_writable(tmp_path):
     """#359: a pre-created world-writable cache root (the poisoning precondition)
diff --git a/cuda_pathfinder/tests/test_find_nvidia_binaries.py b/cuda_pathfinder/tests/test_find_nvidia_binaries.py
index b4cb907dde8..26608c03b6d 100644
--- a/cuda_pathfinder/tests/test_find_nvidia_binaries.py
+++ b/cuda_pathfinder/tests/test_find_nvidia_binaries.py
@@ -372,7 +372,6 @@ def test_caching_per_utility():
         assert nvdisasm1 != nvcc1
 
 
-@pytest.mark.agent_authored(model="claude-opus-4-8")
 def test_resolve_in_trusted_dirs_returns_absolute_path(tmp_path, monkeypatch, mocker):
     """#374: a match found under a relative search dir must be absolutized.
 

From 1c5eb3eab7369a294a25ac49c98cfcd2401662ff Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Tue, 21 Jul 2026 14:01:12 -0700
Subject: [PATCH 04/14] Remove weak coredump buffer-lifetime test

It didn't actually guard the two lifetime fixes: the driver copies the path
during the set call (so the #379 latent UAF can't trigger), and a missing free
leaks silently (so #381 wouldn't fail the test either).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 cuda_bindings/tests/test_cuda.py | 34 --------------------------------
 1 file changed, 34 deletions(-)

diff --git a/cuda_bindings/tests/test_cuda.py b/cuda_bindings/tests/test_cuda.py
index e0ae2d4d6d9..192ad0f72fe 100644
--- a/cuda_bindings/tests/test_cuda.py
+++ b/cuda_bindings/tests/test_cuda.py
@@ -554,40 +554,6 @@ def test_cuda_coredump_attr():
     assert attr_list[3] is True
 
 
-@pytest.mark.skipif(
-    driverVersionLessThan(12010)
-    or not supportsCudaAPI("cuCoredumpSetAttributeGlobal")
-    or not supportsCudaAPI("cuCoredumpGetAttributeGlobal"),
-    reason="Coredump API not present",
-)
-def test_cuda_coredump_attr_buffer_lifetime():
-    """Regression for the _HelperCUcoredumpSettings lifetime fixes.
-
-    #379: the setter aliases the caller's bytes into a ``char*``; the helper
-    must retain the object so the borrowed pointer cannot outlive it. Build the
-    value inline and force a GC before/after the call so a missing retain would
-    surface as a corrupted/empty path.
-
-    #381: the getter allocates a 1024-byte heap buffer that ``__dealloc__`` must
-    free. Loop many getter calls and force GC so a broken free (leak or
-    double-free/crash) is exercised.
-    """
-    import gc
-
-    path = ("/tmp/" + "cuda_python_coredump_" + str(0xC0FFEE)).encode("ascii")
-    (err,) = cuda.cuCoredumpSetAttributeGlobal(cuda.CUcoredumpSettings.CU_COREDUMP_FILE, path)
-    assert err == cuda.CUresult.CUDA_SUCCESS
-    del path
-    gc.collect()
-
-    for _ in range(64):
-        err, value = cuda.cuCoredumpGetAttributeGlobal(cuda.CUcoredumpSettings.CU_COREDUMP_FILE)
-        assert err == cuda.CUresult.CUDA_SUCCESS
-        assert value == b"/tmp/cuda_python_coredump_" + str(0xC0FFEE).encode("ascii")
-        del value
-        gc.collect()
-
-
 def test_get_error_name_and_string():
     err, device = cuda.cuDeviceGet(0)
     _, s = cuda.cuGetErrorString(err)

From dc8bb9f05e911a0b4ab3392c6aeadb09ab3f75b1 Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Tue, 21 Jul 2026 14:08:37 -0700
Subject: [PATCH 05/14] Silence bandit S103 on the intentionally world-writable
 test dir

The test pre-creates a 0o777 cache dir to verify the loader tightens it to
0o700; the permissive mask is the point of the test.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 cuda_core/tests/test_program_cache.py | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/cuda_core/tests/test_program_cache.py b/cuda_core/tests/test_program_cache.py
index 722b55d24d5..2be669f49dc 100644
--- a/cuda_core/tests/test_program_cache.py
+++ b/cuda_core/tests/test_program_cache.py
@@ -2567,7 +2567,8 @@ def test_program_cache_dir_tightens_preexisting_world_writable(tmp_path):
 
     root = tmp_path / "pc"
     root.mkdir()
-    os.chmod(root, 0o777)
+    # Simulate the attack precondition: a pre-existing world-writable cache dir.
+    os.chmod(root, 0o777)  # noqa: S103
 
     FileStreamProgramCache(path=root)
 

From c1d2cce70f49dc39d4f2dcac5e1e083f9f581a3c Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Tue, 21 Jul 2026 15:34:35 -0700
Subject: [PATCH 06/14] test: expect absolutized binary path on Windows

The abspath change makes find_nvidia_binary_utility return a drive-qualified
path on Windows (C:\... ), so the three search-order tests must compare against
os.path.abspath(expected). No-op on Linux; fixes the Windows CI failures.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 cuda_pathfinder/tests/test_find_nvidia_binaries.py | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/cuda_pathfinder/tests/test_find_nvidia_binaries.py b/cuda_pathfinder/tests/test_find_nvidia_binaries.py
index 26608c03b6d..6b7906e1996 100644
--- a/cuda_pathfinder/tests/test_find_nvidia_binaries.py
+++ b/cuda_pathfinder/tests/test_find_nvidia_binaries.py
@@ -152,7 +152,7 @@ def test_find_binary_first_matching_dir_wins(monkeypatch, mocker):
     result = find_nvidia_binary_utility("nvcc")
 
     # Conda comes before CUDA_HOME, so the Conda hit wins and CUDA_HOME is never probed.
-    assert result == conda_nvcc
+    assert result == os.path.abspath(conda_nvcc)
     assert checked == [os.path.join(site_dir, "nvcc"), conda_nvcc]
 
 
@@ -173,7 +173,7 @@ def test_find_binary_ctk_root_canary_fallback(monkeypatch, mocker):
 
     result = find_nvidia_binary_utility("nvcc")
 
-    assert result == ctk_nvcc
+    assert result == os.path.abspath(ctk_nvcc)
     canary_mock.assert_called_once_with()
     # No earlier trusted dirs existed, so the only probe is the canary bin dir.
     assert checked == [ctk_nvcc]
@@ -218,7 +218,7 @@ def test_find_binary_canary_not_consulted_when_found_earlier(monkeypatch, mocker
 
     result = find_nvidia_binary_utility("nvcc")
 
-    assert result == conda_nvcc
+    assert result == os.path.abspath(conda_nvcc)
     canary_mock.assert_not_called()
 
 

From e911dd744ebe31430f73e513a869c3e0ff44a2c1 Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Tue, 21 Jul 2026 15:34:36 -0700
Subject: [PATCH 07/14] Raise on out-of-range c_int/c_byte kernel arguments
 instead of truncating

The c_int/c_byte param feeders silently narrowed an out-of-range Python int
(e.g. 2**32+5 -> 5). Range-check in the feeder and raise OverflowError; declare
feed() as 'except? -1' so Cython propagates the exception instead of swallowing
it. Addresses Glasswing V9.1 (leofang chose the raise option over matching
ctypes' silent wrap).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 .../cuda/bindings/_lib/param_packer.h         | 24 ++++++++++++--
 .../cuda/bindings/_lib/param_packer.pxd       |  2 +-
 cuda_bindings/tests/test_kernelParams.py      | 31 +++++++++++++++++++
 3 files changed, 54 insertions(+), 3 deletions(-)

diff --git a/cuda_bindings/cuda/bindings/_lib/param_packer.h b/cuda_bindings/cuda/bindings/_lib/param_packer.h
index 160ef5f7c92..f0435d4c76c 100644
--- a/cuda_bindings/cuda/bindings/_lib/param_packer.h
+++ b/cuda_bindings/cuda/bindings/_lib/param_packer.h
@@ -7,6 +7,8 @@
 #include <functional>
 #include <stdexcept>
 #include <string>
+#include <climits>
+#include <cstdint>
 
 static PyObject* ctypes_module = nullptr;
 
@@ -69,7 +71,16 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
         {
             m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int
             {
-                *((int*)ptr) = (int)PyLong_AsLong(value);
+                long v = PyLong_AsLong(value);
+                if (v == -1 && PyErr_Occurred())
+                    return -1;  // value out of C long range; propagate the error
+                if (v < INT_MIN || v > INT_MAX)
+                {
+                    PyErr_Format(PyExc_OverflowError,
+                        "Python int %ld is out of range for a c_int (32-bit) kernel argument", v);
+                    return -1;
+                }
+                *((int*)ptr) = (int)v;
                 return sizeof(int);
             };
             return;
@@ -89,7 +100,16 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
         {
             m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int
             {
-                *((int8_t*)ptr) = (int8_t)PyLong_AsLong(value);
+                long v = PyLong_AsLong(value);
+                if (v == -1 && PyErr_Occurred())
+                    return -1;  // value out of C long range; propagate the error
+                if (v < INT8_MIN || v > INT8_MAX)
+                {
+                    PyErr_Format(PyExc_OverflowError,
+                        "Python int %ld is out of range for a c_byte (8-bit) kernel argument", v);
+                    return -1;
+                }
+                *((int8_t*)ptr) = (int8_t)v;
                 return sizeof(int8_t);
             };
             return;
diff --git a/cuda_bindings/cuda/bindings/_lib/param_packer.pxd b/cuda_bindings/cuda/bindings/_lib/param_packer.pxd
index 1c0ad690be4..d1f84059db1 100644
--- a/cuda_bindings/cuda/bindings/_lib/param_packer.pxd
+++ b/cuda_bindings/cuda/bindings/_lib/param_packer.pxd
@@ -4,4 +4,4 @@
 # Include "param_packer.h" so its contents get compiled into every
 # Cython extension module that depends on param_packer.pxd.
 cdef extern from "param_packer.h":
-    int feed(void* ptr, object o, object ct)
+    int feed(void* ptr, object o, object ct) except? -1
diff --git a/cuda_bindings/tests/test_kernelParams.py b/cuda_bindings/tests/test_kernelParams.py
index 555d6a7284c..1e48d2d2fd6 100644
--- a/cuda_bindings/tests/test_kernelParams.py
+++ b/cuda_bindings/tests/test_kernelParams.py
@@ -800,3 +800,34 @@ def __init__(self, address, typestr):
     ASSERT_DRV(err)
     (err,) = cuda.cuModuleUnload(module)
     ASSERT_DRV(err)
+
+
+def test_kernelParams_c_int_out_of_range_raises(device):
+    # #363: an out-of-range Python int for a c_int / c_byte kernel argument must
+    # raise instead of being silently truncated to fit the declared width.
+    kernelString = """\
+    extern "C" __global__ void take_int(int i) {}
+    """
+    module = common_nvrtc(kernelString, device)
+    err, kernel = cuda.cuModuleGetFunction(module, b"take_int")
+    ASSERT_DRV(err)
+    err, stream = cuda.cuStreamCreate(0)
+    ASSERT_DRV(err)
+
+    # An in-range value still packs and launches fine.
+    (err,) = cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((5,), (ctypes.c_int,)), 0)
+    ASSERT_DRV(err)
+
+    # Out-of-range values now raise OverflowError during packing (previously the
+    # high bits were silently dropped, so the kernel saw a different value).
+    with pytest.raises(OverflowError):
+        cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((2**32 + 5,), (ctypes.c_int,)), 0)
+    with pytest.raises(OverflowError):
+        cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((200,), (ctypes.c_byte,)), 0)
+
+    (err,) = cuda.cuStreamSynchronize(stream)
+    ASSERT_DRV(err)
+    (err,) = cuda.cuStreamDestroy(stream)
+    ASSERT_DRV(err)
+    (err,) = cuda.cuModuleUnload(module)
+    ASSERT_DRV(err)

From 2f5be6539e9b466d7a5e733097e27d1a94695ed2 Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Tue, 21 Jul 2026 15:39:24 -0700
Subject: [PATCH 08/14] param_packer: use PyLong_AsInt for c_int range check on
 Python 3.13+

Per Leo's review: on 3.13+ PyLong_AsInt converts and range-checks in one call
(raising OverflowError itself), so gate the manual INT_MIN/INT_MAX check to
pre-3.13 only. c_byte keeps the manual check (no 8-bit CPython equivalent).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 cuda_bindings/cuda/bindings/_lib/param_packer.h | 9 +++++++++
 1 file changed, 9 insertions(+)

diff --git a/cuda_bindings/cuda/bindings/_lib/param_packer.h b/cuda_bindings/cuda/bindings/_lib/param_packer.h
index f0435d4c76c..f92bfb63cc9 100644
--- a/cuda_bindings/cuda/bindings/_lib/param_packer.h
+++ b/cuda_bindings/cuda/bindings/_lib/param_packer.h
@@ -71,6 +71,14 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
         {
             m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int
             {
+#if PY_VERSION_HEX >= 0x030D0000
+                // Python 3.13+: PyLong_AsInt range-checks and raises
+                // OverflowError itself, so let CPython own the bounds check.
+                int iv = PyLong_AsInt(value);
+                if (iv == -1 && PyErr_Occurred())
+                    return -1;
+                *((int*)ptr) = iv;
+#else
                 long v = PyLong_AsLong(value);
                 if (v == -1 && PyErr_Occurred())
                     return -1;  // value out of C long range; propagate the error
@@ -81,6 +89,7 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
                     return -1;
                 }
                 *((int*)ptr) = (int)v;
+#endif
                 return sizeof(int);
             };
             return;

From 8d6ee671e3cd3cb5967967c172b3e616f6f55cac Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Tue, 21 Jul 2026 15:49:42 -0700
Subject: [PATCH 09/14] param_packer: unify c_int/c_byte range check via
 PyLong_AsLongAndOverflow

Per Leo's follow-up: use a single code path on all supported Pythons instead of
version-gating PyLong_AsInt. PyLong_AsLongAndOverflow flags out-of-long values
via its overflow out-param (no exception set), then we bounds-check the 32-bit
(c_int) / 8-bit (c_byte) target range and raise OverflowError. Drops the
PY_VERSION_HEX gate.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 .../cuda/bindings/_lib/param_packer.h         | 38 +++++++++----------
 1 file changed, 17 insertions(+), 21 deletions(-)

diff --git a/cuda_bindings/cuda/bindings/_lib/param_packer.h b/cuda_bindings/cuda/bindings/_lib/param_packer.h
index f92bfb63cc9..7161094a447 100644
--- a/cuda_bindings/cuda/bindings/_lib/param_packer.h
+++ b/cuda_bindings/cuda/bindings/_lib/param_packer.h
@@ -71,25 +71,20 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
         {
             m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int
             {
-#if PY_VERSION_HEX >= 0x030D0000
-                // Python 3.13+: PyLong_AsInt range-checks and raises
-                // OverflowError itself, so let CPython own the bounds check.
-                int iv = PyLong_AsInt(value);
-                if (iv == -1 && PyErr_Occurred())
-                    return -1;
-                *((int*)ptr) = iv;
-#else
-                long v = PyLong_AsLong(value);
-                if (v == -1 && PyErr_Occurred())
-                    return -1;  // value out of C long range; propagate the error
-                if (v < INT_MIN || v > INT_MAX)
+                // One code path across all supported Pythons: AsLongAndOverflow
+                // flags values outside C long via `overflow` (without setting an
+                // exception), then we bounds-check the 32-bit int range.
+                int overflow = 0;
+                long v = PyLong_AsLongAndOverflow(value, &overflow);
+                if (overflow == 0 && v == -1 && PyErr_Occurred())
+                    return -1;  // non-overflow conversion error; exception already set
+                if (overflow != 0 || v < INT_MIN || v > INT_MAX)
                 {
-                    PyErr_Format(PyExc_OverflowError,
-                        "Python int %ld is out of range for a c_int (32-bit) kernel argument", v);
+                    PyErr_SetString(PyExc_OverflowError,
+                        "Python int is out of range for a c_int (32-bit) kernel argument");
                     return -1;
                 }
                 *((int*)ptr) = (int)v;
-#endif
                 return sizeof(int);
             };
             return;
@@ -109,13 +104,14 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
         {
             m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int
             {
-                long v = PyLong_AsLong(value);
-                if (v == -1 && PyErr_Occurred())
-                    return -1;  // value out of C long range; propagate the error
-                if (v < INT8_MIN || v > INT8_MAX)
+                int overflow = 0;
+                long v = PyLong_AsLongAndOverflow(value, &overflow);
+                if (overflow == 0 && v == -1 && PyErr_Occurred())
+                    return -1;  // non-overflow conversion error; exception already set
+                if (overflow != 0 || v < INT8_MIN || v > INT8_MAX)
                 {
-                    PyErr_Format(PyExc_OverflowError,
-                        "Python int %ld is out of range for a c_byte (8-bit) kernel argument", v);
+                    PyErr_SetString(PyExc_OverflowError,
+                        "Python int is out of range for a c_byte (8-bit) kernel argument");
                     return -1;
                 }
                 *((int8_t*)ptr) = (int8_t)v;

From 1d7420608a579612d656461cb268a2d5f4fbb71f Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Tue, 21 Jul 2026 16:00:05 -0700
Subject: [PATCH 10/14] Split kernel-arg overflow change (#363) into its own PR
 (#2402)

The out-of-range c_int/c_byte -> OverflowError change is a distinct behavior
change; moved to a standalone PR so it can be reviewed separately from this
hardening batch. No functional change to the remaining hardening work.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 .../cuda/bindings/_lib/param_packer.h         | 29 ++---------------
 .../cuda/bindings/_lib/param_packer.pxd       |  2 +-
 cuda_bindings/tests/test_kernelParams.py      | 31 -------------------
 3 files changed, 3 insertions(+), 59 deletions(-)

diff --git a/cuda_bindings/cuda/bindings/_lib/param_packer.h b/cuda_bindings/cuda/bindings/_lib/param_packer.h
index 7161094a447..160ef5f7c92 100644
--- a/cuda_bindings/cuda/bindings/_lib/param_packer.h
+++ b/cuda_bindings/cuda/bindings/_lib/param_packer.h
@@ -7,8 +7,6 @@
 #include <functional>
 #include <stdexcept>
 #include <string>
-#include <climits>
-#include <cstdint>
 
 static PyObject* ctypes_module = nullptr;
 
@@ -71,20 +69,7 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
         {
             m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int
             {
-                // One code path across all supported Pythons: AsLongAndOverflow
-                // flags values outside C long via `overflow` (without setting an
-                // exception), then we bounds-check the 32-bit int range.
-                int overflow = 0;
-                long v = PyLong_AsLongAndOverflow(value, &overflow);
-                if (overflow == 0 && v == -1 && PyErr_Occurred())
-                    return -1;  // non-overflow conversion error; exception already set
-                if (overflow != 0 || v < INT_MIN || v > INT_MAX)
-                {
-                    PyErr_SetString(PyExc_OverflowError,
-                        "Python int is out of range for a c_int (32-bit) kernel argument");
-                    return -1;
-                }
-                *((int*)ptr) = (int)v;
+                *((int*)ptr) = (int)PyLong_AsLong(value);
                 return sizeof(int);
             };
             return;
@@ -104,17 +89,7 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
         {
             m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int
             {
-                int overflow = 0;
-                long v = PyLong_AsLongAndOverflow(value, &overflow);
-                if (overflow == 0 && v == -1 && PyErr_Occurred())
-                    return -1;  // non-overflow conversion error; exception already set
-                if (overflow != 0 || v < INT8_MIN || v > INT8_MAX)
-                {
-                    PyErr_SetString(PyExc_OverflowError,
-                        "Python int is out of range for a c_byte (8-bit) kernel argument");
-                    return -1;
-                }
-                *((int8_t*)ptr) = (int8_t)v;
+                *((int8_t*)ptr) = (int8_t)PyLong_AsLong(value);
                 return sizeof(int8_t);
             };
             return;
diff --git a/cuda_bindings/cuda/bindings/_lib/param_packer.pxd b/cuda_bindings/cuda/bindings/_lib/param_packer.pxd
index d1f84059db1..1c0ad690be4 100644
--- a/cuda_bindings/cuda/bindings/_lib/param_packer.pxd
+++ b/cuda_bindings/cuda/bindings/_lib/param_packer.pxd
@@ -4,4 +4,4 @@
 # Include "param_packer.h" so its contents get compiled into every
 # Cython extension module that depends on param_packer.pxd.
 cdef extern from "param_packer.h":
-    int feed(void* ptr, object o, object ct) except? -1
+    int feed(void* ptr, object o, object ct)
diff --git a/cuda_bindings/tests/test_kernelParams.py b/cuda_bindings/tests/test_kernelParams.py
index 1e48d2d2fd6..555d6a7284c 100644
--- a/cuda_bindings/tests/test_kernelParams.py
+++ b/cuda_bindings/tests/test_kernelParams.py
@@ -800,34 +800,3 @@ def __init__(self, address, typestr):
     ASSERT_DRV(err)
     (err,) = cuda.cuModuleUnload(module)
     ASSERT_DRV(err)
-
-
-def test_kernelParams_c_int_out_of_range_raises(device):
-    # #363: an out-of-range Python int for a c_int / c_byte kernel argument must
-    # raise instead of being silently truncated to fit the declared width.
-    kernelString = """\
-    extern "C" __global__ void take_int(int i) {}
-    """
-    module = common_nvrtc(kernelString, device)
-    err, kernel = cuda.cuModuleGetFunction(module, b"take_int")
-    ASSERT_DRV(err)
-    err, stream = cuda.cuStreamCreate(0)
-    ASSERT_DRV(err)
-
-    # An in-range value still packs and launches fine.
-    (err,) = cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((5,), (ctypes.c_int,)), 0)
-    ASSERT_DRV(err)
-
-    # Out-of-range values now raise OverflowError during packing (previously the
-    # high bits were silently dropped, so the kernel saw a different value).
-    with pytest.raises(OverflowError):
-        cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((2**32 + 5,), (ctypes.c_int,)), 0)
-    with pytest.raises(OverflowError):
-        cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((200,), (ctypes.c_byte,)), 0)
-
-    (err,) = cuda.cuStreamSynchronize(stream)
-    ASSERT_DRV(err)
-    (err,) = cuda.cuStreamDestroy(stream)
-    ASSERT_DRV(err)
-    (err,) = cuda.cuModuleUnload(module)
-    ASSERT_DRV(err)

From 6fd6b6cb45fc610ed174712b02342f2256537804 Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Wed, 22 Jul 2026 08:58:51 -0700
Subject: [PATCH 11/14] Address review: scope cache perms to tmp, revert cybind
 docstring

Per @leofang's review on PR #2399:
- _file_stream.py: only the tmp/ staging dir is created owner-only (0o700).
  root/ and entries/ now inherit the umask / any pre-existing permissions so a
  deliberately shared kernel cache (e.g. group-writable on a compute cluster)
  keeps working. Dropped the post-mkdir chmod that would have re-tightened an
  existing shared directory. 0o700 in mkdir mode needs no chmod: umask only
  clears bits.
- _internal/utils.pyx: revert the get_buffer_pointer docstring change; that is
  cybind code and is out of scope for this PR.
- Tests updated to match: assert only tmp is 0o700, and that a pre-existing
  0o777 shared root is used as-is.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 .../cuda/bindings/_internal/utils.pyx         |  8 +-----
 .../core/utils/_program_cache/_file_stream.py | 18 ++++++------
 cuda_core/tests/test_program_cache.py         | 28 +++++++++----------
 3 files changed, 24 insertions(+), 30 deletions(-)

diff --git a/cuda_bindings/cuda/bindings/_internal/utils.pyx b/cuda_bindings/cuda/bindings/_internal/utils.pyx
index ec415f29dcc..a5931f1f080 100644
--- a/cuda_bindings/cuda/bindings/_internal/utils.pyx
+++ b/cuda_bindings/cuda/bindings/_internal/utils.pyx
@@ -42,13 +42,7 @@ cdef bint is_nested_sequence(data):
 
 
 cdef void* get_buffer_pointer(buf, Py_ssize_t size, readonly=True) except*:
-    """Return a raw pointer into ``buf``.
-
-    The caller must keep ``buf`` alive AND not resize it while the pointer is in
-    use. The buffer export is released before this returns, so resizing a mutable
-    buffer (e.g. ``bytearray``) during a ``nogil`` call is a use-after-free.
-    Prefer immutable ``bytes``.
-    """
+    """The caller must ensure ``buf`` is alive when the returned pointer is in use."""
     cdef void* bufPtr
     cdef int flags = cpython.PyBUF_ANY_CONTIGUOUS
     if not readonly:
diff --git a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py b/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
index cc3d877fe4c..c700ca28f92 100644
--- a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
+++ b/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
@@ -397,16 +397,16 @@ def __init__(
         self._entries = self._root / _ENTRIES_SUBDIR
         self._tmp = self._root / _TMP_SUBDIR
         self._max_size_bytes = max_size_bytes
-        # Cache holds compiled device code loaded via cuLibraryLoadData, so keep
-        # it owner-only (0o700) so other local users can't read or replace it.
-        # mkdir's mode is umask-masked and exist_ok keeps an existing dir's
-        # perms, so re-chmod on POSIX.
-        self._root.mkdir(parents=True, exist_ok=True, mode=0o700)
-        self._entries.mkdir(exist_ok=True, mode=0o700)
+        self._root.mkdir(parents=True, exist_ok=True)
+        self._entries.mkdir(exist_ok=True)
+        # ``tmp`` stages in-flight writes before they are atomically renamed into
+        # ``entries``. Create it owner-only (0o700) to keep other local users from
+        # reading or replacing compiled device code mid-write. ``root``/``entries``
+        # intentionally inherit the umask / any pre-existing permissions so a
+        # deliberately shared cache directory keeps working (see PR #2399 review).
+        # mkdir's mode is not chmod'd afterward, so an existing ``tmp`` is left
+        # as-is rather than re-tightened.
         self._tmp.mkdir(exist_ok=True, mode=0o700)
-        if os.name != "nt":
-            for d in (self._root, self._entries, self._tmp):
-                os.chmod(d, 0o700)
         # Opportunistic startup sweep of orphaned temp files left by any
         # crashed writers. Age-based so concurrent in-flight writes from
         # other processes are preserved.
diff --git a/cuda_core/tests/test_program_cache.py b/cuda_core/tests/test_program_cache.py
index 2be669f49dc..aa368a08c6d 100644
--- a/cuda_core/tests/test_program_cache.py
+++ b/cuda_core/tests/test_program_cache.py
@@ -2539,11 +2539,11 @@ def reader(tid: int) -> None:
 
 
 @pytest.mark.skipif(os.name == "nt", reason="POSIX permission bits only")
-def test_program_cache_dir_created_owner_only(tmp_path):
-    """#359/#375: the on-disk program cache stores raw executable device code,
-    so its directory tree must be created owner-only (0o700), independent of the
-    inherited umask, so other local principals can neither read cached kernels
-    nor plant a binary this process would later load."""
+def test_program_cache_tmp_dir_created_owner_only(tmp_path):
+    """#359/#375: ``tmp`` stages in-flight compiled device code before the atomic
+    rename into ``entries``, so it must be created owner-only (0o700) regardless of
+    the inherited umask. ``root``/``entries`` intentionally inherit the umask to keep
+    deliberately shared caches working (PR #2399 review), so only ``tmp`` is asserted."""
     import stat
 
     from cuda.core.utils._program_cache._file_stream import FileStreamProgramCache
@@ -2551,25 +2551,25 @@ def test_program_cache_dir_created_owner_only(tmp_path):
     root = tmp_path / "pc"
     FileStreamProgramCache(path=root)
 
-    for d in (root, root / "entries", root / "tmp"):
-        mode = stat.S_IMODE(os.stat(d).st_mode)
-        assert mode == 0o700, f"{d} has mode {oct(mode)}"
+    mode = stat.S_IMODE(os.stat(root / "tmp").st_mode)
+    assert mode == 0o700, f"tmp has mode {oct(mode)}"
 
 
 @pytest.mark.skipif(os.name == "nt", reason="POSIX permission bits only")
-def test_program_cache_dir_tightens_preexisting_world_writable(tmp_path):
-    """#359: a pre-created world-writable cache root (the poisoning precondition)
-    must be re-tightened to 0o700; mkdir(exist_ok=True) alone would keep the
-    attacker-supplied permissions."""
+def test_program_cache_preexisting_shared_root_used_as_is(tmp_path):
+    """PR #2399 review: a deliberately shared cache root (e.g. group-writable on a
+    compute cluster) must be used as-is, not re-tightened. Only ``tmp`` is forced
+    owner-only; ``root`` keeps whatever permissions it was created with."""
     import stat
 
     from cuda.core.utils._program_cache._file_stream import FileStreamProgramCache
 
     root = tmp_path / "pc"
     root.mkdir()
-    # Simulate the attack precondition: a pre-existing world-writable cache dir.
+    # Simulate an intentionally shared cache directory.
     os.chmod(root, 0o777)  # noqa: S103
 
     FileStreamProgramCache(path=root)
 
-    assert stat.S_IMODE(os.stat(root).st_mode) == 0o700
+    assert stat.S_IMODE(os.stat(root).st_mode) == 0o777
+    assert stat.S_IMODE(os.stat(root / "tmp").st_mode) == 0o700

From aee8b16387042a7ae0074f369e46bc24cf7d2e44 Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Wed, 22 Jul 2026 10:15:19 -0700
Subject: [PATCH 12/14] Document accepted security trade-off for shared program
 caches

Record the residual risk from narrowing the cache-permission hardening
(glasswing #359/#375) per PR #2399 review: committed entry files stay 0o600
(contents owner-only via mkstemp+os.replace), tmp/ is owner-only to protect
in-flight writes, but root/entries inherit the umask so shared caches keep
working. In a group-writable shared cache a group member could replace a
cached kernel; tamper-proofing that path needs load-time integrity checks and
is out of scope here.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 .../core/utils/_program_cache/_file_stream.py | 22 +++++++++++++------
 1 file changed, 15 insertions(+), 7 deletions(-)

diff --git a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py b/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
index c700ca28f92..867f7802b8c 100644
--- a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
+++ b/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
@@ -399,13 +399,21 @@ def __init__(
         self._max_size_bytes = max_size_bytes
         self._root.mkdir(parents=True, exist_ok=True)
         self._entries.mkdir(exist_ok=True)
-        # ``tmp`` stages in-flight writes before they are atomically renamed into
-        # ``entries``. Create it owner-only (0o700) to keep other local users from
-        # reading or replacing compiled device code mid-write. ``root``/``entries``
-        # intentionally inherit the umask / any pre-existing permissions so a
-        # deliberately shared cache directory keeps working (see PR #2399 review).
-        # mkdir's mode is not chmod'd afterward, so an existing ``tmp`` is left
-        # as-is rather than re-tightened.
+        # Security scope (glasswing #359/#375, narrowed per PR #2399 review):
+        #  - Confidentiality: committed entries are written via ``mkstemp`` (0o600)
+        #    and ``os.replace`` preserves that mode, so cached kernel *contents*
+        #    stay owner-readable even though ``entries/`` itself inherits the umask.
+        #    Other local users can list entry filenames (key hashes) but not read
+        #    the compiled code.
+        #  - ``tmp`` stages in-flight writes before the atomic rename; create it
+        #    owner-only (0o700) so no one can read or swap device code mid-write.
+        #  - ``root``/``entries`` intentionally inherit the umask / any pre-existing
+        #    permissions so a deliberately shared cache (e.g. group-writable on a
+        #    cluster) keeps working. Accepted trade-off: in that shared setup a
+        #    group member with write access to ``entries/`` could replace a cached
+        #    kernel this process later loads. Tamper-proofing shared caches is out
+        #    of scope for permissions alone; it needs load-time integrity checks.
+        # mkdir's mode is not chmod'd afterward, so an existing dir is used as-is.
         self._tmp.mkdir(exist_ok=True, mode=0o700)
         # Opportunistic startup sweep of orphaned temp files left by any
         # crashed writers. Age-based so concurrent in-flight writes from

From 644ae82c5a7c98bd8b8ee1d0011015e44bf72c74 Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Wed, 22 Jul 2026 10:18:50 -0700
Subject: [PATCH 13/14] Simplify cache permission comment; drop internal tool
 name

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 .../core/utils/_program_cache/_file_stream.py | 25 ++++++++-----------
 1 file changed, 10 insertions(+), 15 deletions(-)

diff --git a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py b/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
index 867f7802b8c..eb71abf5446 100644
--- a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
+++ b/cuda_core/cuda/core/utils/_program_cache/_file_stream.py
@@ -397,23 +397,18 @@ def __init__(
         self._entries = self._root / _ENTRIES_SUBDIR
         self._tmp = self._root / _TMP_SUBDIR
         self._max_size_bytes = max_size_bytes
+        # Permissions (see PR #2399):
+        # root/ and entries/ use default permissions so a shared cache (e.g. one
+        # a group shares on a cluster) keeps working. The cached files themselves
+        # are still private: each is written to tmp/ as owner-only and moved into
+        # entries/, which keeps its permissions. tmp/ is made owner-only so no one
+        # can read or swap a file while it's being written. We don't chmod, so an
+        # existing directory is left as-is.
+        # Trade-off: if a group deliberately shares a writable entries/, a member
+        # could replace a cached file. Blocking that needs a check at load time,
+        # not just permissions, and is out of scope here.
         self._root.mkdir(parents=True, exist_ok=True)
         self._entries.mkdir(exist_ok=True)
-        # Security scope (glasswing #359/#375, narrowed per PR #2399 review):
-        #  - Confidentiality: committed entries are written via ``mkstemp`` (0o600)
-        #    and ``os.replace`` preserves that mode, so cached kernel *contents*
-        #    stay owner-readable even though ``entries/`` itself inherits the umask.
-        #    Other local users can list entry filenames (key hashes) but not read
-        #    the compiled code.
-        #  - ``tmp`` stages in-flight writes before the atomic rename; create it
-        #    owner-only (0o700) so no one can read or swap device code mid-write.
-        #  - ``root``/``entries`` intentionally inherit the umask / any pre-existing
-        #    permissions so a deliberately shared cache (e.g. group-writable on a
-        #    cluster) keeps working. Accepted trade-off: in that shared setup a
-        #    group member with write access to ``entries/`` could replace a cached
-        #    kernel this process later loads. Tamper-proofing shared caches is out
-        #    of scope for permissions alone; it needs load-time integrity checks.
-        # mkdir's mode is not chmod'd afterward, so an existing dir is used as-is.
         self._tmp.mkdir(exist_ok=True, mode=0o700)
         # Opportunistic startup sweep of orphaned temp files left by any
         # crashed writers. Age-based so concurrent in-flight writes from

From 8623e745c4932b12c263fabb01a0891c59ff5cbd Mon Sep 17 00:00:00 2001
From: Rob Parolin <rparolin@nvidia.com>
Date: Wed, 22 Jul 2026 10:20:21 -0700
Subject: [PATCH 14/14] Strip internal issue numbers from cache permission test
 docstring

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---
 cuda_core/tests/test_program_cache.py | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/cuda_core/tests/test_program_cache.py b/cuda_core/tests/test_program_cache.py
index aa368a08c6d..a8d3fc85f7e 100644
--- a/cuda_core/tests/test_program_cache.py
+++ b/cuda_core/tests/test_program_cache.py
@@ -2540,10 +2540,10 @@ def reader(tid: int) -> None:
 
 @pytest.mark.skipif(os.name == "nt", reason="POSIX permission bits only")
 def test_program_cache_tmp_dir_created_owner_only(tmp_path):
-    """#359/#375: ``tmp`` stages in-flight compiled device code before the atomic
-    rename into ``entries``, so it must be created owner-only (0o700) regardless of
-    the inherited umask. ``root``/``entries`` intentionally inherit the umask to keep
-    deliberately shared caches working (PR #2399 review), so only ``tmp`` is asserted."""
+    """``tmp`` stages in-flight compiled device code before the atomic rename into
+    ``entries``, so it must be created owner-only (0o700) regardless of the inherited
+    umask. ``root``/``entries`` intentionally inherit the umask to keep deliberately
+    shared caches working (PR #2399 review), so only ``tmp`` is asserted."""
     import stat
 
     from cuda.core.utils._program_cache._file_stream import FileStreamProgramCache