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

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 (empty)

13 

14import abc (empty)

15import collections.abc (empty)

16import hashlib (empty)

17from typing import Any, Callable, Sequence (empty)

18 

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 (empty)

27from cuda.core._utils.cuda_utils import ( (empty)

28 driver as _driver, 

29) 

30from cuda.core._utils.cuda_utils import ( (empty)

31 handle_return as _handle_return, 

32) 

33from cuda.core._utils.cuda_utils import ( (empty)

34 nvrtc as _nvrtc, 

35) 

36from cuda.core._utils.validators import check_str_enum, format_or_list (empty)

37from cuda.core.typing import SourceCodeType (empty)

38 

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

40_KEY_SCHEMA_VERSION = 2 (empty)

41 

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

43 

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

45# SUPPORTED_TARGETS matrix in _program.pyx. 

46_SUPPORTED_TARGETS_BY_CODE_TYPE = { (empty)

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

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

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

50} 

51 

52 

53# Map each ProgramOptions field that reaches the Linker via 

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

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

56# ``_prepare_nvjitlink_options`` and ``_prepare_driver_options`` in 

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

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

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

60# cache unnecessarily. Collapsing inputs through these gates means 

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

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

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

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

65# field-name list to keep in sync. 

66def _gate_presence(v: Any) -> bool: (empty)

67 return v is not None 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

68 

69 

70def _gate_truthy(v: Any) -> bool: (empty)

71 return bool(v) 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

72 

73 

74def _gate_is_true(v: Any) -> bool: (empty)

75 return v is True 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

76 

77 

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

79 return None if v is None else bool(v) 34 ctx1mnopqrsFbtuvwxyzABCGHIJKLMghcdefDE

80 

81 

82def _gate_identity(v: Any) -> Any: (empty)

83 return v 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

84 

85 

86def _gate_ptxas_options(v: Any) -> Any: (empty)

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

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

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

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

91 # 

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

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

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

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

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

97 # ordering-sensitive case. 

98 if v is None: 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

99 return None 35 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMcdDE

100 if isinstance(v, str): 6 ctx1ghcdef

101 return ("-Xptxas=" + v,) 2 ctx1ef

102 if isinstance(v, collections.abc.Sequence): 6 ctx1ghcdef

103 if len(v) == 0: 6 ctx1ghcdef

104 return None 3 ctx1gcd

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

106 return v 

107 

108 

109_LINKER_FIELD_GATES = { (empty)

110 "name": _gate_identity, 

111 "arch": _gate_identity, 

112 "max_register_count": _gate_identity, 

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

114 "link_time_optimization": _gate_truthy, 

115 "debug": _gate_truthy, 

116 "lineinfo": _gate_truthy, 

117 "ftz": _gate_tristate_bool, 

118 "prec_div": _gate_tristate_bool, 

119 "prec_sqrt": _gate_tristate_bool, 

120 "fma": _gate_tristate_bool, 

121 "split_compile": _gate_identity, 

122 "ptxas_options": _gate_ptxas_options, 

123 "no_cache": _gate_is_true, 

124} 

125 

126 

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

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

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

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

131# cause cache misses from otherwise-equivalent configurations. 

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

133 

134 

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

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

137 

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

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

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

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

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

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

144 that case, to stay conservative. 

145 """ 

146 parts = [] 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

147 driver_ignored = use_driver_linker is True 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

148 for name, gate in _LINKER_FIELD_GATES.items(): 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

149 if driver_ignored and name in _DRIVER_IGNORED_LINKER_FIELDS: 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

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

151 continue 5 ctx1aijkl

152 gated = gate(getattr(options, name, None)) 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

153 parts.append(f"{name}={gated!r}".encode()) 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

154 return parts 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

155 

156 

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

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

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

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

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

162_DRIVER_LINKER_UNSUPPORTED_FIELDS = ("time", "ptxas_options", "split_compile") (empty)

163 

164 

165def _driver_version() -> int: (empty)

166 return int(_handle_return(_driver.cuDriverGetVersion())) 5 ctx1aijkl

167 

168 

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

170 major, minor = _handle_return(_nvrtc.nvrtcVersion()) 44 ctx1YOPQRST]^Z01234X_5678V9WU!#$%'()*+,-./:;=?@[

171 return int(major), int(minor) 44 ctx1YOPQRST]^Z01234X_5678V9WU!#$%'()*+,-./:;=?@[

172 

173 

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

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

176 

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

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

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

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

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

182 key). 

183 

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

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

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

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

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

189 

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

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

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

193 path the linker actually uses. 

194 """ 

195 import sys 31 ctx1mnopqrsabtuvwxyzAijklBCghcdefDE

196 

197 if use_driver: 31 ctx1mnopqrsabtuvwxyzAijklBCghcdefDE

198 return ("driver", str(_driver_version())) 5 ctx1aijkl

199 nvjitlink = sys.modules.get("cuda.bindings.nvjitlink") 26 ctx1mnopqrsbtuvwxyzABCghcdefDE

200 if nvjitlink is None: 26 ctx1mnopqrsbtuvwxyzABCghcdefDE

201 from cuda.bindings import nvjitlink as _nvjitlink 

202 

203 nvjitlink = _nvjitlink 

204 

205 return ("nvJitLink", str(nvjitlink.version())) 26 ctx1mnopqrsbtuvwxyzABCghcdefDE

206 

207 

208def _nvvm_fingerprint() -> str: (empty)

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

210 

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

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

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

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

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

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

217 

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

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

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

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

222 """ 

223 from cuda.core._program import _get_nvvm_module 10 ctx2{ } ~ abbbcbdb| ` eb

224 

225 module = _get_nvvm_module() 10 ctx2{ } ~ abbbcbdb| ` eb

226 lib_major, lib_minor = module.version() # type: ignore[attr-defined] 10 ctx2{ } ~ abbbcbdb| ` eb

227 major, minor, debug_major, debug_minor = module.ir_version() # type: ignore[attr-defined] 10 ctx2{ } ~ abbbcbdb| ` eb

228 return f"lib={lib_major}.{lib_minor};ir={major}.{minor}.{debug_major}.{debug_minor}" 10 ctx2{ } ~ abbbcbdb| ` eb

229 

230 

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

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

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

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

235_EXTERNAL_CONTENT_OPTIONS = ( (empty)

236 "include_path", 

237 "pre_include", 

238 "pch", 

239 "use_pch", 

240 "pch_dir", 

241) 

242 

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

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

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

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

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

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

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

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

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

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

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

254_SIDE_EFFECT_OPTIONS = ("create_pch", "time", "fdevice_time_trace") (empty)

255 

256 

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

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

259_BOOLEAN_OPTION_FIELDS = frozenset({"pch"}) (empty)

260 

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

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

263# is silently ignored at compile time. 

264_STR_OR_SEQUENCE_OPTION_FIELDS = frozenset({"include_path", "pre_include"}) (empty)

265 

266 

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

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

269 

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

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

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

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

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

275 trigger a cache-time guard. 

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

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

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

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

280 """ 

281 value = getattr(options, name, None) 78 ctx2Y O P Q R S T ] ^ Z 0 1 2 3 4 X _ 5 6 7 8 V 9 W U ! # $ % ' ( lbmbgbnbqbobzbtbubvbrbsbwbxbXbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) * + , - . / : ; = pbEbWb? @ [

282 if value is None: 78 ctx2Y O P Q R S T ] ^ Z 0 1 2 3 4 X _ 5 6 7 8 V 9 W U ! # $ % ' ( lbmbgbnbqbobzbtbubvbrbsbwbxbXbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) * + , - . / : ; = pbEbWb? @ [

283 return False 78 ctx2Y O P Q R S T ] ^ Z 0 1 2 3 4 X _ 5 6 7 8 V 9 W U ! # $ % ' ( lbmbgbnbqbobzbtbubvbrbsbwbxbXbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) * + , - . / : ; = pbEbWb? @ [

284 if name in _BOOLEAN_OPTION_FIELDS: 26 ctx2O P Q R S T qbobzbtbubvbrbsbwbxbObPbQbRbSbTbUbVbpbWb

285 return bool(value) 1 ctx2zb

286 if name in _STR_OR_SEQUENCE_OPTION_FIELDS: 25 ctx2O P Q R S T qbobtbubvbrbsbwbxbObPbQbRbSbTbUbVbpbWb

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

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

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

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

291 if isinstance(value, str): 11 ctx2O P Q R S T qbobrbsbpb

292 return True 3 ctx2qbrbsb

293 if isinstance(value, collections.abc.Sequence): 8 ctx2O P Q R S T obpb

294 return len(value) > 0 6 ctx2O P Q S obpb

295 return False 2 ctx1RT

296 return True 14 ctx2tbubvbwbxbObPbQbRbSbTbUbVbWb

297 

298 

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

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

301 

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

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

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

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

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

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

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

309 persistent within either. 

310 """ 

311 update(f"{label}_probe_failed", type(exc).__name__.encode()) 3 ctx1abV

312 

313 

314class _KeyBackend(abc.ABC): (empty)

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

316 

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

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

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

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

321 backend, and assembles the digest. 

322 """ 

323 

324 @abc.abstractmethod (empty)

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

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

327 

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

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

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

331 """ 

332 

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

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

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

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

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

338 if isinstance(code, (bytes, bytearray)): 3 ctx2mbnbyb

339 raise TypeError( 2 ctx2nbyb

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

341 ) 

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

343 

344 @abc.abstractmethod (empty)

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

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

347 

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

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

350 backend does not consume them. 

351 

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

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

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

355 empty tuple means the backend supports them but the caller 

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

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

358 """ 

359 return None 50 ctx2{ m n o p q r s } a F b ~ t abu v w x bby z cbdbfb| ` A i j k l B C G H I J K L M g h c d e f D ebE

360 

361 @abc.abstractmethod (empty)

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

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

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

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

366 digest is stable across processes hitting the same failure 

367 mode. 

368 """ 

369 

370 def hash_extra_payload(self, options: ProgramOptions, update: Callable[[str, bytes], None]) -> None: # noqa: B027 84 ctx1NYOPQRSTmnopq]^Z01rs234X_5678aFbVtuvw9xWU!yz#$%'(AijklBCGHIJKLMghcdef)D*+E,-./:;=?@[

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

372 ``use_libdevice``). Default: nothing. 

373 """ 

374 

375 

376class _NvrtcBackend(_KeyBackend): (empty)

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

378 # Side-effect options are NVRTC-specific: 

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

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

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

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

383 # the NVRTC path. 

384 side_effects = [name for name in _SIDE_EFFECT_OPTIONS if _option_is_set(options, name)] 78 ctx2Y O P Q R S T ] ^ Z 0 1 2 3 4 X _ 5 6 7 8 V 9 W U ! # $ % ' ( lbmbgbnbqbobzbtbubvbrbsbwbxbXbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) * + , - . / : ; = pbEbWb? @ [

385 if side_effects: 78 ctx2Y O P Q R S T ] ^ Z 0 1 2 3 4 X _ 5 6 7 8 V 9 W U ! # $ % ' ( lbmbgbnbqbobzbtbubvbrbsbwbxbXbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) * + , - . / : ; = pbEbWb? @ [

386 raise ValueError( 9 ctx2ObPbQbRbSbTbUbVbWb

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

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

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

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

391 ) 

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

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

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

395 if getattr(options, "extra_sources", None) is not None: 69 ctx2Y O P Q R S T ] ^ Z 0 1 2 3 4 X _ 5 6 7 8 V 9 W U ! # $ % ' ( lbmbgbnbqbobzbtbubvbrbsbwbxbXbhbibjbkbAbBbCbDb) * + , - . / : ; = pbEb? @ [

396 raise ValueError( 1 ctx2Xb

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

398 ) 

399 if extra_digest is None: 68 ctx2Y O P Q R S T ] ^ Z 0 1 2 3 4 X _ 5 6 7 8 V 9 W U ! # $ % ' ( lbmbgbnbqbobzbtbubvbrbsbwbxbhbibjbkbAbBbCbDb) * + , - . / : ; = pbEb? @ [

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

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

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

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

404 if external: 65 ctx2Y O P Q R S T Z 0 1 2 3 4 X 5 6 7 8 V 9 W U ! # $ % ' ( lbmbgbnbqbobzbtbubvbrbsbwbxbhbibjbkbAbBbCbDb) * + , - . / : ; = pbEb? @ [

405 raise ValueError( 11 ctx2qbobzbtbubvbrbsbwbxbpb

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

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

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

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

410 ) 

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

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

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

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

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

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

417 options_name = getattr(options, "name", None) 54 ctx2Y O P Q R S T Z 0 1 2 3 4 X 5 6 7 8 V 9 W U ! # $ % ' ( lbmbgbnbhbibjbkbAbBbCbDb) * + , - . / : ; = Eb? @ [

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

419 raise ValueError( 5 ctx2AbBbCbDbEb

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

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

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

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

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

425 ) 

426 

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

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

429 # compile-time flag surface for NVRTC. 

430 return options.as_bytes("nvrtc", target_type) 50 ctx2Y O P Q R S T ] ^ Z 0 1 2 3 4 X _ 5 6 7 8 V 9 W U ! # $ % ' ( lbgbhbibjbkb) * + , - . / : ; = ? @ [

431 

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

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

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

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

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

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

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

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

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

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

442 # can't handle. 

443 def _tag(n: Any) -> bytes: 49 ctx2Y O P Q R S T ] ^ Z 0 1 2 3 4 X _ 5 6 7 8 V 9 W U ! # $ % ' ( gbhbibjbkb) * + , - . / : ; = ? @ [

444 if isinstance(n, bytes): 7 ctx2W U gbhbibjbkb

445 return b"b:" + n 1 ctx1U

446 if isinstance(n, str): 7 ctx2W U gbhbibjbkb

447 return b"s:" + n.encode("utf-8") 7 ctx2W U gbhbibjbkb

448 if isinstance(n, bytearray): 5 ctx2gbhbibjbkb

449 raise TypeError( 1 ctx2gb

450 "name_expressions elements must be str or bytes; " 

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

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

453 ) 

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

455 

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

457 

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

459 try: 44 ctx1YOPQRST]^Z01234X_5678V9WU!#$%'()*+,-./:;=?@[

460 major, minor = _nvrtc_version() 44 ctx1YOPQRST]^Z01234X_5678V9WU!#$%'()*+,-./:;=?@[

461 except Exception as exc: 1 ctx1V

462 _hash_probe_failure(update, "nvrtc", exc) 1 ctx1V

463 return 1 ctx1V

464 update("nvrtc", f"{major}.{minor}".encode("ascii")) 44 ctx1YOPQRST]^Z01234X_5678V9WU!#$%'()*+,-./:;=?@[

465 

466 

467_DECISION_UNSET = object() (empty)

468 

469 

470class _LinkerBackend(_KeyBackend): (empty)

471 def __init__(self) -> None: (empty)

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

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

474 # ``validate``, ``option_fingerprint``, and ``hash_version_probe`` 

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

476 # otherwise mints a key whose option fingerprint and version 

477 # probe disagree on which linker is in use. 

478 self._cached_decision = _DECISION_UNSET 50 ctx2m n o p q r s a F b t u v w x y z A i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbybYbD E

479 self._cached_decision_exc: BaseException | None = None 50 ctx2m n o p q r s a F b t u v w x y z A i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbybYbD E

480 

481 def _decide_driver(self) -> bool | None: (empty)

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

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

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

485 the digest instead of a backend label). 

486 """ 

487 if self._cached_decision is _DECISION_UNSET: 49 ctx2m n o p q r s a F b t u v w x y z A i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbybD E

488 try: 49 ctx2m n o p q r s a F b t u v w x y z A i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbybD E

489 from cuda.core._linker import _decide_nvjitlink_or_driver 49 ctx2m n o p q r s a F b t u v w x y z A i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbybD E

490 

491 self._cached_decision = _decide_nvjitlink_or_driver() 49 ctx2m n o p q r s a F b t u v w x y z A i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbybD E

492 except Exception as exc: 

493 self._cached_decision = None 

494 self._cached_decision_exc = exc 

495 return self._cached_decision # type: ignore[return-value] 49 ctx2m n o p q r s a F b t u v w x y z A i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbybD E

496 

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

498 if getattr(options, "extra_sources", None) is not None: 50 ctx2m n o p q r s a F b t u v w x y z A i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbybYbD E

499 raise ValueError( 1 ctx2Yb

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

501 ) 

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

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

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

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

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

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

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

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

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

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

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

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

514 if self._decide_driver() is True: 49 ctx2m n o p q r s a F b t u v w x y z A i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbybD E

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

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

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

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

519 unsupported = [ 14 ctx2a i j k l FbGbHbIbJbKbLbMbNb

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

521 ] 

522 if unsupported: 14 ctx2a i j k l FbGbHbIbJbKbLbMbNb

523 raise ValueError( 9 ctx2FbGbHbIbJbKbLbMbNb

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

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

526 f"configuration before producing an ObjectCode." 

527 ) 

528 

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

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

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

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

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

534 # force spurious cache misses on PTX. 

535 return _linker_option_fingerprint(options, use_driver_linker=self._decide_driver()) 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

536 

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

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

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

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

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

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

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

544 use_driver = self._decide_driver() 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

545 if use_driver is None: 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

546 assert self._cached_decision_exc is not None 

547 _hash_probe_failure(update, "linker", self._cached_decision_exc) 

548 return 

549 try: 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

550 lb_name, lb_version = _linker_backend_and_version(use_driver) 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

551 except Exception as exc: 2 ctx1ab

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

553 return 2 ctx1ab

554 update("linker_backend", lb_name.encode("ascii")) 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

555 update("linker_version", lb_version.encode("ascii")) 39 ctx1mnopqrsaFbtuvwxyzAijklBCGHIJKLMghcdefDE

556 

557 

558class _NvvmBackend(_KeyBackend): (empty)

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

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

561 if isinstance(code, str): 11 ctx2{ } ~ abbbcbdbfb| ` eb

562 return code.encode("utf-8") 11 ctx2{ } ~ abbbcbdbfb| ` eb

563 if isinstance(code, (bytes, bytearray)): 1 ctx1{

564 return bytes(code) 1 ctx1{

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

566 

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

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

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

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

571 # CUDA_HOME / libdevice upgrade changes the linked output without 

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

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

574 # libdevice bytes (or simply disable use_libdevice for 

575 # caching-sensitive workflows). 

576 if extra_digest is None and getattr(options, "use_libdevice", None): 11 ctx2{ } ~ abbbcbdbfb| ` eb

577 raise ValueError( 1 ctx1`

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

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

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

581 "extra_digest that fingerprints the libdevice file you intend " 

582 "to link against, or disable use_libdevice." 

583 ) 

584 

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

586 return options.as_bytes("nvvm", target_type) 11 ctx2{ } ~ abbbcbdbfb| ` eb

587 

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

589 try: 11 ctx2{ } ~ abbbcbdbfb| ` eb

590 fp = _nvvm_fingerprint() 11 ctx2{ } ~ abbbcbdbfb| ` eb

591 except Exception as exc: 

592 _hash_probe_failure(update, "nvvm", exc) 

593 return 

594 update("nvvm", fp.encode("ascii")) 11 ctx2{ } ~ abbbcbdbfb| ` eb

595 

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

597 extra_sources = getattr(options, "extra_sources", None) 11 ctx2{ } ~ abbbcbdbfb| ` eb

598 if extra_sources: 11 ctx2{ } ~ abbbcbdbfb| ` eb

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

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

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

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

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

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

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

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

607 # caller order is the safe default. 

608 update("extra_sources_count", str(len(extra_sources)).encode("ascii")) 

609 for item in extra_sources: 

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

611 if isinstance(item, (tuple, list)) and len(item) == 2: 

612 name, src = item 

613 update("extra_source_name", str(name).encode("utf-8")) 

614 if isinstance(src, str): 

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

616 elif isinstance(src, (bytes, bytearray)): 

617 update("extra_source_code", bytes(src)) 

618 else: 

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

620 else: 

621 # Fallback for unexpected format. 

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

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

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

625 if getattr(options, "use_libdevice", None): 11 ctx2{ } ~ abbbcbdbfb| ` eb

626 update("use_libdevice", b"1") 2 ctx1|`

627 

628 

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

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

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

632_BACKENDS_BY_CODE_TYPE: dict[str, type[_KeyBackend]] = { (empty)

633 "c++": _NvrtcBackend, 

634 "ptx": _LinkerBackend, 

635 "nvvm": _NvvmBackend, 

636} 

637 

638 

639def make_program_cache_key( (empty)

640 *, 

641 code: str | bytes, 

642 code_type: str, 

643 options: ProgramOptions, 

644 target_type: str, 

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

646 extra_digest: bytes | None = None, 

647) -> bytes: 

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

649 

650 Parameters 

651 ---------- 

652 code: 

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

654 code_type: 

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

656 options: 

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

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

659 device). 

660 target_type: 

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

662 name_expressions: 

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

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

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

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

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

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

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

670 extra_digest: 

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

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

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

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

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

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

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

678 

679 Returns 

680 ------- 

681 bytes 

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

683 

684 Raises 

685 ------ 

686 ValueError 

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

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

689 effect would not occur. 

690 ValueError 

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

692 whose compilation effect depends on external file content that the 

693 key cannot otherwise observe. 

694 

695 Examples 

696 -------- 

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

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

699 which derives the key, returns the cached 

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

701 result on miss:: 

702 

703 from cuda.core import Program, ProgramOptions 

704 from cuda.core.utils import FileStreamProgramCache 

705 

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

707 options = ProgramOptions(arch="sm_80") 

708 

709 with FileStreamProgramCache() as cache: 

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

711 

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

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

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

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

716 

717 from cuda.core import ObjectCode 

718 from cuda.core.utils import FileStreamProgramCache, make_program_cache_key 

719 

720 with FileStreamProgramCache() as cache: 

721 key = make_program_cache_key( 

722 code=source, 

723 code_type="c++", 

724 options=options, 

725 target_type="cubin", 

726 extra_digest=fingerprint_headers(options.include_path), 

727 ) 

728 data = cache.get(key) 

729 if data is None: 

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

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

732 else: 

733 obj = ObjectCode.from_cubin(data) 

734 

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

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

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

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

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

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

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

742 symbol explicitly. 

743 

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

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

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

747 which NVRTC uses for relative-include resolution) require 

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

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

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

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

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

753 cases. 

754 """ 

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

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

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

758 # the lowercase form. 

759 code_type = code_type.lower() if isinstance(code_type, str) else code_type 143 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` lbA i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbmbZb0b2b1bgbnbybqbobzbtbubvbrbsbwbxbXbYbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) D * eb+ E , - . / : ; = pbEbWb? @ [

760 target_type = target_type.lower() if isinstance(target_type, str) else target_type 143 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` lbA i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbmbZb0b2b1bgbnbybqbobzbtbubvbrbsbwbxbXbYbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) D * eb+ E , - . / : ; = pbEbWb? @ [

761 check_str_enum(code_type, SourceCodeType) 143 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` lbA i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbmbZb0b2b1bgbnbybqbobzbtbubvbrbsbwbxbXbYbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) D * eb+ E , - . / : ; = pbEbWb? @ [

762 if target_type not in _VALID_TARGET_TYPES: 142 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` lbA i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbmbZb0b1bgbnbybqbobzbtbubvbrbsbwbxbXbYbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) D * eb+ E , - . / : ; = pbEbWb? @ [

763 raise ValueError( 1 ctx21b

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

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

766 ) 

767 supported_for_code = _SUPPORTED_TARGETS_BY_CODE_TYPE[code_type] 141 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` lbA i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbmbZb0bgbnbybqbobzbtbubvbrbsbwbxbXbYbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) D * eb+ E , - . / : ; = pbEbWb? @ [

768 if target_type not in supported_for_code: 141 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` lbA i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbmbZb0bgbnbybqbobzbtbubvbrbsbwbxbXbYbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) D * eb+ E , - . / : ; = pbEbWb? @ [

769 raise ValueError( 2 ctx2Zb0b

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

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

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

773 ) 

774 

775 backend = _BACKENDS_BY_CODE_TYPE[code_type]() 139 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` lbA i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbmbgbnbybqbobzbtbubvbrbsbwbxbXbYbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) D * eb+ E , - . / : ; = pbEbWb? @ [

776 backend.validate(options, target_type, extra_digest) 139 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` lbA i j k l B C G H I J K L M g h c d e f FbGbHbIbJbKbLbMbNbmbgbnbybqbobzbtbubvbrbsbwbxbXbYbhbibjbkbAbBbCbDbObPbQbRbSbTbUbVb) D * eb+ E , - . / : ; = pbEbWb? @ [

777 

778 code_bytes = backend.encode_code(code, code_type) 103 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` lbA i j k l B C G H I J K L M g h c d e f mbgbnbybhbibjbkb) D * eb+ E , - . / : ; = ? @ [

779 option_bytes = backend.option_fingerprint(options, target_type) 100 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` lbA i j k l B C G H I J K L M g h c d e f gbhbibjbkb) D * eb+ E , - . / : ; = ? @ [

780 name_tags = backend.encode_name_expressions(name_expressions) 99 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f gbhbibjbkb) D * eb+ E , - . / : ; = ? @ [

781 

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

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

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

785 hasher = hashlib.sha256(usedforsecurity=False) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

786 

787 def _update(label: str, payload: bytes) -> None: 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

788 hasher.update(label.encode("ascii")) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

789 hasher.update(len(payload).to_bytes(8, "big")) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

790 hasher.update(payload) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

791 

792 _update("schema", str(_KEY_SCHEMA_VERSION).encode("ascii")) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

793 backend.hash_version_probe(_update) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

794 _update("code_type", code_type.encode("ascii")) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

795 _update("target_type", target_type.encode("ascii")) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

796 _update("code", code_bytes) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

797 _update("option_count", str(len(option_bytes)).encode("ascii")) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

798 for opt in option_bytes: 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

799 _update("option", bytes(opt)) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

800 if name_tags is not None: 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

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

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

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

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

805 _update("names_count", str(len(name_tags)).encode("ascii")) 44 ctx1YOPQRST]^Z01234X_5678V9WU!#$%'()*+,-./:;=?@[

806 for n in name_tags: 44 ctx1YOPQRST]^Z01234X_5678V9WU!#$%'()*+,-./:;=?@[

807 _update("name", n) 2 ctx1WU

808 backend.hash_extra_payload(options, _update) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

809 

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

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

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

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

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

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

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

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

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

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

820 options_name = getattr(options, "name", None) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

821 if options_name is not None: 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

822 _update("options_name", str(options_name).encode("utf-8")) 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

823 

824 if extra_digest is not None: 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [

825 _update("extra_digest", bytes(extra_digest)) 6 ctx1]^X_|`

826 

827 return hasher.digest() 94 ctx2Y { O P Q R S T m n o p q ] ^ Z 0 1 r s 2 3 4 X _ 5 6 7 8 } a F b V ~ t abu v w 9 x W U ! bby z # $ % ' ( cbdbfb| ` A i j k l B C G H I J K L M g h c d e f ) D * eb+ E , - . / : ; = ? @ [