Coverage for cuda/core/utils/_program_cache/_keys.py: 90.20%
245 statements
« prev ^ index » next coverage.py v7.15.2, created at 2026-07-19 01:12 +0000
1# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
2#
3# SPDX-License-Identifier: Apache-2.0
5"""Cache-key construction.
7A backend-strategy hierarchy (:class:`_KeyBackend`) owns the per-code-type
8guard / fingerprint / version-probe logic; :func:`make_program_cache_key`
9dispatches to the right backend and assembles the digest.
10"""
12from __future__ import annotations
14import abc
15import collections.abc
16import hashlib
17from typing import Any, Callable, Sequence
19# Mutual-dependency contract: this module imports ProgramOptions from
20# cuda.core._program at module level, and cuda.core._program imports
21# ProgramCacheResource / make_program_cache_key from cuda.core.utils
22# only via deferred imports inside ``Program.compile``. Adding a
23# top-level ``from cuda.core.utils import ...`` to _program.pyx would
24# turn this into a real import cycle -- keep the import in _program.pyx
25# deferred (or import the symbols from the leaf submodule directly).
26from cuda.core._program import ProgramOptions
27from cuda.core._utils.cuda_utils import (
28 driver as _driver,
29)
30from cuda.core._utils.cuda_utils import (
31 handle_return as _handle_return,
32)
33from cuda.core._utils.cuda_utils import (
34 nvrtc as _nvrtc,
35)
36from cuda.core._utils.validators import check_str_enum, format_or_list
37from cuda.core.typing import SourceCodeType
39# Bump when the key schema changes in a way that invalidates existing caches.
40_KEY_SCHEMA_VERSION = 2
42_VALID_TARGET_TYPES = frozenset({"ptx", "cubin", "ltoir"})
44# code_type -> allowed target_type set, mirroring Program.compile's
45# SUPPORTED_TARGETS matrix in _program.pyx.
46_SUPPORTED_TARGETS_BY_CODE_TYPE = {
47 "c++": frozenset({"ptx", "cubin", "ltoir"}),
48 "ptx": frozenset({"cubin", "ptx"}),
49 "nvvm": frozenset({"ptx", "ltoir"}),
50}
53# Map each ProgramOptions field that reaches the Linker via
54# _translate_program_options (see cuda_core/cuda/core/_program.pyx) to the
55# gate the Linker uses to turn it into a flag (see
56# ``_prepare_nvjitlink_options`` and ``_prepare_driver_options`` in
57# _linker.pyx). All other fields on ProgramOptions are NVRTC-only and must
58# NOT perturb a PTX cache key: a PTX compile with a shared ProgramOptions
59# that happens to set include_path/pch/frandom_seed would otherwise miss the
60# cache unnecessarily. Collapsing inputs through these gates means
61# semantically-equivalent configurations (``debug=False`` vs ``None``,
62# ``time=True`` vs ``time="path"``) hash to the same cache key instead of
63# forcing spurious misses. Single source of truth: every reader iterates
64# this dict, so adding a field here is enough -- there is no parallel
65# field-name list to keep in sync.
66def _gate_presence(v: Any) -> bool:
67 return v is not None 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
70def _gate_truthy(v: Any) -> bool:
71 return bool(v) 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
74def _gate_is_true(v: Any) -> bool:
75 return v is True 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
78def _gate_tristate_bool(v: Any) -> bool | None:
79 return None if v is None else bool(v) 1mnopqrsDbtuvwxyzAEFGHIJKghcdefBC
82def _gate_identity(v: Any) -> Any:
83 return v 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
86def _gate_ptxas_options(v: Any) -> Any:
87 # ``_prepare_nvjitlink_options`` emits one ``-Xptxas=<s>`` per element, and
88 # treats ``str`` as a single-element sequence. Canonicalize to a tuple so
89 # ``"-v"`` / ``["-v"]`` / ``("-v",)`` all hash the same. An empty sequence
90 # emits no flags, so collapse it to ``None`` too.
91 #
92 # Order is preserved on purpose: ptxas accepts ordering-sensitive flag
93 # pairs (e.g. ``-O2`` after ``-O3`` lowers the active level), so
94 # ``["-v", "-O2"]`` and ``["-O2", "-v"]`` are not guaranteed to produce
95 # identical bytes. We accept the spurious miss when callers reorder
96 # flags; treating order as semantic keeps the cache safe in the
97 # ordering-sensitive case.
98 if v is None: 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
99 return None 1mnopqrsaDbtuvwxyijklzAEFGHIJKcdBC
100 if isinstance(v, str): 1ghcdef
101 return ("-Xptxas=" + v,) 1ef
102 if isinstance(v, collections.abc.Sequence): 1ghcdef
103 if len(v) == 0: 1ghcdef
104 return None 1gcd
105 return tuple(f"-Xptxas={s}" for s in v) 1hef
106 return v
109_LINKER_FIELD_GATES = {
110 "name": _gate_identity,
111 "arch": _gate_identity,
112 "max_register_count": _gate_identity,
113 "time": _gate_presence, # linker emits ``-time`` iff value is not None
114 "link_time_optimization": _gate_truthy,
115 "debug": _gate_truthy,
116 "lineinfo": _gate_truthy,
117 "ftz": _gate_tristate_bool,
118 "prec_div": _gate_tristate_bool,
119 "prec_sqrt": _gate_tristate_bool,
120 "fma": _gate_tristate_bool,
121 "split_compile": _gate_identity,
122 "ptxas_options": _gate_ptxas_options,
123 "no_cache": _gate_is_true,
124}
127# LinkerOptions fields the ``cuLink`` driver backend silently ignores
128# (emits only a DeprecationWarning; no actual flag reaches the compiler).
129# When the driver backend is active, collapse them to a single sentinel in
130# the fingerprint so nvJitLink<->driver parity of ``ObjectCode`` doesn't
131# cause cache misses from otherwise-equivalent configurations.
132_DRIVER_IGNORED_LINKER_FIELDS = frozenset({"ftz", "prec_div", "prec_sqrt", "fma"})
135def _linker_option_fingerprint(options: ProgramOptions, *, use_driver_linker: bool | None) -> list[bytes]:
136 """Backend-aware fingerprint of ProgramOptions fields consumed by the Linker.
138 Each field passes through the gate the Linker itself uses so equivalent
139 inputs (e.g. ``debug=False`` / ``None``) hash to the same bytes. When
140 the driver (cuLink) linker backend is in use, fields it silently
141 ignores collapse to one sentinel so those options don't perturb the
142 key on driver-backed hosts either. ``use_driver_linker=None`` means we
143 couldn't probe the backend; we don't collapse driver-ignored fields in
144 that case, to stay conservative.
145 """
146 parts = [] 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
147 driver_ignored = use_driver_linker is True 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
148 for name, gate in _LINKER_FIELD_GATES.items(): 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
149 if driver_ignored and name in _DRIVER_IGNORED_LINKER_FIELDS: 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
150 parts.append(f"{name}=<driver-ignored>".encode()) 1aijkl
151 continue 1aijkl
152 gated = gate(getattr(options, name, None)) 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
153 parts.append(f"{name}={gated!r}".encode()) 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
154 return parts 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
157# ProgramOptions fields that map to LinkerOptions fields the cuLink (driver)
158# backend rejects outright (see _prepare_driver_options in _linker.pyx).
159# ``split_compile_extended`` exists on LinkerOptions but is not exposed via
160# ProgramOptions / _translate_program_options, so it cannot reach the driver
161# linker from the cache path and is omitted here.
162_DRIVER_LINKER_UNSUPPORTED_FIELDS = ("time", "ptxas_options", "split_compile")
165def _driver_version() -> int:
166 return int(_handle_return(_driver.cuDriverGetVersion())) 1aijkl
169def _nvrtc_version() -> tuple[int, int]:
170 major, minor = _handle_return(_nvrtc.nvrtcVersion()) 1WMNOPQR@[XYZ012V]3456T7US89!#$%'()*+,-./:;=?
171 return int(major), int(minor) 1WMNOPQR@[XYZ012V]3456T7US89!#$%'()*+,-./:;=?
174def _linker_backend_and_version(use_driver: bool) -> tuple[str, str]:
175 """Return ``(backend, version)`` for the linker used on PTX inputs.
177 ``use_driver`` is the result of ``_decide_nvjitlink_or_driver()`` and
178 must be passed in so a single ``make_program_cache_key`` call shares
179 one probe across :meth:`_LinkerBackend.validate`,
180 :meth:`option_fingerprint`, and :meth:`hash_version_probe` (otherwise
181 a transient probe flap could write inconsistent fields into the same
182 key).
184 Raises any underlying probe exception. ``make_program_cache_key`` catches
185 and mixes the exception's class name into the digest, so the same probe
186 failure produces the same key across processes -- the cache stays
187 persistent in broken environments, while never sharing a key with a
188 working probe (``_probe_failed`` label vs. ``driver``/``nvrtc``/...).
190 nvJitLink version lookup goes through ``sys.modules`` first so we hit the
191 same module ``_decide_nvjitlink_or_driver()`` already loaded. That keeps
192 fingerprinting aligned with whichever ``cuda.bindings.nvjitlink`` import
193 path the linker actually uses.
194 """
195 import sys 1mnopqrsabtuvwxyijklzAghcdefBC
197 if use_driver: 1mnopqrsabtuvwxyijklzAghcdefBC
198 return ("driver", str(_driver_version())) 1aijkl
199 nvjitlink = sys.modules.get("cuda.bindings.nvjitlink") 1mnopqrsbtuvwxyzAghcdefBC
200 if nvjitlink is None: 1mnopqrsbtuvwxyzAghcdefBC
201 from cuda.bindings import nvjitlink as _nvjitlink
203 nvjitlink = _nvjitlink
205 return ("nvJitLink", str(nvjitlink.version())) 1mnopqrsbtuvwxyzAghcdefBC
208def _nvvm_fingerprint() -> str:
209 """Stable identifier for the loaded NVVM toolchain.
211 Combines the libNVVM library version (``module.version()``) with the IR
212 version reported by ``module.ir_version()``. The library version is the
213 primary invalidation lever: a libNVVM patch upgrade can change codegen
214 while keeping the same IR major/minor, so keying only on the IR pair
215 would silently reuse stale entries. Paired with cuda-core, the IR pair
216 adds defence in depth without making the key any less stable.
218 Both calls go through ``_get_nvvm_module()`` so this fingerprint follows
219 the same availability / cuda-bindings-version gate that real NVVM
220 compilation does -- if NVVM is unusable at compile time, the probe
221 fails the same way and ``_probe`` mixes the failure label into the key.
222 """
223 from cuda.core._program import _get_nvvm_module 2_ { | } ~ abbb` ^ cb
225 module = _get_nvvm_module() 2_ { | } ~ abbb` ^ cb
226 lib_major, lib_minor = module.version() # type: ignore[attr-defined] 2_ { | } ~ abbb` ^ cb
227 major, minor, debug_major, debug_minor = module.ir_version() # type: ignore[attr-defined] 2_ { | } ~ abbb` ^ cb
228 return f"lib={lib_major}.{lib_minor};ir={major}.{minor}.{debug_major}.{debug_minor}" 2_ { | } ~ abbb` ^ cb
231# ProgramOptions fields that reference external files whose *contents* the
232# cache key cannot observe without reading the filesystem. Callers that set
233# any of these must supply an ``extra_digest`` covering the dependency surface
234# (e.g. a hash over all reachable headers / PCH bytes).
235_EXTERNAL_CONTENT_OPTIONS = (
236 "include_path",
237 "pre_include",
238 "pch",
239 "use_pch",
240 "pch_dir",
241)
243# ProgramOptions fields whose compilation effect is not captured in the
244# returned ``ObjectCode`` -- they produce a filesystem artifact as a side
245# effect. A cache hit skips compilation, so that artifact would never be
246# written. Reject these outright: the persistent cache is for pure ObjectCode
247# reuse, not for replaying compile-time side effects.
248# * create_pch -- writes a PCH file (NVRTC).
249# * time -- writes NVRTC timing info to a file.
250# * fdevice_time_trace -- writes a device-compilation time trace file (NVRTC).
251# These are all NVRTC-specific; the Linker's ``-time`` logs to the info log
252# (not a file) and NVVM explicitly rejects all three at compile time. The
253# side-effect guard is therefore gated on ``backend == "nvrtc"`` below.
254_SIDE_EFFECT_OPTIONS = ("create_pch", "time", "fdevice_time_trace")
257# ProgramOptions fields gated by plain truthiness in ``_program.pyx`` (the
258# compiler writes the flag only when the value is truthy).
259_BOOLEAN_OPTION_FIELDS = frozenset({"pch"})
261# Fields whose compiler emission requires ``isinstance(value, str)`` or a
262# non-empty sequence; anything else (``False``, ``int``, ``None``, ``[]``)
263# is silently ignored at compile time.
264_STR_OR_SEQUENCE_OPTION_FIELDS = frozenset({"include_path", "pre_include"})
267def _option_is_set(options: ProgramOptions, name: str) -> bool:
268 """Match how ``_program.pyx`` gates option emission, per field shape.
270 - Boolean flags (``pch``): truthy only.
271 - str-or-sequence fields (``include_path``, ``pre_include``): ``str``
272 (including empty) or a non-empty ``collections.abc.Sequence`` (list,
273 tuple, range, user subclass, ...); everything else (``False``, ``int``,
274 empty sequence, ``None``) is ignored by the compiler and must not
275 trigger a cache-time guard.
276 - Path/string-shaped fields (``create_pch``, ``time``,
277 ``fdevice_time_trace``, ``use_pch``, ``pch_dir``): ``is not None`` --
278 the compiler emits ``--flag=<value>`` for any non-None value, so
279 ``False`` / ``""`` / ``0`` must still count as set.
280 """
281 value = getattr(options, name, None) 2W M N O P Q R @ [ X Y Z 0 1 2 V ] 3 4 5 6 T 7 U S 8 9 ! # $ % jbkbeblbobmbxbrbsbtbpbqbubvbVbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' ( ) * + , - . / : nbCbUb; = ?
282 if value is None: 2W M N O P Q R @ [ X Y Z 0 1 2 V ] 3 4 5 6 T 7 U S 8 9 ! # $ % jbkbeblbobmbxbrbsbtbpbqbubvbVbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' ( ) * + , - . / : nbCbUb; = ?
283 return False 2W M N O P Q R @ [ X Y Z 0 1 2 V ] 3 4 5 6 T 7 U S 8 9 ! # $ % jbkbeblbobmbxbrbsbtbpbqbubvbVbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' ( ) * + , - . / : nbCbUb; = ?
284 if name in _BOOLEAN_OPTION_FIELDS: 2M N O P Q R obmbxbrbsbtbpbqbubvbMbNbObPbQbRbSbTbnbUb
285 return bool(value) 2xb
286 if name in _STR_OR_SEQUENCE_OPTION_FIELDS: 2M N O P Q R obmbrbsbtbpbqbubvbMbNbObPbQbRbSbTbnbUb
287 # Mirror ``_prepare_nvrtc_options_impl``: it checks ``isinstance(v, str)``
288 # first, then ``is_sequence(v)`` (which is ``isinstance(v, Sequence)``).
289 # We therefore accept any ``collections.abc.Sequence`` (range, deque,
290 # user subclass, etc.), not just list/tuple.
291 if isinstance(value, str): 2M N O P Q R obmbpbqbnb
292 return True 2obpbqb
293 if isinstance(value, collections.abc.Sequence): 2M N O P Q R mbnb
294 return len(value) > 0 2M N O Q mbnb
295 return False 1PR
296 return True 2rbsbtbubvbMbNbObPbQbRbSbTbUb
299def _hash_probe_failure(update: Callable[[str, bytes], None], label: str, exc: BaseException) -> None:
300 """Mix a probe failure into the digest under a stable, content-free label.
302 Hashing only the exception's CLASS NAME (not its message) keeps the
303 digest stable across repeated calls within one process (e.g. NVVM's
304 loader reports different messages on first vs. cached-failure attempts)
305 AND across processes that hit the same failure mode. The
306 ``_probe_failed`` label differs from every backend's success label, so a
307 broken environment never collides with a working one -- the cache
308 "fails closed" between broken and working environments while staying
309 persistent within either.
310 """
311 update(f"{label}_probe_failed", type(exc).__name__.encode()) 1abT
314class _KeyBackend(abc.ABC):
315 """Strategy for deriving the cache key for one ``Program`` ``code_type``.
317 Each subclass owns the backend-specific guard logic, code coercion,
318 option fingerprinting, name-expression handling, version probing, and
319 extra-payload hashing. The orchestrator :func:`make_program_cache_key`
320 validates the code_type / target_type pair, dispatches to the right
321 backend, and assembles the digest.
322 """
324 @abc.abstractmethod
325 def validate(self, options: ProgramOptions, target_type: str, extra_digest: bytes | None) -> None:
326 """Reject inputs the cache cannot key safely.
328 Raises ``ValueError`` for options that have compile-time side
329 effects, options that pull in external file content the cache
330 can't observe, or any other backend-specific invariants.
331 """
333 def encode_code(self, code: object, code_type: str) -> bytes:
334 """Coerce ``code`` to bytes. Default rejects bytes-like input
335 (only NVVM accepts it; ``Program()`` does the same)."""
336 if isinstance(code, str): 2W M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 a D b T t u 7 v U S 8 w x 9 ! # $ % jby i j k l z A E F G H I J K g h c d e f kbeblbwbfbgbhbib' B ( ) C * + , - . / : ; = ?
337 return code.encode("utf-8") 2W M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 a D b T t u 7 v U S 8 w x 9 ! # $ % jby i j k l z A E F G H I J K g h c d e f ebfbgbhbib' B ( ) C * + , - . / : ; = ?
338 if isinstance(code, (bytes, bytearray)): 2kblbwb
339 raise TypeError( 2lbwb
340 f"code must be str for code_type={code_type!r}; bytes/bytearray are only accepted for code_type='nvvm'."
341 )
342 raise TypeError(f"code must be str or bytes, got {type(code).__name__}") 2kb
344 @abc.abstractmethod
345 def option_fingerprint(self, options: ProgramOptions, target_type: str) -> list[bytes]:
346 """Fingerprint of the ``ProgramOptions`` fields that reach the compiler."""
348 def encode_name_expressions(self, name_expressions: Sequence[Any]) -> tuple[bytes, ...] | None: # noqa: ARG002
349 """Sorted, type-tagged name expressions, or ``None`` if the
350 backend does not consume them.
352 ``None`` means the orchestrator emits no ``names_count`` /
353 ``name`` entries at all (a backend that ignores
354 ``name_expressions`` should never have them perturb its key). An
355 empty tuple means the backend supports them but the caller
356 passed none -- the orchestrator still emits ``names_count=0`` so
357 the schema is stable across "absent" and "empty".
358 """
359 return None 2_ m n o p q r s { a D b | t } u v ~ w x abbbdb` ^ y i j k l z A E F G H I J K g h c d e f B cbC
361 @abc.abstractmethod
362 def hash_version_probe(self, update: Callable[[str, bytes], None]) -> None:
363 """Mix the runtime/compiler version probe into the digest via
364 ``update(label, payload)``. On probe failure, mix
365 ``_hash_probe_failure(update, "<label>", exc)`` instead so the
366 digest is stable across processes hitting the same failure
367 mode.
368 """
370 def hash_extra_payload(self, options: ProgramOptions, update: Callable[[str, bytes], None]) -> None: # noqa: B027 1LWMNOPQRmnopq@[XYZrs012V]3456aDbTtu7vUS8wx9!#$%yijklzAEFGHIJKghcdef'B()C*+,-./:;=?
371 """Mix backend-specific extras (e.g. NVVM ``extra_sources`` /
372 ``use_libdevice``). Default: nothing.
373 """
376class _NvrtcBackend(_KeyBackend):
377 def validate(self, options: ProgramOptions, target_type: str, extra_digest: bytes | None) -> None: # noqa: ARG002
378 # Side-effect options are NVRTC-specific:
379 # ``time``/``fdevice_time_trace`` write artifacts via NVRTC,
380 # ``create_pch`` writes via NVRTC. The Linker's ``-time`` logs to
381 # the info log (not a file), and NVVM explicitly rejects all three
382 # at compile time, so the side-effect guard is meaningful only for
383 # the NVRTC path.
384 side_effects = [name for name in _SIDE_EFFECT_OPTIONS if _option_is_set(options, name)] 2W M N O P Q R @ [ X Y Z 0 1 2 V ] 3 4 5 6 T 7 U S 8 9 ! # $ % jbkbeblbobmbxbrbsbtbpbqbubvbVbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' ( ) * + , - . / : nbCbUb; = ?
385 if side_effects: 2W M N O P Q R @ [ X Y Z 0 1 2 V ] 3 4 5 6 T 7 U S 8 9 ! # $ % jbkbeblbobmbxbrbsbtbpbqbubvbVbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' ( ) * + , - . / : nbCbUb; = ?
386 raise ValueError( 2MbNbObPbQbRbSbTbUb
387 f"make_program_cache_key() refuses to build a key for options that "
388 f"have compile-time side effects ({', '.join(side_effects)}); a "
389 f"cache hit skips compilation, so the side effect would not occur. "
390 f"Disable the option, or compile directly without the cache."
391 )
392 # ``extra_sources`` is NVVM-only -- ``Program`` raises for non-NVVM
393 # backends (_program.pyx). Reject here so callers get the same
394 # error from the cache-key path as from a real compile.
395 if getattr(options, "extra_sources", None) is not None: 2W M N O P Q R @ [ X Y Z 0 1 2 V ] 3 4 5 6 T 7 U S 8 9 ! # $ % jbkbeblbobmbxbrbsbtbpbqbubvbVbfbgbhbibybzbAbBb' ( ) * + , - . / : nbCb; = ?
396 raise ValueError( 2Vb
397 "extra_sources is only valid for code_type='nvvm'; Program() rejects it for code_type='c++'."
398 )
399 if extra_digest is None: 2W M N O P Q R @ [ X Y Z 0 1 2 V ] 3 4 5 6 T 7 U S 8 9 ! # $ % jbkbeblbobmbxbrbsbtbpbqbubvbfbgbhbibybzbAbBb' ( ) * + , - . / : nbCb; = ?
400 # ``Program.compile`` for PTX inputs runs
401 # ``_translate_program_options``, which drops these entirely;
402 # NVVM rejects them. Only NVRTC reads the external content.
403 external = [name for name in _EXTERNAL_CONTENT_OPTIONS if _option_is_set(options, name)] 2W M N O P Q R X Y Z 0 1 2 V 3 4 5 6 T 7 U S 8 9 ! # $ % jbkbeblbobmbxbrbsbtbpbqbubvbfbgbhbibybzbAbBb' ( ) * + , - . / : nbCb; = ?
404 if external: 2W M N O P Q R X Y Z 0 1 2 V 3 4 5 6 T 7 U S 8 9 ! # $ % jbkbeblbobmbxbrbsbtbpbqbubvbfbgbhbibybzbAbBb' ( ) * + , - . / : nbCb; = ?
405 raise ValueError( 2obmbxbrbsbtbpbqbubvbnb
406 f"make_program_cache_key() refuses to build a key for options that "
407 f"pull in external file content ({', '.join(external)}) without an "
408 f"extra_digest; compute a digest over the header/PCH bytes the "
409 f"compile will read and pass it as extra_digest=..."
410 )
411 # NVRTC uses ``options.name`` as the source filename and
412 # resolves quoted ``#include "x.h"`` directives relative to
413 # the directory component of that name. The directory's
414 # contents are external to anything else the key observes,
415 # so a name with a directory component requires the same
416 # ``extra_digest`` treatment as ``include_path`` etc.
417 options_name = getattr(options, "name", None) 2W M N O P Q R X Y Z 0 1 2 V 3 4 5 6 T 7 U S 8 9 ! # $ % jbkbeblbfbgbhbibybzbAbBb' ( ) * + , - . / : Cb; = ?
418 if isinstance(options_name, str) and ("/" in options_name or "\\" in options_name): 2W M N O P Q R X Y Z 0 1 2 V 3 4 5 6 T 7 U S 8 9 ! # $ % jbkbeblbfbgbhbibybzbAbBb' ( ) * + , - . / : Cb; = ?
419 raise ValueError( 2ybzbAbBbCb
420 f"make_program_cache_key() refuses to build a key for options.name="
421 f"{options_name!r} (NVRTC source-filename with a directory "
422 f"component) without an extra_digest; NVRTC resolves quoted "
423 f"#include directives relative to that directory, so a digest "
424 f"covering the headers it may pull in must be supplied."
425 )
427 def option_fingerprint(self, options: ProgramOptions, target_type: str) -> list[bytes]:
428 # ``ProgramOptions.as_bytes("nvrtc", ...)`` gives the real
429 # compile-time flag surface for NVRTC.
430 return options.as_bytes("nvrtc", target_type) 2W M N O P Q R @ [ X Y Z 0 1 2 V ] 3 4 5 6 T 7 U S 8 9 ! # $ % jbebfbgbhbib' ( ) * + , - . / : ; = ?
432 def encode_name_expressions(self, name_expressions: Sequence[Any]) -> tuple[bytes, ...]:
433 # ``"foo"`` and ``b"foo"`` get distinct tags because
434 # ``Program.compile`` records the original Python object as the
435 # ``ObjectCode.symbol_mapping`` key, so a cached ObjectCode whose
436 # mapping-key type differs from what the caller's later
437 # ``get_kernel`` passes would silently miss. Reject ``bytearray``
438 # because ``Program.compile`` also uses the raw element as a dict
439 # key -- bytearray is unhashable, so a cache miss would compile
440 # then crash in ``symbol_mapping[n] = ...``. Accepting it here
441 # would let the cache serve hits for inputs the uncached path
442 # can't handle.
443 def _tag(n: Any) -> bytes: 2W M N O P Q R @ [ X Y Z 0 1 2 V ] 3 4 5 6 T 7 U S 8 9 ! # $ % ebfbgbhbib' ( ) * + , - . / : ; = ?
444 if isinstance(n, bytes): 2U S ebfbgbhbib
445 return b"b:" + n 1S
446 if isinstance(n, str): 2U S ebfbgbhbib
447 return b"s:" + n.encode("utf-8") 2U S ebfbgbhbib
448 if isinstance(n, bytearray): 2ebfbgbhbib
449 raise TypeError( 2eb
450 "name_expressions elements must be str or bytes; "
451 "bytearray is not accepted because Program.compile uses "
452 "each element as a dict key and bytearray is unhashable."
453 )
454 raise TypeError(f"name_expressions elements must be str or bytes; got {type(n).__name__}") 2fbgbhbib
456 return tuple(sorted(_tag(n) for n in name_expressions)) 2W M N O P Q R @ [ X Y Z 0 1 2 V ] 3 4 5 6 T 7 U S 8 9 ! # $ % ebfbgbhbib' ( ) * + , - . / : ; = ?
458 def hash_version_probe(self, update: Callable[[str, bytes], None]) -> None:
459 try: 1WMNOPQR@[XYZ012V]3456T7US89!#$%'()*+,-./:;=?
460 major, minor = _nvrtc_version() 1WMNOPQR@[XYZ012V]3456T7US89!#$%'()*+,-./:;=?
461 except Exception as exc: 1T
462 _hash_probe_failure(update, "nvrtc", exc) 1T
463 return 1T
464 update("nvrtc", f"{major}.{minor}".encode("ascii")) 1WMNOPQR@[XYZ012V]3456T7US89!#$%'()*+,-./:;=?
467_DECISION_UNSET = object()
470class _LinkerBackend(_KeyBackend):
471 def __init__(self) -> None:
472 # Cache the linker-backend decision (and any probe failure) for
473 # the duration of one ``make_program_cache_key`` call so
474 # ``validate``, ``option_fingerprint``, and ``hash_version_probe``
475 # all see the same answer; a transient probe flap mid-call
476 # otherwise mints a key whose option fingerprint and version
477 # probe disagree on which linker is in use.
478 self._cached_decision = _DECISION_UNSET 2m n o p q r s a D b t u v w x y i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbwbWbB C
479 self._cached_decision_exc: BaseException | None = None 2m n o p q r s a D b t u v w x y i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbwbWbB C
481 def _decide_driver(self) -> bool | None:
482 """``True`` if the cuLink driver linker will be used, ``False`` if
483 nvJitLink, ``None`` if the probe failed (in which case
484 :meth:`hash_version_probe` mixes a ``_probe_failed`` taint into
485 the digest instead of a backend label).
486 """
487 if self._cached_decision is _DECISION_UNSET: 2m n o p q r s a D b t u v w x y i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbwbB C
488 try: 2m n o p q r s a D b t u v w x y i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbwbB C
489 from cuda.core._linker import _decide_nvjitlink_or_driver 2m n o p q r s a D b t u v w x y i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbwbB C
491 self._cached_decision = _decide_nvjitlink_or_driver() 2m n o p q r s a D b t u v w x y i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbwbB C
492 except Exception as exc:
493 self._cached_decision = None
494 self._cached_decision_exc = exc
495 return self._cached_decision # type: ignore[return-value] 2m n o p q r s a D b t u v w x y i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbwbB C
497 def validate(self, options: ProgramOptions, target_type: str, extra_digest: bytes | None) -> None: # noqa: ARG002
498 if getattr(options, "extra_sources", None) is not None: 2m n o p q r s a D b t u v w x y i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbwbWbB C
499 raise ValueError( 2Wb
500 "extra_sources is only valid for code_type='nvvm'; Program() rejects it for code_type='ptx'."
501 )
502 # PTX compiles go through the Linker. When the driver (cuLink)
503 # backend is selected (nvJitLink unavailable), ``Program.compile``
504 # rejects a subset of options that nvJitLink would accept; reject
505 # them here too so we never store a key for a compilation that
506 # can't succeed in this environment. If the probe fails we can't
507 # tell which backend will run, so skip -- the failed-probe taint
508 # in ``hash_version_probe`` already poisons the key.
509 if self._decide_driver() is True: 2m n o p q r s a D b t u v w x y i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbwbB C
510 # Mirror ``_prepare_driver_options``'s exact gate: ``is not
511 # None`` for these fields, so ``time=False`` or
512 # ``ptxas_options=[]`` is still a rejection. Do NOT use the
513 # truthiness-based ``_option_is_set`` helper here.
514 unsupported = [ 2a i j k l DbEbFbGbHbIbJbKbLb
515 name for name in _DRIVER_LINKER_UNSUPPORTED_FIELDS if getattr(options, name, None) is not None
516 ]
517 if unsupported: 2a i j k l DbEbFbGbHbIbJbKbLb
518 raise ValueError( 2DbEbFbGbHbIbJbKbLb
519 f"the cuLink driver linker does not support these options: "
520 f"{', '.join(unsupported)}; Program.compile() would reject this "
521 f"configuration before producing an ObjectCode."
522 )
524 def option_fingerprint(self, options: ProgramOptions, target_type: str) -> list[bytes]: # noqa: ARG002
525 # For PTX inputs the Linker reads only a subset of ProgramOptions
526 # (see ``_translate_program_options`` in _program.pyx); fingerprint
527 # just those fields so shared ProgramOptions carrying NVRTC-only
528 # flags (``include_path``, ``pch_*``, ``frandom_seed``, ...) don't
529 # force spurious cache misses on PTX.
530 return _linker_option_fingerprint(options, use_driver_linker=self._decide_driver()) 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
532 def hash_version_probe(self, update: Callable[[str, bytes], None]) -> None:
533 # Only cuLink (driver-backed linker) goes through the CUDA driver
534 # for codegen. nvJitLink is a separate library, so a driver
535 # upgrade under it does not change the compiled bytes -- skip the
536 # driver version there. ``_linker_backend_and_version`` already
537 # returns the driver version when the driver backend is active,
538 # so the bytes are still in the digest via ``linker_version``.
539 use_driver = self._decide_driver() 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
540 if use_driver is None: 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
541 assert self._cached_decision_exc is not None
542 _hash_probe_failure(update, "linker", self._cached_decision_exc)
543 return
544 try: 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
545 lb_name, lb_version = _linker_backend_and_version(use_driver) 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
546 except Exception as exc: 1ab
547 _hash_probe_failure(update, "linker", exc) 1ab
548 return 1ab
549 update("linker_backend", lb_name.encode("ascii")) 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
550 update("linker_version", lb_version.encode("ascii")) 1mnopqrsaDbtuvwxyijklzAEFGHIJKghcdefBC
553class _NvvmBackend(_KeyBackend):
554 def encode_code(self, code: object, code_type: str) -> bytes: # noqa: ARG002
555 # NVVM accepts both str and bytes (matching ``Program()``).
556 if isinstance(code, str): 2_ { | } ~ abbbdb` ^ cb
557 return code.encode("utf-8") 2_ { | } ~ abbbdb` ^ cb
558 if isinstance(code, (bytes, bytearray)): 1_
559 return bytes(code) 1_
560 raise TypeError(f"code must be str or bytes, got {type(code).__name__}")
562 def validate(self, options: ProgramOptions, target_type: str, extra_digest: bytes | None) -> None: # noqa: ARG002
563 # NVVM with ``use_libdevice=True`` reads external libdevice
564 # bitcode at compile time (see Program_init in _program.pyx). The
565 # file is resolved from the active toolkit, so a changed
566 # CUDA_HOME / libdevice upgrade changes the linked output without
567 # touching any key input the cache can observe. Require the
568 # caller to supply an ``extra_digest`` that fingerprints the
569 # libdevice bytes (or simply disable use_libdevice for
570 # caching-sensitive workflows).
571 if extra_digest is None and getattr(options, "use_libdevice", None): 2_ { | } ~ abbbdb` ^ cb
572 raise ValueError( 1^
573 "make_program_cache_key() refuses to build an NVVM key with "
574 "use_libdevice=True and no extra_digest: the linked libdevice "
575 "bitcode can change out from under a cached ObjectCode. Pass an "
576 "extra_digest that fingerprints the libdevice file you intend "
577 "to link against, or disable use_libdevice."
578 )
580 def option_fingerprint(self, options: ProgramOptions, target_type: str) -> list[bytes]:
581 return options.as_bytes("nvvm", target_type) 2_ { | } ~ abbbdb` ^ cb
583 def hash_version_probe(self, update: Callable[[str, bytes], None]) -> None:
584 try: 2_ { | } ~ abbbdb` ^ cb
585 fp = _nvvm_fingerprint() 2_ { | } ~ abbbdb` ^ cb
586 except Exception as exc:
587 _hash_probe_failure(update, "nvvm", exc)
588 return
589 update("nvvm", fp.encode("ascii")) 2_ { | } ~ abbbdb` ^ cb
591 def hash_extra_payload(self, options: ProgramOptions, update: Callable[[str, bytes], None]) -> None:
592 extra_sources = getattr(options, "extra_sources", None) 2_ { | } ~ abbbdb` ^ cb
593 if extra_sources: 2_ { | } ~ abbbdb` ^ cb
594 # ``extra_sources`` is hashed in caller-provided order on purpose.
595 # NVVM module linking is order-dependent in the general case
596 # (overlapping symbols, weak definitions, definition order can
597 # change which body wins), so canonicalising by sorting on the
598 # source name would produce the same key for two compiles whose
599 # outputs may legitimately differ. If a future test proves the
600 # relevant input subset is order-insensitive, sorting can be
601 # introduced under that proof; absent that proof, preserving
602 # caller order is the safe default.
603 update("extra_sources_count", str(len(extra_sources)).encode("ascii"))
604 for item in extra_sources:
605 # ``extra_sources`` is a sequence of (name, source) tuples.
606 if isinstance(item, (tuple, list)) and len(item) == 2:
607 name, src = item
608 update("extra_source_name", str(name).encode("utf-8"))
609 if isinstance(src, str):
610 update("extra_source_code", src.encode("utf-8"))
611 elif isinstance(src, (bytes, bytearray)):
612 update("extra_source_code", bytes(src))
613 else:
614 update("extra_source_code", str(src).encode("utf-8"))
615 else:
616 # Fallback for unexpected format.
617 update("extra_source", str(item).encode("utf-8"))
618 # ``use_libdevice`` is gated on truthiness to match Program_init's
619 # gate -- ``False`` and ``None`` collapse to the same key.
620 if getattr(options, "use_libdevice", None): 2_ { | } ~ abbbdb` ^ cb
621 update("use_libdevice", b"1") 1`^
624# Class registry keyed by code_type. ``make_program_cache_key`` instantiates
625# fresh per call so backends like ``_LinkerBackend`` can cache per-call probe
626# results on ``self`` without leaking that state across calls.
627_BACKENDS_BY_CODE_TYPE: dict[str, type[_KeyBackend]] = {
628 "c++": _NvrtcBackend,
629 "ptx": _LinkerBackend,
630 "nvvm": _NvvmBackend,
631}
634def make_program_cache_key(
635 *,
636 code: str | bytes,
637 code_type: str,
638 options: ProgramOptions,
639 target_type: str,
640 name_expressions: Sequence[str | bytes | bytearray] = (),
641 extra_digest: bytes | None = None,
642) -> bytes:
643 """Build a stable cache key from compile inputs.
645 Parameters
646 ----------
647 code:
648 Source text. ``str`` is encoded as UTF-8.
649 code_type:
650 One of ``"c++"``, ``"ptx"``, ``"nvvm"``.
651 options:
652 A :class:`cuda.core.ProgramOptions`. Its ``arch`` must be set (the
653 default ``ProgramOptions.__post_init__`` populates it from the current
654 device).
655 target_type:
656 One of ``"ptx"``, ``"cubin"``, ``"ltoir"``.
657 name_expressions:
658 Optional iterable of mangled-name lookups. Order is not significant.
659 Elements may be ``str`` or ``bytes``; ``"foo"`` and ``b"foo"`` produce
660 distinct keys because ``Program.compile`` records the original Python
661 object as the ``ObjectCode.symbol_mapping`` key, and ``get_kernel``
662 lookups must use the same type the cache key recorded. ``bytearray``
663 is rejected because ``Program.compile`` stores each element as a
664 dict key and ``bytearray`` is unhashable.
665 extra_digest:
666 Caller-supplied bytes mixed into the key. Required whenever
667 :class:`cuda.core.ProgramOptions` sets any option that pulls in
668 external file content (``include_path``, ``pre_include``, ``pch``,
669 ``use_pch``, ``pch_dir``) -- the cache cannot read those files on
670 the caller's behalf, so the caller must fingerprint the header /
671 PCH surface and pass it here. Callers may pass this for other
672 inputs too (embedded kernels, generated sources, etc.).
674 Returns
675 -------
676 bytes
677 An opaque bytes digest suitable for use as a cache key.
679 Raises
680 ------
681 ValueError
682 If ``options`` sets an option with compile-time side effects (such
683 as ``create_pch``) -- a cache hit skips compilation, so the side
684 effect would not occur.
685 ValueError
686 If ``extra_digest`` is ``None`` while ``options`` sets any option
687 whose compilation effect depends on external file content that the
688 key cannot otherwise observe.
690 Examples
691 --------
692 For most workflows you should not call ``make_program_cache_key``
693 yourself -- pass ``cache=`` to :meth:`cuda.core.Program.compile`,
694 which derives the key, returns the cached
695 :class:`~cuda.core.ObjectCode` on hit, and stores the compile
696 result on miss::
698 from cuda.core import Program, ProgramOptions
699 from cuda.core.utils import FileStreamProgramCache
701 source = 'extern "C" __global__ void k(int *a){ *a = 1; }'
702 options = ProgramOptions(arch="sm_80")
704 with FileStreamProgramCache() as cache:
705 obj = Program(source, "c++", options=options).compile("cubin", cache=cache)
707 Call ``make_program_cache_key`` directly when the compile inputs
708 require an ``extra_digest`` (the cache cannot read external file
709 content on the caller's behalf) -- ``Program.compile(cache=...)``
710 refuses those inputs with a ``ValueError`` pointing here::
712 from cuda.core import ObjectCode
713 from cuda.core.utils import FileStreamProgramCache, make_program_cache_key
715 with FileStreamProgramCache() as cache:
716 key = make_program_cache_key(
717 code=source,
718 code_type="c++",
719 options=options,
720 target_type="cubin",
721 extra_digest=fingerprint_headers(options.include_path),
722 )
723 data = cache.get(key)
724 if data is None:
725 obj = Program(source, "c++", options=options).compile("cubin")
726 cache[key] = obj # extracts bytes(obj.code)
727 else:
728 obj = ObjectCode.from_cubin(data)
730 The cache stores raw binary bytes -- cubin / PTX / LTO-IR with no
731 pickle, JSON, or framing -- so entry files are directly consumable
732 by external NVIDIA tools (``cuobjdump``, ``nvdisasm``, ...). Note
733 that an :class:`~cuda.core.ObjectCode` round-tripped through the
734 cache loses ``symbol_mapping``: callers that compile with
735 ``name_expressions`` and rely on ``get_kernel(name_expression)``
736 after a cache hit must either compile fresh or look up the mangled
737 symbol explicitly.
739 Options that read external files (``include_path``, ``pre_include``,
740 ``pch``, ``use_pch``, ``pch_dir``; ``use_libdevice=True`` on the NVVM
741 path; and on NVRTC, an ``options.name`` with a directory component,
742 which NVRTC uses for relative-include resolution) require
743 ``extra_digest`` -- fingerprint the bytes the compiler will pull in
744 and pass that digest so changes to those files force a cache miss.
745 Options that have compile-time side effects (``create_pch``,
746 ``time``, ``fdevice_time_trace``) cannot be cached and raise
747 ``ValueError``; compile directly, or disable the flag, for those
748 cases.
749 """
750 # Mirror Program.compile (_program.pyx lowercases code_type at Program
751 # init and target_type at the top of compile); a caller that passes
752 # "PTX" or "C++" must get the same routing and the same cache key as
753 # the lowercase form.
754 code_type = code_type.lower() if isinstance(code_type, str) else code_type 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ jby i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbkbXbYb0bZbeblbwbobmbxbrbsbtbpbqbubvbVbWbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' B ( cb) C * + , - . / : nbCbUb; = ?
755 target_type = target_type.lower() if isinstance(target_type, str) else target_type 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ jby i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbkbXbYb0bZbeblbwbobmbxbrbsbtbpbqbubvbVbWbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' B ( cb) C * + , - . / : nbCbUb; = ?
756 check_str_enum(code_type, SourceCodeType) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ jby i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbkbXbYb0bZbeblbwbobmbxbrbsbtbpbqbubvbVbWbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' B ( cb) C * + , - . / : nbCbUb; = ?
757 if target_type not in _VALID_TARGET_TYPES: 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ jby i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbkbXbYbZbeblbwbobmbxbrbsbtbpbqbubvbVbWbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' B ( cb) C * + , - . / : nbCbUb; = ?
758 raise ValueError( 2Zb
759 f"target_type={target_type!r} is not supported by the program cache "
760 f"(must be {format_or_list(sorted(_VALID_TARGET_TYPES))})"
761 )
762 supported_for_code = _SUPPORTED_TARGETS_BY_CODE_TYPE[code_type] 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ jby i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbkbXbYbeblbwbobmbxbrbsbtbpbqbubvbVbWbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' B ( cb) C * + , - . / : nbCbUb; = ?
763 if target_type not in supported_for_code: 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ jby i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbkbXbYbeblbwbobmbxbrbsbtbpbqbubvbVbWbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' B ( cb) C * + , - . / : nbCbUb; = ?
764 raise ValueError( 2XbYb
765 f"target_type={target_type!r} is not valid for code_type={code_type!r}"
766 f" (supported: {sorted(supported_for_code)}). Program.compile() rejects"
767 f" this combination, so caching a key for it is meaningless."
768 )
770 backend = _BACKENDS_BY_CODE_TYPE[code_type]() 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ jby i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbkbeblbwbobmbxbrbsbtbpbqbubvbVbWbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' B ( cb) C * + , - . / : nbCbUb; = ?
771 backend.validate(options, target_type, extra_digest) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ jby i j k l z A E F G H I J K g h c d e f DbEbFbGbHbIbJbKbLbkbeblbwbobmbxbrbsbtbpbqbubvbVbWbfbgbhbibybzbAbBbMbNbObPbQbRbSbTb' B ( cb) C * + , - . / : nbCbUb; = ?
773 code_bytes = backend.encode_code(code, code_type) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ jby i j k l z A E F G H I J K g h c d e f kbeblbwbfbgbhbib' B ( cb) C * + , - . / : ; = ?
774 option_bytes = backend.option_fingerprint(options, target_type) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ jby i j k l z A E F G H I J K g h c d e f ebfbgbhbib' B ( cb) C * + , - . / : ; = ?
775 name_tags = backend.encode_name_expressions(name_expressions) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ebfbgbhbib' B ( cb) C * + , - . / : ; = ?
777 # IMPORTANT: Must use a FIPS-approved hash algorithm (SHA-2 family).
778 # FIPS-enforcing systems can disable non-approved hashlib algorithms
779 # (for example blake2b) at the OpenSSL level. See #2043.
780 hasher = hashlib.sha256(usedforsecurity=False) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
782 def _update(label: str, payload: bytes) -> None: 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
783 hasher.update(label.encode("ascii")) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
784 hasher.update(len(payload).to_bytes(8, "big")) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
785 hasher.update(payload) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
787 _update("schema", str(_KEY_SCHEMA_VERSION).encode("ascii")) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
788 backend.hash_version_probe(_update) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
789 _update("code_type", code_type.encode("ascii")) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
790 _update("target_type", target_type.encode("ascii")) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
791 _update("code", code_bytes) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
792 _update("option_count", str(len(option_bytes)).encode("ascii")) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
793 for opt in option_bytes: 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
794 _update("option", bytes(opt)) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
795 if name_tags is not None: 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
796 # ``encode_name_expressions`` returns ``None`` from backends that
797 # ignore name_expressions and a (possibly-empty) tuple from those
798 # that consume them. Hashing ``names_count=0`` for the latter
799 # keeps the schema stable across "absent" and "empty" inputs.
800 _update("names_count", str(len(name_tags)).encode("ascii")) 1WMNOPQR@[XYZ012V]3456T7US89!#$%'()*+,-./:;=?
801 for n in name_tags: 1WMNOPQR@[XYZ012V]3456T7US89!#$%'()*+,-./:;=?
802 _update("name", n) 1US
803 backend.hash_extra_payload(options, _update) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
805 # ``Program.compile()`` propagates ``options.name`` onto the returned
806 # ObjectCode, so two compiles identical in everything but name produce
807 # ObjectCodes that differ in their public ``name`` attribute. The key
808 # must reflect that or a cache hit could hand back an entry with the
809 # wrong name. Universal across backends. PTX additionally hashes
810 # ``name`` via ``_linker_option_fingerprint`` (the linker reads it),
811 # so for the linker path the value is mixed in twice under
812 # different labels. The redundancy is harmless -- distinct labels
813 # mean it cannot collide -- and the universal hash here keeps the
814 # ``options.name`` invariant in one place rather than per-backend.
815 options_name = getattr(options, "name", None) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
816 if options_name is not None: 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
817 _update("options_name", str(options_name).encode("utf-8")) 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
819 if extra_digest is not None: 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?
820 _update("extra_digest", bytes(extra_digest)) 1@[V]`^
822 return hasher.digest() 2W _ M N O P Q R m n o p q @ [ X Y Z r s 0 1 2 V ] 3 4 5 6 { a D b T | t } u 7 v U S 8 ~ w x 9 ! # $ % abbbdb` ^ y i j k l z A E F G H I J K g h c d e f ' B ( cb) C * + , - . / : ; = ?