Coverage for cuda/core/utils/_program_cache/_keys.py: 72.29%

1# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. 

2# 

3# SPDX-License-Identifier: Apache-2.0 

4 

5"""Cache-key construction. 

6 

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""" 

11 

12from __future__ import annotations 

13 

14import abc 

15import collections.abc 

16import hashlib 

17from collections.abc import Callable, Sequence 

18from typing import Any 

19 

20# Mutual-dependency contract: this module imports ProgramOptions from 

21# cuda.core._program at module level, and cuda.core._program imports 

22# ProgramCacheResource / make_program_cache_key from cuda.core.utils 

23# only via deferred imports inside ``Program.compile``. Adding a 

24# top-level ``from cuda.core.utils import ...`` to _program.pyx would 

25# turn this into a real import cycle -- keep the import in _program.pyx 

26# deferred (or import the symbols from the leaf submodule directly). 

27from cuda.core._program import ProgramOptions 

28from cuda.core._utils.cuda_utils import ( 

29 driver as _driver, 

30) 

31from cuda.core._utils.cuda_utils import ( 

32 handle_return as _handle_return, 

33) 

34from cuda.core._utils.cuda_utils import ( 

35 nvrtc as _nvrtc, 

36) 

37from cuda.core._utils.validators import check_str_enum, format_or_list 

38from cuda.core.typing import SourceCodeType 

39 

40# Bump when the key schema changes in a way that invalidates existing caches. 

41_KEY_SCHEMA_VERSION = 2 

42 

43_VALID_TARGET_TYPES = frozenset({"ptx", "cubin", "ltoir"}) 

44 

45# code_type -> allowed target_type set, mirroring Program.compile's 

46# SUPPORTED_TARGETS matrix in _program.pyx. 

47_SUPPORTED_TARGETS_BY_CODE_TYPE = { 

48 "c++": frozenset({"ptx", "cubin", "ltoir"}), 

49 "ptx": frozenset({"cubin", "ptx"}), 

50 "nvvm": frozenset({"ptx", "ltoir"}), 

51} 

52 

53 

54# Map each ProgramOptions field that reaches the Linker via 

55# _translate_program_options (see cuda_core/cuda/core/_program.pyx) to the 

56# gate the Linker uses to turn it into a flag (see 

57# ``_prepare_nvjitlink_options`` and ``_prepare_driver_options`` in 

58# _linker.pyx). All other fields on ProgramOptions are NVRTC-only and must 

59# NOT perturb a PTX cache key: a PTX compile with a shared ProgramOptions 

60# that happens to set include_path/pch/frandom_seed would otherwise miss the 

61# cache unnecessarily. Collapsing inputs through these gates means 

62# semantically-equivalent configurations (``debug=False`` vs ``None``, 

63# ``time=True`` vs ``time="path"``) hash to the same cache key instead of 

64# forcing spurious misses. Single source of truth: every reader iterates 

65# this dict, so adding a field here is enough -- there is no parallel 

66# field-name list to keep in sync. 

67def _gate_presence(v: Any) -> bool: 

68 return v is not None 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

69 

70 

71def _gate_truthy(v: Any) -> bool: 

72 return bool(v) 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

73 

74 

75def _gate_is_true(v: Any) -> bool: 

76 return v is True 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

77 

78 

79def _gate_tristate_bool(v: Any) -> bool | None: 

80 return None if v is None else bool(v) 36 ctx1opqrstujHbvwxyzcABCDEIJKLMNOhidefgFG

81 

82 

83def _gate_identity(v: Any) -> Any: 

84 return v 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

85 

86 

87def _gate_ptxas_options(v: Any) -> Any: 

88 # ``_prepare_nvjitlink_options`` emits one ``-Xptxas=<s>`` per element, and 

89 # treats ``str`` as a single-element sequence. Canonicalize to a tuple so 

90 # ``"-v"`` / ``["-v"]`` / ``("-v",)`` all hash the same. An empty sequence 

91 # emits no flags, so collapse it to ``None`` too. 

92 # 

93 # Order is preserved on purpose: ptxas accepts ordering-sensitive flag 

94 # pairs (e.g. ``-O2`` after ``-O3`` lowers the active level), so 

95 # ``["-v", "-O2"]`` and ``["-O2", "-v"]`` are not guaranteed to produce 

96 # identical bytes. We accept the spurious miss when callers reorder 

97 # flags; treating order as semantic keeps the cache safe in the 

98 # ordering-sensitive case. 

99 if v is None: 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

100 return None 37 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOdeFG

101 if isinstance(v, str): 7 ctx1chidefg

102 return ("-Xptxas=" + v,) 3 ctx1cfg

103 if isinstance(v, collections.abc.Sequence): 6 ctx1hidefg

104 if len(v) == 0: 6 ctx1hidefg

105 return None 3 ctx1hde

106 return tuple(f"-Xptxas={s}" for s in v) 3 ctx1ifg

107 return v 

108 

109 

110_LINKER_FIELD_GATES = { 

111 "name": _gate_identity, 

112 "arch": _gate_identity, 

113 "max_register_count": _gate_identity, 

114 "time": _gate_presence, # linker emits ``-time`` iff value is not None 

115 "link_time_optimization": _gate_truthy, 

116 "debug": _gate_truthy, 

117 "lineinfo": _gate_truthy, 

118 "ftz": _gate_tristate_bool, 

119 "prec_div": _gate_tristate_bool, 

120 "prec_sqrt": _gate_tristate_bool, 

121 "fma": _gate_tristate_bool, 

122 "split_compile": _gate_identity, 

123 "ptxas_options": _gate_ptxas_options, 

124 "no_cache": _gate_is_true, 

125} 

126 

127 

128# LinkerOptions fields the ``cuLink`` driver backend silently ignores 

129# (emits only a DeprecationWarning; no actual flag reaches the compiler). 

130# When the driver backend is active, collapse them to a single sentinel in 

131# the fingerprint so nvJitLink<->driver parity of ``ObjectCode`` doesn't 

132# cause cache misses from otherwise-equivalent configurations. 

133_DRIVER_IGNORED_LINKER_FIELDS = frozenset({"ftz", "prec_div", "prec_sqrt", "fma"}) 

134 

135 

136def _linker_option_fingerprint(options: ProgramOptions, *, use_driver_linker: bool | None) -> list[bytes]: 

137 """Backend-aware fingerprint of ProgramOptions fields consumed by the Linker. 

138 

139 Each field passes through the gate the Linker itself uses so equivalent 

140 inputs (e.g. ``debug=False`` / ``None``) hash to the same bytes. When 

141 the driver (cuLink) linker backend is in use, fields it silently 

142 ignores collapse to one sentinel so those options don't perturb the 

143 key on driver-backed hosts either. ``use_driver_linker=None`` means we 

144 couldn't probe the backend; we don't collapse driver-ignored fields in 

145 that case, to stay conservative. 

146 """ 

147 parts = [] 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

148 driver_ignored = use_driver_linker is True 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

149 for name, gate in _LINKER_FIELD_GATES.items(): 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

150 if driver_ignored and name in _DRIVER_IGNORED_LINKER_FIELDS: 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

151 parts.append(f"{name}=<driver-ignored>".encode()) 5 ctx1aklmn

152 continue 5 ctx1aklmn

153 gated = gate(getattr(options, name, None)) 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

154 parts.append(f"{name}={gated!r}".encode()) 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

155 return parts 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

156 

157 

158# ProgramOptions fields that map to LinkerOptions fields the cuLink (driver) 

159# backend rejects outright (see _prepare_driver_options in _linker.pyx). 

160# ``split_compile_extended`` exists on LinkerOptions but is not exposed via 

161# ProgramOptions / _translate_program_options, so it cannot reach the driver 

162# linker from the cache path and is omitted here. 

163_DRIVER_LINKER_UNSUPPORTED_FIELDS = ("time", "ptxas_options", "split_compile") 

164 

165 

166def _driver_version() -> int: 

167 return int(_handle_return(_driver.cuDriverGetVersion())) 5 ctx1aklmn

168 

169 

170def _nvrtc_version() -> tuple[int, int]: 

171 major, minor = _handle_return(_nvrtc.nvrtcVersion()) 44 ctx1ZPQRSTU_`012345Y{6789W!XV#$%'()*+,-./:;=?@[]

172 return int(major), int(minor) 44 ctx1ZPQRSTU_`012345Y{6789W!XV#$%'()*+,-./:;=?@[]

173 

174 

175def _linker_backend_and_version(use_driver: bool) -> tuple[str, str]: 

176 """Return ``(backend, version)`` for the linker used on PTX inputs. 

177 

178 ``use_driver`` is the result of ``_decide_nvjitlink_or_driver()`` and 

179 must be passed in so a single ``make_program_cache_key`` call shares 

180 one probe across :meth:`_LinkerBackend.validate`, 

181 :meth:`option_fingerprint`, and :meth:`hash_version_probe` (otherwise 

182 a transient probe flap could write inconsistent fields into the same 

183 key). 

184 

185 Raises any underlying probe exception. ``make_program_cache_key`` catches 

186 and mixes the exception's class name into the digest, so the same probe 

187 failure produces the same key across processes -- the cache stays 

188 persistent in broken environments, while never sharing a key with a 

189 working probe (``_probe_failed`` label vs. ``driver``/``nvrtc``/...). 

190 

191 nvJitLink version lookup goes through ``sys.modules`` first so we hit the 

192 same module ``_decide_nvjitlink_or_driver()`` already loaded. That keeps 

193 fingerprinting aligned with whichever ``cuda.bindings.nvjitlink`` import 

194 path the linker actually uses. 

195 """ 

196 import sys 32 ctx1opqrstujabvwxyzABCklmnDEhidefgFG

197 

198 if use_driver: 32 ctx1opqrstujabvwxyzABCklmnDEhidefgFG

199 return ("driver", str(_driver_version())) 5 ctx1aklmn

200 nvjitlink = sys.modules.get("cuda.bindings.nvjitlink") 27 ctx1opqrstujbvwxyzABCDEhidefgFG

201 if nvjitlink is None: 27 ctx1opqrstujbvwxyzABCDEhidefgFG

202 from cuda.bindings import nvjitlink as _nvjitlink 1 ctx1j

203 

204 nvjitlink = _nvjitlink 1 ctx1j

205 

206 return ("nvJitLink", str(nvjitlink.version())) 27 ctx1opqrstujbvwxyzABCDEhidefgFG

207 

208 

209def _nvvm_fingerprint() -> str: 

210 """Stable identifier for the loaded NVVM toolchain. 

211 

212 Combines the libNVVM library version (``module.version()``) with the IR 

213 version reported by ``module.ir_version()``. The library version is the 

214 primary invalidation lever: a libNVVM patch upgrade can change codegen 

215 while keeping the same IR major/minor, so keying only on the IR pair 

216 would silently reuse stale entries. Paired with cuda-core, the IR pair 

217 adds defence in depth without making the key any less stable. 

218 

219 Both calls go through ``_get_nvvm_module()`` so this fingerprint follows 

220 the same availability / cuda-bindings-version gate that real NVVM 

221 compilation does -- if NVVM is unusable at compile time, the probe 

222 fails the same way and ``_probe`` mixes the failure label into the key. 

223 """ 

224 from cuda.core._program import _get_nvvm_module 12 ctx2~ cbdbebfbbb^ gbhbab} ib

225 

226 module = _get_nvvm_module() 12 ctx2~ cbdbebfbbb^ gbhbab} ib

227 lib_major, lib_minor = module.version() # type: ignore[attr-defined] 12 ctx2~ cbdbebfbbb^ gbhbab} ib

228 major, minor, debug_major, debug_minor = module.ir_version() # type: ignore[attr-defined] 12 ctx2~ cbdbebfbbb^ gbhbab} ib

229 return f"lib={lib_major}.{lib_minor};ir={major}.{minor}.{debug_major}.{debug_minor}" 12 ctx2~ cbdbebfbbb^ gbhbab} ib

230 

231 

232# ProgramOptions fields that reference external files whose *contents* the 

233# cache key cannot observe without reading the filesystem. Callers that set 

234# any of these must supply an ``extra_digest`` covering the dependency surface 

235# (e.g. a hash over all reachable headers / PCH bytes). 

236_EXTERNAL_CONTENT_OPTIONS = ( 

237 "include_path", 

238 "pre_include", 

239 "pch", 

240 "use_pch", 

241 "pch_dir", 

242) 

243 

244# ProgramOptions fields whose compilation effect is not captured in the 

245# returned ``ObjectCode`` -- they produce a filesystem artifact as a side 

246# effect. A cache hit skips compilation, so that artifact would never be 

247# written. Reject these outright: the persistent cache is for pure ObjectCode 

248# reuse, not for replaying compile-time side effects. 

249# * create_pch -- writes a PCH file (NVRTC). 

250# * time -- writes NVRTC timing info to a file. 

251# * fdevice_time_trace -- writes a device-compilation time trace file (NVRTC). 

252# These are all NVRTC-specific; the Linker's ``-time`` logs to the info log 

253# (not a file) and NVVM explicitly rejects all three at compile time. The 

254# side-effect guard is therefore gated on ``backend == "nvrtc"`` below. 

255_SIDE_EFFECT_OPTIONS = ("create_pch", "time", "fdevice_time_trace") 

256 

257 

258# ProgramOptions fields gated by plain truthiness in ``_program.pyx`` (the 

259# compiler writes the flag only when the value is truthy). 

260_BOOLEAN_OPTION_FIELDS = frozenset({"pch"}) 

261 

262# Fields whose compiler emission requires ``isinstance(value, str)`` or a 

263# non-empty sequence; anything else (``False``, ``int``, ``None``, ``[]``) 

264# is silently ignored at compile time. 

265_STR_OR_SEQUENCE_OPTION_FIELDS = frozenset({"include_path", "pre_include"}) 

266 

267 

268def _option_is_set(options: ProgramOptions, name: str) -> bool: 

269 """Match how ``_program.pyx`` gates option emission, per field shape. 

270 

271 - Boolean flags (``pch``): truthy only. 

272 - str-or-sequence fields (``include_path``, ``pre_include``): ``str`` 

273 (including empty) or a non-empty ``collections.abc.Sequence`` (list, 

274 tuple, range, user subclass, ...); everything else (``False``, ``int``, 

275 empty sequence, ``None``) is ignored by the compiler and must not 

276 trigger a cache-time guard. 

277 - Path/string-shaped fields (``create_pch``, ``time``, 

278 ``fdevice_time_trace``, ``use_pch``, ``pch_dir``): ``is not None`` -- 

279 the compiler emits ``--flag=<value>`` for any non-None value, so 

280 ``False`` / ``""`` / ``0`` must still count as set. 

281 """ 

282 value = getattr(options, name, None) 78 ctx2Z P Q R S T U _ ` 0 1 2 3 4 5 Y { 6 7 8 9 W ! X V # $ % ' ( ) qbrblbsbvbtbEbybzbAbwbxbBbCb2bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* + , - . / : ; = ? ubJb1b@ [ ]

283 if value is None: 78 ctx2Z P Q R S T U _ ` 0 1 2 3 4 5 Y { 6 7 8 9 W ! X V # $ % ' ( ) qbrblbsbvbtbEbybzbAbwbxbBbCb2bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* + , - . / : ; = ? ubJb1b@ [ ]

284 return False 78 ctx2Z P Q R S T U _ ` 0 1 2 3 4 5 Y { 6 7 8 9 W ! X V # $ % ' ( ) qbrblbsbvbtbEbybzbAbwbxbBbCb2bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* + , - . / : ; = ? ubJb1b@ [ ]

285 if name in _BOOLEAN_OPTION_FIELDS: 26 ctx2P Q R S T U vbtbEbybzbAbwbxbBbCbTbUbVbWbXbYbZb0bub1b

286 return bool(value) 1 ctx2Eb

287 if name in _STR_OR_SEQUENCE_OPTION_FIELDS: 25 ctx2P Q R S T U vbtbybzbAbwbxbBbCbTbUbVbWbXbYbZb0bub1b

288 # Mirror ``_prepare_nvrtc_options_impl``: it checks ``isinstance(v, str)`` 

289 # first, then ``is_sequence(v)`` (which is ``isinstance(v, Sequence)``). 

290 # We therefore accept any ``collections.abc.Sequence`` (range, deque, 

291 # user subclass, etc.), not just list/tuple. 

292 if isinstance(value, str): 11 ctx2P Q R S T U vbtbwbxbub

293 return True 3 ctx2vbwbxb

294 if isinstance(value, collections.abc.Sequence): 8 ctx2P Q R S T U tbub

295 return len(value) > 0 6 ctx2P Q R T tbub

296 return False 2 ctx1SU

297 return True 14 ctx2ybzbAbBbCbTbUbVbWbXbYbZb0b1b

298 

299 

300def _hash_probe_failure(update: Callable[[str, bytes], None], label: str, exc: BaseException) -> None: 

301 """Mix a probe failure into the digest under a stable, content-free label. 

302 

303 Hashing only the exception's CLASS NAME (not its message) keeps the 

304 digest stable across repeated calls within one process (e.g. NVVM's 

305 loader reports different messages on first vs. cached-failure attempts) 

306 AND across processes that hit the same failure mode. The 

307 ``_probe_failed`` label differs from every backend's success label, so a 

308 broken environment never collides with a working one -- the cache 

309 "fails closed" between broken and working environments while staying 

310 persistent within either. 

311 """ 

312 update(f"{label}_probe_failed", type(exc).__name__.encode()) 5 ctx2a b W c jb

313 

314 

315class _KeyBackend(abc.ABC): 

316 """Strategy for deriving the cache key for one ``Program`` ``code_type``. 

317 

318 Each subclass owns the backend-specific guard logic, code coercion, 

319 option fingerprinting, name-expression handling, version probing, and 

320 extra-payload hashing. The orchestrator :func:`make_program_cache_key` 

321 validates the code_type / target_type pair, dispatches to the right 

322 backend, and assembles the digest. 

323 """ 

324 

325 @abc.abstractmethod 

326 def validate(self, options: ProgramOptions, target_type: str, extra_digest: bytes | None) -> None: 

327 """Reject inputs the cache cannot key safely. 

328 

329 Raises ``ValueError`` for options that have compile-time side 

330 effects, options that pull in external file content the cache 

331 can't observe, or any other backend-specific invariants. 

332 """ 

333 

334 def encode_code(self, code: object, code_type: str) -> bytes: 

335 """Coerce ``code`` to bytes. Default rejects bytes-like input 

336 (only NVVM accepts it; ``Program()`` does the same).""" 

337 if isinstance(code, str): 94 ctx2Z P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 a H b W v w x y ! z c X V # A B $ % ' ( ) qbC k l m n D E I J K L M N O h i d e f g rblbsbDbmbnbobpb* F + , G - . / : ; = ? @ [ ]

338 return code.encode("utf-8") 91 ctx2Z P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 a H b W v w x y ! z c X V # A B $ % ' ( ) qbC k l m n D E I J K L M N O h i d e f g lbmbnbobpb* F + , G - . / : ; = ? @ [ ]

339 if isinstance(code, (bytes, bytearray)): 3 ctx2rbsbDb

340 raise TypeError( 2 ctx2sbDb

341 f"code must be str for code_type={code_type!r}; bytes/bytearray are only accepted for code_type='nvvm'." 

342 ) 

343 raise TypeError(f"code must be str or bytes, got {type(code).__name__}") 1 ctx2rb

344 

345 @abc.abstractmethod 

346 def option_fingerprint(self, options: ProgramOptions, target_type: str) -> list[bytes]: 

347 """Fingerprint of the ``ProgramOptions`` fields that reach the compiler.""" 

348 

349 def encode_name_expressions(self, name_expressions: Sequence[Any]) -> tuple[bytes, ...] | None: # noqa: ARG002 

350 """Sorted, type-tagged name expressions, or ``None`` if the 

351 backend does not consume them. 

352 

353 ``None`` means the orchestrator emits no ``names_count`` / 

354 ``name`` entries at all (a backend that ignores 

355 ``name_expressions`` should never have them perturb its key). An 

356 empty tuple means the backend supports them but the caller 

357 passed none -- the orchestrator still emits ``names_count=0`` so 

358 the schema is stable across "absent" and "empty". 

359 """ 

360 return None 56 ctx2~ o p q r s t u j cba H b dbv ebw x y z c fbA B bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g F ibG

361 

362 @abc.abstractmethod 

363 def hash_version_probe(self, update: Callable[[str, bytes], None]) -> None: 

364 """Mix the runtime/compiler version probe into the digest via 

365 ``update(label, payload)``. On probe failure, mix 

366 ``_hash_probe_failure(update, "<label>", exc)`` instead so the 

367 digest is stable across processes hitting the same failure 

368 mode. 

369 """ 

370 

371 def hash_extra_payload(self, options: ProgramOptions, update: Callable[[str, bytes], None]) -> None: # noqa: B027 85 ctx1ZPQRSTUopqrs_`012tu3j45Y{6789aHbWvwxy!zcXV#AB$%'()CklmnDEIJKLMNOhidefg*F+,G-./:;=?@[]

372 """Mix backend-specific extras (e.g. NVVM ``extra_sources`` / 

373 ``use_libdevice``). Default: nothing. 

374 """ 

375 

376 

377class _NvrtcBackend(_KeyBackend): 

378 def validate(self, options: ProgramOptions, target_type: str, extra_digest: bytes | None) -> None: # noqa: ARG002 

379 # Side-effect options are NVRTC-specific: 

380 # ``time``/``fdevice_time_trace`` write artifacts via NVRTC, 

381 # ``create_pch`` writes via NVRTC. The Linker's ``-time`` logs to 

382 # the info log (not a file), and NVVM explicitly rejects all three 

383 # at compile time, so the side-effect guard is meaningful only for 

384 # the NVRTC path. 

385 side_effects = [name for name in _SIDE_EFFECT_OPTIONS if _option_is_set(options, name)] 78 ctx2Z P Q R S T U _ ` 0 1 2 3 4 5 Y { 6 7 8 9 W ! X V # $ % ' ( ) qbrblbsbvbtbEbybzbAbwbxbBbCb2bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* + , - . / : ; = ? ubJb1b@ [ ]

386 if side_effects: 78 ctx2Z P Q R S T U _ ` 0 1 2 3 4 5 Y { 6 7 8 9 W ! X V # $ % ' ( ) qbrblbsbvbtbEbybzbAbwbxbBbCb2bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* + , - . / : ; = ? ubJb1b@ [ ]

387 raise ValueError( 9 ctx2TbUbVbWbXbYbZb0b1b

388 f"make_program_cache_key() refuses to build a key for options that " 

389 f"have compile-time side effects ({', '.join(side_effects)}); a " 

390 f"cache hit skips compilation, so the side effect would not occur. " 

391 f"Disable the option, or compile directly without the cache." 

392 ) 

393 # ``extra_sources`` is NVVM-only -- ``Program`` raises for non-NVVM 

394 # backends (_program.pyx). Reject here so callers get the same 

395 # error from the cache-key path as from a real compile. 

396 if getattr(options, "extra_sources", None) is not None: 69 ctx2Z P Q R S T U _ ` 0 1 2 3 4 5 Y { 6 7 8 9 W ! X V # $ % ' ( ) qbrblbsbvbtbEbybzbAbwbxbBbCb2bmbnbobpbFbGbHbIb* + , - . / : ; = ? ubJb@ [ ]

397 raise ValueError( 1 ctx22b

398 "extra_sources is only valid for code_type='nvvm'; Program() rejects it for code_type='c++'." 

399 ) 

400 if extra_digest is None: 68 ctx2Z P Q R S T U _ ` 0 1 2 3 4 5 Y { 6 7 8 9 W ! X V # $ % ' ( ) qbrblbsbvbtbEbybzbAbwbxbBbCbmbnbobpbFbGbHbIb* + , - . / : ; = ? ubJb@ [ ]

401 # ``Program.compile`` for PTX inputs runs 

402 # ``_translate_program_options``, which drops these entirely; 

403 # NVVM rejects them. Only NVRTC reads the external content. 

404 external = [name for name in _EXTERNAL_CONTENT_OPTIONS if _option_is_set(options, name)] 65 ctx2Z P Q R S T U 0 1 2 3 4 5 Y 6 7 8 9 W ! X V # $ % ' ( ) qbrblbsbvbtbEbybzbAbwbxbBbCbmbnbobpbFbGbHbIb* + , - . / : ; = ? ubJb@ [ ]

405 if external: 65 ctx2Z P Q R S T U 0 1 2 3 4 5 Y 6 7 8 9 W ! X V # $ % ' ( ) qbrblbsbvbtbEbybzbAbwbxbBbCbmbnbobpbFbGbHbIb* + , - . / : ; = ? ubJb@ [ ]

406 raise ValueError( 11 ctx2vbtbEbybzbAbwbxbBbCbub

407 f"make_program_cache_key() refuses to build a key for options that " 

408 f"pull in external file content ({', '.join(external)}) without an " 

409 f"extra_digest; compute a digest over the header/PCH bytes the " 

410 f"compile will read and pass it as extra_digest=..." 

411 ) 

412 # NVRTC uses ``options.name`` as the source filename and 

413 # resolves quoted ``#include "x.h"`` directives relative to 

414 # the directory component of that name. The directory's 

415 # contents are external to anything else the key observes, 

416 # so a name with a directory component requires the same 

417 # ``extra_digest`` treatment as ``include_path`` etc. 

418 options_name = getattr(options, "name", None) 54 ctx2Z P Q R S T U 0 1 2 3 4 5 Y 6 7 8 9 W ! X V # $ % ' ( ) qbrblbsbmbnbobpbFbGbHbIb* + , - . / : ; = ? Jb@ [ ]

419 if isinstance(options_name, str) and ("/" in options_name or "\\" in options_name): 54 ctx2Z P Q R S T U 0 1 2 3 4 5 Y 6 7 8 9 W ! X V # $ % ' ( ) qbrblbsbmbnbobpbFbGbHbIb* + , - . / : ; = ? Jb@ [ ]

420 raise ValueError( 5 ctx2FbGbHbIbJb

421 f"make_program_cache_key() refuses to build a key for options.name=" 

422 f"{options_name!r} (NVRTC source-filename with a directory " 

423 f"component) without an extra_digest; NVRTC resolves quoted " 

424 f"#include directives relative to that directory, so a digest " 

425 f"covering the headers it may pull in must be supplied." 

426 ) 

427 

428 def option_fingerprint(self, options: ProgramOptions, target_type: str) -> list[bytes]: 

429 # ``ProgramOptions.as_bytes("nvrtc", ...)`` gives the real 

430 # compile-time flag surface for NVRTC. 

431 return options.as_bytes("nvrtc", target_type) 50 ctx2Z P Q R S T U _ ` 0 1 2 3 4 5 Y { 6 7 8 9 W ! X V # $ % ' ( ) qblbmbnbobpb* + , - . / : ; = ? @ [ ]

432 

433 def encode_name_expressions(self, name_expressions: Sequence[Any]) -> tuple[bytes, ...]: 

434 # ``"foo"`` and ``b"foo"`` get distinct tags because 

435 # ``Program.compile`` records the original Python object as the 

436 # ``ObjectCode.symbol_mapping`` key, so a cached ObjectCode whose 

437 # mapping-key type differs from what the caller's later 

438 # ``get_kernel`` passes would silently miss. Reject ``bytearray`` 

439 # because ``Program.compile`` also uses the raw element as a dict 

440 # key -- bytearray is unhashable, so a cache miss would compile 

441 # then crash in ``symbol_mapping[n] = ...``. Accepting it here 

442 # would let the cache serve hits for inputs the uncached path 

443 # can't handle. 

444 def _tag(n: Any) -> bytes: 49 ctx2Z P Q R S T U _ ` 0 1 2 3 4 5 Y { 6 7 8 9 W ! X V # $ % ' ( ) lbmbnbobpb* + , - . / : ; = ? @ [ ]

445 if isinstance(n, bytes): 7 ctx2X V lbmbnbobpb

446 return b"b:" + n 1 ctx1V

447 if isinstance(n, str): 7 ctx2X V lbmbnbobpb

448 return b"s:" + n.encode("utf-8") 7 ctx2X V lbmbnbobpb

449 if isinstance(n, bytearray): 5 ctx2lbmbnbobpb

450 raise TypeError( 1 ctx2lb

451 "name_expressions elements must be str or bytes; " 

452 "bytearray is not accepted because Program.compile uses " 

453 "each element as a dict key and bytearray is unhashable." 

454 ) 

455 raise TypeError(f"name_expressions elements must be str or bytes; got {type(n).__name__}") 4 ctx2mbnbobpb

456 

457 return tuple(sorted(_tag(n) for n in name_expressions)) 49 ctx2Z P Q R S T U _ ` 0 1 2 3 4 5 Y { 6 7 8 9 W ! X V # $ % ' ( ) lbmbnbobpb* + , - . / : ; = ? @ [ ]

458 

459 def hash_version_probe(self, update: Callable[[str, bytes], None]) -> None: 

460 try: 44 ctx1ZPQRSTU_`012345Y{6789W!XV#$%'()*+,-./:;=?@[]

461 major, minor = _nvrtc_version() 44 ctx1ZPQRSTU_`012345Y{6789W!XV#$%'()*+,-./:;=?@[]

462 except Exception as exc: 1 ctx1W

463 _hash_probe_failure(update, "nvrtc", exc) 1 ctx1W

464 return 1 ctx1W

465 update("nvrtc", f"{major}.{minor}".encode("ascii")) 44 ctx1ZPQRSTU_`012345Y{6789W!XV#$%'()*+,-./:;=?@[]

466 

467 

468_DECISION_UNSET = object() 

469 

470 

471class _LinkerBackend(_KeyBackend): 

472 def __init__(self) -> None: 

473 # Cache the linker-backend decision (and any probe failure) for 

474 # the duration of one ``make_program_cache_key`` call so 

475 # ``validate``, ``option_fingerprint``, and ``hash_version_probe`` 

476 # all see the same answer; a transient probe flap mid-call 

477 # otherwise mints a key whose option fingerprint and version 

478 # probe disagree on which linker is in use. 

479 self._cached_decision = _DECISION_UNSET 52 ctx2o p q r s t u j a H b v w x y z c A B C k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbDb4bF G

480 self._cached_decision_exc: BaseException | None = None 52 ctx2o p q r s t u j a H b v w x y z c A B C k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbDb4bF G

481 

482 def _decide_driver(self) -> bool | None: 

483 """``True`` if the cuLink driver linker will be used, ``False`` if 

484 nvJitLink, ``None`` if the probe failed (in which case 

485 :meth:`hash_version_probe` mixes a ``_probe_failed`` taint into 

486 the digest instead of a backend label). 

487 """ 

488 if self._cached_decision is _DECISION_UNSET: 51 ctx2o p q r s t u j a H b v w x y z c A B C k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbDbF G

489 try: 51 ctx2o p q r s t u j a H b v w x y z c A B C k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbDbF G

490 from cuda.core._linker import _decide_nvjitlink_or_driver 51 ctx2o p q r s t u j a H b v w x y z c A B C k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbDbF G

491 

492 self._cached_decision = _decide_nvjitlink_or_driver() 51 ctx2o p q r s t u j a H b v w x y z c A B C k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbDbF G

493 except Exception as exc: 1 ctx1c

494 self._cached_decision = None 1 ctx1c

495 self._cached_decision_exc = exc 1 ctx1c

496 return self._cached_decision # type: ignore[return-value] 51 ctx2o p q r s t u j a H b v w x y z c A B C k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbDbF G

497 

498 def validate(self, options: ProgramOptions, target_type: str, extra_digest: bytes | None) -> None: # noqa: ARG002 

499 if getattr(options, "extra_sources", None) is not None: 52 ctx2o p q r s t u j a H b v w x y z c A B C k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbDb4bF G

500 raise ValueError( 1 ctx24b

501 "extra_sources is only valid for code_type='nvvm'; Program() rejects it for code_type='ptx'." 

502 ) 

503 # ``numba_debug`` is deliberately not rejected here and is absent from 

504 # ``_LINKER_FIELD_GATES``: for PTX inputs the linker ignores it (with a 

505 # warning from ``_translate_program_options``), so it cannot change the 

506 # generated code and must not perturb the key. Two PTX compiles that 

507 # differ only in ``numba_debug`` are the same compile. 

508 # PTX compiles go through the Linker. When the driver (cuLink) 

509 # backend is selected (nvJitLink unavailable), ``Program.compile`` 

510 # rejects a subset of options that nvJitLink would accept; reject 

511 # them here too so we never store a key for a compilation that 

512 # can't succeed in this environment. If the probe fails we can't 

513 # tell which backend will run, so skip -- the failed-probe taint 

514 # in ``hash_version_probe`` already poisons the key. 

515 if self._decide_driver() is True: 51 ctx2o p q r s t u j a H b v w x y z c A B C k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbDbF G

516 # Mirror ``_prepare_driver_options``'s exact gate: ``is not 

517 # None`` for these fields, so ``time=False`` or 

518 # ``ptxas_options=[]`` is still a rejection. Do NOT use the 

519 # truthiness-based ``_option_is_set`` helper here. 

520 unsupported = [ 14 ctx2a k l m n KbLbMbNbObPbQbRbSb

521 name for name in _DRIVER_LINKER_UNSUPPORTED_FIELDS if getattr(options, name, None) is not None 

522 ] 

523 if unsupported: 14 ctx2a k l m n KbLbMbNbObPbQbRbSb

524 raise ValueError( 9 ctx2KbLbMbNbObPbQbRbSb

525 f"the cuLink driver linker does not support these options: " 

526 f"{', '.join(unsupported)}; Program.compile() would reject this " 

527 f"configuration before producing an ObjectCode." 

528 ) 

529 

530 def option_fingerprint(self, options: ProgramOptions, target_type: str) -> list[bytes]: # noqa: ARG002 

531 # For PTX inputs the Linker reads only a subset of ProgramOptions 

532 # (see ``_translate_program_options`` in _program.pyx); fingerprint 

533 # just those fields so shared ProgramOptions carrying NVRTC-only 

534 # flags (``include_path``, ``pch_*``, ``frandom_seed``, ...) don't 

535 # force spurious cache misses on PTX. 

536 return _linker_option_fingerprint(options, use_driver_linker=self._decide_driver()) 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

537 

538 def hash_version_probe(self, update: Callable[[str, bytes], None]) -> None: 

539 # Only cuLink (driver-backed linker) goes through the CUDA driver 

540 # for codegen. nvJitLink is a separate library, so a driver 

541 # upgrade under it does not change the compiled bytes -- skip the 

542 # driver version there. ``_linker_backend_and_version`` already 

543 # returns the driver version when the driver backend is active, 

544 # so the bytes are still in the digest via ``linker_version``. 

545 use_driver = self._decide_driver() 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

546 if use_driver is None: 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

547 if self._cached_decision_exc is None: 1 ctx1c

548 raise RuntimeError("internal cuda.core error, please report: linker probe failed without an exception") 

549 _hash_probe_failure(update, "linker", self._cached_decision_exc) 1 ctx1c

550 return 1 ctx1c

551 try: 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

552 lb_name, lb_version = _linker_backend_and_version(use_driver) 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

553 except Exception as exc: 2 ctx1ab

554 _hash_probe_failure(update, "linker", exc) 2 ctx1ab

555 return 2 ctx1ab

556 update("linker_backend", lb_name.encode("ascii")) 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

557 update("linker_version", lb_version.encode("ascii")) 41 ctx1opqrstujaHbvwxyzcABCklmnDEIJKLMNOhidefgFG

558 

559 

560class _NvvmBackend(_KeyBackend): 

561 def encode_code(self, code: object, code_type: str) -> bytes: # noqa: ARG002 

562 # NVVM accepts both str and bytes (matching ``Program()``). 

563 if isinstance(code, str): 16 ctx2~ cbdbebfbbb^ | gbhbkbjb3bab} ib

564 return code.encode("utf-8") 14 ctx2~ cbdbebfb^ | gbhbkbjbab} ib

565 if isinstance(code, (bytes, bytearray)): 3 ctx2~ bb3b

566 return bytes(code) 2 ctx2~ bb

567 raise TypeError(f"code must be str or bytes, got {type(code).__name__}") 1 ctx23b

568 

569 def validate(self, options: ProgramOptions, target_type: str, extra_digest: bytes | None) -> None: # noqa: ARG002 

570 # NVVM with ``use_libdevice=True`` reads external libdevice 

571 # bitcode at compile time (see Program_init in _program.pyx). The 

572 # file is resolved from the active toolkit, so a changed 

573 # CUDA_HOME / libdevice upgrade changes the linked output without 

574 # touching any key input the cache can observe. Require the 

575 # caller to supply an ``extra_digest`` that fingerprints the 

576 # libdevice bytes (or simply disable use_libdevice for 

577 # caching-sensitive workflows). 

578 if extra_digest is None and getattr(options, "use_libdevice", None): 16 ctx2~ cbdbebfbbb^ | gbhbkbjb3bab} ib

579 raise ValueError( 1 ctx1}

580 "make_program_cache_key() refuses to build an NVVM key with " 

581 "use_libdevice=True and no extra_digest: the linked libdevice " 

582 "bitcode can change out from under a cached ObjectCode. Pass an " 

583 "extra_digest that fingerprints the libdevice file you intend " 

584 "to link against, or disable use_libdevice." 

585 ) 

586 

587 def option_fingerprint(self, options: ProgramOptions, target_type: str) -> list[bytes]: 

588 return options.as_bytes("nvvm", target_type) 15 ctx2~ cbdbebfbbb^ | gbhbkbjbab} ib

589 

590 def hash_version_probe(self, update: Callable[[str, bytes], None]) -> None: 

591 try: 15 ctx2~ cbdbebfbbb^ | gbhbkbjbab} ib

592 fp = _nvvm_fingerprint() 15 ctx2~ cbdbebfbbb^ | gbhbkbjbab} ib

593 except Exception as exc: 1 ctx2jb

594 _hash_probe_failure(update, "nvvm", exc) 1 ctx2jb

595 return 1 ctx2jb

596 update("nvvm", fp.encode("ascii")) 15 ctx2~ cbdbebfbbb^ | gbhbkbjbab} ib

597 

598 def hash_extra_payload(self, options: ProgramOptions, update: Callable[[str, bytes], None]) -> None: 

599 extra_sources = getattr(options, "extra_sources", None) 15 ctx2~ cbdbebfbbb^ | gbhbkbjbab} ib

600 if extra_sources: 15 ctx2~ cbdbebfbbb^ | gbhbkbjbab} ib

601 # ``extra_sources`` is hashed in caller-provided order on purpose. 

602 # NVVM module linking is order-dependent in the general case 

603 # (overlapping symbols, weak definitions, definition order can 

604 # change which body wins), so canonicalising by sorting on the 

605 # source name would produce the same key for two compiles whose 

606 # outputs may legitimately differ. If a future test proves the 

607 # relevant input subset is order-insensitive, sorting can be 

608 # introduced under that proof; absent that proof, preserving 

609 # caller order is the safe default. 

610 update("extra_sources_count", str(len(extra_sources)).encode("ascii")) 2 ctx1^|

611 for item in extra_sources: 2 ctx1^|

612 # ``extra_sources`` is a sequence of (name, source) tuples. 

613 if isinstance(item, (tuple, list)) and len(item) == 2: 2 ctx1^|

614 name, src = item 2 ctx1^|

615 update("extra_source_name", str(name).encode("utf-8")) 2 ctx1^|

616 match src: 2 ctx1^|

617 case str(): 2 ctx1^|

618 update("extra_source_code", src.encode("utf-8")) 2 ctx1^|

619 case bytes() | bytearray(): 1 ctx1|

620 update("extra_source_code", bytes(src)) 1 ctx1|

621 case _: 

622 update("extra_source_code", str(src).encode("utf-8")) 

623 else: 

624 # Fallback for unexpected format. 

625 update("extra_source", str(item).encode("utf-8")) 

626 # ``use_libdevice`` is gated on truthiness to match Program_init's 

627 # gate -- ``False`` and ``None`` collapse to the same key. 

628 if getattr(options, "use_libdevice", None): 15 ctx2~ cbdbebfbbb^ | gbhbkbjbab} ib

629 update("use_libdevice", b"1") 2 ctx2ab}

630 

631 

632# Class registry keyed by code_type. ``make_program_cache_key`` instantiates 

633# fresh per call so backends like ``_LinkerBackend`` can cache per-call probe 

634# results on ``self`` without leaking that state across calls. 

635_BACKENDS_BY_CODE_TYPE: dict[str, type[_KeyBackend]] = { 

636 "c++": _NvrtcBackend, 

637 "ptx": _LinkerBackend, 

638 "nvvm": _NvvmBackend, 

639} 

640 

641 

642def make_program_cache_key( 

643 *, 

644 code: str | bytes, 

645 code_type: str, 

646 options: ProgramOptions, 

647 target_type: str, 

648 name_expressions: Sequence[str | bytes | bytearray] = (), 

649 extra_digest: bytes | None = None, 

650) -> bytes: 

651 """Build a stable cache key from compile inputs. 

652 

653 Parameters 

654 ---------- 

655 code: 

656 Source text. ``str`` is encoded as UTF-8. 

657 code_type: 

658 One of ``"c++"``, ``"ptx"``, ``"nvvm"``. 

659 options: 

660 A :class:`cuda.core.ProgramOptions`. Its ``arch`` must be set (the 

661 default ``ProgramOptions.__post_init__`` populates it from the current 

662 device). 

663 target_type: 

664 One of ``"ptx"``, ``"cubin"``, ``"ltoir"``. 

665 name_expressions: 

666 Optional iterable of mangled-name lookups. Order is not significant. 

667 Elements may be ``str`` or ``bytes``; ``"foo"`` and ``b"foo"`` produce 

668 distinct keys because ``Program.compile`` records the original Python 

669 object as the ``ObjectCode.symbol_mapping`` key, and ``get_kernel`` 

670 lookups must use the same type the cache key recorded. ``bytearray`` 

671 is rejected because ``Program.compile`` stores each element as a 

672 dict key and ``bytearray`` is unhashable. 

673 extra_digest: 

674 Caller-supplied bytes mixed into the key. Required whenever 

675 :class:`cuda.core.ProgramOptions` sets any option that pulls in 

676 external file content (``include_path``, ``pre_include``, ``pch``, 

677 ``use_pch``, ``pch_dir``) -- the cache cannot read those files on 

678 the caller's behalf, so the caller must fingerprint the header / 

679 PCH surface and pass it here. Callers may pass this for other 

680 inputs too (embedded kernels, generated sources, etc.). 

681 

682 Returns 

683 ------- 

684 bytes 

685 An opaque bytes digest suitable for use as a cache key. 

686 

687 Raises 

688 ------ 

689 ValueError 

690 If ``options`` sets an option with compile-time side effects (such 

691 as ``create_pch``) -- a cache hit skips compilation, so the side 

692 effect would not occur. 

693 ValueError 

694 If ``extra_digest`` is ``None`` while ``options`` sets any option 

695 whose compilation effect depends on external file content that the 

696 key cannot otherwise observe. 

697 

698 Examples 

699 -------- 

700 For most workflows you should not call ``make_program_cache_key`` 

701 yourself -- pass ``cache=`` to :meth:`cuda.core.Program.compile`, 

702 which derives the key, returns the cached 

703 :class:`~cuda.core.ObjectCode` on hit, and stores the compile 

704 result on miss:: 

705 

706 from cuda.core import Program, ProgramOptions 

707 from cuda.core.utils import FileStreamProgramCache 

708 

709 source = 'extern "C" __global__ void k(int *a){ *a = 1; }' 

710 options = ProgramOptions(arch="sm_80") 

711 

712 with FileStreamProgramCache() as cache: 

713 obj = Program(source, "c++", options=options).compile("cubin", cache=cache) 

714 

715 Call ``make_program_cache_key`` directly when the compile inputs 

716 require an ``extra_digest`` (the cache cannot read external file 

717 content on the caller's behalf) -- ``Program.compile(cache=...)`` 

718 refuses those inputs with a ``ValueError`` pointing here:: 

719 

720 from cuda.core import ObjectCode 

721 from cuda.core.utils import FileStreamProgramCache, make_program_cache_key 

722 

723 with FileStreamProgramCache() as cache: 

724 key = make_program_cache_key( 

725 code=source, 

726 code_type="c++", 

727 options=options, 

728 target_type="cubin", 

729 extra_digest=fingerprint_headers(options.include_path), 

730 ) 

731 data = cache.get(key) 

732 if data is None: 

733 obj = Program(source, "c++", options=options).compile("cubin") 

734 cache[key] = obj # extracts bytes(obj.code) 

735 else: 

736 obj = ObjectCode.from_cubin(data) 

737 

738 The cache stores raw binary bytes -- cubin / PTX / LTO-IR with no 

739 pickle, JSON, or framing -- so entry files are directly consumable 

740 by external NVIDIA tools (``cuobjdump``, ``nvdisasm``, ...). Note 

741 that an :class:`~cuda.core.ObjectCode` round-tripped through the 

742 cache loses ``symbol_mapping``: callers that compile with 

743 ``name_expressions`` and rely on ``get_kernel(name_expression)`` 

744 after a cache hit must either compile fresh or look up the mangled 

745 symbol explicitly. 

746 

747 Options that read external files (``include_path``, ``pre_include``, 

748 ``pch``, ``use_pch``, ``pch_dir``; ``use_libdevice=True`` on the NVVM 

749 path; and on NVRTC, an ``options.name`` with a directory component, 

750 which NVRTC uses for relative-include resolution) require 

751 ``extra_digest`` -- fingerprint the bytes the compiler will pull in 

752 and pass that digest so changes to those files force a cache miss. 

753 Options that have compile-time side effects (``create_pch``, 

754 ``time``, ``fdevice_time_trace``) cannot be cached and raise 

755 ``ValueError``; compile directly, or disable the flag, for those 

756 cases. 

757 """ 

758 # Mirror Program.compile (_program.pyx lowercases code_type at Program 

759 # init and target_type at the top of compile); a caller that passes 

760 # "PTX" or "C++" must get the same routing and the same cache key as 

761 # the lowercase form. 

762 code_type = code_type.lower() if isinstance(code_type, str) else code_type 150 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjb3bab} qbC k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbrb5b6b8b7blbsbDbvbtbEbybzbAbwbxbBbCb2b4bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* F + ib, G - . / : ; = ? ubJb1b@ [ ]

763 target_type = target_type.lower() if isinstance(target_type, str) else target_type 150 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjb3bab} qbC k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbrb5b6b8b7blbsbDbvbtbEbybzbAbwbxbBbCb2b4bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* F + ib, G - . / : ; = ? ubJb1b@ [ ]

764 check_str_enum(code_type, SourceCodeType) 150 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjb3bab} qbC k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbrb5b6b8b7blbsbDbvbtbEbybzbAbwbxbBbCb2b4bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* F + ib, G - . / : ; = ? ubJb1b@ [ ]

765 if target_type not in _VALID_TARGET_TYPES: 149 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjb3bab} qbC k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbrb5b6b7blbsbDbvbtbEbybzbAbwbxbBbCb2b4bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* F + ib, G - . / : ; = ? ubJb1b@ [ ]

766 raise ValueError( 1 ctx27b

767 f"target_type={target_type!r} is not supported by the program cache " 

768 f"(must be {format_or_list(sorted(_VALID_TARGET_TYPES))})" 

769 ) 

770 supported_for_code = _SUPPORTED_TARGETS_BY_CODE_TYPE[code_type] 148 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjb3bab} qbC k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbrb5b6blbsbDbvbtbEbybzbAbwbxbBbCb2b4bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* F + ib, G - . / : ; = ? ubJb1b@ [ ]

771 if target_type not in supported_for_code: 148 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjb3bab} qbC k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbrb5b6blbsbDbvbtbEbybzbAbwbxbBbCb2b4bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* F + ib, G - . / : ; = ? ubJb1b@ [ ]

772 raise ValueError( 2 ctx25b6b

773 f"target_type={target_type!r} is not valid for code_type={code_type!r}" 

774 f" (supported: {sorted(supported_for_code)}). Program.compile() rejects" 

775 f" this combination, so caching a key for it is meaningless." 

776 ) 

777 

778 backend = _BACKENDS_BY_CODE_TYPE[code_type]() 146 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjb3bab} qbC k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbrblbsbDbvbtbEbybzbAbwbxbBbCb2b4bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* F + ib, G - . / : ; = ? ubJb1b@ [ ]

779 backend.validate(options, target_type, extra_digest) 146 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjb3bab} qbC k l m n D E I J K L M N O h i d e f g KbLbMbNbObPbQbRbSbrblbsbDbvbtbEbybzbAbwbxbBbCb2b4bmbnbobpbFbGbHbIbTbUbVbWbXbYbZb0b* F + ib, G - . / : ; = ? ubJb1b@ [ ]

780 

781 code_bytes = backend.encode_code(code, code_type) 110 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjb3bab} qbC k l m n D E I J K L M N O h i d e f g rblbsbDbmbnbobpb* F + ib, G - . / : ; = ? @ [ ]

782 option_bytes = backend.option_fingerprint(options, target_type) 106 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} qbC k l m n D E I J K L M N O h i d e f g lbmbnbobpb* F + ib, G - . / : ; = ? @ [ ]

783 name_tags = backend.encode_name_expressions(name_expressions) 105 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g lbmbnbobpb* F + ib, G - . / : ; = ? @ [ ]

784 

785 # IMPORTANT: Must use a FIPS-approved hash algorithm (SHA-2 family). 

786 # FIPS-enforcing systems can disable non-approved hashlib algorithms 

787 # (for example blake2b) at the OpenSSL level. See #2043. 

788 hasher = hashlib.sha256(usedforsecurity=False) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

789 

790 def _update(label: str, payload: bytes) -> None: 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

791 hasher.update(label.encode("ascii")) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

792 hasher.update(len(payload).to_bytes(8, "big")) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

793 hasher.update(payload) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

794 

795 _update("schema", str(_KEY_SCHEMA_VERSION).encode("ascii")) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

796 backend.hash_version_probe(_update) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

797 _update("code_type", code_type.encode("ascii")) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

798 _update("target_type", target_type.encode("ascii")) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

799 _update("code", code_bytes) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

800 _update("option_count", str(len(option_bytes)).encode("ascii")) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

801 for opt in option_bytes: 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

802 _update("option", bytes(opt)) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

803 if name_tags is not None: 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

804 # ``encode_name_expressions`` returns ``None`` from backends that 

805 # ignore name_expressions and a (possibly-empty) tuple from those 

806 # that consume them. Hashing ``names_count=0`` for the latter 

807 # keeps the schema stable across "absent" and "empty" inputs. 

808 _update("names_count", str(len(name_tags)).encode("ascii")) 44 ctx1ZPQRSTU_`012345Y{6789W!XV#$%'()*+,-./:;=?@[]

809 for n in name_tags: 44 ctx1ZPQRSTU_`012345Y{6789W!XV#$%'()*+,-./:;=?@[]

810 _update("name", n) 2 ctx1XV

811 backend.hash_extra_payload(options, _update) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

812 

813 # ``Program.compile()`` propagates ``options.name`` onto the returned 

814 # ObjectCode, so two compiles identical in everything but name produce 

815 # ObjectCodes that differ in their public ``name`` attribute. The key 

816 # must reflect that or a cache hit could hand back an entry with the 

817 # wrong name. Universal across backends. PTX additionally hashes 

818 # ``name`` via ``_linker_option_fingerprint`` (the linker reads it), 

819 # so for the linker path the value is mixed in twice under 

820 # different labels. The redundancy is harmless -- distinct labels 

821 # mean it cannot collide -- and the universal hash here keeps the 

822 # ``options.name`` invariant in one place rather than per-backend. 

823 options_name = getattr(options, "name", None) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

824 if options_name is not None: 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

825 _update("options_name", str(options_name).encode("utf-8")) 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

826 

827 if extra_digest is not None: 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]

828 _update("extra_digest", bytes(extra_digest)) 6 ctx2_ ` Y { ab}

829 

830 return hasher.digest() 100 ctx2Z ~ P Q R S T U o p q r s _ ` 0 1 2 t u 3 j 4 5 Y { 6 7 8 9 cba H b W dbv ebw x y ! z c X V # fbA B $ % ' ( ) bb^ | gbhbkbjbab} C k l m n D E I J K L M N O h i d e f g * F + ib, G - . / : ; = ? @ [ ]