Coverage for cuda/bindings/_v2/nvrtc.pyx: 46.98%

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

2# 

3# SPDX-License-Identifier: Apache-2.0 

4# 

5# This code was automatically generated across versions from 12.9.0 to 13.3.0. Do not modify it directly. 

6# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=a36c7e54cf29166832dd9aebc1fa71cc3649498794a2846e707396419caebe10 

7  

8  

9# <<<< PREAMBLE CONTENT >>>> 

10  

11cimport cpython as _cyb_cpython 

12cimport cpython.buffer as _cyb_cpython_buffer 

13from cython cimport view as _cyb_view 

14from libc.stdint cimport intptr_t 

15from libc.stdlib cimport ( 

16 calloc as _cyb_calloc, 

17 free as _cyb_free, 

18 malloc as _cyb_malloc, 

19) 

20from libc.string cimport ( 

21 memcmp as _cyb_memcmp, 

22 memcpy as _cyb_memcpy, 

23) 

24  

25from cuda.bindings._internal._fast_enum import FastEnum as _cyb_FastEnum (empty)

26  

27import numpy as _numpy (empty)

28  

29cdef _cyb___getbuffer(object self, _cyb_cpython.Py_buffer *buffer, void *ptr, int size, bint readonly): (empty)

30 buffer.buf = <char *>ptr 

31 buffer.format = 'b' 

32 buffer.internal = NULL 

33 buffer.itemsize = 1 

34 buffer.len = size 

35 buffer.ndim = 1 

36 buffer.obj = self 

37 buffer.readonly = readonly 

38 buffer.shape = &buffer.len 

39 buffer.strides = &buffer.itemsize 

40 buffer.suboffsets = NULL 

41  

42cdef _cyb_from_buffer(buffer, size, lowpp_type): (empty)

43 cdef _cyb_cpython.Py_buffer view 

44 if _cyb_cpython.PyObject_GetBuffer(buffer, &view, _cyb_cpython_buffer.PyBUF_SIMPLE) != 0: 

45 raise TypeError("buffer argument does not support the buffer protocol") 

46 try: 

47 if view.itemsize != 1: 

48 raise ValueError("buffer itemsize must be 1 byte") 

49 if view.len != size: 

50 raise ValueError(f"buffer length must be {size} bytes") 

51 return lowpp_type.from_ptr(<intptr_t><void *>view.buf, not view.readonly, buffer) 

52 finally: 

53 _cyb_cpython.PyBuffer_Release(&view) 

54  

55cdef _cyb_from_data(data, dtype_name, expected_dtype, lowpp_type): (empty)

56 # _numpy.recarray is a subclass of _numpy.ndarray, so implicitly handled here. 

57 if isinstance(data, lowpp_type): 

58 return data 

59 if not isinstance(data, _numpy.ndarray): 

60 raise TypeError("data argument must be a NumPy ndarray") 

61 if data.size != 1: 

62 raise ValueError("data array must have a size of 1") 

63 if data.dtype != expected_dtype: 

64 raise ValueError(f"data array must be of dtype {dtype_name}") 

65 return lowpp_type.from_ptr(data.ctypes.data, not data.flags.writeable, data) 

66  

67  

68# <<<< END OF PREAMBLE CONTENT >>>> 

69  

70cimport cython # NOQA 

71from libcpp.vector cimport vector 

72  

73from cuda.bindings._internal._fast_enum import FastEnum as _FastEnum (empty)

74  

75  

76############################################################################### 

77# Enum 

78############################################################################### 

79  

80class Result(_cyb_FastEnum): (empty)

81 """ 

82 The enumerated type `nvrtcResult` defines API call result codes. NVRTC 

83 API functions return `nvrtcResult` to indicate the call result. 

84  

85 See `nvrtcResult`. 

86 """ 

87 SUCCESS = NVRTC_SUCCESS (empty)

88 ERROR_OUT_OF_MEMORY = NVRTC_ERROR_OUT_OF_MEMORY (empty)

89 ERROR_PROGRAM_CREATION_FAILURE = NVRTC_ERROR_PROGRAM_CREATION_FAILURE (empty)

90 ERROR_INVALID_INPUT = NVRTC_ERROR_INVALID_INPUT (empty)

91 ERROR_INVALID_PROGRAM = NVRTC_ERROR_INVALID_PROGRAM (empty)

92 ERROR_INVALID_OPTION = NVRTC_ERROR_INVALID_OPTION (empty)

93 ERROR_COMPILATION = NVRTC_ERROR_COMPILATION (empty)

94 ERROR_BUILTIN_OPERATION_FAILURE = NVRTC_ERROR_BUILTIN_OPERATION_FAILURE (empty)

95 ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION (empty)

96 ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION (empty)

97 ERROR_NAME_EXPRESSION_NOT_VALID = NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID (empty)

98 ERROR_INTERNAL_ERROR = NVRTC_ERROR_INTERNAL_ERROR (empty)

99 ERROR_TIME_FILE_WRITE_FAILED = NVRTC_ERROR_TIME_FILE_WRITE_FAILED (empty)

100 ERROR_NO_PCH_CREATE_ATTEMPTED = NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED (empty)

101 ERROR_PCH_CREATE_HEAP_EXHAUSTED = NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED (empty)

102 ERROR_PCH_CREATE = NVRTC_ERROR_PCH_CREATE (empty)

103 ERROR_CANCELLED = NVRTC_ERROR_CANCELLED (empty)

104 ERROR_TIME_TRACE_FILE_WRITE_FAILED = NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED (empty)

105 ERROR_BUSY = NVRTC_ERROR_BUSY (empty)

106  

107  

108class InstallHeadersFlag(_FastEnum): (empty)

109 """Flags for :func:`install_bundled_headers`.""" 

110 SKIP_IF_EXISTS = ( 

111 0x0, (empty)

112 "Skip installation if version marker exists and version matches. " 

113 "This is the default behavior when flags=0." 

114 ) 

115 FORCE_OVERWRITE = ( 

116 0x1, (empty)

117 "Clear existing directory contents before installation. " 

118 "Guarantees consistency by removing any existing files first." 

119 ) 

120 NO_WAIT = ( 

121 0x2, (empty)

122 "Return NVRTC_ERROR_BUSY immediately if installation is in progress " 

123 "by another process, instead of waiting for the lock. " 

124 "Can be combined with FORCE_OVERWRITE using bitwise OR. " 

125 "Do not wait for installation to complete." 

126 ) 

127  

128  

129############################################################################### 

130# Error handling 

131############################################################################### 

132  

133  

134class NvrtcError(Exception): (empty)

135 def __init__(self, status): (empty)

136 self.status = status 1 ctx1C

137 s = get_error_string(status) 1 ctx1C

138 super(NvrtcError, self).__init__(s) 1 ctx1C

139  

140 def __reduce__(self): (empty)

141 return (type(self), (self.status,)) 

142  

143class OutOfMemoryError(NvrtcError): (empty)

144 pass (empty)

145class ProgramCreationFailureError(NvrtcError): (empty)

146 pass (empty)

147class InvalidInputError(NvrtcError): (empty)

148 pass (empty)

149class InvalidProgramError(NvrtcError): (empty)

150 pass (empty)

151class InvalidOptionError(NvrtcError): (empty)

152 pass (empty)

153class CompilationError(NvrtcError): (empty)

154 pass (empty)

155class BuiltinOperationFailureError(NvrtcError): (empty)

156 pass (empty)

157class NoNameExpressionsAfterCompilationError(NvrtcError): (empty)

158 pass (empty)

159class NoLoweredNamesBeforeCompilationError(NvrtcError): (empty)

160 pass (empty)

161class NameExpressionNotValidError(NvrtcError): (empty)

162 pass (empty)

163class InternalErrorError(NvrtcError): (empty)

164 pass (empty)

165class TimeFileWriteFailedError(NvrtcError): (empty)

166 pass (empty)

167class NoPchCreateAttemptedError(NvrtcError): (empty)

168 pass (empty)

169class PchCreateHeapExhaustedError(NvrtcError): (empty)

170 pass (empty)

171class PchCreateError(NvrtcError): (empty)

172 pass (empty)

173class CancelledError(NvrtcError): (empty)

174 pass (empty)

175class TimeTraceFileWriteFailedError(NvrtcError): (empty)

176 pass (empty)

177class BusyError(NvrtcError): (empty)

178 pass (empty)

179cdef object _nvrtc_error_factory(int status): (empty)

180 cdef object pystatus = status 1 ctx1C

181 if status == 1: 1 ctx1C

182 return OutOfMemoryError(pystatus) 

183 elif status == 2: 

184 return ProgramCreationFailureError(pystatus) 

185 elif status == 3: 

186 return InvalidInputError(pystatus) 

187 elif status == 4: 

188 return InvalidProgramError(pystatus) 1 ctx1C

189 elif status == 5: 

190 return InvalidOptionError(pystatus) 

191 elif status == 6: 

192 return CompilationError(pystatus) 

193 elif status == 7: 

194 return BuiltinOperationFailureError(pystatus) 

195 elif status == 8: 

196 return NoNameExpressionsAfterCompilationError(pystatus) 

197 elif status == 9: 

198 return NoLoweredNamesBeforeCompilationError(pystatus) 

199 elif status == 10: 

200 return NameExpressionNotValidError(pystatus) 

201 elif status == 11: 

202 return InternalErrorError(pystatus) 

203 elif status == 12: 

204 return TimeFileWriteFailedError(pystatus) 

205 elif status == 13: 

206 return NoPchCreateAttemptedError(pystatus) 

207 elif status == 14: 

208 return PchCreateHeapExhaustedError(pystatus) 

209 elif status == 15: 

210 return PchCreateError(pystatus) 

211 elif status == 16: 

212 return CancelledError(pystatus) 

213 elif status == 17: 

214 return TimeTraceFileWriteFailedError(pystatus) 

215 elif status == 18: 

216 return BusyError(pystatus) 

217 return NvrtcError(status) 

218  

219InternalError = InternalErrorError (empty)

220  

221  

222@cython.profile(False) (empty)

223cdef int check_status(int status) except 1 nogil: 

224 if status != 0: 30 ctx1abcdefghijklmnopqrstuvwxyzABCD

225 with gil: 1 ctx1C

226 raise _nvrtc_error_factory(status) 1 ctx1C

227 return status 29 ctx1abcdefghijklmnopqrstuvwxyzABD

228  

229  

230############################################################################### 

231# POD definitions 

232############################################################################### 

233  

234cdef _get_bundled_headers_info_dtype_offsets(): (empty)

235 cdef nvrtcBundledHeadersInfo pod 

236 return _numpy.dtype({ (empty)

237 'names': ['available', 'compressed_size', 'uncompressed_size', 'cuda_version_major', 'cuda_version_minor', 'num_files'], (empty)

238 'formats': [_numpy.int32, _numpy.uint64, _numpy.uint64, _numpy.int32, _numpy.int32, _numpy.uint32], (empty)

239 'offsets': [ (empty)

240 (<intptr_t>&(pod.available)) - (<intptr_t>&pod), (empty)

241 (<intptr_t>&(pod.compressedSize)) - (<intptr_t>&pod), (empty)

242 (<intptr_t>&(pod.uncompressedSize)) - (<intptr_t>&pod), (empty)

243 (<intptr_t>&(pod.cudaVersionMajor)) - (<intptr_t>&pod), (empty)

244 (<intptr_t>&(pod.cudaVersionMinor)) - (<intptr_t>&pod), (empty)

245 (<intptr_t>&(pod.numFiles)) - (<intptr_t>&pod), (empty)

246 ], 

247 'itemsize': sizeof(nvrtcBundledHeadersInfo), (empty)

248 }) 

249  

250bundled_headers_info_dtype = _get_bundled_headers_info_dtype_offsets() (empty)

251  

252cdef class BundledHeadersInfo: 

253 """Empty-initialize an instance of `nvrtcBundledHeadersInfo`. 

254  

255  

256 .. seealso:: `nvrtcBundledHeadersInfo` 

257 """ 

258 cdef: 

259 nvrtcBundledHeadersInfo *_ptr 

260 object _owner 

261 bint _owned 

262 bint _readonly 

263  

264 def __init__(self): 

265 self._ptr = <nvrtcBundledHeadersInfo *>_cyb_calloc(1, sizeof(nvrtcBundledHeadersInfo)) 

266 if self._ptr == NULL: 

267 raise MemoryError("Error allocating BundledHeadersInfo") 

268 self._owner = None 

269 self._owned = True 

270 self._readonly = False 

271  

272 def __dealloc__(self): 

273 cdef nvrtcBundledHeadersInfo *ptr 

274 if self._owned and self._ptr != NULL: 

275 ptr = self._ptr 

276 self._ptr = NULL 

277 _cyb_free(ptr) 

278  

279 def __repr__(self): 

280 return f"<{__name__}.BundledHeadersInfo object at {hex(id(self))}>" 

281  

282 @property 

283 def ptr(self): 

284 """Get the pointer address to the data as Python :class:`int`.""" 

285 return <intptr_t>(self._ptr) 

286  

287 cdef intptr_t _get_ptr(self): (empty)

288 return <intptr_t>(self._ptr) 

289  

290 def __int__(self): 

291 return <intptr_t>(self._ptr) 

292  

293 def __eq__(self, other): 

294 cdef BundledHeadersInfo other_ 

295 if not isinstance(other, BundledHeadersInfo): 

296 return False 

297 other_ = other 

298 return (_cyb_memcmp(<void *><intptr_t>(self._ptr), <void *><intptr_t>(other_._ptr), sizeof(nvrtcBundledHeadersInfo)) == 0) 

299  

300 def __getbuffer__(self, _cyb_cpython.Py_buffer *buffer, int flags): 

301 _cyb___getbuffer(self, buffer, <void *>self._ptr, sizeof(nvrtcBundledHeadersInfo), self._readonly) 

302  

303 def __releasebuffer__(self, Py_buffer *buffer): 

304 pass 

305  

306 def __setitem__(self, key, val): 

307 if key == 0 and isinstance(val, _numpy.ndarray): 

308 self._ptr = <nvrtcBundledHeadersInfo *>_cyb_malloc(sizeof(nvrtcBundledHeadersInfo)) 

309 if self._ptr == NULL: 

310 raise MemoryError("Error allocating BundledHeadersInfo") 

311 _cyb_memcpy(<void*>self._ptr, <void*><intptr_t>val.ctypes.data, sizeof(nvrtcBundledHeadersInfo)) 

312 self._owner = None 

313 self._owned = True 

314 self._readonly = not val.flags.writeable 

315 else: 

316 setattr(self, key, val) 

317  

318 @property 

319 def available(self): 

320 """int: Non-zero if bundled headers are available""" 

321 return self._ptr[0].available 

322  

323 @available.setter 

324 def available(self, val): 

325 if self._readonly: 

326 raise ValueError("This BundledHeadersInfo instance is read-only") 

327 self._ptr[0].available = val 

328  

329 @property 

330 def compressed_size(self): 

331 """int: Size of compressed archive in bytes""" 

332 return self._ptr[0].compressedSize 

333  

334 @compressed_size.setter 

335 def compressed_size(self, val): 

336 if self._readonly: 

337 raise ValueError("This BundledHeadersInfo instance is read-only") 

338 self._ptr[0].compressedSize = val 

339  

340 @property 

341 def uncompressed_size(self): 

342 """int: Estimated size when extracted in bytes""" 

343 return self._ptr[0].uncompressedSize 

344  

345 @uncompressed_size.setter 

346 def uncompressed_size(self, val): 

347 if self._readonly: 

348 raise ValueError("This BundledHeadersInfo instance is read-only") 

349 self._ptr[0].uncompressedSize = val 

350  

351 @property 

352 def cuda_version_major(self): 

353 """int: CUDA major version of bundled headers""" 

354 return self._ptr[0].cudaVersionMajor 

355  

356 @cuda_version_major.setter 

357 def cuda_version_major(self, val): 

358 if self._readonly: 

359 raise ValueError("This BundledHeadersInfo instance is read-only") 

360 self._ptr[0].cudaVersionMajor = val 

361  

362 @property 

363 def cuda_version_minor(self): 

364 """int: CUDA minor version of bundled headers""" 

365 return self._ptr[0].cudaVersionMinor 

366  

367 @cuda_version_minor.setter 

368 def cuda_version_minor(self, val): 

369 if self._readonly: 

370 raise ValueError("This BundledHeadersInfo instance is read-only") 

371 self._ptr[0].cudaVersionMinor = val 

372  

373 @property 

374 def num_files(self): 

375 """int: Number of header files in the bundle""" 

376 return self._ptr[0].numFiles 

377  

378 @num_files.setter 

379 def num_files(self, val): 

380 if self._readonly: 

381 raise ValueError("This BundledHeadersInfo instance is read-only") 

382 self._ptr[0].numFiles = val 

383  

384 @staticmethod (empty)

385 def from_buffer(buffer): 

386 """Create an BundledHeadersInfo instance with the memory from the given buffer.""" 

387 return _cyb_from_buffer(buffer, sizeof(nvrtcBundledHeadersInfo), BundledHeadersInfo) 

388  

389 @staticmethod (empty)

390 def from_data(data): 

391 """Create an BundledHeadersInfo instance wrapping the given NumPy array. 

392  

393 Args: 

394 data (_numpy.ndarray): a single-element array of dtype `bundled_headers_info_dtype` holding the data. 

395 """ 

396 return _cyb_from_data(data, "bundled_headers_info_dtype", bundled_headers_info_dtype, BundledHeadersInfo) 

397  

398 @staticmethod (empty)

399 def from_ptr(intptr_t ptr, bint readonly=False, object owner=None): 

400 """Create an BundledHeadersInfo instance wrapping the given pointer. 

401  

402 Args: 

403 ptr (intptr_t): pointer address as Python :class:`int` to the data. 

404 owner (object): The Python object that owns the pointer. If not provided, data will be copied. 

405 readonly (bool): whether the data is read-only (to the user). default is `False`. 

406 """ 

407 if ptr == 0: 

408 raise ValueError("ptr must not be null (0)") 

409 cdef BundledHeadersInfo obj = BundledHeadersInfo.__new__(BundledHeadersInfo) 

410 if owner is None: 

411 obj._ptr = <nvrtcBundledHeadersInfo *>_cyb_malloc(sizeof(nvrtcBundledHeadersInfo)) 

412 if obj._ptr == NULL: 

413 raise MemoryError("Error allocating BundledHeadersInfo") 

414 _cyb_memcpy(<void*>(obj._ptr), <void*>ptr, sizeof(nvrtcBundledHeadersInfo)) 

415 obj._owner = None 

416 obj._owned = True 

417 else: 

418 obj._ptr = <nvrtcBundledHeadersInfo *>ptr 

419 obj._owner = owner 

420 obj._owned = False 

421 obj._readonly = readonly 

422 return obj 

423  

424  

425############################################################################### 

426# Wrapper functions 

427############################################################################### 

428  

429cpdef intptr_t create_program(bytes src, name, headers=None, include_names=None) except? 0: (empty)

430 """nvrtcCreateProgram creates an instance of nvrtcProgram with the given input parameters. 

431  

432 Args: 

433 src (bytes): CUDA program source. 

434 name (bytes | None): CUDA program name. ``None`` or ``""`` causes ``"default_program"`` 

435 to be used. 

436 headers (list[bytes] | None): Sources of the headers. ``None`` is treated as an empty 

437 list (no headers). 

438 include_names (list[bytes] | None): Name of each header by which it can be included in 

439 the CUDA program source. Must have the same length as *headers*. 

440  

441 Returns: 

442 intptr_t: Opaque handle to the created program. Pass to :func:`destroy_program` when done. 

443  

444 .. seealso:: `nvrtcCreateProgram` 

445 """ 

446 if headers is None: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

447 headers = [] 27 ctx1bcdefghijklmnopqrstuvwxyzAB

448 if include_names is None: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

449 include_names = [] 27 ctx1bcdefghijklmnopqrstuvwxyzAB

450 if len(headers) != len(include_names): 27 ctx1bcdefghijklmnopqrstuvwxyzAB

451 raise ValueError( 

452 f"headers and include_names must have the same length " 

453 f"({len(headers)} != {len(include_names)})" 

454 ) 

455 cdef int num_headers = len(headers) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

456 cdef Program prog 

457 cdef const char* c_src = src 27 ctx1bcdefghijklmnopqrstuvwxyzAB

458 cdef bytes _name = b"" if name is None else name 27 ctx1bcdefghijklmnopqrstuvwxyzAB

459 cdef const char* c_name = _name 27 ctx1bcdefghijklmnopqrstuvwxyzAB

460 cdef vector[const char*] cy_headers = headers 27 ctx1bcdefghijklmnopqrstuvwxyzAB

461 cdef vector[const char*] cy_include_names = include_names 27 ctx1bcdefghijklmnopqrstuvwxyzAB

462 cdef const char** hdr_data = NULL 27 ctx1bcdefghijklmnopqrstuvwxyzAB

463 cdef const char** inc_data = NULL 27 ctx1bcdefghijklmnopqrstuvwxyzAB

464 if num_headers: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

465 hdr_data = cy_headers.data() 

466 inc_data = cy_include_names.data() 

467 with nogil: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

468 __status__ = nvrtcCreateProgram(&prog, c_src, c_name, num_headers, hdr_data, inc_data) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

469 check_status(__status__) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

470 return <intptr_t>prog 27 ctx1bcdefghijklmnopqrstuvwxyzAB

471  

472  

473cpdef compile_program(intptr_t prog, options=None): (empty)

474 """nvrtcCompileProgram compiles the given program. 

475  

476 It supports compile options listed in Supported Compile Options. 

477  

478 Args: 

479 prog (intptr_t): CUDA Runtime Compilation program. 

480 options (list[bytes] | None): Compiler options as a list of byte strings. 

481 May be ``None`` or an empty list for no options. 

482  

483 .. seealso:: `nvrtcCompileProgram` 

484 """ 

485 if options is None: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

486 options = [] 

487 cdef int num_options = len(options) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

488 cdef vector[const char*] cy_options = options 27 ctx1bcdefghijklmnopqrstuvwxyzAB

489 cdef const char** opt_data = NULL 27 ctx1bcdefghijklmnopqrstuvwxyzAB

490 if num_options: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

491 opt_data = cy_options.data() 27 ctx1bcdefghijklmnopqrstuvwxyzAB

492 with nogil: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

493 __status__ = nvrtcCompileProgram(<Program>prog, num_options, opt_data) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

494 check_status(__status__) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

495  

496  

497cpdef set_flow_callback(intptr_t prog, intptr_t callback, intptr_t payload): (empty)

498 """nvrtcSetFlowCallback registers a callback that the compiler invokes during :func:`compile_program`. 

499  

500 The callback signature must be ``int callback(void *param1, void *param2)``. 

501 The compiler passes *payload* as *param1* and ``NULL`` as *param2* (reserved). 

502 Return 1 to cancel compilation, 0 to continue; the callback must return 

503 consistently, be thread-safe, and must not call any NVRTC/libnvvm/PTX APIs. 

504  

505 Pass *callback* as a raw C function-pointer integer (e.g. via ``ctypes.cast``). 

506 Pass 0 for *callback* to clear a previously registered callback. 

507  

508 Args: 

509 prog (intptr_t): CUDA Runtime Compilation program. 

510 callback (intptr_t): C function pointer ``int (*)(void*, void*)`` cast to an integer, 

511 or 0 to clear. 

512 payload (intptr_t): Opaque pointer passed to the callback as its first argument. 

513  

514 .. seealso:: `nvrtcSetFlowCallback` 

515 """ 

516 with nogil: 

517 __status__ = nvrtcSetFlowCallback(<Program>prog, <void*>callback, <void*>payload) 

518 check_status(__status__) 

519  

520  

521cpdef bytes get_lowered_name(intptr_t prog, bytes name_expression): (empty)

522 """nvrtcGetLoweredName extracts the lowered (mangled) name for a ``__global__`` function or ``__device__``/``__constant__`` variable. 

523  

524 The memory containing the name is released when the program is destroyed by 

525 :func:`destroy_program`. The identical *name_expression* must have been previously 

526 provided to :func:`add_name_expression`. 

527  

528 Args: 

529 prog (intptr_t): CUDA Runtime Compilation program. 

530 name_expression (bytes): Constant expression denoting the address of a 

531 ``__global__`` function or ``__device__``/``__constant__`` variable. 

532  

533 Returns: 

534 bytes: C string containing the lowered (mangled) name, or ``None``. 

535  

536 .. seealso:: `nvrtcGetLoweredName` 

537 """ 

538 cdef const char* c_name_expression = name_expression 1 ctx1C

539 cdef const char* lowered_name = NULL 1 ctx1C

540 with nogil: 1 ctx1C

541 __status__ = nvrtcGetLoweredName(<Program>prog, c_name_expression, &lowered_name) 1 ctx1C

542 check_status(__status__) 1 ctx1C

543 return <bytes>lowered_name if lowered_name != NULL else None 

544  

545  

546cpdef install_bundled_headers(bytes install_path, unsigned int flags): (empty)

547 """nvrtcInstallBundledHeaders extracts CUDA headers bundled with NVRTC to a specified directory. 

548  

549 NVRTC bundles a set of CUDA Toolkit headers and CCCL within libnvrtc-builtins. 

550 After extraction, compile kernels by passing ``-I<installPath>`` and 

551 ``-I<installPath>/cccl`` to :func:`compile_program`. A version marker file 

552 (``.nvrtc_headers_version``) is created to track the installed version. 

553 The function is thread-safe and process-safe; concurrent calls are serialized 

554 using file locking. 

555  

556 Args: 

557 install_path (bytes): Path where headers should be extracted (UTF-8 encoded). 

558 The directory is created if it does not exist. 

559 flags (unsigned int): Bitwise OR of :class:`InstallHeadersFlag` values (or 0 

560 for the default ``SKIP_IF_EXISTS`` behaviour). 

561  

562 Returns: 

563 bytes | None: Detailed error message on failure, or ``None`` on success. 

564  

565 .. seealso:: `nvrtcInstallBundledHeaders` 

566 """ 

567 cdef const char* c_install_path = install_path 

568 cdef const char* error_log = NULL 

569 with nogil: 

570 __status__ = nvrtcInstallBundledHeaders(c_install_path, flags, &error_log) 

571 check_status(__status__) 

572 return <bytes>error_log if error_log != NULL else None 

573  

574  

575cpdef tuple get_bundled_headers_info(): (empty)

576 """nvrtcGetBundledHeadersInfo queries information about the bundled headers without extracting them. 

577  

578 Allows users to determine if bundled headers are available and get size estimates 

579 before calling :func:`install_bundled_headers`. 

580  

581 Returns: 

582 tuple[BundledHeadersInfo, bytes | None]: Header information struct and an optional 

583 detailed error message (``None`` on success). 

584  

585 .. seealso:: `nvrtcGetBundledHeadersInfo` 

586 """ 

587 cdef BundledHeadersInfo info = BundledHeadersInfo() 

588 cdef nvrtcBundledHeadersInfo* c_info = <nvrtcBundledHeadersInfo*><intptr_t>(info._get_ptr()) 

589 cdef const char* error_log = NULL 

590 with nogil: 

591 __status__ = nvrtcGetBundledHeadersInfo(c_info, &error_log) 

592 check_status(__status__) 

593 return info, (<bytes>error_log if error_log != NULL else None) 

594  

595  

596cpdef remove_bundled_headers(bytes install_path): (empty)

597 """nvrtcRemoveBundledHeaders removes previously installed bundled headers. 

598  

599 Recursively removes all files and subdirectories within the installation 

600 directory to help manage disk space. 

601  

602 .. note:: This removes ALL contents of the specified directory, not just files 

603 installed by NVRTC. Use with caution. 

604  

605 Args: 

606 install_path (bytes): Path where headers were previously installed; must be 

607 the same path used with :func:`install_bundled_headers`. 

608  

609 Returns: 

610 bytes | None: Detailed error message on failure, or ``None`` on success. 

611  

612 .. seealso:: `nvrtcRemoveBundledHeaders` 

613 """ 

614 cdef const char* c_install_path = install_path 

615 cdef const char* error_log = NULL 

616 with nogil: 

617 __status__ = nvrtcRemoveBundledHeaders(c_install_path, &error_log) 

618 check_status(__status__) 

619 return <bytes>error_log if error_log != NULL else None 

620  

621  

622cpdef str get_error_string(int result): (empty)

623 """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"``. 

624  

625 Args: 

626 result (Result): CUDA Runtime Compilation API result code. 

627  

628 .. seealso:: `nvrtcGetErrorString` 

629 """ 

630 cdef const char *_output_cstr_ 

631 cdef bytes _output_ 

632 with nogil: 1 ctx1C

633 _output_cstr_ = nvrtcGetErrorString(<_Result>result) 1 ctx1C

634 _output_ = _output_cstr_ 1 ctx1C

635 return _output_.decode() 1 ctx1C

636  

637  

638cpdef tuple version(): (empty)

639 """nvrtcVersion sets the output parameters ``major`` and ``minor`` with the CUDA Runtime Compilation version number. 

640  

641 Returns: 

642 A 2-tuple containing: 

643  

644 - int: CUDA Runtime Compilation major version number. 

645 - int: CUDA Runtime Compilation minor version number. 

646  

647 .. seealso:: `nvrtcVersion` 

648 """ 

649 cdef int major 

650 cdef int minor 

651 with nogil: 10 ctx1abcdefghij

652 __status__ = nvrtcVersion(&major, &minor) 10 ctx1abcdefghij

653 check_status(__status__) 10 ctx1abcdefghij

654 return (major, minor) 10 ctx1abcdefghij

655  

656  

657cpdef int get_num_supported_archs() except? -1: (empty)

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

659  

660 Returns: 

661 int: number of supported architectures. 

662  

663 .. seealso:: `nvrtcGetNumSupportedArchs` 

664 """ 

665 cdef int num_archs 

666 with nogil: 

667 __status__ = nvrtcGetNumSupportedArchs(&num_archs) 

668 check_status(__status__) 

669 return num_archs 

670  

671  

672cpdef object get_supported_archs(): (empty)

673 """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``. 

674  

675 Returns: 

676 int: sorted array of supported architectures. 

677  

678 .. seealso:: `nvrtcGetSupportedArchs` 

679 """ 

680 cdef int numArchs 

681 with nogil: 1 ctx1D

682 __status__ = nvrtcGetNumSupportedArchs(&numArchs) 1 ctx1D

683 check_status(__status__) 1 ctx1D

684 if numArchs == 0: 1 ctx1D

685 return _cyb_view.array(shape=(1,), itemsize=sizeof(int), format="i", mode="c")[:0] 

686 cdef _cyb_view.array supported_archs = _cyb_view.array(shape=(numArchs,), itemsize=sizeof(int), format="i", mode="c") 1 ctx1D

687 cdef int *supported_archs_ptr = <int *>(supported_archs.data) 1 ctx1D

688 with nogil: 1 ctx1D

689 __status__ = nvrtcGetSupportedArchs(supported_archs_ptr) 1 ctx1D

690 check_status(__status__) 1 ctx1D

691 return supported_archs 1 ctx1D

692  

693  

694cpdef destroy_program(intptr_t prog): (empty)

695 """nvrtcDestroyProgram destroys the given program. 

696  

697 Args: 

698 prog (intptr_t): CUDA Runtime Compilation program. 

699  

700 .. seealso:: `nvrtcDestroyProgram` 

701 """ 

702 cdef Program _prog_ = <Program>prog 

703 with nogil: 

704 __status__ = nvrtcDestroyProgram(&_prog_) 

705 check_status(__status__) 

706  

707  

708cpdef size_t get_ptx_size(intptr_t prog) except? 0: (empty)

709 """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``). 

710  

711 Args: 

712 prog (intptr_t): CUDA Runtime Compilation program. 

713  

714 Returns: 

715 size_t: Size of the generated PTX (including the trailing 

716 ``NULL``). 

717  

718 .. seealso:: `nvrtcGetPTXSize` 

719 """ 

720 cdef size_t ptx_size_ret 

721 with nogil: 

722 __status__ = nvrtcGetPTXSize(<Program>prog, &ptx_size_ret) 

723 check_status(__status__) 

724 return ptx_size_ret 

725  

726  

727cpdef bytes get_ptx(intptr_t prog): (empty)

728 """nvrtcGetPTX stores the PTX generated by the previous compilation of ``prog`` in the memory pointed by ``ptx``. 

729  

730 Args: 

731 prog (intptr_t): CUDA Runtime Compilation program. 

732  

733 Returns: 

734 char: Compiled result. 

735  

736 .. seealso:: `nvrtcGetPTX` 

737 """ 

738 cdef size_t ptxSizeRet 

739 with nogil: 

740 __status__ = nvrtcGetPTXSize(<Program>prog, &ptxSizeRet) 

741 check_status(__status__) 

742 if ptxSizeRet == 0: 

743 return b"" 

744 cdef bytes _ptx_ = bytes(ptxSizeRet) 

745 cdef char* ptx = _ptx_ 

746 with nogil: 

747 __status__ = nvrtcGetPTX(<Program>prog, ptx) 

748 check_status(__status__) 

749 return _ptx_ 

750  

751  

752cpdef size_t get_cubin_size(intptr_t prog) except? 0: (empty)

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

754  

755 Args: 

756 prog (intptr_t): CUDA Runtime Compilation program. 

757  

758 Returns: 

759 size_t: Size of the generated cubin. 

760  

761 .. seealso:: `nvrtcGetCUBINSize` 

762 """ 

763 cdef size_t cubin_size_ret 

764 with nogil: 

765 __status__ = nvrtcGetCUBINSize(<Program>prog, &cubin_size_ret) 

766 check_status(__status__) 

767 return cubin_size_ret 

768  

769  

770cpdef bytes get_cubin(intptr_t prog): (empty)

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

772  

773 Args: 

774 prog (intptr_t): CUDA Runtime Compilation program. 

775  

776 Returns: 

777 char: Compiled and assembled result. 

778  

779 .. seealso:: `nvrtcGetCUBIN` 

780 """ 

781 cdef size_t cubinSizeRet 

782 with nogil: 15 ctx1bcdefghijopqrst

783 __status__ = nvrtcGetCUBINSize(<Program>prog, &cubinSizeRet) 15 ctx1bcdefghijopqrst

784 check_status(__status__) 15 ctx1bcdefghijopqrst

785 if cubinSizeRet == 0: 15 ctx1bcdefghijopqrst

786 return b"" 

787 cdef bytes _cubin_ = bytes(cubinSizeRet) 15 ctx1bcdefghijopqrst

788 cdef char* cubin = _cubin_ 15 ctx1bcdefghijopqrst

789 with nogil: 15 ctx1bcdefghijopqrst

790 __status__ = nvrtcGetCUBIN(<Program>prog, cubin) 15 ctx1bcdefghijopqrst

791 check_status(__status__) 15 ctx1bcdefghijopqrst

792 return _cubin_ 15 ctx1bcdefghijopqrst

793  

794  

795cpdef size_t get_ltoir_size(intptr_t prog) except? 0: (empty)

796 """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``. 

797  

798 Args: 

799 prog (intptr_t): CUDA Runtime Compilation program. 

800  

801 Returns: 

802 size_t: Size of the generated LTO IR. 

803  

804 .. seealso:: `nvrtcGetLTOIRSize` 

805 """ 

806 cdef size_t ltoir_size_ret 

807 with nogil: 

808 __status__ = nvrtcGetLTOIRSize(<Program>prog, &ltoir_size_ret) 

809 check_status(__status__) 

810 return ltoir_size_ret 

811  

812  

813cpdef bytes get_ltoir(intptr_t prog): (empty)

814 """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``. 

815  

816 Args: 

817 prog (intptr_t): CUDA Runtime Compilation program. 

818  

819 Returns: 

820 char: Compiled result. 

821  

822 .. seealso:: `nvrtcGetLTOIR` 

823 """ 

824 cdef size_t LTOIRSizeRet 

825 with nogil: 12 ctx1klmnuvwxyzAB

826 __status__ = nvrtcGetLTOIRSize(<Program>prog, &LTOIRSizeRet) 12 ctx1klmnuvwxyzAB

827 check_status(__status__) 12 ctx1klmnuvwxyzAB

828 if LTOIRSizeRet == 0: 12 ctx1klmnuvwxyzAB

829 return b"" 

830 cdef bytes _ltoir_ = bytes(LTOIRSizeRet) 12 ctx1klmnuvwxyzAB

831 cdef char* ltoir = _ltoir_ 12 ctx1klmnuvwxyzAB

832 with nogil: 12 ctx1klmnuvwxyzAB

833 __status__ = nvrtcGetLTOIR(<Program>prog, ltoir) 12 ctx1klmnuvwxyzAB

834 check_status(__status__) 12 ctx1klmnuvwxyzAB

835 return _ltoir_ 12 ctx1klmnuvwxyzAB

836  

837  

838cpdef size_t get_optix_ir_size(intptr_t prog) except? 0: (empty)

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

840  

841 Args: 

842 prog (intptr_t): CUDA Runtime Compilation program. 

843  

844 Returns: 

845 size_t: Size of the generated LTO IR. 

846  

847 .. seealso:: `nvrtcGetOptiXIRSize` 

848 """ 

849 cdef size_t optixir_size_ret 

850 with nogil: 

851 __status__ = nvrtcGetOptiXIRSize(<Program>prog, &optixir_size_ret) 

852 check_status(__status__) 

853 return optixir_size_ret 

854  

855  

856cpdef bytes get_optix_ir(intptr_t prog): (empty)

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

858  

859 Args: 

860 prog (intptr_t): CUDA Runtime Compilation program. 

861  

862 Returns: 

863 char: Optix IR Compiled result. 

864  

865 .. seealso:: `nvrtcGetOptiXIR` 

866 """ 

867 cdef size_t optixirSizeRet 

868 with nogil: 

869 __status__ = nvrtcGetOptiXIRSize(<Program>prog, &optixirSizeRet) 

870 check_status(__status__) 

871 if optixirSizeRet == 0: 

872 return b"" 

873 cdef bytes _optixir_ = bytes(optixirSizeRet) 

874 cdef char* optixir = _optixir_ 

875 with nogil: 

876 __status__ = nvrtcGetOptiXIR(<Program>prog, optixir) 

877 check_status(__status__) 

878 return _optixir_ 

879  

880  

881cpdef size_t get_program_log_size(intptr_t prog) except? 0: (empty)

882 """nvrtcGetProgramLogSize sets ``log_size_ret`` with the size of the log generated by the previous compilation of ``prog`` (including the trailing ``NULL``). 

883  

884 Args: 

885 prog (intptr_t): CUDA Runtime Compilation program. 

886  

887 Returns: 

888 size_t: Size of the compilation log (including the trailing 

889 ``NULL``). 

890  

891 .. seealso:: `nvrtcGetProgramLogSize` 

892 """ 

893 cdef size_t log_size_ret 

894 with nogil: 

895 __status__ = nvrtcGetProgramLogSize(<Program>prog, &log_size_ret) 

896 check_status(__status__) 

897 return log_size_ret 

898  

899  

900cpdef bytes get_program_log(intptr_t prog): (empty)

901 """nvrtcGetProgramLog stores the log generated by the previous compilation of ``prog`` in the memory pointed by ``log``. 

902  

903 Args: 

904 prog (intptr_t): CUDA Runtime Compilation program. 

905  

906 Returns: 

907 char: Compilation log. 

908  

909 .. seealso:: `nvrtcGetProgramLog` 

910 """ 

911 cdef size_t logSizeRet 

912 with nogil: 9 ctx1bcdefghij

913 __status__ = nvrtcGetProgramLogSize(<Program>prog, &logSizeRet) 9 ctx1bcdefghij

914 check_status(__status__) 9 ctx1bcdefghij

915 if logSizeRet == 0: 9 ctx1bcdefghij

916 return b"" 

917 cdef bytes _log_ = bytes(logSizeRet) 9 ctx1bcdefghij

918 cdef char* log = _log_ 9 ctx1bcdefghij

919 with nogil: 9 ctx1bcdefghij

920 __status__ = nvrtcGetProgramLog(<Program>prog, log) 9 ctx1bcdefghij

921 check_status(__status__) 9 ctx1bcdefghij

922 return _log_ 9 ctx1bcdefghij

923  

924  

925cpdef add_name_expression(intptr_t prog, name_expression): (empty)

926 """nvrtcAddNameExpression notes the given name expression denoting the address of a global function or device/__constant__ variable. 

927  

928 Args: 

929 prog (intptr_t): CUDA Runtime Compilation program. 

930 name_expression (str): constant expression denoting the 

931 address of a global function or device/__constant__ 

932 variable. 

933  

934 .. seealso:: `nvrtcAddNameExpression` 

935 """ 

936 if not isinstance(name_expression, str): 

937 raise TypeError("name_expression must be a Python str") 

938 cdef bytes _temp_name_expression_ = (<str>name_expression).encode() 

939 cdef char* _name_expression_ = _temp_name_expression_ 

940 with nogil: 

941 __status__ = nvrtcAddNameExpression(<Program>prog, <const char* const>_name_expression_) 

942 check_status(__status__) 

943  

944  

945cpdef size_t get_pch_heap_size() except? 0: (empty)

946 """retrieve the current size of the PCH Heap. 

947  

948 Returns: 

949 size_t: pointer to location where the size of the PCH Heap 

950 will be stored. 

951  

952 .. seealso:: `nvrtcGetPCHHeapSize` 

953 """ 

954 cdef size_t ret 

955 with nogil: 

956 __status__ = nvrtcGetPCHHeapSize(&ret) 

957 check_status(__status__) 

958 return ret 

959  

960  

961cpdef set_pch_heap_size(size_t size): (empty)

962 """set the size of the PCH Heap. 

963  

964 Args: 

965 size (size_t): requested size of the PCH Heap, in bytes. 

966  

967 .. seealso:: `nvrtcSetPCHHeapSize` 

968 """ 

969 with nogil: 

970 __status__ = nvrtcSetPCHHeapSize(size) 

971 check_status(__status__) 

972  

973  

974cpdef int get_pch_create_status(intptr_t prog) except? -1: (empty)

975 """returns the PCH creation status. 

976  

977 Args: 

978 prog (intptr_t): CUDA Runtime Compilation program. 

979  

980 .. seealso:: `nvrtcGetPCHCreateStatus` 

981 """ 

982 cdef int ret 

983 with nogil: 

984 ret = <int>nvrtcGetPCHCreateStatus(<Program>prog) 

985 return ret 

986  

987  

988cpdef size_t get_pch_heap_size_required(intptr_t prog) except? 0: (empty)

989 """retrieve the required size of the PCH heap required to compile the given program. 

990  

991 Args: 

992 prog (intptr_t): CUDA Runtime Compilation program. 

993  

994 Returns: 

995 size_t: pointer to location where the required size of the PCH 

996 Heap will be stored. 

997  

998 .. seealso:: `nvrtcGetPCHHeapSizeRequired` 

999 """ 

1000 cdef size_t size 

1001 with nogil: 

1002 __status__ = nvrtcGetPCHHeapSizeRequired(<Program>prog, &size) 

1003 check_status(__status__) 

1004 return size 

1005  

1006  

1007cpdef size_t get_tile_ir_size(intptr_t prog) except? 0: (empty)

1008 """nvrtcGetTileIRSize sets the value of ``tile_ir_size_ret`` with the size of the cuda_tile IR generated by the previous compilation of ``prog``. 

1009  

1010 Args: 

1011 prog (intptr_t): CUDA Runtime Compilation program. 

1012  

1013 Returns: 

1014 size_t: Size of the generated cuda_tile IR. 

1015  

1016 .. seealso:: `nvrtcGetTileIRSize` 

1017 """ 

1018 cdef size_t tile_ir_size_ret 

1019 with nogil: 

1020 __status__ = nvrtcGetTileIRSize(<Program>prog, &tile_ir_size_ret) 

1021 check_status(__status__) 

1022 return tile_ir_size_ret 

1023  

1024  

1025cpdef bytes get_tile_ir(intptr_t prog): (empty)

1026 """nvrtcGettile_ir stores the cuda_tile IR generated by the previous compilation of ``prog`` in the memory pointed by ``tile_ir``. 

1027  

1028 Args: 

1029 prog (intptr_t): CUDA Runtime Compilation program. 

1030  

1031 Returns: 

1032 char: Generated cuda_tile IR. 

1033  

1034 .. seealso:: `nvrtcGetTileIR` 

1035 """ 

1036 cdef size_t TileIRSizeRet 

1037 with nogil: 

1038 __status__ = nvrtcGetTileIRSize(<Program>prog, &TileIRSizeRet) 

1039 check_status(__status__) 

1040 if TileIRSizeRet == 0: 

1041 return b"" 

1042 cdef bytes _tile_ir_ = bytes(TileIRSizeRet) 

1043 cdef char* tile_ir = _tile_ir_ 

1044 with nogil: 

1045 __status__ = nvrtcGetTileIR(<Program>prog, tile_ir) 

1046 check_status(__status__) 

1047 return _tile_ir_ 

1048  

1049  

1050del _cyb_FastEnum (empty)