Coverage for cuda / core / _program.pyx: 81.73%
624 statements
« prev ^ index » next coverage.py v7.13.4, created at 2026-03-08 01:07 +0000
1# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
2#
3# SPDX-License-Identifier: Apache-2.0
4"""Compilation machinery for CUDA programs.
6This module provides :class:`Program` for compiling source code into
7:class:`~cuda.core.ObjectCode`, with :class:`ProgramOptions` for configuration.
8"""
10from __future__ import annotations
12from dataclasses import dataclass
13import threading
14from warnings import warn
16from cuda.bindings import driver, nvrtc
17from cuda.pathfinder import optional_cuda_import
19from libcpp.vector cimport vector
21from ._resource_handles cimport (
22 as_cu,
23 as_py,
24 create_nvrtc_program_handle,
25 create_nvvm_program_handle,
26)
27from cuda.bindings cimport cynvrtc, cynvvm
28from cuda.core._utils.cuda_utils cimport HANDLE_RETURN_NVRTC, HANDLE_RETURN_NVVM
29from cuda.core._device import Device
30from cuda.core._linker import Linker, LinkerHandleT, LinkerOptions
31from cuda.core._module import ObjectCode
32from cuda.core._utils.clear_error_support import assert_type
33from cuda.core._utils.cuda_utils import (
34 CUDAError,
35 _handle_boolean_option,
36 check_or_create_options,
37 get_binding_version,
38 handle_return,
39 is_nested_sequence,
40 is_sequence,
41)
43__all__ = ["Program", "ProgramOptions"]
45ProgramHandleT = nvrtc.nvrtcProgram | int | LinkerHandleT
46"""Type alias for program handle types across different backends.
48The ``int`` type covers NVVM handles, which don't have a wrapper class.
49"""
52# =============================================================================
53# Principal Class
54# =============================================================================
57cdef class Program:
58 """Represent a compilation machinery to process programs into
59 :class:`~cuda.core.ObjectCode`.
61 This object provides a unified interface to multiple underlying
62 compiler libraries. Compilation support is enabled for a wide
63 range of code types and compilation types.
65 Parameters
66 ----------
67 code : str | bytes | bytearray
68 The source code to compile. For C++ and PTX, must be a string.
69 For NVVM IR, can be str, bytes, or bytearray.
70 code_type : str
71 The type of source code. Must be one of ``"c++"``, ``"ptx"``, or ``"nvvm"``.
72 options : :class:`ProgramOptions`, optional
73 Options to customize the compilation process.
74 """
76 def __init__(self, code: str | bytes | bytearray, code_type: str, options: ProgramOptions | None = None):
77 Program_init(self, code, code_type, options) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-d.c/cb dc/da ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f +d'd%d,c)d:d.d(d3d4d5d-c6d7d8d9d!d#d
79 def close(self):
80 """Destroy this program."""
81 if self._linker: 2.c/cb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 'd3d4d5d6d7d8d9d!d#d
82 self._linker.close() 23d4d5d6d7d8d9d!d#d
83 # Reset handles - the C++ shared_ptr destructor handles cleanup
84 self._h_nvrtc.reset() 2.c/cb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 'd3d4d5d6d7d8d9d!d#d
85 self._h_nvvm.reset() 2.c/cb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 'd3d4d5d6d7d8d9d!d#d
87 def compile(
88 self, target_type: str, name_expressions: tuple | list = (), logs = None
89 ) -> ObjectCode:
90 """Compile the program to the specified target type.
92 Parameters
93 ----------
94 target_type : str
95 The compilation target. Must be one of ``"ptx"``, ``"cubin"``, or ``"ltoir"``.
96 name_expressions : tuple | list, optional
97 Sequence of name expressions to make accessible in the compiled code.
98 Used for template instantiation and similar cases.
99 logs : object, optional
100 Object with a ``write`` method to receive compilation logs.
102 Returns
103 -------
104 :class:`~cuda.core.ObjectCode`
105 The compiled object code.
106 """
107 return Program_compile(self, target_type, name_expressions, logs) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTd.c/cb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f %d,c3d4d5d-c6d7d8d9d!d#d
109 @property
110 def pch_status(self) -> str | None:
111 """PCH creation outcome from the most recent :meth:`compile` call.
113 Possible values:
115 * ``"created"`` — PCH file was written successfully.
116 * ``"not_attempted"`` — PCH creation was not attempted (e.g. the
117 compiler decided not to, or automatic PCH processing skipped it).
118 * ``"failed"`` — an error prevented PCH creation.
119 * ``None`` — PCH was not requested, the program has not been
120 compiled yet, the backend is not NVRTC (e.g. PTX or NVVM),
121 or the NVRTC bindings are too old to report status.
123 When ``create_pch`` is set in :class:`ProgramOptions` and the PCH
124 heap is too small, :meth:`compile` automatically resizes the heap
125 and retries, so ``"created"`` should be the common outcome.
127 .. note::
129 PCH is only supported for ``code_type="c++"`` programs that
130 use the NVRTC backend. For PTX and NVVM programs this property
131 always returns ``None``.
132 """
133 return self._pch_status 2:cb dc
135 @property
136 def backend(self) -> str:
137 """Return this Program instance's underlying backend."""
138 return self._backend 2.c/ca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f +d(d3d4d5d6d7d8d9d!d#d
140 @property
141 def handle(self) -> ProgramHandleT:
142 """Return the underlying handle object.
144 .. note::
146 The type of the returned object depends on the backend.
148 .. caution::
150 This handle is a Python object. To get the memory address of the underlying C
151 handle, call ``int(Program.handle)``.
152 """
153 if self._backend == "NVRTC": 2+c)d(d
154 return as_py(self._h_nvrtc) 2)d(d
155 elif self._backend == "NVVM": 2+c
156 return as_py(self._h_nvvm) # returns int (NVVM uses raw integers) 2+c
157 else:
158 return self._linker.handle
160 def __repr__(self) -> str:
161 return f"<Program backend='{self._backend}'>" 2*d$d+c
164# =============================================================================
165# Other Public Classes
166# =============================================================================
169@dataclass
170class ProgramOptions:
171 """Customizable options for configuring :class:`Program`.
173 Attributes
174 ----------
175 name : str, optional
176 Name of the program. If the compilation succeeds, the name is passed down to the generated `ObjectCode`.
177 arch : str, optional
178 Pass the SM architecture value, such as ``sm_<CC>`` (for generating CUBIN) or
179 ``compute_<CC>`` (for generating PTX). If not provided, the current device's architecture
180 will be used.
181 relocatable_device_code : bool, optional
182 Enable (disable) the generation of relocatable device code.
183 Default: False
184 extensible_whole_program : bool, optional
185 Do extensible whole program compilation of device code.
186 Default: False
187 debug : bool, optional
188 Generate debug information. If --dopt is not specified, then turns off all optimizations.
189 Default: False
190 lineinfo: bool, optional
191 Generate line-number information.
192 Default: False
193 device_code_optimize : bool, optional
194 Enable device code optimization. When specified along with '-G', enables limited debug information generation
195 for optimized device code.
196 Default: None
197 ptxas_options : Union[str, list[str]], optional
198 Specify one or more options directly to ptxas, the PTX optimizing assembler. Options should be strings.
199 For example ["-v", "-O2"].
200 Default: None
201 max_register_count : int, optional
202 Specify the maximum amount of registers that GPU functions can use.
203 Default: None
204 ftz : bool, optional
205 When performing single-precision floating-point operations, flush denormal values to zero or preserve denormal
206 values.
207 Default: False
208 prec_sqrt : bool, optional
209 For single-precision floating-point square root, use IEEE round-to-nearest mode or use a faster approximation.
210 Default: True
211 prec_div : bool, optional
212 For single-precision floating-point division and reciprocals, use IEEE round-to-nearest mode or use a faster
213 approximation.
214 Default: True
215 fma : bool, optional
216 Enables (disables) the contraction of floating-point multiplies and adds/subtracts into floating-point
217 multiply-add operations.
218 Default: True
219 use_fast_math : bool, optional
220 Make use of fast math operations.
221 Default: False
222 extra_device_vectorization : bool, optional
223 Enables more aggressive device code vectorization in the NVVM optimizer.
224 Default: False
225 link_time_optimization : bool, optional
226 Generate intermediate code for later link-time optimization.
227 Default: False
228 gen_opt_lto : bool, optional
229 Run the optimizer passes before generating the LTO IR.
230 Default: False
231 define_macro : Union[str, tuple[str, str], list[Union[str, tuple[str, str]]]], optional
232 Predefine a macro. Can be either a string, in which case that macro will be set to 1, a 2 element tuple of
233 strings, in which case the first element is defined as the second, or a list of strings or tuples.
234 Default: None
235 undefine_macro : Union[str, list[str]], optional
236 Cancel any previous definition of a macro, or list of macros.
237 Default: None
238 include_path : Union[str, list[str]], optional
239 Add the directory or directories to the list of directories to be searched for headers.
240 Default: None
241 pre_include : Union[str, list[str]], optional
242 Preinclude one or more headers during preprocessing. Can be either a string or a list of strings.
243 Default: None
244 no_source_include : bool, optional
245 Disable the default behavior of adding the directory of each input source to the include path.
246 Default: False
247 std : str, optional
248 Set language dialect to C++03, C++11, C++14, C++17 or C++20.
249 Default: c++17
250 builtin_move_forward : bool, optional
251 Provide builtin definitions of std::move and std::forward.
252 Default: True
253 builtin_initializer_list : bool, optional
254 Provide builtin definitions of std::initializer_list class and member functions.
255 Default: True
256 disable_warnings : bool, optional
257 Inhibit all warning messages.
258 Default: False
259 restrict : bool, optional
260 Programmer assertion that all kernel pointer parameters are restrict pointers.
261 Default: False
262 device_as_default_execution_space : bool, optional
263 Treat entities with no execution space annotation as __device__ entities.
264 Default: False
265 device_int128 : bool, optional
266 Allow the __int128 type in device code.
267 Default: False
268 optimization_info : str, optional
269 Provide optimization reports for the specified kind of optimization.
270 Default: None
271 no_display_error_number : bool, optional
272 Disable the display of a diagnostic number for warning messages.
273 Default: False
274 diag_error : Union[int, list[int]], optional
275 Emit error for a specified diagnostic message number or comma separated list of numbers.
276 Default: None
277 diag_suppress : Union[int, list[int]], optional
278 Suppress a specified diagnostic message number or comma separated list of numbers.
279 Default: None
280 diag_warn : Union[int, list[int]], optional
281 Emit warning for a specified diagnostic message number or comma separated lis of numbers.
282 Default: None
283 brief_diagnostics : bool, optional
284 Disable or enable showing source line and column info in a diagnostic.
285 Default: False
286 time : str, optional
287 Generate a CSV table with the time taken by each compilation phase.
288 Default: None
289 split_compile : int, optional
290 Perform compiler optimizations in parallel.
291 Default: 1
292 fdevice_syntax_only : bool, optional
293 Ends device compilation after front-end syntax checking.
294 Default: False
295 minimal : bool, optional
296 Omit certain language features to reduce compile time for small programs.
297 Default: False
298 no_cache : bool, optional
299 Disable compiler caching.
300 Default: False
301 fdevice_time_trace : str, optional
302 Generate time trace JSON for profiling compilation (NVRTC only).
303 Default: None
304 device_float128 : bool, optional
305 Allow __float128 type in device code (NVRTC only).
306 Default: False
307 frandom_seed : str, optional
308 Set random seed for randomized optimizations (NVRTC only).
309 Default: None
310 ofast_compile : str, optional
311 Fast compilation mode: "0", "min", "mid", or "max" (NVRTC only).
312 Default: None
313 pch : bool, optional
314 Use default precompiled header (NVRTC only, CUDA 12.8+).
315 Default: False
316 create_pch : str, optional
317 Create precompiled header file (NVRTC only, CUDA 12.8+).
318 Default: None
319 use_pch : str, optional
320 Use specific precompiled header file (NVRTC only, CUDA 12.8+).
321 Default: None
322 pch_dir : str, optional
323 PCH directory location (NVRTC only, CUDA 12.8+).
324 Default: None
325 pch_verbose : bool, optional
326 Verbose PCH output (NVRTC only, CUDA 12.8+).
327 Default: False
328 pch_messages : bool, optional
329 Control PCH diagnostic messages (NVRTC only, CUDA 12.8+).
330 Default: False
331 instantiate_templates_in_pch : bool, optional
332 Control template instantiation in PCH (NVRTC only, CUDA 12.8+).
333 Default: False
334 extra_sources : list of 2-tuples or tuple of 2-tuples, optional
335 Additional NVVM IR modules to compile together with the main program, specified as
336 ``((name1, source1), (name2, source2), ...)``. Each name is a string identifier used
337 in diagnostic messages. Each source can be a string (textual LLVM IR) or bytes/bytearray
338 (LLVM bitcode). Only supported for the NVVM backend.
339 Default: None
340 use_libdevice : bool, optional
341 Load NVIDIA's `libdevice <https://docs.nvidia.com/cuda/libdevice-users-guide/>`_
342 math builtins library. Only supported for the NVVM backend.
343 Default: False
344 """
346 name: str | None = "default_program"
347 arch: str | None = None
348 relocatable_device_code: bool | None = None
349 extensible_whole_program: bool | None = None
350 debug: bool | None = None
351 lineinfo: bool | None = None
352 device_code_optimize: bool | None = None
353 ptxas_options: str | list[str] | tuple[str] | None = None
354 max_register_count: int | None = None
355 ftz: bool | None = None
356 prec_sqrt: bool | None = None
357 prec_div: bool | None = None
358 fma: bool | None = None
359 use_fast_math: bool | None = None
360 extra_device_vectorization: bool | None = None
361 link_time_optimization: bool | None = None
362 gen_opt_lto: bool | None = None
363 define_macro: str | tuple[str, str] | list[str | tuple[str, str]] | tuple[str | tuple[str, str], ...] | None = None
364 undefine_macro: str | list[str] | tuple[str] | None = None
365 include_path: str | list[str] | tuple[str] | None = None
366 pre_include: str | list[str] | tuple[str] | None = None
367 no_source_include: bool | None = None
368 std: str | None = None
369 builtin_move_forward: bool | None = None
370 builtin_initializer_list: bool | None = None
371 disable_warnings: bool | None = None
372 restrict: bool | None = None
373 device_as_default_execution_space: bool | None = None
374 device_int128: bool | None = None
375 optimization_info: str | None = None
376 no_display_error_number: bool | None = None
377 diag_error: int | list[int] | tuple[int] | None = None
378 diag_suppress: int | list[int] | tuple[int] | None = None
379 diag_warn: int | list[int] | tuple[int] | None = None
380 brief_diagnostics: bool | None = None
381 time: str | None = None
382 split_compile: int | None = None
383 fdevice_syntax_only: bool | None = None
384 minimal: bool | None = None
385 no_cache: bool | None = None
386 fdevice_time_trace: str | None = None
387 device_float128: bool | None = None
388 frandom_seed: str | None = None
389 ofast_compile: str | None = None
390 pch: bool | None = None
391 create_pch: str | None = None
392 use_pch: str | None = None
393 pch_dir: str | None = None
394 pch_verbose: bool | None = None
395 pch_messages: bool | None = None
396 instantiate_templates_in_pch: bool | None = None
397 extra_sources: list[tuple[str, str | bytes | bytearray]] | tuple[tuple[str, str | bytes | bytearray], ...] | None = None
398 use_libdevice: bool | None = None # For libdevice execution
399 numba_debug: bool | None = None # Custom option for Numba debugging
401 def __post_init__(self):
402 self._name = self.name.encode() 2:cW X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-d.c/cb dc/da ~b2d+cc f +d'd%d,c)d:d.d(d@dZd1d0d-c
403 # Set arch to default if not provided
404 if self.arch is None: 2:cW X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-d.c/cb dc/da ~b2d+cc f +d'd%d,c)d:d.d(d@dZd1d0d-c
405 self.arch = f"sm_{Device().arch}" 2:c_cEcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dc/da ~b2d+cc f +d'd%d,c)d:d.d(d-c
407 def _prepare_nvrtc_options(self) -> list[bytes]:
408 return _prepare_nvrtc_options_impl(self) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
410 def _prepare_nvvm_options(self, as_bytes: bool = True) -> list[bytes] | list[str]:
411 return _prepare_nvvm_options_impl(self, as_bytes) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
413 def as_bytes(self, backend: str, target_type: str | None = None) -> list[bytes]:
414 """Convert program options to bytes format for the specified backend.
416 This method transforms the program options into a format suitable for the
417 specified compiler backend. Different backends may use different option names
418 and formats even for the same conceptual options.
420 Parameters
421 ----------
422 backend : str
423 The compiler backend to prepare options for. Must be either "nvrtc" or "nvvm".
424 target_type : str, optional
425 The compilation target type (e.g., "ptx", "cubin", "ltoir"). Some backends
426 require additional options based on the target type.
428 Returns
429 -------
430 list[bytes]
431 List of option strings encoded as bytes.
433 Raises
434 ------
435 ValueError
436 If an unknown backend is specified.
437 CUDAError
438 If an option incompatible with the specified backend is set.
440 Examples
441 --------
442 >>> options = ProgramOptions(arch="sm_80", debug=True)
443 >>> nvrtc_options = options.as_bytes("nvrtc")
444 """
445 backend = backend.lower() 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTd.c/cb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f ,c@dZd1d0d-c
446 if backend == "nvrtc": 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTd.c/cb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f ,c@dZd1d0d-c
447 return self._prepare_nvrtc_options() 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
448 elif backend == "nvvm": 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f @d1d0d
449 options = self._prepare_nvvm_options(as_bytes=True) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
450 if target_type == "ltoir" and b"-gen-lto" not in options: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 0d
451 options.append(b"-gen-lto") 2.c/c=c?c@cCc[cDc]cc
452 return options 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 0d
453 else:
454 raise ValueError(f"Unknown backend '{backend}'. Must be one of: 'nvrtc', 'nvvm'") 2@d
456 def __repr__(self):
457 return f"ProgramOptions(name={self.name!r}, arch={self.arch!r})"
460# =============================================================================
461# Private Classes and Helper Functions
462# =============================================================================
464# Module-level state for NVVM lazy loading
465_nvvm_module = None
466_nvvm_import_attempted = False
469def _get_nvvm_module():
470 """Get the NVVM module, importing it lazily with availability checks."""
471 global _nvvm_module, _nvvm_import_attempted
473 if _nvvm_import_attempted: 2:c$d;d=d?d.c/c2d[d+c=cUd?cVd@cWdCc#c[c]dXdDc$c]c^dYdc f
474 if _nvvm_module is None: 2:c$d.c/c2d[d+c=cUd?cVd@cWdCc#c[c]dXdDc$c]c^dYdc f
475 raise RuntimeError("NVVM module is not available (previous import attempt failed)")
476 return _nvvm_module 2:c$d.c/c2d[d+c=cUd?cVd@cWdCc#c[c]dXdDc$c]c^dYdc f
478 _nvvm_import_attempted = True 2:c;d=d?d
480 try: 2:c;d=d?d
481 version = get_binding_version() 2:c;d=d?d
482 if version < (12, 9): 2:c;d=d?d
483 raise RuntimeError(
484 f"NVVM bindings require cuda-bindings >= 12.9.0, but found {version[0]}.{version[1]}.x. "
485 "Please update cuda-bindings to use NVVM features."
486 )
488 nvvm = optional_cuda_import( 2:c;d=d?d
489 "cuda.bindings.nvvm",
490 probe_function=lambda module: module.version(), # probe triggers libnvvm load 2:c;d=d?d
491 )
492 if nvvm is None: 2:c;d=d
493 raise RuntimeError( 2;d=d
494 "NVVM support is unavailable: cuda.bindings.nvvm is missing or libnvvm cannot be loaded."
495 )
497 _nvvm_module = nvvm
498 return _nvvm_module
500 except RuntimeError: 2;d=d?d
501 _nvvm_module = None 2;d=d
502 raise 2;d=d
504def _find_libdevice_path():
505 """Find libdevice*.bc for NVVM compilation using cuda.pathfinder."""
506 from cuda.pathfinder import find_bitcode_lib
507 return find_bitcode_lib("device")
512cdef inline bint _process_define_macro_inner(list options, object macro) except? -1:
513 """Process a single define macro, returning True if successful."""
514 if isinstance(macro, str): 1Rhde
515 options.append(f"--define-macro={macro}") 1R
516 return True 1R
517 if isinstance(macro, tuple): 1hde
518 if len(macro) != 2 or any(not isinstance(val, str) for val in macro): 1hde
519 raise RuntimeError(f"Expected define_macro tuple[str, str], got {macro}")
520 options.append(f"--define-macro={macro[0]}={macro[1]}") 1hde
521 return True 1hde
522 return False 1de
525cdef inline void _process_define_macro(list options, object macro) except *:
526 """Process define_macro option which can be str, tuple, or list thereof."""
527 union_type = "Union[str, tuple[str, str]]" 1Rhde
528 if _process_define_macro_inner(options, macro): 1Rhde
529 return 1Rh
530 if is_nested_sequence(macro): 1de
531 for seq_macro in macro: 1de
532 if not _process_define_macro_inner(options, seq_macro): 1de
533 raise RuntimeError(f"Expected define_macro {union_type}, got {seq_macro}")
534 return 1de
535 raise RuntimeError(f"Expected define_macro {union_type}, list[{union_type}], got {macro}")
538cpdef bint _can_load_generated_ptx() except? -1:
539 """Check if the driver can load PTX generated by the current NVRTC version."""
540 driver_ver = handle_return(driver.cuDriverGetVersion()) 2EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c_dO `dP {d|db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
541 nvrtc_major, nvrtc_minor = handle_return(nvrtc.nvrtcVersion()) 2EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c_dO `dP {d|db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
542 return nvrtc_major * 1000 + nvrtc_minor * 10 <= driver_ver 2EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c_dO `dP {d|db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
545cdef inline object _translate_program_options(object options):
546 """Translate ProgramOptions to LinkerOptions for PTX compilation."""
547 return LinkerOptions( 23d4d5d6d7d8d9d!d#d
548 name=options.name, 23d4d5d6d7d8d9d!d#d
549 arch=options.arch, 23d4d5d6d7d8d9d!d#d
550 max_register_count=options.max_register_count, 23d4d5d6d7d8d9d!d#d
551 time=options.time, 23d4d5d6d7d8d9d!d#d
552 link_time_optimization=options.link_time_optimization, 23d4d5d6d7d8d9d!d#d
553 debug=options.debug, 23d4d5d6d7d8d9d!d#d
554 lineinfo=options.lineinfo, 23d4d5d6d7d8d9d!d#d
555 ftz=options.ftz, 23d4d5d6d7d8d9d!d#d
556 prec_div=options.prec_div, 23d4d5d6d7d8d9d!d#d
557 prec_sqrt=options.prec_sqrt, 23d4d5d6d7d8d9d!d#d
558 fma=options.fma, 23d4d5d6d7d8d9d!d#d
559 split_compile=options.split_compile, 23d4d5d6d7d8d9d!d#d
560 ptxas_options=options.ptxas_options, 23d4d5d6d7d8d9d!d#d
561 no_cache=options.no_cache, 23d4d5d6d7d8d9d!d#d
562 )
565cdef inline int Program_init(Program self, object code, str code_type, object options) except -1:
566 """Initialize a Program instance."""
567 cdef cynvrtc.nvrtcProgram nvrtc_prog
568 cdef cynvvm.nvvmProgram nvvm_prog
569 cdef bytes code_bytes
570 cdef const char* code_ptr
571 cdef const char* name_ptr
572 cdef size_t code_len
573 cdef bytes module_bytes
574 cdef const char* module_ptr
575 cdef size_t module_len
577 self._options = options = check_or_create_options(ProgramOptions, options, "Program options") 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-d.c/cb dc/da ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f +d'd%d,c)d:d.d(d3d4d5d-c6d7d8d9d!d#d
578 code_type = code_type.lower() 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-d.c/cb dc/da ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f +d'd%d,c)d:d.d(d3d4d5d-c6d7d8d9d!d#d
579 self._compile_lock = threading.Lock() 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-d.c/cb dc/da ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f +d'd%d,c)d:d.d(d3d4d5d-c6d7d8d9d!d#d
580 self._use_libdevice = False 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-d.c/cb dc/da ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f +d'd%d,c)d:d.d(d3d4d5d-c6d7d8d9d!d#d
581 self._libdevice_added = False 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-d.c/cb dc/da ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f +d'd%d,c)d:d.d(d3d4d5d-c6d7d8d9d!d#d
583 self._pch_status = None 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-d.c/cb dc/da ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f +d'd%d,c)d:d.d(d3d4d5d-c6d7d8d9d!d#d
585 if code_type == "c++": 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-d.c/cb dc/da ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f +d'd%d,c)d:d.d(d3d4d5d-c6d7d8d9d!d#d
586 assert_type(code, str) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dc/da ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+d'd%d,c)d:d(d-c
587 if options.extra_sources is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dc/da ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+d'd%d,c)d(d-c
588 raise ValueError("extra_sources is not supported by the NVRTC backend (C++ code_type)") 2/d
590 # TODO: support pre-loaded headers & include names
591 code_bytes = code.encode() 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+d'd%d,c)d(d-c
592 code_ptr = <const char*>code_bytes 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+d'd%d,c)d(d-c
593 name_ptr = <const char*>options._name 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+d'd%d,c)d(d-c
595 with nogil: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+d'd%d,c)d(d-c
596 HANDLE_RETURN_NVRTC(NULL, cynvrtc.nvrtcCreateProgram( 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+d'd%d,c)d(d-c
597 &nvrtc_prog, code_ptr, name_ptr, 0, NULL, NULL))
598 self._h_nvrtc = create_nvrtc_program_handle(nvrtc_prog) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+d'd%d,c)d(d-c
599 self._nvrtc_code = code_bytes 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+d'd%d,c)d(d-c
600 self._backend = "NVRTC" 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+d'd%d,c)d(d-c
601 self._linker = None 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-db dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+d'd%d,c)d(d-c
603 elif code_type == "ptx": 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f .d3d4d5d6d7d8d9d!d#d
604 assert_type(code, str) 23d4d5d6d7d8d9d!d#d
605 if options.extra_sources is not None: 23d4d5d6d7d8d9d!d#d
606 raise ValueError("extra_sources is not supported by the PTX backend.")
607 self._linker = Linker( 23d4d5d6d7d8d9d!d#d
608 ObjectCode._init(code.encode(), code_type), options=_translate_program_options(options) 23d4d5d6d7d8d9d!d#d
609 )
610 self._backend = self._linker.backend 23d4d5d6d7d8d9d!d#d
612 elif code_type == "nvvm": 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f .d
613 _get_nvvm_module() # Validate NVVM availability 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
614 if isinstance(code, str): 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
615 code = code.encode("utf-8") 2$d2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
616 elif not isinstance(code, (bytes, bytearray)): 2.c/c
617 raise TypeError("NVVM IR code must be provided as str, bytes, or bytearray")
619 code_ptr = <const char*>(<bytes>code) 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
620 name_ptr = <const char*>options._name 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
621 code_len = len(code) 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
623 with nogil: 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
624 HANDLE_RETURN_NVVM(NULL, cynvvm.nvvmCreateProgram(&nvvm_prog)) 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
625 self._h_nvvm = create_nvvm_program_handle(nvvm_prog) # RAII from here 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
626 with nogil: 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
627 HANDLE_RETURN_NVVM(nvvm_prog, cynvvm.nvvmAddModuleToProgram(nvvm_prog, code_ptr, code_len, name_ptr)) 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
629 # Add extra modules if provided
630 if options.extra_sources is not None: 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
631 if not is_sequence(options.extra_sources): 1cf
632 raise TypeError(
633 "extra_sources must be a sequence of 2-tuples: ((name1, source1), (name2, source2), ...)"
634 )
635 for i, module in enumerate(options.extra_sources): 1cf
636 if not isinstance(module, tuple) or len(module) != 2: 1cf
637 raise TypeError(
638 f"Each extra module must be a 2-tuple (name, source)"
639 f", got {type(module).__name__} at index {i}"
640 )
642 module_name, module_source = module 1cf
644 if not isinstance(module_name, str): 1cf
645 raise TypeError(f"Module name at index {i} must be a string, got {type(module_name).__name__}")
647 if isinstance(module_source, str): 1cf
648 # Textual LLVM IR - encode to UTF-8 bytes
649 module_source = module_source.encode("utf-8") 1cf
650 elif not isinstance(module_source, (bytes, bytearray)):
651 raise TypeError(
652 f"Module source at index {i} must be str (textual LLVM IR), bytes (textual LLVM IR or bitcode), "
653 f"or bytearray, got {type(module_source).__name__}"
654 )
656 if len(module_source) == 0: 1cf
657 raise ValueError(f"Module source for '{module_name}' (index {i}) cannot be empty")
659 # Add the module using NVVM API
660 module_bytes = module_source if isinstance(module_source, bytes) else bytes(module_source) 1cf
661 module_ptr = <const char*>module_bytes 1cf
662 module_len = len(module_bytes) 1cf
663 module_name_bytes = module_name.encode() 1cf
664 module_name_ptr = <const char*>module_name_bytes 1cf
666 with nogil: 1cf
667 HANDLE_RETURN_NVVM(nvvm_prog, cynvvm.nvvmAddModuleToProgram( 1cf
668 nvvm_prog, module_ptr, module_len, module_name_ptr))
670 # Store use_libdevice flag
671 if options.use_libdevice: 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
672 self._use_libdevice = True
674 self._backend = "NVVM" 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
675 self._linker = None 2$d.c/c2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
677 else:
678 supported_code_types = ("c++", "ptx", "nvvm") 2.d
679 assert code_type not in supported_code_types, f"{code_type=}" 2.d
680 if options.use_libdevice: 2.d
681 raise ValueError("use_libdevice is only supported by the NVVM backend")
682 raise RuntimeError(f"Unsupported {code_type=} ({supported_code_types=})") 2.d
684 return 0 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFd*d$dGdHdIdJdKdLdMdNdOdPdQdRd,dSdTd-d.c/cb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f +d'd%d,c)d(d3d4d5d-c6d7d8d9d!d#d
687cdef object _nvrtc_compile_and_extract(
688 cynvrtc.nvrtcProgram prog, str target_type, object name_expressions,
689 object logs, list options_list, str name,
690):
691 """Run nvrtcCompileProgram on *prog* and extract the output.
693 This is the inner compile+extract loop, factored out so the PCH
694 auto-retry path can call it on a fresh program handle.
695 """
696 cdef size_t output_size = 0 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
697 cdef size_t logsize = 0 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
698 cdef vector[const char*] options_vec
699 cdef char* data_ptr = NULL 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
700 cdef bytes name_bytes
701 cdef const char* name_ptr = NULL 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
702 cdef const char* lowered_name = NULL 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
703 cdef dict symbol_mapping = {} 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
705 # Add name expressions before compilation
706 if name_expressions: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
707 for n in name_expressions: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbi j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhbibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b
708 name_bytes = n.encode() if isinstance(n, str) else n 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbi j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhbibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b
709 name_ptr = <const char*>name_bytes 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbi j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhbibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b
710 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcAddNameExpression(prog, name_ptr)) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbi j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhbibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b
712 # Build options array
713 options_vec.resize(len(options_list)) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
714 for i in range(len(options_list)): 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
715 options_vec[i] = <const char*>(<bytes>options_list[i]) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
717 # Compile
718 with nogil: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
719 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcCompileProgram(prog, <int>options_vec.size(), options_vec.data())) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
721 # Get compiled output based on target type
722 if target_type == "ptx": 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
723 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcGetPTXSize(prog, &output_size)) 2EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9cO P b dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
724 data = bytearray(output_size) 2EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9cO P b dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
725 data_ptr = <char*>(<bytearray>data) 2EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9cO P b dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
726 with nogil: 2EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9cO P b dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
727 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcGetPTX(prog, data_ptr)) 2EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9cO P b dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
728 elif target_type == "cubin": 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ;cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTd
729 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcGetCUBINSize(prog, &output_size)) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbxbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdGdHdIdJdKdLdMdNdOdPdQdRdSdTd
730 data = bytearray(output_size) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbxbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdGdHdIdJdKdLdMdNdOdPdQdRdSdTd
731 data_ptr = <char*>(<bytearray>data) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbxbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdGdHdIdJdKdLdMdNdOdPdQdRdSdTd
732 with nogil: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbxbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdGdHdIdJdKdLdMdNdOdPdQdRdSdTd
733 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcGetCUBIN(prog, data_ptr)) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbxbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdGdHdIdJdKdLdMdNdOdPdQdRdSdTd
734 else: # ltoir
735 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcGetLTOIRSize(prog, &output_size)) 2;cU V Fd
736 data = bytearray(output_size) 2;cU V Fd
737 data_ptr = <char*>(<bytearray>data) 2;cU V Fd
738 with nogil: 2;cU V Fd
739 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcGetLTOIR(prog, data_ptr)) 2;cU V Fd
741 # Get lowered names after compilation
742 if name_expressions: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
743 for n in name_expressions: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbi j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhbibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b
744 name_bytes = n.encode() if isinstance(n, str) else n 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbi j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhbibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b
745 name_ptr = <const char*>name_bytes 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbi j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhbibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b
746 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcGetLoweredName(prog, name_ptr, &lowered_name)) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbi j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhbibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b
747 symbol_mapping[n] = lowered_name if lowered_name != NULL else None 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbi j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N ebfbgbhbibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b
749 # Get compilation log if requested
750 if logs is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
751 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcGetProgramLogSize(prog, &logsize))
752 if logsize > 1:
753 log = bytearray(logsize)
754 data_ptr = <char*>(<bytearray>log)
755 with nogil:
756 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcGetProgramLog(prog, data_ptr))
757 logs.write(log.decode("utf-8", errors="backslashreplace"))
759 return ObjectCode._init(bytes(data), target_type, symbol_mapping=symbol_mapping, name=name) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
762cdef int _nvrtc_pch_apis_cached = -1 # -1 = unchecked
764cdef bint _has_nvrtc_pch_apis():
765 global _nvrtc_pch_apis_cached
766 if _nvrtc_pch_apis_cached < 0: 1bag
767 _nvrtc_pch_apis_cached = hasattr(nvrtc, "nvrtcGetPCHCreateStatus") 1bg
768 return _nvrtc_pch_apis_cached 1bag
771cdef str _PCH_STATUS_CREATED = "created"
772cdef str _PCH_STATUS_NOT_ATTEMPTED = "not_attempted"
773cdef str _PCH_STATUS_FAILED = "failed"
776cdef str _read_pch_status(cynvrtc.nvrtcProgram prog):
777 """Query nvrtcGetPCHCreateStatus and translate to a high-level string."""
778 cdef cynvrtc.nvrtcResult err
779 with nogil: 1bag
780 err = cynvrtc.nvrtcGetPCHCreateStatus(prog) 1bag
781 if err == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 1bag
782 return _PCH_STATUS_CREATED 1bag
783 if err == cynvrtc.nvrtcResult.NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED:
784 return None # sentinel: caller should auto-retry
785 if err == cynvrtc.nvrtcResult.NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED:
786 return _PCH_STATUS_NOT_ATTEMPTED
787 return _PCH_STATUS_FAILED
790cdef object Program_compile_nvrtc(Program self, str target_type, object name_expressions, object logs):
791 """Compile using NVRTC backend and return ObjectCode."""
792 cdef cynvrtc.nvrtcProgram prog = as_cu(self._h_nvrtc) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
793 cdef list options_list = self._options.as_bytes("nvrtc", target_type) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
795 result = _nvrtc_compile_and_extract( 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
796 prog, target_type, name_expressions, logs, options_list, self._options.name, 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
797 )
799 cdef bint pch_creation_possible = self._options.create_pch or self._options.pch 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
800 if not pch_creation_possible or not _has_nvrtc_pch_apis(): 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
801 self._pch_status = None 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTddca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcAcBcmc,c-c
802 return result 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTddca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcAcBcmc,c-c
804 try: 1bag
805 status = _read_pch_status(prog) 1bag
806 except RuntimeError as e:
807 raise RuntimeError(
808 "PCH was requested but the runtime libnvrtc does not support "
809 "PCH APIs. Update to CUDA toolkit 12.8 or newer."
810 ) from e
812 if status is not None: 1bag
813 self._pch_status = status 1bag
814 return result 1bag
816 # Heap exhausted — auto-resize and retry with a fresh program
817 cdef size_t required = 0
818 with nogil:
819 HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcGetPCHHeapSizeRequired(prog, &required))
820 HANDLE_RETURN_NVRTC(NULL, cynvrtc.nvrtcSetPCHHeapSize(required))
822 cdef cynvrtc.nvrtcProgram retry_prog
823 cdef const char* code_ptr = <const char*>self._nvrtc_code
824 cdef const char* name_ptr = <const char*>self._options._name
825 with nogil:
826 HANDLE_RETURN_NVRTC(NULL, cynvrtc.nvrtcCreateProgram(
827 &retry_prog, code_ptr, name_ptr, 0, NULL, NULL))
828 self._h_nvrtc = create_nvrtc_program_handle(retry_prog)
830 result = _nvrtc_compile_and_extract(
831 retry_prog, target_type, name_expressions, logs, options_list, self._options.name,
832 )
834 status = _read_pch_status(retry_prog)
835 self._pch_status = status if status is not None else _PCH_STATUS_FAILED
836 return result
839cdef object Program_compile_nvvm(Program self, str target_type, object logs):
840 """Compile using NVVM backend and return ObjectCode."""
841 cdef cynvvm.nvvmProgram prog = as_cu(self._h_nvvm) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
842 cdef size_t output_size = 0 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
843 cdef size_t logsize = 0 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
844 cdef vector[const char*] options_vec
845 cdef char* data_ptr = NULL 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
846 cdef bytes libdevice_bytes
847 cdef const char* libdevice_ptr
848 cdef size_t libdevice_len
849 # Build options array
850 options_list = self._options.as_bytes("nvvm", target_type) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
851 options_vec.resize(len(options_list)) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
852 for i in range(len(options_list)): 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
853 options_vec[i] = <const char*>(<bytes>options_list[i]) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
855 # Serialize NVVM program mutation/use per Program instance.
856 with self._compile_lock: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
857 with nogil: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
858 HANDLE_RETURN_NVVM(prog, cynvvm.nvvmVerifyProgram(prog, <int>options_vec.size(), options_vec.data())) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
860 # Load libdevice if requested - following numba-cuda.
861 if self._use_libdevice and not self._libdevice_added: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
862 libdevice_path = _find_libdevice_path()
863 with open(libdevice_path, "rb") as f:
864 libdevice_bytes = f.read()
865 libdevice_ptr = <const char*>libdevice_bytes
866 libdevice_len = len(libdevice_bytes)
867 with nogil:
868 HANDLE_RETURN_NVVM(prog, cynvvm.nvvmLazyAddModuleToProgram(
869 prog, libdevice_ptr, libdevice_len, NULL))
870 self._libdevice_added = True
872 with nogil: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
873 HANDLE_RETURN_NVVM(prog, cynvvm.nvvmCompileProgram(prog, <int>options_vec.size(), options_vec.data())) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
875 HANDLE_RETURN_NVVM(prog, cynvvm.nvvmGetCompiledResultSize(prog, &output_size)) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
876 data = bytearray(output_size) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
877 data_ptr = <char*>(<bytearray>data) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
878 with nogil: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
879 HANDLE_RETURN_NVVM(prog, cynvvm.nvvmGetCompiledResult(prog, data_ptr)) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
881 # Get compilation log if requested
882 if logs is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
883 HANDLE_RETURN_NVVM(prog, cynvvm.nvvmGetProgramLogSize(prog, &logsize))
884 if logsize > 1:
885 log = bytearray(logsize)
886 data_ptr = <char*>(<bytearray>log)
887 with nogil:
888 HANDLE_RETURN_NVVM(prog, cynvvm.nvvmGetProgramLog(prog, data_ptr))
889 logs.write(log.decode("utf-8", errors="backslashreplace"))
891 return ObjectCode._init(bytes(data), target_type, name=self._options.name) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
893# Supported target types per backend
894cdef dict SUPPORTED_TARGETS = {
895 "NVRTC": ("ptx", "cubin", "ltoir"),
896 "NVVM": ("ptx", "ltoir"),
897 "nvJitLink": ("cubin", "ptx"),
898 "driver": ("cubin", "ptx"),
899}
902cdef object Program_compile(Program self, str target_type, object name_expressions, object logs):
903 """Compile the program to the specified target type."""
904 # Validate target_type for this backend
905 supported = SUPPORTED_TARGETS.get(self._backend) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTd.c/cb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f %d,c3d4d5d-c6d7d8d9d!d#d
906 if supported is None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTd.c/cb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f %d,c3d4d5d-c6d7d8d9d!d#d
907 raise ValueError(f'Unknown backend="{self._backend}"')
908 if target_type not in supported: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTd.c/cb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc2d+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f %d,c3d4d5d-c6d7d8d9d!d#d
909 raise ValueError( 22d%d
910 f'Unsupported target_type="{target_type}" for {self._backend} ' 22d%d
911 f'(supported: {", ".join(repr(t) for t in supported)})' 22d%d
912 )
914 if self._backend == "NVRTC": 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTd.c/cb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f ,c3d4d5d-c6d7d8d9d!d#d
915 if target_type == "ptx" and not _can_load_generated_ptx(): 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
916 warn( 2EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9cO P b dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
917 "The CUDA driver version is older than the backend version. "
918 "The generated ptx will not be loadable by the current driver.",
919 stacklevel=2,
920 category=RuntimeWarning, 2EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9cO P b dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
921 )
922 return Program_compile_nvrtc(self, target_type, name_expressions, logs) 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,c-c
924 elif self._backend == "NVVM": 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 3d4d5d6d7d8d9d!d#d
925 return Program_compile_nvvm(self, target_type, logs) 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f
927 else:
928 return self._linker.link(target_type) 23d4d5d6d7d8d9d!d#d
931cdef inline list _prepare_nvrtc_options_impl(object opts):
932 """Build NVRTC-specific compiler options."""
933 options = [f"-arch={opts.arch}"] 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
934 if opts.relocatable_device_code is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
935 options.append(f"--relocatable-device-code={_handle_boolean_option(opts.relocatable_device_code)}") 2EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9cgc
936 if opts.extensible_whole_program is not None and opts.extensible_whole_program: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
937 options.append("--extensible-whole-program") 2sc
938 if opts.debug is not None and opts.debug: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
939 options.append("--device-debug") 2fcZd
940 if opts.lineinfo is not None and opts.lineinfo: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
941 options.append("--generate-line-info") 2Zd
942 if opts.device_code_optimize is not None and opts.device_code_optimize: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
943 options.append("--dopt=on") 2fc
944 if opts.ptxas_options is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
945 opt_name = "--ptxas-options" 1ST
946 if isinstance(opts.ptxas_options, str): 1ST
947 options.append(f"{opt_name}={opts.ptxas_options}")
948 elif is_sequence(opts.ptxas_options): 1ST
949 for opt_value in opts.ptxas_options: 1ST
950 options.append(f"{opt_name}={opt_value}") 1ST
951 if opts.max_register_count is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
952 options.append(f"--maxrregcount={opts.max_register_count}") 2gc
953 if opts.ftz is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
954 options.append(f"--ftz={_handle_boolean_option(opts.ftz)}") 2ecZd
955 if opts.prec_sqrt is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
956 options.append(f"--prec-sqrt={_handle_boolean_option(opts.prec_sqrt)}") 2ec
957 if opts.prec_div is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
958 options.append(f"--prec-div={_handle_boolean_option(opts.prec_div)}") 2ec
959 if opts.fma is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
960 options.append(f"--fmad={_handle_boolean_option(opts.fma)}") 2hc
961 if opts.use_fast_math is not None and opts.use_fast_math: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
962 options.append("--use_fast_math") 2hc
963 if opts.extra_device_vectorization is not None and opts.extra_device_vectorization: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
964 options.append("--extra-device-vectorization") 2nc
965 if opts.link_time_optimization is not None and opts.link_time_optimization: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
966 options.append("--dlink-time-opt") 2;cU V oc
967 if opts.gen_opt_lto is not None and opts.gen_opt_lto: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
968 options.append("--gen-opt-lto") 2uc
969 if opts.define_macro is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
970 _process_define_macro(options, opts.define_macro) 1Rhde
971 if opts.undefine_macro is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
972 if isinstance(opts.undefine_macro, str): 2acbc
973 options.append(f"--undefine-macro={opts.undefine_macro}") 2bc
974 elif is_sequence(opts.undefine_macro): 2ac
975 for macro in opts.undefine_macro: 2ac
976 options.append(f"--undefine-macro={macro}") 2ac
977 if opts.include_path is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
978 if isinstance(opts.include_path, str): 2%c'c(c)c*ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N bc
979 options.append(f"--include-path={opts.include_path}") 2bc
980 elif is_sequence(opts.include_path): 2%c'c(c)c*ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N
981 for path in opts.include_path: 2%c'c(c)c*ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N
982 options.append(f"--include-path={path}") 2%c'c(c)c*ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N
983 if opts.pre_include is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
984 if isinstance(opts.pre_include, str):
985 options.append(f"--pre-include={opts.pre_include}")
986 elif is_sequence(opts.pre_include):
987 for header in opts.pre_include:
988 options.append(f"--pre-include={header}")
989 if opts.no_source_include is not None and opts.no_source_include: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
990 options.append("--no-source-include") 2wc
991 if opts.std is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
992 options.append(f"--std={opts.std}") 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N S T
993 if opts.builtin_move_forward is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
994 options.append(f"--builtin-move-forward={_handle_boolean_option(opts.builtin_move_forward)}") 2rc
995 if opts.builtin_initializer_list is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
996 options.append(f"--builtin-initializer-list={_handle_boolean_option(opts.builtin_initializer_list)}") 2ic
997 if opts.disable_warnings is not None and opts.disable_warnings: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
998 options.append("--disable-warnings") 2ic
999 if opts.restrict is not None and opts.restrict: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1000 options.append("--restrict") 2jc
1001 if opts.device_as_default_execution_space is not None and opts.device_as_default_execution_space: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1002 options.append("--device-as-default-execution-space") 2jc
1003 if opts.device_int128 is not None and opts.device_int128: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1004 options.append("--device-int128") 2kc
1005 if opts.device_float128 is not None and opts.device_float128: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1006 options.append("--device-float128")
1007 if opts.optimization_info is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1008 options.append(f"--optimization-info={opts.optimization_info}") 2kc
1009 if opts.no_display_error_number is not None and opts.no_display_error_number: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1010 options.append("--no-display-error-number") 2pc
1011 if opts.diag_error is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1012 if isinstance(opts.diag_error, int): 2ccQ
1013 options.append(f"--diag-error={opts.diag_error}") 2cc
1014 elif is_sequence(opts.diag_error): 1Q
1015 for error in opts.diag_error: 1Q
1016 options.append(f"--diag-error={error}") 1Q
1017 if opts.diag_suppress is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1018 if isinstance(opts.diag_suppress, int): 2ccQ
1019 options.append(f"--diag-suppress={opts.diag_suppress}") 2cc
1020 elif is_sequence(opts.diag_suppress): 1Q
1021 for suppress in opts.diag_suppress: 1Q
1022 options.append(f"--diag-suppress={suppress}") 1Q
1023 if opts.diag_warn is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1024 if isinstance(opts.diag_warn, int): 2lc
1025 options.append(f"--diag-warn={opts.diag_warn}") 2lc
1026 elif is_sequence(opts.diag_warn):
1027 for w in opts.diag_warn:
1028 options.append(f"--diag-warn={w}")
1029 if opts.brief_diagnostics is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1030 options.append(f"--brief-diagnostics={_handle_boolean_option(opts.brief_diagnostics)}") 2qc
1031 if opts.time is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1032 options.append(f"--time={opts.time}")
1033 if opts.split_compile is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1034 options.append(f"--split-compile={opts.split_compile}")
1035 if opts.fdevice_syntax_only is not None and opts.fdevice_syntax_only: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1036 options.append("--fdevice-syntax-only") 2tc
1037 if opts.minimal is not None and opts.minimal: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1038 options.append("--minimal") 2vc
1039 if opts.no_cache is not None and opts.no_cache: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1040 options.append("--no-cache") 2xc
1041 if opts.fdevice_time_trace is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1042 options.append(f"--fdevice-time-trace={opts.fdevice_time_trace}") 2~b
1043 if opts.frandom_seed is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1044 options.append(f"--frandom-seed={opts.frandom_seed}") 2yc
1045 if opts.ofast_compile is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1046 options.append(f"--Ofast-compile={opts.ofast_compile}") 2zc
1047 # PCH options (CUDA 12.8+)
1048 if opts.pch is not None and opts.pch: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1049 options.append("--pch") 1g
1050 if opts.create_pch is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1051 options.append(f"--create-pch={opts.create_pch}") 1ba
1052 if opts.use_pch is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1053 options.append(f"--use-pch={opts.use_pch}") 1a
1054 if opts.pch_dir is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1055 options.append(f"--pch-dir={opts.pch_dir}")
1056 if opts.pch_verbose is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1057 options.append(f"--pch-verbose={_handle_boolean_option(opts.pch_verbose)}") 2Ac
1058 if opts.pch_messages is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1059 options.append(f"--pch-messages={_handle_boolean_option(opts.pch_messages)}") 2Bc
1060 if opts.instantiate_templates_in_pch is not None: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1061 options.append( 2mc
1062 f"--instantiate-templates-in-pch={_handle_boolean_option(opts.instantiate_templates_in_pch)}" 2mc
1063 )
1064 if opts.numba_debug: 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1065 options.append("--numba-debug")
1066 return [o.encode() for o in options] 2W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb%c'c(c^c)c*c_ci j k l m n o p q r s t u v w x y z A B C D E F G H I J K L M N EcFcGcHcIcJcKcLcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c;c9cebfbgbhb`cibjbkblbmbnbobpbqbrbsbtbubvbwbU V O P xbybzbAbBbCbDbEbFbGbHbIbJbKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b(b)b*b+b,b-b.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b{c|c}c~cadbdcdddedfdgdhdidjdkdldmdndodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdb dca ~b!cfcgcechcncocR h d e acbcicjckcpcccQ lcS T qcrcsctcucvcwcxcyczcg AcBcmc,cZd-c
1069cdef inline object _prepare_nvvm_options_impl(object opts, bint as_bytes):
1070 """Build NVVM-specific compiler options."""
1071 options = [] 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1073 # Options supported by NVVM
1074 assert opts.arch is not None 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1075 arch = opts.arch 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1076 if arch.startswith("sm_"): 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1077 arch = f"compute_{arch[3:]}" 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1078 options.append(f"-arch={arch}") 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1079 if opts.debug is not None and opts.debug: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1080 options.append("-g") 20d
1081 if opts.device_code_optimize is False: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1082 options.append("-opt=0") 2=cUd?cVd[cXd
1083 elif opts.device_code_optimize is True: 2.c/c+c@cWdCc#cDc$c]cYdc f 1d0d
1084 options.append("-opt=3") 2Cc#cDc$c0d
1085 # NVVM uses 0/1 instead of true/false for boolean options
1086 if opts.ftz is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1087 options.append(f"-ftz={'1' if opts.ftz else '0'}") 2Cc#cDc$c0d
1088 if opts.prec_sqrt is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1089 options.append(f"-prec-sqrt={'1' if opts.prec_sqrt else '0'}") 2Cc#cDc$c
1090 if opts.prec_div is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1091 options.append(f"-prec-div={'1' if opts.prec_div else '0'}") 2Cc#cDc$c
1092 if opts.fma is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1093 options.append(f"-fma={'1' if opts.fma else '0'}") 2Cc#cDc$c
1095 # Check for unsupported options and raise error if they are set
1096 unsupported = [] 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1097 if opts.relocatable_device_code is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1098 unsupported.append("relocatable_device_code")
1099 if opts.extensible_whole_program is not None and opts.extensible_whole_program: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1100 unsupported.append("extensible_whole_program")
1101 if opts.lineinfo is not None and opts.lineinfo: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1102 unsupported.append("lineinfo") 21d
1103 if opts.ptxas_options is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1104 unsupported.append("ptxas_options")
1105 if opts.max_register_count is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1106 unsupported.append("max_register_count")
1107 if opts.use_fast_math is not None and opts.use_fast_math: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1108 unsupported.append("use_fast_math")
1109 if opts.extra_device_vectorization is not None and opts.extra_device_vectorization: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1110 unsupported.append("extra_device_vectorization")
1111 if opts.gen_opt_lto is not None and opts.gen_opt_lto: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1112 unsupported.append("gen_opt_lto")
1113 if opts.define_macro is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1114 unsupported.append("define_macro")
1115 if opts.undefine_macro is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1116 unsupported.append("undefine_macro")
1117 if opts.include_path is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1118 unsupported.append("include_path")
1119 if opts.pre_include is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1120 unsupported.append("pre_include")
1121 if opts.no_source_include is not None and opts.no_source_include: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1122 unsupported.append("no_source_include")
1123 if opts.std is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1124 unsupported.append("std")
1125 if opts.builtin_move_forward is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1126 unsupported.append("builtin_move_forward")
1127 if opts.builtin_initializer_list is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1128 unsupported.append("builtin_initializer_list")
1129 if opts.disable_warnings is not None and opts.disable_warnings: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1130 unsupported.append("disable_warnings")
1131 if opts.restrict is not None and opts.restrict: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1132 unsupported.append("restrict")
1133 if opts.device_as_default_execution_space is not None and opts.device_as_default_execution_space: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1134 unsupported.append("device_as_default_execution_space")
1135 if opts.device_int128 is not None and opts.device_int128: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1136 unsupported.append("device_int128")
1137 if opts.optimization_info is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1138 unsupported.append("optimization_info")
1139 if opts.no_display_error_number is not None and opts.no_display_error_number: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1140 unsupported.append("no_display_error_number")
1141 if opts.diag_error is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1142 unsupported.append("diag_error")
1143 if opts.diag_suppress is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1144 unsupported.append("diag_suppress")
1145 if opts.diag_warn is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1146 unsupported.append("diag_warn")
1147 if opts.brief_diagnostics is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1148 unsupported.append("brief_diagnostics")
1149 if opts.time is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1150 unsupported.append("time")
1151 if opts.split_compile is not None: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1152 unsupported.append("split_compile")
1153 if opts.fdevice_syntax_only is not None and opts.fdevice_syntax_only: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1154 unsupported.append("fdevice_syntax_only")
1155 if opts.minimal is not None and opts.minimal: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1156 unsupported.append("minimal")
1157 if opts.numba_debug is not None and opts.numba_debug: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1158 unsupported.append("numba_debug")
1159 if unsupported: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 1d0d
1160 raise CUDAError(f"The following options are not supported by NVVM backend: {', '.join(unsupported)}") 21d
1162 if as_bytes: 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 0d
1163 return [o.encode() for o in options] 2.c/c+c=cUd?cVd@cWdCc#c[cXdDc$c]cYdc f 0d
1164 else:
1165 return options