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