Coverage for cuda/bindings/_v2/nvrtc.pyx: 57.75%
445 statements
« prev ^ index » next coverage.py v7.17.2, created at 2026-11-05 01:11 +0000
1# SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
2#
3# SPDX-License-Identifier: Apache-2.0
4#
5# This code was automatically generated across versions from 12.9.0 to 13.4.1. Do not modify it directly.
6# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=19c58d80d518d7503530eddd3838b601f741b4e53ee1d46aa8ab4f8d62731795
9# <<<< PREAMBLE CONTENT >>>>
11cimport cpython as _cyb_cpython
12cimport cpython.buffer as _cyb_cpython_buffer
13from cpython.buffer cimport (
14 PyBUF_SIMPLE as _cyb_PyBUF_SIMPLE,
15 PyBuffer_Release as _cyb_PyBuffer_Release,
16 PyObject_GetBuffer as _cyb_PyObject_GetBuffer,
17 Py_buffer as _cyb_Py_buffer,
18)
19from cython cimport view as _cyb_view
20from libc.stdint cimport intptr_t
21from libc.string cimport (
22 memcmp as _cyb_memcmp,
23 memcpy as _cyb_memcpy,
24 memmove as _cyb_memmove,
25)
27from cuda.bindings._internal._fast_enum import FastEnum as _cyb_FastEnum
29import numpy as _numpy
31cdef _cyb___getbuffer(object self, _cyb_cpython.Py_buffer *buffer, void *ptr, int size, bint readonly):
32 buffer.buf = <char *>ptr
33 buffer.format = 'b'
34 buffer.internal = NULL
35 buffer.itemsize = 1
36 buffer.len = size
37 buffer.ndim = 1
38 buffer.obj = self
39 buffer.readonly = readonly
40 buffer.shape = &buffer.len
41 buffer.strides = &buffer.itemsize
42 buffer.suboffsets = NULL
44cdef _cyb_from_buffer(buffer, size, lowpp_type):
45 cdef _cyb_cpython.Py_buffer view
46 if _cyb_cpython.PyObject_GetBuffer(buffer, &view, _cyb_cpython_buffer.PyBUF_SIMPLE) != 0:
47 raise TypeError("buffer argument does not support the buffer protocol")
48 try:
49 if view.itemsize != 1:
50 raise ValueError("buffer itemsize must be 1 byte")
51 if view.len != size:
52 raise ValueError(f"buffer length must be {size} bytes")
53 return lowpp_type.from_ptr(<intptr_t><void *>view.buf, view.readonly != 0, buffer)
54 finally:
55 _cyb_cpython.PyBuffer_Release(&view)
57cdef _cyb_from_data(data, dtype_name, expected_dtype, lowpp_type):
58 # _numpy.recarray is a subclass of _numpy.ndarray, so implicitly handled here.
59 if isinstance(data, lowpp_type):
60 return data
61 if not isinstance(data, _numpy.ndarray):
62 raise TypeError("data argument must be a NumPy ndarray")
63 if data.size != 1:
64 raise ValueError("data array must have a size of 1")
65 if data.dtype != expected_dtype:
66 raise ValueError(f"data array must be of dtype {dtype_name}")
67 return lowpp_type.from_ptr(data.ctypes.data, not data.flags.writeable, data)
70# <<<< END OF PREAMBLE CONTENT >>>>
72cimport cython # NOQA
73from libcpp.vector cimport vector
75from cuda.bindings._internal._fast_enum import FastEnum as _FastEnum
78###############################################################################
79# Enum
80###############################################################################
82class Result(_cyb_FastEnum):
83 """
84 The enumerated type `nvrtcResult` defines API call result codes. NVRTC
85 API functions return `nvrtcResult` to indicate the call result.
87 See `nvrtcResult`.
88 """
89 SUCCESS = NVRTC_SUCCESS
90 ERROR_OUT_OF_MEMORY = NVRTC_ERROR_OUT_OF_MEMORY
91 ERROR_PROGRAM_CREATION_FAILURE = NVRTC_ERROR_PROGRAM_CREATION_FAILURE
92 ERROR_INVALID_INPUT = NVRTC_ERROR_INVALID_INPUT
93 ERROR_INVALID_PROGRAM = NVRTC_ERROR_INVALID_PROGRAM
94 ERROR_INVALID_OPTION = NVRTC_ERROR_INVALID_OPTION
95 ERROR_COMPILATION = NVRTC_ERROR_COMPILATION
96 ERROR_BUILTIN_OPERATION_FAILURE = NVRTC_ERROR_BUILTIN_OPERATION_FAILURE
97 ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION
98 ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION
99 ERROR_NAME_EXPRESSION_NOT_VALID = NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID
100 ERROR_INTERNAL_ERROR = NVRTC_ERROR_INTERNAL_ERROR
101 ERROR_TIME_FILE_WRITE_FAILED = NVRTC_ERROR_TIME_FILE_WRITE_FAILED
102 ERROR_NO_PCH_CREATE_ATTEMPTED = NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED
103 ERROR_PCH_CREATE_HEAP_EXHAUSTED = NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED
104 ERROR_PCH_CREATE = NVRTC_ERROR_PCH_CREATE
105 ERROR_CANCELLED = NVRTC_ERROR_CANCELLED
106 ERROR_TIME_TRACE_FILE_WRITE_FAILED = NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED
107 ERROR_BUSY = NVRTC_ERROR_BUSY
110class InstallHeadersFlag(_FastEnum):
111 """Flags for :func:`install_bundled_headers`."""
112 SKIP_IF_EXISTS = (
113 0x0,
114 "Skip installation if version marker exists and version matches. "
115 "This is the default behavior when flags=0."
116 )
117 FORCE_OVERWRITE = (
118 0x1,
119 "Clear existing directory contents before installation. "
120 "Guarantees consistency by removing any existing files first."
121 )
122 NO_WAIT = (
123 0x2,
124 "Return NVRTC_ERROR_BUSY immediately if installation is in progress "
125 "by another process, instead of waiting for the lock. "
126 "Can be combined with FORCE_OVERWRITE using bitwise OR. "
127 "Do not wait for installation to complete."
128 )
131###############################################################################
132# Error handling
133###############################################################################
136class NvrtcError(Exception):
137 def __init__(self, status):
138 self.status = status 1E
139 s = get_error_string(status) 1E
140 super(NvrtcError, self).__init__(s) 1E
142 def __reduce__(self):
143 return (type(self), (self.status,))
145class OutOfMemoryError(NvrtcError):
146 pass
147class ProgramCreationFailureError(NvrtcError):
148 pass
149class InvalidInputError(NvrtcError):
150 pass
151class InvalidProgramError(NvrtcError):
152 pass
153class InvalidOptionError(NvrtcError):
154 pass
155class CompilationError(NvrtcError):
156 pass
157class BuiltinOperationFailureError(NvrtcError):
158 pass
159class NoNameExpressionsAfterCompilationError(NvrtcError):
160 pass
161class NoLoweredNamesBeforeCompilationError(NvrtcError):
162 pass
163class NameExpressionNotValidError(NvrtcError):
164 pass
165class InternalError(NvrtcError):
166 pass
167class TimeFileWriteFailedError(NvrtcError):
168 pass
169class NoPchCreateAttemptedError(NvrtcError):
170 pass
171class PchCreateHeapExhaustedError(NvrtcError):
172 pass
173class PchCreateError(NvrtcError):
174 pass
175class CancelledError(NvrtcError):
176 pass
177class TimeTraceFileWriteFailedError(NvrtcError):
178 pass
179class BusyError(NvrtcError):
180 pass
181cdef object _nvrtc_error_factory(int status):
182 cdef object pystatus = status 1E
183 if status == 1: 1E
184 return OutOfMemoryError(pystatus)
185 elif status == 2:
186 return ProgramCreationFailureError(pystatus)
187 elif status == 3:
188 return InvalidInputError(pystatus)
189 elif status == 4:
190 return InvalidProgramError(pystatus) 1E
191 elif status == 5:
192 return InvalidOptionError(pystatus)
193 elif status == 6:
194 return CompilationError(pystatus)
195 elif status == 7:
196 return BuiltinOperationFailureError(pystatus)
197 elif status == 8:
198 return NoNameExpressionsAfterCompilationError(pystatus)
199 elif status == 9:
200 return NoLoweredNamesBeforeCompilationError(pystatus)
201 elif status == 10:
202 return NameExpressionNotValidError(pystatus)
203 elif status == 11:
204 return InternalError(pystatus)
205 elif status == 12:
206 return TimeFileWriteFailedError(pystatus)
207 elif status == 13:
208 return NoPchCreateAttemptedError(pystatus)
209 elif status == 14:
210 return PchCreateHeapExhaustedError(pystatus)
211 elif status == 15:
212 return PchCreateError(pystatus)
213 elif status == 16:
214 return CancelledError(pystatus)
215 elif status == 17:
216 return TimeTraceFileWriteFailedError(pystatus)
217 elif status == 18:
218 return BusyError(pystatus)
219 return NvrtcError(status)
222@cython.profile(False)
223cdef int check_status(int status) except 1 nogil:
224 if status != 0: 1abcdefghijkmnopqrstuvwxyzABCDFElHG
225 with gil: 1E
226 raise _nvrtc_error_factory(status) 1E
227 return status 1abcdefghijkmnopqrstuvwxyzABCDFlHG
230###############################################################################
231# POD definitions
232###############################################################################
234cdef _get_bundled_headers_info_dtype_offsets():
235 cdef nvrtcBundledHeadersInfo pod
236 return _numpy.dtype({
237 'names': ['available', 'compressed_size', 'uncompressed_size', 'cuda_version_major', 'cuda_version_minor', 'num_files'],
238 'formats': [_numpy.int32, _numpy.uint64, _numpy.uint64, _numpy.int32, _numpy.int32, _numpy.uint32],
239 'offsets': [
240 (<intptr_t>&(pod.available)) - (<intptr_t>&pod),
241 (<intptr_t>&(pod.compressedSize)) - (<intptr_t>&pod),
242 (<intptr_t>&(pod.uncompressedSize)) - (<intptr_t>&pod),
243 (<intptr_t>&(pod.cudaVersionMajor)) - (<intptr_t>&pod),
244 (<intptr_t>&(pod.cudaVersionMinor)) - (<intptr_t>&pod),
245 (<intptr_t>&(pod.numFiles)) - (<intptr_t>&pod),
246 ],
247 'itemsize': sizeof(nvrtcBundledHeadersInfo),
248 })
250bundled_headers_info_dtype = _get_bundled_headers_info_dtype_offsets()
252cdef class BundledHeadersInfo:
253 """Empty-initialize an instance of `nvrtcBundledHeadersInfo`.
256 .. seealso:: `nvrtcBundledHeadersInfo`
257 """
258 cdef:
259 nvrtcBundledHeadersInfo _data
260 nvrtcBundledHeadersInfo *_ptr
261 object _owner
262 bint _readonly
264 def __init__(self):
265 self._ptr = &self._data 1F
266 self._owner = None 1F
267 self._readonly = False 1F
269 def __repr__(self):
270 return f"<{__name__}.BundledHeadersInfo object at {hex(id(self))}>"
272 @property
273 def ptr(self):
274 """Get the pointer address to the data as Python :class:`int`."""
275 return <intptr_t>(self._ptr)
277 cdef intptr_t _get_ptr(self):
278 return <intptr_t>(self._ptr) 1F
280 def __int__(self):
281 return <intptr_t>(self._ptr)
283 def __eq__(self, other):
284 cdef BundledHeadersInfo other_
285 if not isinstance(other, BundledHeadersInfo):
286 return False
287 other_ = other
288 return (_cyb_memcmp(<void *><intptr_t>(self._ptr), <void *><intptr_t>(other_._ptr), sizeof(nvrtcBundledHeadersInfo)) == 0)
290 def __getbuffer__(self, _cyb_cpython.Py_buffer *buffer, int flags):
291 _cyb___getbuffer(self, buffer, <void *>self._ptr, sizeof(nvrtcBundledHeadersInfo), self._readonly)
293 def __releasebuffer__(self, Py_buffer *buffer):
294 pass
296 def __setitem__(self, key, val):
297 cdef _cyb_Py_buffer view
298 if key == 0:
299 if self._readonly:
300 raise ValueError("This BundledHeadersInfo instance is read-only")
301 _cyb_PyObject_GetBuffer(val, &view, _cyb_PyBUF_SIMPLE)
302 try:
303 if <size_t>view.len < sizeof(nvrtcBundledHeadersInfo):
304 raise ValueError(
305 "source buffer too small: expected at least %d bytes, got %d" % (sizeof(nvrtcBundledHeadersInfo), view.len)
306 )
307 _cyb_memmove(<void*>&self._data, view.buf, sizeof(nvrtcBundledHeadersInfo))
308 self._ptr = &self._data
309 self._owner = None
310 self._readonly = view.readonly != 0
311 finally:
312 _cyb_PyBuffer_Release(&view)
313 else:
314 setattr(self, key, val)
316 @property
317 def available(self):
318 """int: Non-zero if bundled headers are available"""
319 return self._ptr[0].available 1F
321 @available.setter
322 def available(self, val):
323 if self._readonly:
324 raise ValueError("This BundledHeadersInfo instance is read-only")
325 self._ptr[0].available = val
327 @property
328 def compressed_size(self):
329 """int: Size of compressed archive in bytes"""
330 return self._ptr[0].compressedSize 1F
332 @compressed_size.setter
333 def compressed_size(self, val):
334 if self._readonly:
335 raise ValueError("This BundledHeadersInfo instance is read-only")
336 self._ptr[0].compressedSize = val
338 @property
339 def uncompressed_size(self):
340 """int: Estimated size when extracted in bytes"""
341 return self._ptr[0].uncompressedSize 1F
343 @uncompressed_size.setter
344 def uncompressed_size(self, val):
345 if self._readonly:
346 raise ValueError("This BundledHeadersInfo instance is read-only")
347 self._ptr[0].uncompressedSize = val
349 @property
350 def cuda_version_major(self):
351 """int: CUDA major version of bundled headers"""
352 return self._ptr[0].cudaVersionMajor 1F
354 @cuda_version_major.setter
355 def cuda_version_major(self, val):
356 if self._readonly:
357 raise ValueError("This BundledHeadersInfo instance is read-only")
358 self._ptr[0].cudaVersionMajor = val
360 @property
361 def cuda_version_minor(self):
362 """int: CUDA minor version of bundled headers"""
363 return self._ptr[0].cudaVersionMinor 1F
365 @cuda_version_minor.setter
366 def cuda_version_minor(self, val):
367 if self._readonly:
368 raise ValueError("This BundledHeadersInfo instance is read-only")
369 self._ptr[0].cudaVersionMinor = val
371 @property
372 def num_files(self):
373 """int: Number of header files in the bundle"""
374 return self._ptr[0].numFiles 1F
376 @num_files.setter
377 def num_files(self, val):
378 if self._readonly:
379 raise ValueError("This BundledHeadersInfo instance is read-only")
380 self._ptr[0].numFiles = val
382 @staticmethod
383 def from_buffer(buffer):
384 """Create an BundledHeadersInfo instance with the memory from the given buffer."""
385 return _cyb_from_buffer(buffer, sizeof(nvrtcBundledHeadersInfo), BundledHeadersInfo)
387 @staticmethod
388 def from_data(data):
389 """Create an BundledHeadersInfo instance wrapping the given NumPy array.
391 Args:
392 data (_numpy.ndarray): a single-element array of dtype `bundled_headers_info_dtype` holding the data.
393 """
394 return _cyb_from_data(data, "bundled_headers_info_dtype", bundled_headers_info_dtype, BundledHeadersInfo)
396 @staticmethod
397 def from_ptr(intptr_t ptr, bint readonly=False, object owner=None):
398 """Create an BundledHeadersInfo instance wrapping the given pointer.
400 Args:
401 ptr (intptr_t): pointer address as Python :class:`int` to the data.
402 owner (object): The Python object that owns the pointer. If not provided, data will be copied.
403 readonly (bool): whether the data is read-only (to the user). default is `False`.
404 """
405 if ptr == 0:
406 raise ValueError("ptr must not be null (0)")
407 cdef BundledHeadersInfo obj = BundledHeadersInfo.__new__(BundledHeadersInfo)
408 if owner is None:
409 _cyb_memcpy(<void*>&obj._data, <void*>ptr, sizeof(nvrtcBundledHeadersInfo))
410 obj._ptr = &obj._data
411 obj._owner = None
412 else:
413 obj._ptr = <nvrtcBundledHeadersInfo *>ptr
414 obj._owner = owner
415 obj._readonly = readonly
416 return obj
419###############################################################################
420# Wrapper functions
421###############################################################################
423cpdef intptr_t create_program(bytes src, name, headers=None, include_names=None) except? 0:
424 """nvrtcCreateProgram creates an instance of nvrtcProgram with the given input parameters.
426 Args:
427 src (bytes): CUDA program source.
428 name (bytes | None): CUDA program name. ``None`` or ``""`` causes ``"default_program"``
429 to be used.
430 headers (list[bytes] | None): Sources of the headers. ``None`` is treated as an empty
431 list (no headers).
432 include_names (list[bytes] | None): Name of each header by which it can be included in
433 the CUDA program source. Must have the same length as *headers*.
435 Returns:
436 intptr_t: Opaque handle to the created program. Pass to :func:`destroy_program` when done.
438 .. seealso:: `nvrtcCreateProgram`
439 """
440 if headers is None: 1bcdefghijkmnopqrstuvwxyzABCDl
441 headers = [] 1bcdefghijkmnopqrstuvwxyzABCD
442 if include_names is None: 1bcdefghijkmnopqrstuvwxyzABCDl
443 include_names = [] 1bcdefghijkmnopqrstuvwxyzABCD
444 if len(headers) != len(include_names): 1bcdefghijkmnopqrstuvwxyzABCDl
445 raise ValueError(
446 f"headers and include_names must have the same length "
447 f"({len(headers)} != {len(include_names)})"
448 )
449 cdef int num_headers = len(headers) 1bcdefghijkmnopqrstuvwxyzABCDl
450 cdef Program prog
451 cdef const char* c_src = src 1bcdefghijkmnopqrstuvwxyzABCDl
452 cdef bytes _name
453 if name is None: 1bcdefghijkmnopqrstuvwxyzABCDl
454 _name = b""
455 else:
456 _name = name 1bcdefghijkmnopqrstuvwxyzABCDl
457 cdef const char* c_name = _name 1bcdefghijkmnopqrstuvwxyzABCDl
458 cdef vector[const char*] cy_headers = headers 1bcdefghijkmnopqrstuvwxyzABCDl
459 cdef vector[const char*] cy_include_names = include_names 1bcdefghijkmnopqrstuvwxyzABCDl
460 cdef const char** hdr_data = NULL 1bcdefghijkmnopqrstuvwxyzABCDl
461 cdef const char** inc_data = NULL 1bcdefghijkmnopqrstuvwxyzABCDl
462 if num_headers: 1bcdefghijkmnopqrstuvwxyzABCDl
463 hdr_data = cy_headers.data()
464 inc_data = cy_include_names.data()
465 with nogil: 1bcdefghijkmnopqrstuvwxyzABCDl
466 __status__ = nvrtcCreateProgram(&prog, c_src, c_name, num_headers, hdr_data, inc_data) 1bcdefghijkmnopqrstuvwxyzABCDl
467 check_status(__status__) 1bcdefghijkmnopqrstuvwxyzABCDl
468 return <intptr_t>prog 1bcdefghijkmnopqrstuvwxyzABCDl
471cpdef compile_program(intptr_t prog, options=None):
472 """nvrtcCompileProgram compiles the given program.
474 It supports compile options listed in Supported Compile Options.
476 Args:
477 prog (intptr_t): CUDA Runtime Compilation program.
478 options (list[bytes] | None): Compiler options as a list of byte strings.
479 May be ``None`` or an empty list for no options.
481 .. seealso:: `nvrtcCompileProgram`
482 """
483 if options is None: 1bcdefghijkmnopqrstuvwxyzABCDl
484 options = []
485 cdef int num_options = len(options) 1bcdefghijkmnopqrstuvwxyzABCDl
486 cdef vector[const char*] cy_options = options 1bcdefghijkmnopqrstuvwxyzABCDl
487 cdef const char** opt_data = NULL 1bcdefghijkmnopqrstuvwxyzABCDl
488 if num_options: 1bcdefghijkmnopqrstuvwxyzABCDl
489 opt_data = cy_options.data() 1bcdefghijkmnopqrstuvwxyzABCD
490 with nogil: 1bcdefghijkmnopqrstuvwxyzABCDl
491 __status__ = nvrtcCompileProgram(<Program>prog, num_options, opt_data) 1bcdefghijkmnopqrstuvwxyzABCDl
492 check_status(__status__) 1bcdefghijkmnopqrstuvwxyzABCDl
495cpdef set_flow_callback(intptr_t prog, intptr_t callback, intptr_t payload):
496 """nvrtcSetFlowCallback registers a callback that the compiler invokes during :func:`compile_program`.
498 The callback signature must be ``int callback(void *param1, void *param2)``.
499 The compiler passes *payload* as *param1* and ``NULL`` as *param2* (reserved).
500 Return 1 to cancel compilation, 0 to continue; the callback must return
501 consistently, be thread-safe, and must not call any NVRTC/libnvvm/PTX APIs.
503 Pass *callback* as a raw C function-pointer integer (e.g. via ``ctypes.cast``).
504 Pass 0 for *callback* to clear a previously registered callback.
506 Args:
507 prog (intptr_t): CUDA Runtime Compilation program.
508 callback (intptr_t): C function pointer ``int (*)(void*, void*)`` cast to an integer,
509 or 0 to clear.
510 payload (intptr_t): Opaque pointer passed to the callback as its first argument.
512 .. seealso:: `nvrtcSetFlowCallback`
513 """
514 with nogil:
515 __status__ = nvrtcSetFlowCallback(<Program>prog, <void*>callback, <void*>payload)
516 check_status(__status__)
519cpdef str get_error_string(int result):
520 """nvrtcGetErrorString is a helper function that returns a string describing the given ``nvrtcResult`` code, e.g., NVRTC_SUCCESS to ``"NVRTC_SUCCESS"``. For unrecognized enumeration values, it returns ``"NVRTC_ERROR unknown"``.
523 Args:
524 result (Result): CUDA Runtime Compilation API result code.
526 .. seealso:: `nvrtcGetErrorString`
527 """
528 cdef const char *_output_cstr_
529 cdef bytes _output_
530 with nogil: 1E
531 _output_cstr_ = nvrtcGetErrorString(<nvrtcResult>result) 1E
532 _output_ = _output_cstr_ 1E
533 return _output_.decode() 1E
536cpdef tuple version():
537 """nvrtcVersion sets the output parameters ``major`` and ``minor`` with the CUDA Runtime Compilation version number.
540 Returns:
541 A 2-tuple containing:
542 - int: CUDA Runtime Compilation major version number.
543 - int: CUDA Runtime Compilation minor version number.
545 .. seealso:: `nvrtcVersion`
546 """
547 cdef int major
548 cdef int minor
549 with nogil: 1abcdefghijk
550 __status__ = nvrtcVersion(&major, &minor) 1abcdefghijk
551 check_status(__status__) 1abcdefghijk
552 return (major, minor) 1abcdefghijk
555cpdef int get_num_supported_archs() except? -1:
556 """nvrtcGetNumSupportedArchs sets the output parameter ``num_archs`` with the number of architectures supported by NVRTC. This can then be used to pass an array to ``nvrtcGetSupportedArchs`` to get the supported architectures.
558 see ``nvrtcGetSupportedArchs``.
560 Returns:
561 int: number of supported architectures.
563 .. seealso:: `nvrtcGetNumSupportedArchs`
564 """
565 cdef int num_archs
566 with nogil:
567 __status__ = nvrtcGetNumSupportedArchs(&num_archs)
568 check_status(__status__)
569 return num_archs
572cpdef object get_supported_archs():
573 """nvrtcGetSupportedArchs populates the array passed via the output parameter ``supported_archs`` with the architectures supported by NVRTC. The array is sorted in the ascending order. The size of the array to be passed can be determined using ``nvrtcGetNumSupportedArchs``.
575 see ``nvrtcGetNumSupportedArchs``.
577 Returns:
578 int: sorted array of supported architectures.
580 .. seealso:: `nvrtcGetSupportedArchs`
581 """
582 cdef int numArchs
583 with nogil: 1H
584 __status__ = nvrtcGetNumSupportedArchs(&numArchs) 1H
585 check_status(__status__) 1H
586 cdef _cyb_view.array _supported_archs_alloc_ = _cyb_view.array(shape=(max(numArchs, 1),), itemsize=sizeof(int), format="i", mode="c") 1H
587 cdef int *supported_archs_ptr = <int *>(_supported_archs_alloc_.data) 1H
588 cdef object supported_archs = _supported_archs_alloc_[:numArchs] 1H
589 if numArchs != 0: 1H
590 with nogil: 1H
591 __status__ = nvrtcGetSupportedArchs(supported_archs_ptr) 1H
592 check_status(__status__) 1H
593 return supported_archs 1H
596cpdef destroy_program(intptr_t prog):
597 """nvrtcDestroyProgram destroys the given program.
600 Args:
601 prog (intptr_t): CUDA Runtime Compilation program.
603 .. seealso:: `nvrtcDestroyProgram`
604 """
605 cdef nvrtcProgram _prog_ = <nvrtcProgram>prog 1l
606 with nogil: 1l
607 __status__ = nvrtcDestroyProgram(&_prog_) 1l
608 check_status(__status__) 1l
611cpdef size_t get_ptx_size(intptr_t prog) except? 0:
612 """nvrtcGetPTXSize sets the value of ``ptx_size_ret`` with the size of the PTX generated by the previous compilation of ``prog`` (including the trailing ``NULL``).
615 Args:
616 prog (intptr_t): CUDA Runtime Compilation program.
618 Returns:
619 size_t: Size of the generated PTX (including the trailing
620 ``NULL``).
622 .. seealso:: `nvrtcGetPTXSize`
623 """
624 cdef size_t ptx_size_ret
625 with nogil:
626 __status__ = nvrtcGetPTXSize(<nvrtcProgram>prog, &ptx_size_ret)
627 check_status(__status__)
628 return ptx_size_ret
631cpdef bytes get_ptx(intptr_t prog):
632 """nvrtcGetPTX stores the PTX generated by the previous compilation of ``prog`` in the memory pointed by ``ptx``.
635 Args:
636 prog (intptr_t): CUDA Runtime Compilation program.
638 Returns:
639 bytes: Compiled result.
641 .. seealso:: `nvrtcGetPTX`
642 """
643 cdef size_t ptxSizeRet
644 with nogil:
645 __status__ = nvrtcGetPTXSize(<nvrtcProgram>prog, &ptxSizeRet)
646 check_status(__status__)
647 cdef bytes _ptx_ = bytes(ptxSizeRet)
648 cdef char* ptx = _ptx_
649 if ptxSizeRet != 0:
650 with nogil:
651 __status__ = nvrtcGetPTX(<nvrtcProgram>prog, ptx)
652 check_status(__status__)
653 return _ptx_
656cpdef size_t get_cubin_size(intptr_t prog) except? 0:
657 """nvrtcGetCUBINSize sets the value of ``cubin_size_ret`` with the size of the cubin generated by the previous compilation of ``prog``. The value of cubin_size_ret is set to 0 if the value specified to ``-arch`` is a virtual architecture instead of an actual architecture.
660 Args:
661 prog (intptr_t): CUDA Runtime Compilation program.
663 Returns:
664 size_t: Size of the generated cubin.
666 .. seealso:: `nvrtcGetCUBINSize`
667 """
668 cdef size_t cubin_size_ret
669 with nogil:
670 __status__ = nvrtcGetCUBINSize(<nvrtcProgram>prog, &cubin_size_ret)
671 check_status(__status__)
672 return cubin_size_ret
675cpdef bytes get_cubin(intptr_t prog):
676 """nvrtcGetCUBIN stores the cubin generated by the previous compilation of ``prog`` in the memory pointed by ``cubin``. No cubin is available if the value specified to ``-arch`` is a virtual architecture instead of an actual architecture. The cubin does not contain code for the Tile functions (``__tile__`` / ``__tile_global__``) or variables (``__tile__``); use :func:`get_tile_ir` to extract the cuda_tile IR generated for Tile code.
679 Args:
680 prog (intptr_t): CUDA Runtime Compilation program.
682 Returns:
683 bytes: Compiled and assembled result.
685 .. seealso:: `nvrtcGetCUBIN`
686 """
687 cdef size_t cubinSizeRet
688 with nogil: 1bcdefghijkqrstuv
689 __status__ = nvrtcGetCUBINSize(<nvrtcProgram>prog, &cubinSizeRet) 1bcdefghijkqrstuv
690 check_status(__status__) 1bcdefghijkqrstuv
691 cdef bytes _cubin_ = bytes(cubinSizeRet) 1bcdefghijkqrstuv
692 cdef char* cubin = _cubin_ 1bcdefghijkqrstuv
693 if cubinSizeRet != 0: 1bcdefghijkqrstuv
694 with nogil: 1bcdefghijkqrstuv
695 __status__ = nvrtcGetCUBIN(<nvrtcProgram>prog, cubin) 1bcdefghijkqrstuv
696 check_status(__status__) 1bcdefghijkqrstuv
697 return _cubin_ 1bcdefghijkqrstuv
700cpdef size_t get_ltoir_size(intptr_t prog) except? 0:
701 """nvrtcGetLTOIRSize sets the value of ``ltoir_size_ret`` with the size of the LTO IR generated by the previous compilation of ``prog``. The value of ltoir_size_ret is set to 0 if the program was not compiled with ``-dlto``.
704 Args:
705 prog (intptr_t): CUDA Runtime Compilation program.
707 Returns:
708 size_t: Size of the generated LTO IR.
710 .. seealso:: `nvrtcGetLTOIRSize`
711 """
712 cdef size_t ltoir_size_ret
713 with nogil:
714 __status__ = nvrtcGetLTOIRSize(<nvrtcProgram>prog, <oir_size_ret)
715 check_status(__status__)
716 return ltoir_size_ret
719cpdef bytes get_ltoir(intptr_t prog):
720 """nvrtcGetltoir stores the LTO IR generated by the previous compilation of ``prog`` in the memory pointed by ``ltoir``. No LTO IR is available if the program was compiled without ``-dlto``.
723 Args:
724 prog (intptr_t): CUDA Runtime Compilation program.
726 Returns:
727 bytes: Compiled result.
729 .. seealso:: `nvrtcGetLTOIR`
730 """
731 cdef size_t LTOIRSizeRet
732 with nogil: 1mnopwxyzABCD
733 __status__ = nvrtcGetLTOIRSize(<nvrtcProgram>prog, <OIRSizeRet) 1mnopwxyzABCD
734 check_status(__status__) 1mnopwxyzABCD
735 cdef bytes _ltoir_ = bytes(LTOIRSizeRet) 1mnopwxyzABCD
736 cdef char* ltoir = _ltoir_ 1mnopwxyzABCD
737 if LTOIRSizeRet != 0: 1mnopwxyzABCD
738 with nogil: 1mnopwxyzABCD
739 __status__ = nvrtcGetLTOIR(<nvrtcProgram>prog, ltoir) 1mnopwxyzABCD
740 check_status(__status__) 1mnopwxyzABCD
741 return _ltoir_ 1mnopwxyzABCD
744cpdef size_t get_optix_ir_size(intptr_t prog) except? 0:
745 """nvrtcGetOptiXIRSize sets the value of ``optixir_size_ret`` with the size of the OptiX IR generated by the previous compilation of ``prog``. The value of nvrtcGetOptiXIRSize is set to 0 if the program was compiled with options incompatible with OptiX IR generation.
748 Args:
749 prog (intptr_t): CUDA Runtime Compilation program.
751 Returns:
752 size_t: Size of the generated LTO IR.
754 .. seealso:: `nvrtcGetOptiXIRSize`
755 """
756 cdef size_t optixir_size_ret
757 with nogil:
758 __status__ = nvrtcGetOptiXIRSize(<nvrtcProgram>prog, &optixir_size_ret)
759 check_status(__status__)
760 return optixir_size_ret
763cpdef bytes get_optix_ir(intptr_t prog):
764 """nvrtcGetOptiXIR stores the OptiX IR generated by the previous compilation of ``prog`` in the memory pointed by ``optixir``. No OptiX IR is available if the program was compiled with options incompatible with OptiX IR generation.
767 Args:
768 prog (intptr_t): CUDA Runtime Compilation program.
770 Returns:
771 bytes: Optix IR Compiled result.
773 .. seealso:: `nvrtcGetOptiXIR`
774 """
775 cdef size_t optixirSizeRet
776 with nogil:
777 __status__ = nvrtcGetOptiXIRSize(<nvrtcProgram>prog, &optixirSizeRet)
778 check_status(__status__)
779 cdef bytes _optixir_ = bytes(optixirSizeRet)
780 cdef char* optixir = _optixir_
781 if optixirSizeRet != 0:
782 with nogil:
783 __status__ = nvrtcGetOptiXIR(<nvrtcProgram>prog, optixir)
784 check_status(__status__)
785 return _optixir_
788cpdef size_t get_program_log_size(intptr_t prog) except? 0:
789 """nvrtcGetProgramLogSize sets ``log_size_ret`` with the size of the log generated by the previous compilation of ``prog`` (including the trailing ``NULL``).
791 Note that compilation log may be generated with warnings and informative
792 messages, even when the compilation of ``prog`` succeeds.
794 Args:
795 prog (intptr_t): CUDA Runtime Compilation program.
797 Returns:
798 size_t: Size of the compilation log (including the trailing
799 ``NULL``).
801 .. seealso:: `nvrtcGetProgramLogSize`
802 """
803 cdef size_t log_size_ret
804 with nogil:
805 __status__ = nvrtcGetProgramLogSize(<nvrtcProgram>prog, &log_size_ret)
806 check_status(__status__)
807 return log_size_ret
810cpdef bytes get_program_log(intptr_t prog):
811 """nvrtcGetProgramLog stores the log generated by the previous compilation of ``prog`` in the memory pointed by ``log``.
814 Args:
815 prog (intptr_t): CUDA Runtime Compilation program.
817 Returns:
818 bytes: Compilation log.
820 .. seealso:: `nvrtcGetProgramLog`
821 """
822 cdef size_t logSizeRet
823 with nogil: 1bcdefghijk
824 __status__ = nvrtcGetProgramLogSize(<nvrtcProgram>prog, &logSizeRet) 1bcdefghijk
825 check_status(__status__) 1bcdefghijk
826 cdef bytes _log_ = bytes(logSizeRet) 1bcdefghijk
827 cdef char* log = _log_ 1bcdefghijk
828 if logSizeRet != 0: 1bcdefghijk
829 with nogil: 1bcdefghijk
830 __status__ = nvrtcGetProgramLog(<nvrtcProgram>prog, log) 1bcdefghijk
831 check_status(__status__) 1bcdefghijk
832 return _log_ 1bcdefghijk
835cpdef add_name_expression(intptr_t prog, name_expression):
836 """nvrtcAddNameExpression notes the given name expression denoting the address of a global function or device/__constant__ variable.
838 The identical name expression string must be provided on a subsequent call to
839 nvrtcGetLoweredName to extract the lowered name.
841 Args:
842 prog (intptr_t): CUDA Runtime Compilation program.
843 name_expression (str): constant expression denoting the address of
844 a global function or device/__constant__ variable.
846 .. seealso:: `nvrtcAddNameExpression`
847 """
848 if not isinstance(name_expression, str): 1l
849 raise TypeError("name_expression must be a Python str")
850 cdef bytes _temp_name_expression_ = (<str>name_expression).encode() 1l
851 cdef char* _name_expression_ = _temp_name_expression_ 1l
852 with nogil: 1l
853 __status__ = nvrtcAddNameExpression(<nvrtcProgram>prog, <const char*>_name_expression_) 1l
854 check_status(__status__) 1l
857cpdef bytes get_lowered_name(intptr_t prog, name_expression):
858 """nvrtcGetLoweredName extracts the lowered (mangled) name for a global function or device/__constant__ variable, and updates \*lowered_name to point to it. The memory containing the name is released when the NVRTC program is destroyed by nvrtcDestroyProgram. The identical name expression must have been previously provided to nvrtcAddNameExpression.
861 Args:
862 prog (intptr_t): CUDA Runtime Compilation program.
863 name_expression (str): constant expression denoting the address of
864 a global function or device/__constant__ variable.
866 Returns:
867 bytes: initialized by the function to point to a C string
868 containing the lowered (mangled) name corresponding to the
869 provided name expression.
871 .. seealso:: `nvrtcGetLoweredName`
872 """
873 if not isinstance(name_expression, str): 1El
874 raise TypeError("name_expression must be a Python str")
875 cdef bytes _temp_name_expression_ = (<str>name_expression).encode() 1El
876 cdef char* _name_expression_ = _temp_name_expression_ 1El
877 cdef const char* lowered_name = NULL 1El
878 with nogil: 1El
879 __status__ = nvrtcGetLoweredName(<nvrtcProgram>prog, <const char*>_name_expression_, &lowered_name) 1El
880 check_status(__status__) 1El
881 return (<bytes>lowered_name) if lowered_name != NULL else b"" 1l
884cpdef size_t get_pch_heap_size() except? 0:
885 """retrieve the current size of the PCH Heap.
888 Returns:
889 size_t: pointer to location where the size of the PCH Heap will be
890 stored.
892 .. seealso:: `nvrtcGetPCHHeapSize`
893 """
894 cdef size_t ret
895 with nogil:
896 __status__ = nvrtcGetPCHHeapSize(&ret)
897 check_status(__status__)
898 return ret
901cpdef set_pch_heap_size(size_t size):
902 """set the size of the PCH Heap.
904 The requested size may be rounded up to a platform dependent alignment (e.g.
905 page size). If the PCH Heap has already been allocated, the heap memory will be
906 freed and a new PCH Heap will be allocated.
908 Args:
909 size (size_t): requested size of the PCH Heap, in bytes.
911 .. seealso:: `nvrtcSetPCHHeapSize`
912 """
913 with nogil:
914 __status__ = nvrtcSetPCHHeapSize(size)
915 check_status(__status__)
918cpdef int get_pch_create_status(intptr_t prog) except? -1:
919 """returns the PCH creation status.
921 NVRTC_SUCCESS indicates that the PCH was successfully created.
922 NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED indicates that no PCH creation was
923 attempted, either because PCH functionality was not requested during the
924 preceding nvrtcCompileProgram call, or automatic PCH processing was requested,
925 and compiler chose not to create a PCH file.
926 NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED indicates that a PCH file could
927 potentially have been created, but the compiler ran out space in the PCH heap.
928 In this scenario, the :func:`get_pch_heap_size_required` can be used to query
929 the required heap size, the heap can be reallocated for this size with
930 :func:`set_pch_heap_size` and PCH creation may be reattempted again invoking
931 :func:`compile_program` with a new NVRTC program instance.
932 NVRTC_ERROR_PCH_CREATE indicates that an error condition prevented the PCH file
933 from being created.
935 Args:
936 prog (intptr_t): CUDA Runtime Compilation program.
938 .. seealso:: `nvrtcGetPCHCreateStatus`
939 """
940 cdef int ret
941 with nogil:
942 ret = <int>nvrtcGetPCHCreateStatus(<nvrtcProgram>prog)
943 return ret
946cpdef size_t get_pch_heap_size_required(intptr_t prog) except? 0:
947 """retrieve the required size of the PCH heap required to compile the given program.
950 Args:
951 prog (intptr_t): CUDA Runtime Compilation program.
953 Returns:
954 size_t: pointer to location where the required size of the PCH
955 Heap will be stored.
957 .. seealso:: `nvrtcGetPCHHeapSizeRequired`
958 """
959 cdef size_t size
960 with nogil:
961 __status__ = nvrtcGetPCHHeapSizeRequired(<nvrtcProgram>prog, &size)
962 check_status(__status__)
963 return size
966cpdef size_t get_tile_ir_size(intptr_t prog) except? 0:
967 """nvrtcGetTileIRSize sets the value of ``tile_ir_size_ret`` with the size of the cuda_tile IR generated by the previous compilation of ``prog``.
970 Args:
971 prog (intptr_t): CUDA Runtime Compilation program.
973 Returns:
974 size_t: Size of the generated cuda_tile IR.
976 .. seealso:: `nvrtcGetTileIRSize`
977 """
978 cdef size_t tile_ir_size_ret
979 with nogil:
980 __status__ = nvrtcGetTileIRSize(<nvrtcProgram>prog, &tile_ir_size_ret)
981 check_status(__status__)
982 return tile_ir_size_ret
985cpdef bytes get_tile_ir(intptr_t prog):
986 """nvrtcGettile_ir stores the cuda_tile IR generated by the previous compilation of ``prog`` in the memory pointed by ``tile_ir``.
989 Args:
990 prog (intptr_t): CUDA Runtime Compilation program.
992 Returns:
993 bytes: Generated cuda_tile IR.
995 .. seealso:: `nvrtcGetTileIR`
996 """
997 cdef size_t TileIRSizeRet
998 with nogil:
999 __status__ = nvrtcGetTileIRSize(<nvrtcProgram>prog, &TileIRSizeRet)
1000 check_status(__status__)
1001 cdef bytes _tile_ir_ = bytes(TileIRSizeRet)
1002 cdef char* tile_ir = _tile_ir_
1003 if TileIRSizeRet != 0:
1004 with nogil:
1005 __status__ = nvrtcGetTileIR(<nvrtcProgram>prog, tile_ir)
1006 check_status(__status__)
1007 return _tile_ir_
1010cpdef bytes install_bundled_headers(install_path, unsigned int flags):
1011 """nvrtcInstallBundledHeaders extracts CUDA headers bundled with NVRTC to a specified directory for use during compilation.
1013 NVRTC bundles a set of CUDA Toolkit headers and CUDA C++ Core Libraries (CCCL)
1014 within libnvrtc-builtins. This function extracts these headers to the specified
1015 directory, allowing NVRTC programs to compile without requiring a separate CUDA
1016 Toolkit installation. The bundled headers match those available in the CUDA
1017 Toolkit plus CCCL libraries.
1018 After extraction, users can compile kernels by passing appropriate include
1019 paths (such as "-I<install_path>" and "-I<install_path>/cccl") to
1020 nvrtcCompileProgram.
1021 A version marker file (.nvrtc_headers_version) is created in the installation
1022 directory to track the installed version.
1023 This function is thread-safe and process-safe. Concurrent calls from multiple
1024 threads or processes will be serialized using file locking. By default, the
1025 function waits for the lock; use NVRTC_INSTALL_HEADERS_NO_WAIT to return
1026 immediately with NVRTC_ERROR_BUSY if another process holds the lock.
1028 Args:
1029 install_path (str): Path where headers should be extracted (UTF-8
1030 encoded). The directory will be created if it doesn't exist.
1031 flags (unsigned int): NVRTC_INSTALL_HEADERS_* flags:.
1033 Returns:
1034 bytes: Optional pointer to receive detailed error message on
1035 failure. If non-NULL, ``*error_log`` will be set to point to a
1036 string describing the error cause. Note: subsequent API calls
1037 from the same thread may overwrite this message. May be NULL
1038 if error details are not needed.
1040 .. note::
1041 Use NVRTC_INSTALL_HEADERS_SKIP_IF_EXISTS to avoid reinstalling if headers
1042 already exist. Use NVRTC_INSTALL_HEADERS_FORCE_OVERWRITE to guarantee
1043 consistency by clearing the directory first.
1044 .. seealso:: `nvrtcInstallBundledHeaders`
1045 """
1046 if not isinstance(install_path, str): 1G
1047 raise TypeError("install_path must be a Python str")
1048 cdef bytes _temp_install_path_ = (<str>install_path).encode() 1G
1049 cdef char* _install_path_ = _temp_install_path_ 1G
1050 cdef const char* error_log = NULL 1G
1051 with nogil: 1G
1052 __status__ = nvrtcInstallBundledHeaders(<const char*>_install_path_, flags, &error_log) 1G
1053 check_status(__status__) 1G
1054 return (<bytes>error_log) if error_log != NULL else b"" 1G
1057cpdef tuple get_bundled_headers_info():
1058 """nvrtcGetBundledHeadersInfo queries information about the bundled headers without extracting them.
1060 This function allows users to determine if bundled headers are available and
1061 get size estimates before calling nvrtcInstallBundledHeaders.
1063 Returns:
1064 A 2-tuple containing:
1065 - nvrtcBundledHeadersInfo: Pointer to structure to receive header
1066 information.
1067 - bytes: Optional pointer to receive detailed error message on
1068 failure. If non-NULL, ``*error_log`` will be set to point to a
1069 string describing the error cause. Note: subsequent API calls
1070 from the same thread may overwrite this message. May be NULL
1071 if error details are not needed.
1073 .. seealso:: `nvrtcGetBundledHeadersInfo`
1074 """
1075 cdef BundledHeadersInfo info_py = BundledHeadersInfo() 1F
1076 cdef nvrtcBundledHeadersInfo *info = <nvrtcBundledHeadersInfo *><intptr_t>(info_py._get_ptr()) 1F
1077 cdef const char* error_log = NULL 1F
1078 with nogil: 1F
1079 __status__ = nvrtcGetBundledHeadersInfo(info, &error_log) 1F
1080 check_status(__status__) 1F
1081 return (info_py, (<bytes>error_log) if error_log != NULL else b"") 1F
1084cpdef bytes remove_bundled_headers(install_path):
1085 """nvrtcRemoveBundledHeaders removes previously installed bundled headers.
1087 This function removes the headers installed by nvrtcInstallBundledHeaders,
1088 helping users manage disk space. It recursively removes all files and
1089 subdirectories within the installation directory.
1091 Args:
1092 install_path (str): Path where headers were previously installed.
1093 Must be the same path used with nvrtcInstallBundledHeaders.
1095 Returns:
1096 bytes: Optional pointer to receive detailed error message on
1097 failure. If non-NULL, ``*error_log`` will be set to point to a
1098 string describing the error cause. Note: subsequent API calls
1099 from the same thread may overwrite this message. May be NULL
1100 if error details are not needed.
1102 .. note::
1103 This function will remove ALL contents of the specified directory, not just
1104 files installed by NVRTC. Use with caution.
1105 .. seealso:: `nvrtcRemoveBundledHeaders`
1106 """
1107 if not isinstance(install_path, str): 1G
1108 raise TypeError("install_path must be a Python str")
1109 cdef bytes _temp_install_path_ = (<str>install_path).encode() 1G
1110 cdef char* _install_path_ = _temp_install_path_ 1G
1111 cdef const char* error_log = NULL 1G
1112 with nogil: 1G
1113 __status__ = nvrtcRemoveBundledHeaders(<const char*>_install_path_, &error_log) 1G
1114 check_status(__status__) 1G
1115 return (<bytes>error_log) if error_log != NULL else b"" 1G
1116del _cyb_FastEnum