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=e5360fc057cdd7b8e28b4d502821443d190e589b200dce1a234494ad6e9abf93 

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.stdlib cimport ( 

15 calloc as _cyb_calloc, 

16 free as _cyb_free, 

17 malloc as _cyb_malloc, 

18) 

19from libc.string cimport ( 

20 memcmp as _cyb_memcmp, 

21 memcpy as _cyb_memcpy, 

22) 

23  

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

25  

26import numpy as _numpy (empty)

27  

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

29 buffer.buf = <char *>ptr 

30 buffer.format = 'b' 

31 buffer.internal = NULL 

32 buffer.itemsize = 1 

33 buffer.len = size 

34 buffer.ndim = 1 

35 buffer.obj = self 

36 buffer.readonly = readonly 

37 buffer.shape = &buffer.len 

38 buffer.strides = &buffer.itemsize 

39 buffer.suboffsets = NULL 

40  

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

42 cdef _cyb_cpython.Py_buffer view 

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

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

45 try: 

46 if view.itemsize != 1: 

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

48 if view.len != size: 

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

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

51 finally: 

52 _cyb_cpython.PyBuffer_Release(&view) 

53  

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

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

56 if isinstance(data, lowpp_type): 

57 return data 

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

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

60 if data.size != 1: 

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

62 if data.dtype != expected_dtype: 

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

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

65  

66  

67# <<<< END OF PREAMBLE CONTENT >>>> 

68  

69cimport cython # NOQA 

70from libcpp.vector cimport vector 

71  

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

73  

74  

75############################################################################### 

76# Enum 

77############################################################################### 

78  

79class Result(_cyb_FastEnum): (empty)

80 """ 

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

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

83  

84 See `nvrtcResult`. 

85 """ 

86 SUCCESS = NVRTC_SUCCESS (empty)

87 ERROR_OUT_OF_MEMORY = NVRTC_ERROR_OUT_OF_MEMORY (empty)

88 ERROR_PROGRAM_CREATION_FAILURE = NVRTC_ERROR_PROGRAM_CREATION_FAILURE (empty)

89 ERROR_INVALID_INPUT = NVRTC_ERROR_INVALID_INPUT (empty)

90 ERROR_INVALID_PROGRAM = NVRTC_ERROR_INVALID_PROGRAM (empty)

91 ERROR_INVALID_OPTION = NVRTC_ERROR_INVALID_OPTION (empty)

92 ERROR_COMPILATION = NVRTC_ERROR_COMPILATION (empty)

93 ERROR_BUILTIN_OPERATION_FAILURE = NVRTC_ERROR_BUILTIN_OPERATION_FAILURE (empty)

94 ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION (empty)

95 ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION (empty)

96 ERROR_NAME_EXPRESSION_NOT_VALID = NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID (empty)

97 ERROR_INTERNAL_ERROR = NVRTC_ERROR_INTERNAL_ERROR (empty)

98 ERROR_TIME_FILE_WRITE_FAILED = NVRTC_ERROR_TIME_FILE_WRITE_FAILED (empty)

99 ERROR_NO_PCH_CREATE_ATTEMPTED = NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED (empty)

100 ERROR_PCH_CREATE_HEAP_EXHAUSTED = NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED (empty)

101 ERROR_PCH_CREATE = NVRTC_ERROR_PCH_CREATE (empty)

102 ERROR_CANCELLED = NVRTC_ERROR_CANCELLED (empty)

103 ERROR_TIME_TRACE_FILE_WRITE_FAILED = NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED (empty)

104 ERROR_BUSY = NVRTC_ERROR_BUSY (empty)

105  

106  

107class InstallHeadersFlag(_FastEnum): (empty)

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

109 SKIP_IF_EXISTS = ( 

110 0x0, (empty)

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

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

113 ) 

114 FORCE_OVERWRITE = ( 

115 0x1, (empty)

116 "Clear existing directory contents before installation. " 

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

118 ) 

119 NO_WAIT = ( 

120 0x2, (empty)

121 "Return NVRTC_ERROR_BUSY immediately if installation is in progress " 

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

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

124 "Do not wait for installation to complete." 

125 ) 

126  

127  

128############################################################################### 

129# Error handling 

130############################################################################### 

131  

132  

133class NvrtcError(Exception): (empty)

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

135 self.status = status 1 ctx1C

136 s = get_error_string(status) 1 ctx1C

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

138  

139 def __reduce__(self): (empty)

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

141  

142class OutOfMemoryError(NvrtcError): (empty)

143 pass (empty)

144class ProgramCreationFailureError(NvrtcError): (empty)

145 pass (empty)

146class InvalidInputError(NvrtcError): (empty)

147 pass (empty)

148class InvalidProgramError(NvrtcError): (empty)

149 pass (empty)

150class InvalidOptionError(NvrtcError): (empty)

151 pass (empty)

152class CompilationError(NvrtcError): (empty)

153 pass (empty)

154class BuiltinOperationFailureError(NvrtcError): (empty)

155 pass (empty)

156class NoNameExpressionsAfterCompilationError(NvrtcError): (empty)

157 pass (empty)

158class NoLoweredNamesBeforeCompilationError(NvrtcError): (empty)

159 pass (empty)

160class NameExpressionNotValidError(NvrtcError): (empty)

161 pass (empty)

162class InternalErrorError(NvrtcError): (empty)

163 pass (empty)

164class TimeFileWriteFailedError(NvrtcError): (empty)

165 pass (empty)

166class NoPchCreateAttemptedError(NvrtcError): (empty)

167 pass (empty)

168class PchCreateHeapExhaustedError(NvrtcError): (empty)

169 pass (empty)

170class PchCreateError(NvrtcError): (empty)

171 pass (empty)

172class CancelledError(NvrtcError): (empty)

173 pass (empty)

174class TimeTraceFileWriteFailedError(NvrtcError): (empty)

175 pass (empty)

176class BusyError(NvrtcError): (empty)

177 pass (empty)

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

179 cdef object pystatus = status 1 ctx1C

180 if status == 1: 1 ctx1C

181 return OutOfMemoryError(pystatus) 

182 elif status == 2: 

183 return ProgramCreationFailureError(pystatus) 

184 elif status == 3: 

185 return InvalidInputError(pystatus) 

186 elif status == 4: 

187 return InvalidProgramError(pystatus) 1 ctx1C

188 elif status == 5: 

189 return InvalidOptionError(pystatus) 

190 elif status == 6: 

191 return CompilationError(pystatus) 

192 elif status == 7: 

193 return BuiltinOperationFailureError(pystatus) 

194 elif status == 8: 

195 return NoNameExpressionsAfterCompilationError(pystatus) 

196 elif status == 9: 

197 return NoLoweredNamesBeforeCompilationError(pystatus) 

198 elif status == 10: 

199 return NameExpressionNotValidError(pystatus) 

200 elif status == 11: 

201 return InternalErrorError(pystatus) 

202 elif status == 12: 

203 return TimeFileWriteFailedError(pystatus) 

204 elif status == 13: 

205 return NoPchCreateAttemptedError(pystatus) 

206 elif status == 14: 

207 return PchCreateHeapExhaustedError(pystatus) 

208 elif status == 15: 

209 return PchCreateError(pystatus) 

210 elif status == 16: 

211 return CancelledError(pystatus) 

212 elif status == 17: 

213 return TimeTraceFileWriteFailedError(pystatus) 

214 elif status == 18: 

215 return BusyError(pystatus) 

216 return NvrtcError(status) 

217  

218InternalError = InternalErrorError (empty)

219  

220  

221@cython.profile(False) (empty)

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

223 if status != 0: 30 ctx1abcdefghijklmnopqrstuvwxyzABCD

224 with gil: 1 ctx1C

225 raise _nvrtc_error_factory(status) 1 ctx1C

226 return status 29 ctx1abcdefghijklmnopqrstuvwxyzABD

227  

228  

229############################################################################### 

230# POD definitions 

231############################################################################### 

232  

233cdef _get_bundled_headers_info_dtype_offsets(): (empty)

234 cdef nvrtcBundledHeadersInfo pod 

235 return _numpy.dtype({ (empty)

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

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

238 'offsets': [ (empty)

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

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

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

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

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

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

245 ], 

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

247 }) 

248  

249bundled_headers_info_dtype = _get_bundled_headers_info_dtype_offsets() (empty)

250  

251cdef class BundledHeadersInfo: 

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

253  

254  

255 .. seealso:: `nvrtcBundledHeadersInfo` 

256 """ 

257 cdef: 

258 nvrtcBundledHeadersInfo *_ptr 

259 object _owner 

260 bint _owned 

261 bint _readonly 

262  

263 def __init__(self): 

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

265 if self._ptr == NULL: 

266 raise MemoryError("Error allocating BundledHeadersInfo") 

267 self._owner = None 

268 self._owned = True 

269 self._readonly = False 

270  

271 def __dealloc__(self): 

272 cdef nvrtcBundledHeadersInfo *ptr 

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

274 ptr = self._ptr 

275 self._ptr = NULL 

276 _cyb_free(ptr) 

277  

278 def __repr__(self): 

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

280  

281 @property 

282 def ptr(self): 

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

284 return <intptr_t>(self._ptr) 

285  

286 cdef intptr_t _get_ptr(self): (empty)

287 return <intptr_t>(self._ptr) 

288  

289 def __int__(self): 

290 return <intptr_t>(self._ptr) 

291  

292 def __eq__(self, other): 

293 cdef BundledHeadersInfo other_ 

294 if not isinstance(other, BundledHeadersInfo): 

295 return False 

296 other_ = other 

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

298  

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

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

301  

302 def __releasebuffer__(self, Py_buffer *buffer): 

303 pass 

304  

305 def __setitem__(self, key, val): 

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

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

308 if self._ptr == NULL: 

309 raise MemoryError("Error allocating BundledHeadersInfo") 

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

311 self._owner = None 

312 self._owned = True 

313 self._readonly = not val.flags.writeable 

314 else: 

315 setattr(self, key, val) 

316  

317 @property 

318 def available(self): 

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

320 return self._ptr[0].available 

321  

322 @available.setter 

323 def available(self, val): 

324 if self._readonly: 

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

326 self._ptr[0].available = val 

327  

328 @property 

329 def compressed_size(self): 

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

331 return self._ptr[0].compressedSize 

332  

333 @compressed_size.setter 

334 def compressed_size(self, val): 

335 if self._readonly: 

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

337 self._ptr[0].compressedSize = val 

338  

339 @property 

340 def uncompressed_size(self): 

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

342 return self._ptr[0].uncompressedSize 

343  

344 @uncompressed_size.setter 

345 def uncompressed_size(self, val): 

346 if self._readonly: 

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

348 self._ptr[0].uncompressedSize = val 

349  

350 @property 

351 def cuda_version_major(self): 

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

353 return self._ptr[0].cudaVersionMajor 

354  

355 @cuda_version_major.setter 

356 def cuda_version_major(self, val): 

357 if self._readonly: 

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

359 self._ptr[0].cudaVersionMajor = val 

360  

361 @property 

362 def cuda_version_minor(self): 

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

364 return self._ptr[0].cudaVersionMinor 

365  

366 @cuda_version_minor.setter 

367 def cuda_version_minor(self, val): 

368 if self._readonly: 

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

370 self._ptr[0].cudaVersionMinor = val 

371  

372 @property 

373 def num_files(self): 

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

375 return self._ptr[0].numFiles 

376  

377 @num_files.setter 

378 def num_files(self, val): 

379 if self._readonly: 

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

381 self._ptr[0].numFiles = val 

382  

383 @staticmethod (empty)

384 def from_buffer(buffer): 

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

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

387  

388 @staticmethod (empty)

389 def from_data(data): 

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

391  

392 Args: 

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

394 """ 

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

396  

397 @staticmethod (empty)

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

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

400  

401 Args: 

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

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

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

405 """ 

406 if ptr == 0: 

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

408 cdef BundledHeadersInfo obj = BundledHeadersInfo.__new__(BundledHeadersInfo) 

409 if owner is None: 

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

411 if obj._ptr == NULL: 

412 raise MemoryError("Error allocating BundledHeadersInfo") 

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

414 obj._owner = None 

415 obj._owned = True 

416 else: 

417 obj._ptr = <nvrtcBundledHeadersInfo *>ptr 

418 obj._owner = owner 

419 obj._owned = False 

420 obj._readonly = readonly 

421 return obj 

422  

423  

424############################################################################### 

425# Wrapper functions 

426############################################################################### 

427  

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

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

430  

431 Args: 

432 src (bytes): CUDA program source. 

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

434 to be used. 

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

436 list (no headers). 

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

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

439  

440 Returns: 

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

442  

443 .. seealso:: `nvrtcCreateProgram` 

444 """ 

445 if headers is None: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

446 headers = [] 27 ctx1bcdefghijklmnopqrstuvwxyzAB

447 if include_names is None: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

448 include_names = [] 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

450 raise ValueError( 

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

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

453 ) 

454 cdef int num_headers = len(headers) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

455 cdef Program prog 

456 cdef const char* c_src = src 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

458 cdef const char* c_name = _name 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

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

461 cdef const char** hdr_data = NULL 27 ctx1bcdefghijklmnopqrstuvwxyzAB

462 cdef const char** inc_data = NULL 27 ctx1bcdefghijklmnopqrstuvwxyzAB

463 if num_headers: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

464 hdr_data = cy_headers.data() 

465 inc_data = cy_include_names.data() 

466 with nogil: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

468 check_status(__status__) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

469 return <intptr_t>prog 27 ctx1bcdefghijklmnopqrstuvwxyzAB

470  

471  

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

473 """nvrtcCompileProgram compiles the given program. 

474  

475 It supports compile options listed in Supported Compile Options. 

476  

477 Args: 

478 prog (intptr_t): CUDA Runtime Compilation program. 

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

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

481  

482 .. seealso:: `nvrtcCompileProgram` 

483 """ 

484 if options is None: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

485 options = [] 

486 cdef int num_options = len(options) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

488 cdef const char** opt_data = NULL 27 ctx1bcdefghijklmnopqrstuvwxyzAB

489 if num_options: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

490 opt_data = cy_options.data() 27 ctx1bcdefghijklmnopqrstuvwxyzAB

491 with nogil: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

493 check_status(__status__) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

494  

495  

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

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

498  

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

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

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

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

503  

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

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

506  

507 Args: 

508 prog (intptr_t): CUDA Runtime Compilation program. 

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

510 or 0 to clear. 

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

512  

513 .. seealso:: `nvrtcSetFlowCallback` 

514 """ 

515 with nogil: 

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

517 check_status(__status__) 

518  

519  

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

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

522  

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

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

525 provided to :func:`add_name_expression`. 

526  

527 Args: 

528 prog (intptr_t): CUDA Runtime Compilation program. 

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

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

531  

532 Returns: 

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

534  

535 .. seealso:: `nvrtcGetLoweredName` 

536 """ 

537 cdef const char* c_name_expression = name_expression 1 ctx1C

538 cdef const char* lowered_name = NULL 1 ctx1C

539 with nogil: 1 ctx1C

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

541 check_status(__status__) 1 ctx1C

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

543  

544  

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

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

547  

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

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

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

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

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

553 using file locking. 

554  

555 Args: 

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

557 The directory is created if it does not exist. 

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

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

560  

561 Returns: 

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

563  

564 .. seealso:: `nvrtcInstallBundledHeaders` 

565 """ 

566 cdef const char* c_install_path = install_path 

567 cdef const char* error_log = NULL 

568 with nogil: 

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

570 check_status(__status__) 

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

572  

573  

574cpdef tuple get_bundled_headers_info(): (empty)

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

576  

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

578 before calling :func:`install_bundled_headers`. 

579  

580 Returns: 

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

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

583  

584 .. seealso:: `nvrtcGetBundledHeadersInfo` 

585 """ 

586 cdef BundledHeadersInfo info = BundledHeadersInfo() 

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

588 cdef const char* error_log = NULL 

589 with nogil: 

590 __status__ = nvrtcGetBundledHeadersInfo(c_info, &error_log) 

591 check_status(__status__) 

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

593  

594  

595cpdef remove_bundled_headers(bytes install_path): (empty)

596 """nvrtcRemoveBundledHeaders removes previously installed bundled headers. 

597  

598 Recursively removes all files and subdirectories within the installation 

599 directory to help manage disk space. 

600  

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

602 installed by NVRTC. Use with caution. 

603  

604 Args: 

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

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

607  

608 Returns: 

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

610  

611 .. seealso:: `nvrtcRemoveBundledHeaders` 

612 """ 

613 cdef const char* c_install_path = install_path 

614 cdef const char* error_log = NULL 

615 with nogil: 

616 __status__ = nvrtcRemoveBundledHeaders(c_install_path, &error_log) 

617 check_status(__status__) 

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

619  

620  

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

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

623  

624 Args: 

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

626  

627 .. seealso:: `nvrtcGetErrorString` 

628 """ 

629 cdef const char *_output_cstr_ 

630 cdef bytes _output_ 

631 with nogil: 1 ctx1C

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

633 _output_ = _output_cstr_ 1 ctx1C

634 return _output_.decode() 1 ctx1C

635  

636  

637cpdef tuple version(): (empty)

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

639  

640 Returns: 

641 A 2-tuple containing: 

642  

643 - int: CUDA Runtime Compilation major version number. 

644 - int: CUDA Runtime Compilation minor version number. 

645  

646 .. seealso:: `nvrtcVersion` 

647 """ 

648 cdef int major 

649 cdef int minor 

650 with nogil: 10 ctx1abcdefghij

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

652 check_status(__status__) 10 ctx1abcdefghij

653 return (major, minor) 10 ctx1abcdefghij

654  

655  

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

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

658  

659 Returns: 

660 int: number of supported architectures. 

661  

662 .. seealso:: `nvrtcGetNumSupportedArchs` 

663 """ 

664 cdef int num_archs 

665 with nogil: 

666 __status__ = nvrtcGetNumSupportedArchs(&num_archs) 

667 check_status(__status__) 

668 return num_archs 

669  

670  

671cpdef object get_supported_archs(): (empty)

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

673  

674 Returns: 

675 int: sorted array of supported architectures. 

676  

677 .. seealso:: `nvrtcGetSupportedArchs` 

678 """ 

679 cdef int numArchs 

680 with nogil: 1 ctx1D

681 __status__ = nvrtcGetNumSupportedArchs(&numArchs) 1 ctx1D

682 check_status(__status__) 1 ctx1D

683 if numArchs == 0: 1 ctx1D

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

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

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

687 with nogil: 1 ctx1D

688 __status__ = nvrtcGetSupportedArchs(supported_archs_ptr) 1 ctx1D

689 check_status(__status__) 1 ctx1D

690 return supported_archs 1 ctx1D

691  

692  

693cpdef destroy_program(intptr_t prog): (empty)

694 """nvrtcDestroyProgram destroys the given program. 

695  

696 Args: 

697 prog (intptr_t): CUDA Runtime Compilation program. 

698  

699 .. seealso:: `nvrtcDestroyProgram` 

700 """ 

701 cdef Program _prog_ = <Program>prog 

702 with nogil: 

703 __status__ = nvrtcDestroyProgram(&_prog_) 

704 check_status(__status__) 

705  

706  

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

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

709  

710 Args: 

711 prog (intptr_t): CUDA Runtime Compilation program. 

712  

713 Returns: 

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

715 ``NULL``). 

716  

717 .. seealso:: `nvrtcGetPTXSize` 

718 """ 

719 cdef size_t ptx_size_ret 

720 with nogil: 

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

722 check_status(__status__) 

723 return ptx_size_ret 

724  

725  

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

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

728  

729 Args: 

730 prog (intptr_t): CUDA Runtime Compilation program. 

731  

732 Returns: 

733 char: Compiled result. 

734  

735 .. seealso:: `nvrtcGetPTX` 

736 """ 

737 cdef size_t ptxSizeRet 

738 with nogil: 

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

740 check_status(__status__) 

741 if ptxSizeRet == 0: 

742 return b"" 

743 cdef bytes _ptx_ = bytes(ptxSizeRet) 

744 cdef char* ptx = _ptx_ 

745 with nogil: 

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

747 check_status(__status__) 

748 return _ptx_ 

749  

750  

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

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

753  

754 Args: 

755 prog (intptr_t): CUDA Runtime Compilation program. 

756  

757 Returns: 

758 size_t: Size of the generated cubin. 

759  

760 .. seealso:: `nvrtcGetCUBINSize` 

761 """ 

762 cdef size_t cubin_size_ret 

763 with nogil: 

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

765 check_status(__status__) 

766 return cubin_size_ret 

767  

768  

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

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

771  

772 Args: 

773 prog (intptr_t): CUDA Runtime Compilation program. 

774  

775 Returns: 

776 char: Compiled and assembled result. 

777  

778 .. seealso:: `nvrtcGetCUBIN` 

779 """ 

780 cdef size_t cubinSizeRet 

781 with nogil: 15 ctx1bcdefghijopqrst

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

783 check_status(__status__) 15 ctx1bcdefghijopqrst

784 if cubinSizeRet == 0: 15 ctx1bcdefghijopqrst

785 return b"" 

786 cdef bytes _cubin_ = bytes(cubinSizeRet) 15 ctx1bcdefghijopqrst

787 cdef char* cubin = _cubin_ 15 ctx1bcdefghijopqrst

788 with nogil: 15 ctx1bcdefghijopqrst

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

790 check_status(__status__) 15 ctx1bcdefghijopqrst

791 return _cubin_ 15 ctx1bcdefghijopqrst

792  

793  

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

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

796  

797 Args: 

798 prog (intptr_t): CUDA Runtime Compilation program. 

799  

800 Returns: 

801 size_t: Size of the generated LTO IR. 

802  

803 .. seealso:: `nvrtcGetLTOIRSize` 

804 """ 

805 cdef size_t ltoir_size_ret 

806 with nogil: 

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

808 check_status(__status__) 

809 return ltoir_size_ret 

810  

811  

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

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

814  

815 Args: 

816 prog (intptr_t): CUDA Runtime Compilation program. 

817  

818 Returns: 

819 char: Compiled result. 

820  

821 .. seealso:: `nvrtcGetLTOIR` 

822 """ 

823 cdef size_t LTOIRSizeRet 

824 with nogil: 12 ctx1klmnuvwxyzAB

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

826 check_status(__status__) 12 ctx1klmnuvwxyzAB

827 if LTOIRSizeRet == 0: 12 ctx1klmnuvwxyzAB

828 return b"" 

829 cdef bytes _ltoir_ = bytes(LTOIRSizeRet) 12 ctx1klmnuvwxyzAB

830 cdef char* ltoir = _ltoir_ 12 ctx1klmnuvwxyzAB

831 with nogil: 12 ctx1klmnuvwxyzAB

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

833 check_status(__status__) 12 ctx1klmnuvwxyzAB

834 return _ltoir_ 12 ctx1klmnuvwxyzAB

835  

836  

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

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

839  

840 Args: 

841 prog (intptr_t): CUDA Runtime Compilation program. 

842  

843 Returns: 

844 size_t: Size of the generated LTO IR. 

845  

846 .. seealso:: `nvrtcGetOptiXIRSize` 

847 """ 

848 cdef size_t optixir_size_ret 

849 with nogil: 

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

851 check_status(__status__) 

852 return optixir_size_ret 

853  

854  

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

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

857  

858 Args: 

859 prog (intptr_t): CUDA Runtime Compilation program. 

860  

861 Returns: 

862 char: Optix IR Compiled result. 

863  

864 .. seealso:: `nvrtcGetOptiXIR` 

865 """ 

866 cdef size_t optixirSizeRet 

867 with nogil: 

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

869 check_status(__status__) 

870 if optixirSizeRet == 0: 

871 return b"" 

872 cdef bytes _optixir_ = bytes(optixirSizeRet) 

873 cdef char* optixir = _optixir_ 

874 with nogil: 

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

876 check_status(__status__) 

877 return _optixir_ 

878  

879  

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

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

882  

883 Args: 

884 prog (intptr_t): CUDA Runtime Compilation program. 

885  

886 Returns: 

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

888 ``NULL``). 

889  

890 .. seealso:: `nvrtcGetProgramLogSize` 

891 """ 

892 cdef size_t log_size_ret 

893 with nogil: 

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

895 check_status(__status__) 

896 return log_size_ret 

897  

898  

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

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

901  

902 Args: 

903 prog (intptr_t): CUDA Runtime Compilation program. 

904  

905 Returns: 

906 char: Compilation log. 

907  

908 .. seealso:: `nvrtcGetProgramLog` 

909 """ 

910 cdef size_t logSizeRet 

911 with nogil: 9 ctx1bcdefghij

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

913 check_status(__status__) 9 ctx1bcdefghij

914 if logSizeRet == 0: 9 ctx1bcdefghij

915 return b"" 

916 cdef bytes _log_ = bytes(logSizeRet) 9 ctx1bcdefghij

917 cdef char* log = _log_ 9 ctx1bcdefghij

918 with nogil: 9 ctx1bcdefghij

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

920 check_status(__status__) 9 ctx1bcdefghij

921 return _log_ 9 ctx1bcdefghij

922  

923  

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

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

926  

927 Args: 

928 prog (intptr_t): CUDA Runtime Compilation program. 

929 name_expression (str): constant expression denoting the 

930 address of a global function or device/__constant__ 

931 variable. 

932  

933 .. seealso:: `nvrtcAddNameExpression` 

934 """ 

935 if not isinstance(name_expression, str): 

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

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

938 cdef char* _name_expression_ = _temp_name_expression_ 

939 with nogil: 

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

941 check_status(__status__) 

942  

943  

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

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

946  

947 Returns: 

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

949 will be stored. 

950  

951 .. seealso:: `nvrtcGetPCHHeapSize` 

952 """ 

953 cdef size_t ret 

954 with nogil: 

955 __status__ = nvrtcGetPCHHeapSize(&ret) 

956 check_status(__status__) 

957 return ret 

958  

959  

960cpdef set_pch_heap_size(size_t size): (empty)

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

962  

963 Args: 

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

965  

966 .. seealso:: `nvrtcSetPCHHeapSize` 

967 """ 

968 with nogil: 

969 __status__ = nvrtcSetPCHHeapSize(size) 

970 check_status(__status__) 

971  

972  

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

974 """returns the PCH creation status. 

975  

976 Args: 

977 prog (intptr_t): CUDA Runtime Compilation program. 

978  

979 .. seealso:: `nvrtcGetPCHCreateStatus` 

980 """ 

981 cdef int ret 

982 with nogil: 

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

984 return ret 

985  

986  

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

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

989  

990 Args: 

991 prog (intptr_t): CUDA Runtime Compilation program. 

992  

993 Returns: 

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

995 Heap will be stored. 

996  

997 .. seealso:: `nvrtcGetPCHHeapSizeRequired` 

998 """ 

999 cdef size_t size 

1000 with nogil: 

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

1002 check_status(__status__) 

1003 return size 

1004  

1005  

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

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

1008  

1009 Args: 

1010 prog (intptr_t): CUDA Runtime Compilation program. 

1011  

1012 Returns: 

1013 size_t: Size of the generated cuda_tile IR. 

1014  

1015 .. seealso:: `nvrtcGetTileIRSize` 

1016 """ 

1017 cdef size_t tile_ir_size_ret 

1018 with nogil: 

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

1020 check_status(__status__) 

1021 return tile_ir_size_ret 

1022  

1023  

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

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

1026  

1027 Args: 

1028 prog (intptr_t): CUDA Runtime Compilation program. 

1029  

1030 Returns: 

1031 char: Generated cuda_tile IR. 

1032  

1033 .. seealso:: `nvrtcGetTileIR` 

1034 """ 

1035 cdef size_t TileIRSizeRet 

1036 with nogil: 

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

1038 check_status(__status__) 

1039 if TileIRSizeRet == 0: 

1040 return b"" 

1041 cdef bytes _tile_ir_ = bytes(TileIRSizeRet) 

1042 cdef char* tile_ir = _tile_ir_ 

1043 with nogil: 

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

1045 check_status(__status__) 

1046 return _tile_ir_ 

1047  

1048  

1049del _cyb_FastEnum (empty)