Add persistent program cache for Program.compile (#1912)

* feat(core.utils): persistent program cache

Add a bytes-in / bytes-out cache abstraction and two backends for
caching compiled CUDA programs across process boundaries.

* ``ProgramCacheResource`` -- abstract base. Concrete backends store
  raw binary bytes keyed by ``bytes`` or ``str``; reads return the
  same payload. ``__setitem__`` accepts ``bytes``, ``bytearray``,
  ``memoryview``, or any :class:`~cuda.core.ObjectCode` (path-backed
  too -- the file is read at write time so the cached entry holds
  the binary content, not a path that could move). Provides default
  ``get``, ``update`` (mapping or pairs), ``close``, and context
  manager. ``__contains__`` is intentionally NOT abstract: the racy
  ``if key in cache; data = cache[key]`` idiom is steered toward
  ``cache.get(key)`` instead.

* ``InMemoryProgramCache`` -- single-process LRU on
  ``collections.OrderedDict`` with ``threading.RLock`` and a
  size-only cap. Reads promote via ``move_to_end``.

* ``FileStreamProgramCache`` -- directory of atomic per-entry files.
  Writes stage to ``tmp/`` then ``os.replace`` into
  ``entries/<2-char>/<blake2b-hex>``; concurrent readers never see a
  torn file. Each entry is the raw compiled binary (no pickle, no
  framing) so files are directly consumable by external NVIDIA
  tools (``cuobjdump``, ``nvdisasm``, ``cuda-gdb``). Eviction is
  true LRU via ``st_atime`` (the read path calls ``os.utime`` to
  bypass ``relatime`` / ``NtfsDisableLastAccessUpdate`` /
  ``noatime``). Stat-guarded prunes refuse to unlink entries
  another process replaced mid-eviction. ``tmp/`` is recreated on
  every write so an external wipe doesn't crash later writes.
  Default cache directory comes from
  ``platformdirs.user_cache_path("cuda-python", appauthor=False,
  opinion=False) / "program-cache"``.

* Windows sharing-violation handling -- ``os.replace``,
  ``path.stat() + read_bytes()``, and ``path.unlink`` all retry on
  winerror 5/32/33 with a bounded backoff (~185 ms). The
  ``_is_windows_sharing_violation`` predicate filters EACCES only
  when ``winerror`` is absent so non-sharing winerrors propagate as
  the real config errors they are. Off-Windows ``PermissionError``
  always propagates.

* ``make_program_cache_key`` -- escape hatch for callers whose
  compile inputs require an ``extra_digest`` (header / PCH content
  fingerprints, NVVM libdevice). Builds a 32-byte blake2b digest
  via a backend-strategy pattern: a ``_KeyBackend`` ABC with
  per-code-type subclasses (``_NvrtcBackend``, ``_LinkerBackend``,
  ``_NvvmBackend``) owns each backend's validation, code coercion,
  option fingerprinting, name-expression handling, version probe,
  and extra-payload hashing. The orchestrator dispatches via
  ``_BACKENDS_BY_CODE_TYPE[code_type]`` and assembles the digest in
  fixed order. Backend gates match ``Program.compile``: rejects
  inputs the real compile would reject (side-effect options,
  external-content options without an ``extra_digest``,
  driver-linker-unsupported options, NVRTC ``options.name`` with a
  directory component). NVVM ``extra_sources`` is hashed in
  caller-provided order because NVVM module linking is
  order-dependent in the general case (overlapping symbols, weak
  definitions); canonicalising would silently change behavior for
  order-dependent inputs.

Adds ``platformdirs >=3.0`` to ``cuda_core/pyproject.toml`` and the
matching pixi manifests.

Tests cover the abstract contract, key-construction matrix
(deterministic, supported-target gates, backend-probe taint, gate
canonicalization, side-effect / external-content / dir-component
guards, schema version mixing), single-process CRUD and LRU,
atomic-write race coverage, atime LRU promotion, stat-guarded
prune / atime touch / clear / size-cap, default-dir resolution via
platformdirs, the ``_is_windows_sharing_violation`` predicate's
truth table including the regression case (non-sharing winerror
plus EACCES propagates), tmp-dir recreation after external wipe,
multiprocess concurrent writers / reader-vs-writer torn-file
safety / size-cap eviction race.

* feat(core): Program.compile(cache=...) convenience wrapper

Adds a ``cache=`` keyword to :meth:`cuda.core.Program.compile` that
threads the persistent cache machinery into the high-level compile
path. With ``cache=None`` (the default) the call is byte-identical
to the un-cached path -- no key derivation, no extra import, no
behavior change.

When a cache is provided, the wrapper derives a key via
:func:`~cuda.core.utils.make_program_cache_key` from the program's
source, options, and target type; checks the cache; on hit,
returns a fresh
``ObjectCode._init(hit_bytes, target_type, name=self._options.name)``;
on miss, runs the underlying compile and stores
``cache[key] = compiled`` (the cache extracts ``bytes(obj.code)``).

Two compile-time guards close obvious footguns:

* ``name_expressions`` plus ``cache=`` raises ``ValueError``.
  NVRTC populates ``ObjectCode.symbol_mapping`` from name-expression
  mangling at compile time, and that mapping isn't carried in the
  binary the cache stores. Without this guard the first call (miss)
  would return an ObjectCode with mappings populated, while every
  subsequent call (hit) would return one without -- silently
  breaking later ``get_kernel(name_expression)`` lookups that work
  on the uncached path. Compiles that need name_expressions should
  run without ``cache=``, or look up mangled symbols by hand from
  the cached ``ObjectCode``.

* Inputs whose compilation effect isn't captured by the key
  (``include_path``, ``pre_include``, ``pch``, ``use_pch``,
  ``pch_dir``, NVVM ``use_libdevice=True``, NVRTC ``options.name``
  with a directory component, side-effect options like
  ``create_pch`` / ``time`` / ``fdevice_time_trace``) propagate the
  ``ValueError`` from ``make_program_cache_key`` -- those callers
  should use ``make_program_cache_key`` directly with an
  ``extra_digest`` covering the external content.

Cache hits also mirror the uncached path's NVRTC-PTX loadability
warning: when ``self._backend == "NVRTC"``, ``target_type ==
"ptx"``, and ``_can_load_generated_ptx()`` returns False, a
``RuntimeWarning`` is emitted before returning the cached bytes.
Loadability is a property of the active driver, not of how the
bytes were produced, so the warning applies equally to cached PTX.

Supporting refactors:

* Unify ``Program``'s source retention into a single ``_code``
  field (was split between ``_code`` for NVVM and a separate
  ``_source`` for c++/ptx). ``_code`` is now always bytes; the
  cache wrapper decodes back to ``str`` for c++/ptx before passing
  to ``make_program_cache_key`` (which only accepts bytes for NVVM).

* Move the actual compile call into a module-level
  ``_program_compile_uncached`` so tests can monkeypatch the seam
  without going through NVRTC. ``Program`` is a ``cdef class``, so
  its methods cannot be reassigned from Python -- the seam has to
  live outside the class.

* The unified ``_code`` field also exposed a pre-existing bug on
  the NVVM path: the C pointer was being recomputed from the
  caller's original ``code`` argument rather than from
  ``self._code``, which crashed for ``bytearray`` inputs that the
  field's bytes coercion handled cleanly. Fixed; regression test
  added in ``test_program.py``.

Tests in ``test_program_compile_cache.py`` cover both halves of the
contract: the wrapper-level miss/hit/error paths against a recording
stub (verifying it's duck-typed and doesn't require subclassing
``ProgramCacheResource``), the rejection paths (name_expressions,
extra_digest-required options, side-effect options, NVRTC
``options.name`` with a directory component), the PTX loadability
warning on cache hit (positive: warns when the driver can't load
the cached PTX; negative: stays quiet otherwise), and a real NVRTC
end-to-end roundtrip using ``FileStreamProgramCache`` across reopen
so the bytes match across processes.

* refactor(core.utils): import program cache eagerly, fix ruff findings

Drop the module-level __getattr__/__dir__ shim that lazily exposed the
program-cache classes. Import them eagerly alongside the memoryview
helpers. Remove the two tests that pinned the lazy-import behaviour.

Also pick up ruff's auto-fixes (import-block blank lines, long-line
reformat) and rename the unused classmethod argument to satisfy ARG005.

* docs(core): nest Program caches under CUDA compilation toolchain

Move the Program caches autosummary out of the trailing
cuda.core.utils block and into a subsection of CUDA compilation
toolchain so the cache classes sit next to Program/Linker. The
subsection switches the current module to cuda.core.utils for the
autosummary and switches back afterwards. make_program_cache_key
moves with the cache classes.

* docs(core.utils): nudge cache callers toward close()

The default close() on ProgramCacheResource is a no-op because the
in-tree backends happen not to hold long-lived state. Future
backends will -- file handles, sockets, db connections -- so the
docstring now spells that out and tells callers to use the context
manager or call close() explicitly so their code is portable across
backends.

* docs(core.utils): teach cache examples with public ObjectCode API

Two docstring examples (in ProgramCacheResource and
make_program_cache_key) reached into the private cuda.core._module
path and used the private ObjectCode._init constructor. Switch to
the public from cuda.core import ObjectCode plus
ObjectCode.from_cubin(...) so users learning the cache API don't
get steered at private surface.

* feat(core): require Program.compile(cache=...) to be a ProgramCacheResource

cache= used to accept any object with the get/__setitem__ duck-typed
shape. Tighten the contract: the wrapper now isinstance-checks against
ProgramCacheResource and raises TypeError up front when callers pass
a plain dict-like. Subclasses get the get/update/close/__enter__/__exit__
defaults from the ABC and a portable interface across backends.

Update the recording test cache to subclass the ABC and add a regression
test that a duck-typed cache is rejected with TypeError.

* test(core): mock driver_version to keep _can_load_generated_ptx as cpdef

The PTX cache-hit warning tests used to monkeypatch
_can_load_generated_ptx itself, which forced the helper to a plain def
so Cython would not early-bind the in-module call past the patch.
Restore it to cpdef bint ... except? -1 and instead pin
driver_version() in the test to (0, 0, 0) (warning expected) or
(999, 0, 0) (no warning) so the cpdef body computes the desired
result. Cython compiles each global lookup inside the helper as a
fresh module-globals fetch, so swapping driver_version on the
module object is enough to steer the comparison.

* docs(core): note cache hits do not write to logs

* fix(core.utils): cache _LinkerBackend driver decision per key call

_LinkerBackend.validate, option_fingerprint, and hash_version_probe each
re-probed _decide_nvjitlink_or_driver(), so a flapping probe could mint
a key whose option fingerprint and version probe disagreed on which
linker is in use.

Cache the decision (and any probe exception) on a per-instance basis,
instantiate _BACKENDS_BY_CODE_TYPE entries fresh per make_program_cache_key
call so the cache lives exactly one call, and thread the decision into
_linker_backend_and_version() instead of letting it probe a third time.

Tests that monkeypatched _linker_backend_and_version now accept the
extra use_driver argument (or *args/**kwargs in the failure-path test).

* fix(core): lowercase target_type in Program.compile and cache key

code_type was already normalised at Program init, but target_type was
checked case-sensitively against {"ptx", "cubin", "ltoir"} in
Program_compile (so compile(target_type="PTX") used to raise) and the
cache key path inherited the same asymmetry from
make_program_cache_key. Lowercase target_type at the top of
Program.compile and at the entry to make_program_cache_key so callers
who pass "PTX" get the same dispatch and the same cache key as "ptx".

* refactor(core.utils): collapse _LINKER_RELEVANT_FIELDS into _LINKER_FIELD_GATES

The names tuple and the gates dict had to be kept in sync by hand: a
field added to one but forgotten in the other would silently slip out
of the PTX fingerprint. Drop the tuple and iterate the dict directly,
so the dict is the single source of truth for which ProgramOptions
fields perturb a PTX cache key.

* fix(core.utils): contextual error when path-backed ObjectCode vanishes

_extract_bytes used to let a bare FileNotFoundError bubble up from
Path(code).read_bytes(), so cache[key] = obj failures pointed only at
the missing path with no hint that the cache was reading a path-backed
ObjectCode. Wrap the FileNotFoundError with a message that names both
the cache operation and the missing file so debugging the case stays
self-explanatory.

* docs(core): clarify cache= vs ProgramOptions.no_cache independence

* test(core.utils): tighten multiprocess cache assertions

The reader test only checked that every read returned non-None; a
half-written file with non-empty bytes would pass. Carry the seeded
payload into the worker, count exact-byte mismatches, and require
zero. The eviction-race test wrote a final uncontested cache[key] =
payload + b"final" after the churner exited, so the post-race
endswith assertion would pass even with a broken stat-guard. Drop the
final write and assert the entry survives carrying a rewriter payload
prefix -- if the stat-guarded eviction path is broken, the in-race
write is the one that vanishes.

* test(core.utils): add InMemoryProgramCache thread test and FileStream overwrite test

InMemoryProgramCache claims thread-safety via the RLock that wraps every
method but had no concurrent-thread coverage. Add a stress test with 4
writers + 4 readers x 200 ops against a size-capped cache that verifies
no exceptions, no deadlocks (RLock reentrance through __setitem__ ->
_evict_to_caps -> popitem), and that internal accounting (_total_bytes,
len(_entries), len(cache)) stays consistent under contention.

FileStreamProgramCache had no overwrite test analogous to
test_inmemory_cache_overwrite_replaces_value_and_updates_size. Add one
that writes a key twice and asserts the second value reads back, len
stays at 1, and exactly one entry file lives on disk -- so a leaked
entry from a botched os.replace would surface here.

* refactor(core.utils): drop dead empty-key branch in _path_for_key

* perf(core.utils): drop redundant stat in FileStreamProgramCache.__len__

* feat(core.utils): reject max_size_bytes=0 in cache backends

max_size_bytes=0 used to slip past the >=0 guard but turned the cache
into a black hole: every write was immediately evicted on its own
size-cap pass. There is no legitimate use for that, so tighten the
guard to >0 (or None for unbounded) and update both backends and the
matching tests.

* docs(core.utils): note ptxas_options order is intentionally significant

* docs(core.utils): document the _keys.py <-> _program.pyx import contract

* docs(core.utils): note options.name double-hash on PTX is intentional

* docs(core.utils): document burst-write over-eviction trade-off

* docs(core): explain why Program.compile cache decode is provably safe

* refactor(core.utils): extract _stat_key for the four stat-guarded paths

The (st_ino, st_size, st_mtime_ns) triple was open-coded in
_touch_atime (fd-based and path-based fallbacks),
_prune_if_stat_unchanged, and _enforce_size_cap. Centralise the
fingerprint as _stat_key(st) so all four readers compare the same
fields and the invariant has one place to read.

* refactor(core.utils): collapse Windows sharing-retry loops

Replace, stat-and-read, and unlink each carried their own copy of the
_REPLACE_RETRY_DELAYS / sleep / try-op / PermissionError loop. Centralise
the loop as _with_sharing_retry(op, on_exhausted=...) and let each
caller plug in its own success-on-success and exhausted-budget
behaviour. Net behaviour is unchanged (exhaustion semantics for each
public helper are preserved via the on_exhausted callback).

* fix(test): walk sharded entries dir in FileStream overwrite assertion

FileStreamProgramCache shards cache files into
entries/<digest[:2]>/<digest[2:]>, so the overwrite test's iterdir
filter on entries/ saw only the digest-prefix subdir (no is_file()
match) and reported 0 entries. Switch to rglob so the assertion
counts actual entry files. CI from the previous push caught this.

* fix(test): collapse nested with into a single contextmanager group (SIM117)

* fix(core.utils): keep path-backed ObjectCode error message Windows-friendly

`{code!r}` doubles every backslash on Windows, so the error string holds
`'C:\\Users\\...\\file'` while `str(src)` only has single
backslashes. Naive `str(path) in str(exc)` checks (and the
`test_filestream_cache_path_backed_object_code_missing_file_message`
assertion) failed on win-64 as a result. Drop the `!r` so the path
appears verbatim; the message still quotes it via the surrounding
sentence punctuation.

* perf(core.utils): track FileStream cache size incrementally

Avoid the per-write directory walk in `_enforce_size_cap` by maintaining
a running byte total. The tracker is seeded once from `_compute_total_size`
at open time, updated on `__setitem__` (net delta from the old entry's
size), `__delitem__` (subtract the unlinked file's size), and `clear`
(re-derive from the post-clear state). When a write doesn't push the
running total above `max_size_bytes`, eviction stays a no-op -- writes
become O(1) instead of O(n) in the cache size.

Cross-process drift (other writers/deleters working on the same root)
self-corrects: any time `_enforce_size_cap` actually runs its scan, it
reseeds `_tracked_size_bytes` from the observed disk total, so
overestimates trigger one extra scan and then settle.

Skipped entirely when `max_size_bytes is None` (no cap, no need for a
tracker). Mutations are guarded by `_size_lock` so multi-threaded
writers in the same process don't interleave the read-modify-write on
the int.

Also switch `_iter_entry_paths` from `Path.iterdir` to `os.scandir`:
the dirent type cache lets `is_dir`/`is_file` answer without a
separate `stat` syscall on filesystems that report it (ext4, NTFS,
...). Behaviourally equivalent (the cache layout never creates symlinks,
and we pass `follow_symlinks=False` to match `pathlib`'s previous
no-symlink-confusion default for our paths).

Tests cover three guarantees: (1) writes that stay under the cap don't
call `_enforce_size_cap`, (2) writes that cross the cap do call it,
and (3) external deletion behind the cache's back is reconciled by the
next eviction pass.

* test(core.utils): drop SUPPORTED_TARGETS cross-check

The test parsed `_program.pyx` to extract `SUPPORTED_TARGETS` and
compare against `_SUPPORTED_TARGETS_BY_CODE_TYPE` in the cache. Source
parsing for a duplication check is not worth the maintenance cost --
a reviewer eyeballing both definitions catches drift just as well, and
the test was already broken once by upstream's StrEnum migration.

* refactor(core.utils): idiomatic with-as on os.scandir, yield from inner

`os.scandir` returns an iterator that holds an OS directory handle
(POSIX dirent fd, Windows FindFirstFileW); the context manager closes
that handle deterministically. Using `with os.scandir(...) as outer`
is the conventional spelling -- splitting the call from the `with` to
land the FileNotFoundError catch on the call alone added one binding
for no behavioural difference. Wrap each `with` in its own
`try/except FileNotFoundError` instead.

The inner loop is also a plain filter-and-yield over the DirEntry
iterator, which is exactly what `yield from <generator-expr>`
expresses.

* refactor(core.utils): extract _iter_tmp_entries and _sum_tmp_sizes

The same scandir-then-stat-for-size loop appeared in three places
(`_compute_total_size`, `_sweep_stale_tmp_files`, `_enforce_size_cap`).

`_iter_tmp_entries` covers the scandir + is_file filter +
context-managed cleanup that all three need.

`_sum_tmp_sizes` covers the size accumulation that both
`_compute_total_size` and `_enforce_size_cap` need.

Net result: each callsite now reads as one obvious line instead of
seven, and a future change to the temp-walk discipline (e.g. switching
to a different scan strategy) lands in one place.

* fix(core.utils): clamp tracked size at zero on __delitem__

`__delitem__` could walk `_tracked_size_bytes` negative under a race
with `_enforce_size_cap`'s reseed: if the eviction scan runs AFTER
this delete unlinks (so its reseed value excludes the deleted entry)
but BEFORE this delete's subtract, the subtract undercounts by
`size`. Repeated under contention, the tracker crosses zero -- and
once negative, the `tracker > cap` check that gates eviction never
fires again, so the cache grows without bound and there is no
self-healing path (the only reseed point is the function that no
longer runs).

Clamp `tracker = max(0, tracker - size)` so the tracker can't enter
the permanently-broken state. Worst case after the race is
undercounting reality, which the next eviction's reseed corrects;
that's the same self-healing path the existing tests already exercise.

Reported by leofang in PR review.

Author: cpcloud

cuda_core/cuda/core/_program.pxd

@@ -17,5 +17,6 @@ cdef class Program:
         object _compile_lock  # Per-instance lock for compile-time mutation
         bint _use_libdevice      # Flag for libdevice loading
         bint _libdevice_added
-        bytes _nvrtc_code       # Source code for NVRTC retry (PCH auto-resize)
+        bytes _code             # Source code as bytes: used for key derivation and NVRTC PCH retry
+        str _code_type          # Normalised code_type ("c++", "ptx", "nvvm")
         str _pch_status         # PCH creation outcome after compile

cuda_core/cuda/core/_program.pyx

@@ -85,7 +85,12 @@ cdef class Program:
         self._h_nvvm.reset()
 
     def compile(
-        self, target_type: ObjectCodeFormatType | str, name_expressions: tuple | list = (), logs = None
+        self,
+        target_type: ObjectCodeFormatType | str,
+        name_expressions: tuple | list = (),
+        logs=None,
+        *,
+        cache: "ProgramCacheResource | None" = None,
     ) -> ObjectCode:
         """Compile the program to the specified target type.
 
@@ -98,13 +103,126 @@ cdef class Program:
             Used for template instantiation and similar cases.
         logs : object, optional
             Object with a ``write`` method to receive compilation logs.
+            On a cache hit no compilation runs and ``logs`` receives
+            nothing -- callers that rely on log output to confirm a
+            compile happened should compile without ``cache=``.
+        cache : :class:`~cuda.core.utils.ProgramCacheResource`, optional
+            If provided, the compiled binary is looked up in ``cache`` via a
+            key derived from the program's code, options, and ``target_type``.
+            On a hit the cached bytes are wrapped in a fresh
+            :class:`~cuda.core.ObjectCode` (with the same ``target_type``
+            and ``ProgramOptions.name``) and returned without re-compiling;
+            on a miss the compile output is stored as raw bytes (the cache
+            extracts ``bytes(object_code.code)``). Passing a non-empty
+            ``name_expressions`` together with ``cache=`` raises
+            ``ValueError``: NVRTC populates
+            ``ObjectCode.symbol_mapping`` at compile time and that mapping
+            is not carried in the binary the cache stores, so cache hits
+            would silently miss ``get_kernel(name_expression)`` lookups.
+            Options that require an ``extra_digest`` (``include_path``,
+            ``pre_include``, ``pch``, ``use_pch``, ``pch_dir``, NVVM
+            ``use_libdevice=True``, or NVRTC ``options.name`` with a
+            directory component) raise ``ValueError`` via
+            :func:`~cuda.core.utils.make_program_cache_key`; for those
+            compiles, use the manual ``make_program_cache_key(...)``
+            pattern directly.
+
+            ``cache=`` is independent of ``ProgramOptions.no_cache``: the
+            former controls this program-level cache (compiled-output
+            reuse across calls), while ``no_cache`` is forwarded to the
+            Linker to disable its in-process JIT cache for cuLink/nvJitLink.
+            Setting ``options.no_cache=True`` does not bypass ``cache=``,
+            and vice-versa.
 
         Returns
         -------
         :class:`~cuda.core.ObjectCode`
             The compiled object code.
         """
-        return Program_compile(self, str(target_type), name_expressions, logs)
+        # Mirror Program_init's code_type normalization so callers can pass
+        # ``ObjectCodeFormatType.PTX`` or ``"PTX"`` and get the same routing
+        # / cache key as the lowercase string. ``Program_compile_nvrtc``
+        # keys on lowercase ``target_type`` and ``make_program_cache_key``
+        # lowercases too.
+        target_type = str(target_type).lower()
+
+        if cache is None:
+            return _program_compile_uncached(self, target_type, name_expressions, logs)
+
+        # Deferred import to avoid a circular import between _program and
+        # cuda.core.utils._program_cache (the cache module already imports
+        # ProgramOptions from this module). Import from the leaf module so
+        # tests that monkeypatch make_program_cache_key via that path
+        # intercept reliably.
+        from cuda.core.utils._program_cache import (
+            ProgramCacheResource,
+            make_program_cache_key,
+        )
+
+        if not isinstance(cache, ProgramCacheResource):
+            raise TypeError(
+                "cache must be an instance of "
+                "cuda.core.utils.ProgramCacheResource; got "
+                f"{type(cache).__name__}"
+            )
+
+        # ``name_expressions`` is incompatible with the cache: NVRTC
+        # populates ``ObjectCode.symbol_mapping`` from name-expression
+        # mangling at compile time, and that mapping isn't carried in
+        # the binary bytes the cache stores. Without this guard the
+        # first call (cache miss) would return an ObjectCode with
+        # symbol_mapping populated, while every subsequent call (hit)
+        # would return one without -- silently breaking later
+        # ``get_kernel(name_expression)`` lookups that work on the
+        # uncached path. Fail loud here instead.
+        if name_expressions:
+            raise ValueError(
+                "Program.compile(cache=...) does not support name_expressions: "
+                "ObjectCode.symbol_mapping is populated by NVRTC at compile "
+                "time and is not preserved across a cache round-trip, so cache "
+                "hits would silently break get_kernel(name_expression) lookups "
+                "that the uncached path supports. Compile without cache= when "
+                "name_expressions are needed, or look up mangled symbols by "
+                "hand from the cached ObjectCode."
+            )
+
+        # ``self._code`` is always stored as bytes (see ``Program_init``),
+        # but ``make_program_cache_key`` only accepts bytes when
+        # ``code_type == "nvvm"`` -- c++/ptx must be ``str``. The bytes
+        # came from ``code.encode()`` on a ``str`` Program_init validated
+        # via ``assert_type(code, str)``, so this round-trip is always
+        # safe; no try/except needed.
+        code_for_key = self._code if self._code_type == "nvvm" else self._code.decode("utf-8")
+
+        key = make_program_cache_key(
+            code=code_for_key,
+            code_type=self._code_type,
+            options=self._options,
+            target_type=target_type,
+        )
+        hit_bytes = cache.get(key)
+        if hit_bytes is not None:
+            # The uncached NVRTC path warns when the active driver can't
+            # load freshly-generated PTX; that loadability is a property
+            # of the driver, not of how the bytes were produced, so the
+            # warning applies equally to cached PTX. Mirror it here so a
+            # cache hit doesn't silently hide an incompatibility that the
+            # uncached call would have surfaced.
+            if (
+                self._backend == "NVRTC"
+                and target_type == "ptx"
+                and not _can_load_generated_ptx()
+            ):
+                warn(
+                    "The CUDA driver version is older than the backend version. "
+                    "The generated ptx will not be loadable by the current driver.",
+                    stacklevel=2,
+                    category=RuntimeWarning,
+                )
+            return ObjectCode._init(hit_bytes, target_type, name=self._options.name)
+        compiled = _program_compile_uncached(self, target_type, name_expressions, logs)
+        cache[key] = compiled
+        return compiled
 
     @property
     def pch_status(self) -> PCHStatusType | None:
@@ -505,6 +623,19 @@ class ProgramOptions:
 # Private Classes and Helper Functions
 # =============================================================================
 
+
+def _program_compile_uncached(program, target_type, name_expressions, logs):
+    """Run ``Program_compile`` without the cache wrapper.
+
+    Module-level Python function so tests can monkeypatch it from
+    ``cuda.core._program`` to avoid invoking NVRTC when exercising the cache
+    wrapper in :meth:`Program.compile`. ``Program`` itself is a ``cdef class``
+    and its methods cannot be reassigned from Python, so the seam must live
+    outside the class.
+    """
+    return Program_compile(program, target_type, name_expressions, logs)
+
+
 # Module-level state for NVVM lazy loading
 _nvvm_module = None
 _nvvm_import_attempted = False
@@ -620,6 +751,7 @@ cdef inline int Program_init(Program self, object code, str code_type, object op
 
     self._options = options = check_or_create_options(ProgramOptions, options, "Program options")
     code_type = code_type.lower()
+    self._code_type = code_type
     self._compile_lock = threading.Lock()
     self._use_libdevice = False
     self._libdevice_added = False
@@ -640,16 +772,18 @@ cdef inline int Program_init(Program self, object code, str code_type, object op
             HANDLE_RETURN_NVRTC(NULL, cynvrtc.nvrtcCreateProgram(
                 &nvrtc_prog, code_ptr, name_ptr, 0, NULL, NULL))
         self._h_nvrtc = create_nvrtc_program_handle(nvrtc_prog)
-        self._nvrtc_code = code_bytes
+        self._code = code_bytes
         self._backend = str(CompilerBackendType.NVRTC)
         self._linker = None
 
     elif code_type == "ptx":
         assert_type(code, str)
         if options.extra_sources is not None:
             raise ValueError("extra_sources is not supported by the PTX backend.")
+        code_bytes = code.encode()
+        self._code = code_bytes
         self._linker = Linker(
-            ObjectCode._init(code.encode(), code_type), options=_translate_program_options(options)
+            ObjectCode._init(code_bytes, code_type), options=_translate_program_options(options)
         )
         self._backend = str(self._linker.backend)
 
@@ -659,10 +793,13 @@ cdef inline int Program_init(Program self, object code, str code_type, object op
             code = code.encode("utf-8")
         elif not isinstance(code, (bytes, bytearray)):
             raise TypeError("NVVM IR code must be provided as str, bytes, or bytearray")
+        self._code = bytes(code)  # Coerce bytearray -> bytes so retention type is stable
 
-        code_ptr = <const char*>(<bytes>code)
+        # Use self._code (strictly bytes) for the C pointer so a bytearray
+        # input doesn't trip the `<bytes>code` cast at runtime.
+        code_ptr = <const char*>self._code
         name_ptr = <const char*>options._name
-        code_len = len(code)
+        code_len = len(self._code)
 
         with nogil:
             HANDLE_RETURN_NVVM(NULL, cynvvm.nvvmCreateProgram(&nvvm_prog))
@@ -828,7 +965,7 @@ cdef object Program_compile_nvrtc(Program self, str target_type, object name_exp
         HANDLE_RETURN_NVRTC(NULL, cynvrtc.nvrtcSetPCHHeapSize(required))
 
     cdef cynvrtc.nvrtcProgram retry_prog
-    cdef const char* code_ptr = <const char*>self._nvrtc_code
+    cdef const char* code_ptr = <const char*>self._code
     cdef const char* name_ptr = <const char*>self._options._name
     with nogil:
         HANDLE_RETURN_NVRTC(NULL, cynvrtc.nvrtcCreateProgram(

cuda_core/cuda/core/utils.py

@@ -0,0 +1,23 @@
+# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+from cuda.core._memoryview import (
+    StridedMemoryView,
+    args_viewable_as_strided_memory,
+)
+from cuda.core.utils._program_cache import (
+    FileStreamProgramCache,
+    InMemoryProgramCache,
+    ProgramCacheResource,
+    make_program_cache_key,
+)
+
+__all__ = [
+    "FileStreamProgramCache",
+    "InMemoryProgramCache",
+    "ProgramCacheResource",
+    "StridedMemoryView",
+    "args_viewable_as_strided_memory",
+    "make_program_cache_key",
+]

cuda_core/cuda/core/utils/__init__.py

@@ -0,0 +1,36 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Persistent program cache for cuda.core.
+
+Public surface:
+
+* :class:`ProgramCacheResource` -- bytes-in / bytes-out ABC.
+* :class:`InMemoryProgramCache` -- thread-safe LRU dict-backed cache.
+* :class:`FileStreamProgramCache` -- atomic, multi-process directory cache.
+* :func:`make_program_cache_key` -- key derivation for arbitrary
+  ``Program`` configurations.
+
+The package is split into submodules by concern. Tests that need to
+monkeypatch internals (Windows flag, version probes, helpers, ...)
+should reach into the owning submodule (e.g.
+``_program_cache._file_stream._IS_WINDOWS``,
+``_program_cache._keys._linker_backend_and_version``) rather than the
+package object: the symbols re-exported here are only convenience
+aliases and don't intercept calls within the submodules.
+"""
+
+from __future__ import annotations
+
+from ._abc import ProgramCacheResource
+from ._file_stream import FileStreamProgramCache
+from ._in_memory import InMemoryProgramCache
+from ._keys import make_program_cache_key
+
+__all__ = [
+    "FileStreamProgramCache",
+    "InMemoryProgramCache",
+    "ProgramCacheResource",
+    "make_program_cache_key",
+]