Add PersistentProgramCache (sqlite + filestream backends)

[cpcloud]

Summary

• Convert cuda.core.utils from a module to a package; expose cache APIs lazily via __getattr__ so from cuda.core.utils import StridedMemoryView stays lightweight.
• Add ProgramCacheResource ABC with bytes | str keys, context manager, pickle-safety warning, and rejection of path-backed ObjectCode at write time.
• Add make_program_cache_key() — blake2b(32) digest with backend-specific gates that mirror Program/Linker:
  • Versions: cuda-core, NVRTC (c++), libNVVM lib+IR (nvvm), linker backend+version (ptx); driver only on the cuLink path.
  • Validates code_type/target_type against Program.compile's SUPPORTED_TARGETS; rejects bytes-like code for non-NVVM and extra_sources for non-NVVM.
  • NVRTC side-effect (create_pch, time, fdevice_time_trace) and external-content (include_path, pre_include, pch, use_pch, pch_dir) options require extra_digest; NVVM use_libdevice=True likewise.
  • PTX (Linker) options pass through per-field gates that match _prepare_nvjitlink_options / _prepare_driver_options; ptxas_options canonicalised across str/list/tuple/empty shapes; driver-linker hard rejections (time, ptxas_options, split_compile) raise at key time; ftz/prec_div/prec_sqrt/fma collapse under driver linker.
  • Failed env probes mix the exception class name into a *_probe_failed label so broken environments never collide with working ones, while staying stable across processes and repeated calls.
• Add SQLiteProgramCache — single-file sqlite3 (WAL + autocommit), LRU eviction, optional size cap, wal_checkpoint(TRUNCATE) + VACUUM after evictions so the cap bounds real on-disk usage. __contains__ is read-only; __len__ validates and prunes corrupt rows. threading.RLock serialises connection use. Schema-mismatch on open drops tables and rebuilds; corrupt / non-SQLite files reinitialise empty; OperationalError (lock/busy) propagates without nuking the file (and closes the partial connection).
• Add FileStreamProgramCache — multi-process via tmp + os.replace. Hash-based filenames so arbitrary-length keys don't overflow filesystem limits. Reader pruning, clear(), and _enforce_size_cap are all stat-guarded (snapshot (ino, size, mtime_ns), refuse unlink on mismatch) so a concurrent writer's os.replace is preserved. Stale temp files swept on open; live temps count toward the size cap. Windows ERROR_SHARING_VIOLATION/ERROR_LOCK_VIOLATION on os.replace are retried with bounded backoff (~185ms) before being treated as a non-fatal cache miss; other PermissionError and all POSIX failures propagate. __len__ also rejects stored_key/path mismatch.

Program.compile(cache=...) integration is out of scope (tracked by #176/#179).

Test plan

• 177 cache tests — single-process CRUD; LRU/size-cap (logical and on-disk); corruption + __len__ pruning; schema-mismatch table-DROP; threaded SQLite (4 writers + 4 readers × 200 ops); cross-process FileStream stress (writer/reader race exercising the stat-guard prune; clear/eviction race injection via generator cleanup); Windows vs POSIX PermissionError narrowing (winerror 32/33 swallow + retry, others propagate; partial-conn close on OperationalError); lazy-import subprocess test; _SUPPORTED_TARGETS_BY_CODE_TYPE parity test that parses _program.pyx via tokenize + ast.literal_eval.
• End-to-end: real CUDA C++ compile → store in cache → reopen → get_kernel on the deserialised ObjectCode, parametrized over both backends.
• CI: clean across all platforms.

Closes #178

[github-actions]

Doc Preview CI
🚀 View preview at https://nvidia.github.io/cuda-python/pr-preview/pr-1912/
https://nvidia.github.io/cuda-python/pr-preview/pr-1912/cuda-core/
https://nvidia.github.io/cuda-python/pr-preview/pr-1912/cuda-bindings/
https://nvidia.github.io/cuda-python/pr-preview/pr-1912/cuda-pathfinder/
Preview will be ready when the GitHub Pages deployment is complete.

[rwgk]

Generated with the help of Cursor GPT-5.4 Extra High Fast

---

High: make_program_cache_key() misses implicit source-directory header dependencies

make_program_cache_key() only forces extra_digest for explicit include/PCH options in cuda_core/cuda/core/utils/_program_cache.py:393 and cuda_core/cuda/core/utils/_program_cache.py:592, but NVRTC also implicitly searches the source file's directory unless no_source_include is set in cuda_core/cuda/core/_program.pyx:1001.

Program passes options.name straight to nvrtcCreateProgram() in cuda_core/cuda/core/_program.pyx:635, while the key builder only hashes that path string in cuda_core/cuda/core/utils/_program_cache.py:778. That means a workflow like options.name="/path/to/kernel.cu" plus #include "local.h" can reuse a stale cached ObjectCode after local.h changes.

The new tests cover explicit include/PCH knobs, but not this default source-directory include path (cuda_core/tests/test_program_cache.py:765).

Medium: FileStreamProgramCache._enforce_size_cap() can over-evict under concurrent capped writers

After the re-stat at cuda_core/cuda/core/utils/_program_cache.py:1515, a concurrent deleter can remove the candidate before path.unlink(). That FileNotFoundError is suppressed at cuda_core/cuda/core/utils/_program_cache.py:1530, but total is not adjusted, so eviction continues and can delete newer entries unnecessarily.

For a backend explicitly documented for multi-process use, that turns ordinary contention into avoidable cache data loss. The current multiprocess coverage exercises concurrent writes and prune races, but not max_size_bytes under concurrency (cuda_core/tests/test_program_cache_multiprocess.py).

Reduced simulation I used locally:

import time
from pathlib import Path

from cuda.core._module import ObjectCode
from cuda.core.utils import FileStreamProgramCache

cache = FileStreamProgramCache("/tmp/cuda_cache_review_race", max_size_bytes=1000)
cache[b"old"] = ObjectCode._init(b"a" * 600, "cubin", name="old")
time.sleep(0.01)

old_path = cache._path_for_key(b"old")
orig_unlink = Path.unlink
state = {"done": False}


def flaky_unlink(self, *args, **kwargs):
    if self == old_path and not state["done"]:
        state["done"] = True
        # Simulate another process deleting the file after stat() but before
        # _enforce_size_cap() updates its bookkeeping.
        orig_unlink(self, *args, **kwargs)
        raise FileNotFoundError(self)
    return orig_unlink(self, *args, **kwargs)


Path.unlink = flaky_unlink
try:
    cache[b"new"] = ObjectCode._init(b"b" * 600, "cubin", name="new")
finally:
    Path.unlink = orig_unlink

remaining = [key for key in (b"old", b"new") if cache.get(key) is not None]
print(remaining)  # []

That produced [] for me: once the first deletion race is swallowed without decrementing total, the loop keeps evicting and drops the fresh entry too.

Low: from cuda.core.utils import * now eagerly imports the cache stack

The package conversion keeps explicit imports like from cuda.core.utils import StridedMemoryView lightweight, but from cuda.core.utils import * walks __all__, resolves the lazy cache symbols, and imports _program_cache (cuda_core/cuda/core/utils/__init__.py:10, cuda_core/cuda/core/utils/__init__.py:32).

I verified that star-import now loads cuda.core.utils._program_cache. That said, this only affects import *, which is already discouraged. I think a short comment explaining that the laziness guarantee is intended for explicit imports, not star-import, seems sufficient here.

[rwgk]

Small readability/maintenance cleanup suggestion:

__all__ and _LAZY_CACHE_ATTRS currently duplicate the same cache-export names, so defining the ordered lazy-export list once and reusing it in __all__ seems a bit easier to scan and reduces the chance that the two drift apart later.

Something along these lines:

_LAZY_CACHE_ATTRS = (
    "FileStreamProgramCache",
    "ProgramCacheResource",
    "SQLiteProgramCache",
    "make_program_cache_key",
)

__all__ = [
    "StridedMemoryView",
    "args_viewable_as_strided_memory",
    *_LAZY_CACHE_ATTRS,
]

Mostly just a readability nit, but I think this makes the relationship between "lazy exports" and "public exports" a little clearer.