Skip to content

Commit 378dc78

Browse files
committed
fixup! add ProgramCacheResource.update + transparency tests + lead-with-high-level docs
* Add ``ProgramCacheResource.update(items)`` -- default ABC method that delegates to ``__setitem__`` so backend coercion (bytes extraction, size-cap enforcement, atomic-write retry) runs per entry. Accepts a Mapping or any iterable of (key, value) pairs. Symmetric with the existing ``clear`` convenience. * Test ``cache.update`` with both Mapping and pairs forms. * New transparency test that puts the SAME bytes into the cache via every supported input form (raw bytes, bytearray, memoryview, bytes-backed ObjectCode, path-backed ObjectCode pointing at a file with those bytes) and asserts every read returns the identical payload AND the on-disk file is the raw bytes. * Reframe the ProgramCacheResource ABC docstring example so it leads with ``program.compile(\"cubin\", cache=cache)`` and treats the manual key + bytes round-trip as the escape hatch for ``extra_digest`` cases. * Same reframe in ``make_program_cache_key``'s Examples block.
1 parent f1fbab7 commit 378dc78

2 files changed

Lines changed: 121 additions & 17 deletions

File tree

cuda_core/cuda/core/utils/_program_cache.py

Lines changed: 74 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -95,30 +95,44 @@ class ProgramCacheResource(abc.ABC):
9595
``bytes``.
9696
9797
The values written are the compiled program bytes themselves --
98-
cubin, PTX, LTO-IR, etc. Callers that compile via
99-
:class:`~cuda.core.Program` typically pass the resulting
100-
:class:`~cuda.core.ObjectCode` directly; the cache extracts
101-
``bytes(object_code.code)`` for storage. Reads return raw bytes so
102-
cache files remain consumable by external NVIDIA tools
103-
(``cuobjdump``, ``nvdisasm``, ``cuda-gdb``, ...). Callers
104-
reconstruct an :class:`~cuda.core.ObjectCode` themselves when they
105-
need one::
98+
cubin, PTX, LTO-IR, etc. Reads return raw bytes so cache files
99+
remain consumable by external NVIDIA tools (``cuobjdump``,
100+
``nvdisasm``, ``cuda-gdb``, ...).
106101
102+
Most callers don't interact with this object directly. The
103+
recommended usage is :meth:`cuda.core.Program.compile`'s ``cache=``
104+
keyword, which derives the key, returns a fresh
105+
:class:`~cuda.core.ObjectCode` on hit, and stores the compile
106+
result on miss::
107+
108+
with FileStreamProgramCache() as cache:
109+
obj = program.compile("cubin", cache=cache)
110+
111+
The escape hatch -- only needed when the compile inputs require an
112+
``extra_digest`` (header / PCH content fingerprints, NVVM
113+
libdevice) -- is to call :func:`make_program_cache_key` yourself
114+
and use the cache as a plain ``bytes`` mapping::
115+
116+
from cuda.core._module import ObjectCode
117+
118+
key = make_program_cache_key(
119+
code=source, code_type="c++", options=options,
120+
target_type="cubin", extra_digest=header_fingerprint(),
121+
)
107122
data = cache.get(key)
108123
if data is None:
109124
obj = program.compile("cubin")
110125
cache[key] = obj # extracts bytes(obj.code)
111-
data = bytes(obj.code)
112126
else:
113127
obj = ObjectCode._init(data, "cubin")
114128
115129
The cache layer does no payload validation; bytes go in and come
116130
back out unchanged. Symbol-mapping metadata that
117131
:class:`~cuda.core.ObjectCode` carries when produced with NVRTC
118132
name expressions is **not** preserved across a cache round-trip --
119-
the binary alone is stored. Callers that need symbol_mapping for
120-
``get_kernel(name_expression)`` should compile fresh, or look the
121-
mangled symbol up by hand.
133+
the binary alone is stored. Callers that need ``symbol_mapping``
134+
for ``get_kernel(name_expression)`` should compile fresh, or look
135+
the mangled symbol up by hand.
122136
"""
123137

124138
@abc.abstractmethod
@@ -170,6 +184,29 @@ def get(self, key: bytes | str, default: bytes | None = None) -> bytes | None:
170184
except KeyError:
171185
return default
172186

187+
def update(
188+
self,
189+
items: (
190+
collections.abc.Mapping[bytes | str, bytes | bytearray | memoryview | ObjectCode]
191+
| collections.abc.Iterable[
192+
tuple[bytes | str, bytes | bytearray | memoryview | ObjectCode]
193+
]
194+
),
195+
/,
196+
) -> None:
197+
"""Bulk ``__setitem__``.
198+
199+
Accepts a mapping or an iterable of ``(key, value)`` pairs. Each
200+
write goes through ``__setitem__`` so backend-specific value
201+
coercion (e.g. extracting bytes from an :class:`~cuda.core.ObjectCode`)
202+
and size-cap enforcement run on every entry. Not transactional --
203+
a failure mid-iteration leaves earlier writes committed.
204+
"""
205+
if isinstance(items, collections.abc.Mapping):
206+
items = items.items()
207+
for key, value in items:
208+
self[key] = value
209+
173210
def close(self) -> None: # noqa: B027
174211
"""Release backend resources. No-op by default."""
175212

@@ -516,18 +553,38 @@ def make_program_cache_key(
516553
517554
Examples
518555
--------
519-
Wiring a cache around :class:`~cuda.core.Program` compile::
556+
For most workflows you should not call ``make_program_cache_key``
557+
yourself -- pass ``cache=`` to :meth:`cuda.core.Program.compile`,
558+
which derives the key, returns the cached
559+
:class:`~cuda.core.ObjectCode` on hit, and stores the compile
560+
result on miss::
520561
521562
from cuda.core import Program, ProgramOptions
522-
from cuda.core._module import ObjectCode
523-
from cuda.core.utils import FileStreamProgramCache, make_program_cache_key
563+
from cuda.core.utils import FileStreamProgramCache
524564
525565
source = 'extern "C" __global__ void k(int *a){ *a = 1; }'
526566
options = ProgramOptions(arch="sm_80")
527567
528-
with FileStreamProgramCache("/var/cache/myapp/cuda") as cache:
568+
with FileStreamProgramCache() as cache:
569+
obj = Program(source, "c++", options=options).compile(
570+
"cubin", cache=cache
571+
)
572+
573+
Call ``make_program_cache_key`` directly when the compile inputs
574+
require an ``extra_digest`` (the cache cannot read external file
575+
content on the caller's behalf) -- ``Program.compile(cache=...)``
576+
refuses those inputs with a ``ValueError`` pointing here::
577+
578+
from cuda.core._module import ObjectCode
579+
from cuda.core.utils import FileStreamProgramCache, make_program_cache_key
580+
581+
with FileStreamProgramCache() as cache:
529582
key = make_program_cache_key(
530-
code=source, code_type="c++", options=options, target_type="cubin"
583+
code=source,
584+
code_type="c++",
585+
options=options,
586+
target_type="cubin",
587+
extra_digest=fingerprint_headers(options.include_path),
531588
)
532589
data = cache.get(key)
533590
if data is None:

cuda_core/tests/test_program_cache.py

Lines changed: 47 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1276,6 +1276,53 @@ def test_filestream_cache_accepts_path_backed_object_code(tmp_path):
12761276
assert cache[b"k"] == b"hello-cubin-bytes"
12771277

12781278

1279+
def test_program_cache_resource_update_accepts_mapping_and_pairs(tmp_path):
1280+
"""``update`` is a default ABC method; it must accept either a Mapping
1281+
or an iterable of (key, value) pairs and dispatch each item through
1282+
``__setitem__`` so backend coercion (bytes extraction, size-cap
1283+
enforcement) still runs."""
1284+
from cuda.core.utils import FileStreamProgramCache
1285+
1286+
with FileStreamProgramCache(tmp_path / "fc-mapping") as cache:
1287+
cache.update({b"a": b"v-a", b"b": b"v-b"})
1288+
assert cache[b"a"] == b"v-a"
1289+
assert cache[b"b"] == b"v-b"
1290+
1291+
with FileStreamProgramCache(tmp_path / "fc-pairs") as cache:
1292+
cache.update([(b"x", b"v-x"), (b"y", b"v-y")])
1293+
assert cache[b"x"] == b"v-x"
1294+
assert cache[b"y"] == b"v-y"
1295+
1296+
1297+
def test_filestream_cache_input_forms_are_byte_equivalent(tmp_path):
1298+
"""Whether the caller writes raw bytes, a bytearray, a memoryview, a
1299+
bytes-backed ObjectCode, or a path-backed ObjectCode pointing at a file
1300+
with the same bytes, the cache content is byte-identical and the on-disk
1301+
file has those exact bytes. Demonstrates the transparency contract:
1302+
callers don't have to normalise their input shape themselves."""
1303+
from cuda.core._module import ObjectCode
1304+
from cuda.core.utils import FileStreamProgramCache
1305+
1306+
payload = b"\x7fELF\x02\x01\x01\x00fake-cubin-bytes"
1307+
src = tmp_path / "src.cubin"
1308+
src.write_bytes(payload)
1309+
1310+
inputs = {
1311+
b"raw-bytes": payload,
1312+
b"bytearray": bytearray(payload),
1313+
b"memoryview": memoryview(payload),
1314+
b"obj-bytes-backed": ObjectCode._init(payload, "cubin", name="x"),
1315+
b"obj-path-backed": ObjectCode.from_cubin(str(src), name="y"),
1316+
}
1317+
1318+
with FileStreamProgramCache(tmp_path / "fc") as cache:
1319+
cache.update(inputs)
1320+
for k in inputs:
1321+
assert cache[k] == payload, f"value for {k!r} round-tripped to a different byte string"
1322+
on_disk = cache._path_for_key(k).read_bytes()
1323+
assert on_disk == payload, f"on-disk file for {k!r} is not the raw payload"
1324+
1325+
12791326
def test_filestream_cache_rejects_negative_size_cap(tmp_path):
12801327
from cuda.core.utils import FileStreamProgramCache
12811328

0 commit comments

Comments
 (0)