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

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.4.1. Do not modify it directly. 

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

7  

8  

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

10  

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) 

26  

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

28  

29import numpy as _numpy (empty)

30  

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

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 

43  

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

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) 

56  

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

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) 

68  

69  

70# <<<< END OF PREAMBLE CONTENT >>>> 

71  

72cimport cython # NOQA 

73from libcpp.vector cimport vector 

74  

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

76  

77  

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

79# Enum 

80############################################################################### 

81  

82class Result(_cyb_FastEnum): (empty)

83 """ 

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

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

86  

87 See `nvrtcResult`. 

88 """ 

89 SUCCESS = NVRTC_SUCCESS (empty)

90 ERROR_OUT_OF_MEMORY = NVRTC_ERROR_OUT_OF_MEMORY (empty)

91 ERROR_PROGRAM_CREATION_FAILURE = NVRTC_ERROR_PROGRAM_CREATION_FAILURE (empty)

92 ERROR_INVALID_INPUT = NVRTC_ERROR_INVALID_INPUT (empty)

93 ERROR_INVALID_PROGRAM = NVRTC_ERROR_INVALID_PROGRAM (empty)

94 ERROR_INVALID_OPTION = NVRTC_ERROR_INVALID_OPTION (empty)

95 ERROR_COMPILATION = NVRTC_ERROR_COMPILATION (empty)

96 ERROR_BUILTIN_OPERATION_FAILURE = NVRTC_ERROR_BUILTIN_OPERATION_FAILURE (empty)

97 ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION (empty)

98 ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION (empty)

99 ERROR_NAME_EXPRESSION_NOT_VALID = NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID (empty)

100 ERROR_INTERNAL_ERROR = NVRTC_ERROR_INTERNAL_ERROR (empty)

101 ERROR_TIME_FILE_WRITE_FAILED = NVRTC_ERROR_TIME_FILE_WRITE_FAILED (empty)

102 ERROR_NO_PCH_CREATE_ATTEMPTED = NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED (empty)

103 ERROR_PCH_CREATE_HEAP_EXHAUSTED = NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED (empty)

104 ERROR_PCH_CREATE = NVRTC_ERROR_PCH_CREATE (empty)

105 ERROR_CANCELLED = NVRTC_ERROR_CANCELLED (empty)

106 ERROR_TIME_TRACE_FILE_WRITE_FAILED = NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED (empty)

107 ERROR_BUSY = NVRTC_ERROR_BUSY (empty)

108  

109  

110class InstallHeadersFlag(_FastEnum): (empty)

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

112 SKIP_IF_EXISTS = ( 

113 0x0, (empty)

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

119 "Clear existing directory contents before installation. " 

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

121 ) 

122 NO_WAIT = ( 

123 0x2, (empty)

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 ) 

129  

130  

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

132# Error handling 

133############################################################################### 

134  

135  

136class NvrtcError(Exception): (empty)

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

138 self.status = status 1 ctx1D

139 s = get_error_string(status) 1 ctx1D

140 super(NvrtcError, self).__init__(s) 1 ctx1D

141  

142 def __reduce__(self): (empty)

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

144  

145class OutOfMemoryError(NvrtcError): (empty)

146 pass (empty)

147class ProgramCreationFailureError(NvrtcError): (empty)

148 pass (empty)

149class InvalidInputError(NvrtcError): (empty)

150 pass (empty)

151class InvalidProgramError(NvrtcError): (empty)

152 pass (empty)

153class InvalidOptionError(NvrtcError): (empty)

154 pass (empty)

155class CompilationError(NvrtcError): (empty)

156 pass (empty)

157class BuiltinOperationFailureError(NvrtcError): (empty)

158 pass (empty)

159class NoNameExpressionsAfterCompilationError(NvrtcError): (empty)

160 pass (empty)

161class NoLoweredNamesBeforeCompilationError(NvrtcError): (empty)

162 pass (empty)

163class NameExpressionNotValidError(NvrtcError): (empty)

164 pass (empty)

165class InternalErrorError(NvrtcError): (empty)

166 pass (empty)

167class TimeFileWriteFailedError(NvrtcError): (empty)

168 pass (empty)

169class NoPchCreateAttemptedError(NvrtcError): (empty)

170 pass (empty)

171class PchCreateHeapExhaustedError(NvrtcError): (empty)

172 pass (empty)

173class PchCreateError(NvrtcError): (empty)

174 pass (empty)

175class CancelledError(NvrtcError): (empty)

176 pass (empty)

177class TimeTraceFileWriteFailedError(NvrtcError): (empty)

178 pass (empty)

179class BusyError(NvrtcError): (empty)

180 pass (empty)

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

182 cdef object pystatus = status 1 ctx1D

183 if status == 1: 1 ctx1D

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) 1 ctx1D

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

220  

221InternalError = InternalErrorError (empty)

222  

223  

224@cython.profile(False) (empty)

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

226 if status != 0: 31 ctx1abcdefghijklmnopqrstuvwxyzABCDE

227 with gil: 1 ctx1D

228 raise _nvrtc_error_factory(status) 1 ctx1D

229 return status 30 ctx1abcdefghijklmnopqrstuvwxyzABCE

230  

231  

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

233# POD definitions 

234############################################################################### 

235  

236cdef _get_bundled_headers_info_dtype_offsets(): (empty)

237 cdef nvrtcBundledHeadersInfo pod 

238 return _numpy.dtype({ (empty)

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

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

241 'offsets': [ (empty)

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

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

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

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

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

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

248 ], 

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

250 }) 

251  

252bundled_headers_info_dtype = _get_bundled_headers_info_dtype_offsets() (empty)

253  

254cdef class BundledHeadersInfo: 

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

256  

257  

258 .. seealso:: `nvrtcBundledHeadersInfo` 

259 """ 

260 cdef: 

261 nvrtcBundledHeadersInfo _data 

262 nvrtcBundledHeadersInfo *_ptr 

263 object _owner 

264 bint _readonly 

265  

266 def __init__(self): 

267 self._ptr = &self._data 1 ctx1C

268 self._owner = None 1 ctx1C

269 self._readonly = False 1 ctx1C

270  

271 def __repr__(self): 

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

273  

274 @property 

275 def ptr(self): 

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

277 return <intptr_t>(self._ptr) 

278  

279 cdef intptr_t _get_ptr(self): (empty)

280 return <intptr_t>(self._ptr) 1 ctx1C

281  

282 def __int__(self): 

283 return <intptr_t>(self._ptr) 

284  

285 def __eq__(self, other): 

286 cdef BundledHeadersInfo other_ 

287 if not isinstance(other, BundledHeadersInfo): 

288 return False 

289 other_ = other 

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

291  

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

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

294  

295 def __releasebuffer__(self, Py_buffer *buffer): 

296 pass 

297  

298 def __setitem__(self, key, val): 

299 cdef _cyb_Py_buffer view 

300 if key == 0: 

301 if self._readonly: 

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

303 _cyb_PyObject_GetBuffer(val, &view, _cyb_PyBUF_SIMPLE) 

304 try: 

305 if <size_t>view.len < sizeof(nvrtcBundledHeadersInfo): 

306 raise ValueError( 

307 "source buffer too small: expected at least %d bytes, got %d" % (sizeof(nvrtcBundledHeadersInfo), view.len) 

308 ) 

309 _cyb_memmove(<void*>&self._data, view.buf, sizeof(nvrtcBundledHeadersInfo)) 

310 self._ptr = &self._data 

311 self._owner = None 

312 self._readonly = view.readonly != 0 

313 finally: 

314 _cyb_PyBuffer_Release(&view) 

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 1 ctx1C

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 1 ctx1C

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 1 ctx1C

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 1 ctx1C

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 1 ctx1C

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 1 ctx1C

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 _cyb_memcpy(<void*>&obj._data, <void*>ptr, sizeof(nvrtcBundledHeadersInfo)) 

412 obj._ptr = &obj._data 

413 obj._owner = None 

414 else: 

415 obj._ptr = <nvrtcBundledHeadersInfo *>ptr 

416 obj._owner = owner 

417 obj._readonly = readonly 

418 return obj 

419  

420  

421############################################################################### 

422# Wrapper functions 

423############################################################################### 

424  

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

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

427  

428 Args: 

429 src (bytes): CUDA program source. 

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

431 to be used. 

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

433 list (no headers). 

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

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

436  

437 Returns: 

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

439  

440 .. seealso:: `nvrtcCreateProgram` 

441 """ 

442 if headers is None: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

443 headers = [] 27 ctx1bcdefghijklmnopqrstuvwxyzAB

444 if include_names is None: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

445 include_names = [] 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

447 raise ValueError( 

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

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

450 ) 

451 cdef int num_headers = len(headers) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

452 cdef Program prog 

453 cdef const char* c_src = src 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

455 cdef const char* c_name = _name 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

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

458 cdef const char** hdr_data = NULL 27 ctx1bcdefghijklmnopqrstuvwxyzAB

459 cdef const char** inc_data = NULL 27 ctx1bcdefghijklmnopqrstuvwxyzAB

460 if num_headers: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

461 hdr_data = cy_headers.data() 

462 inc_data = cy_include_names.data() 

463 with nogil: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

465 check_status(__status__) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

466 return <intptr_t>prog 27 ctx1bcdefghijklmnopqrstuvwxyzAB

467  

468  

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

470 """nvrtcCompileProgram compiles the given program. 

471  

472 It supports compile options listed in Supported Compile Options. 

473  

474 Args: 

475 prog (intptr_t): CUDA Runtime Compilation program. 

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

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

478  

479 .. seealso:: `nvrtcCompileProgram` 

480 """ 

481 if options is None: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

482 options = [] 

483 cdef int num_options = len(options) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

485 cdef const char** opt_data = NULL 27 ctx1bcdefghijklmnopqrstuvwxyzAB

486 if num_options: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

487 opt_data = cy_options.data() 27 ctx1bcdefghijklmnopqrstuvwxyzAB

488 with nogil: 27 ctx1bcdefghijklmnopqrstuvwxyzAB

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

490 check_status(__status__) 27 ctx1bcdefghijklmnopqrstuvwxyzAB

491  

492  

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

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

495  

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

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

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

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

500  

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

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

503  

504 Args: 

505 prog (intptr_t): CUDA Runtime Compilation program. 

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

507 or 0 to clear. 

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

509  

510 .. seealso:: `nvrtcSetFlowCallback` 

511 """ 

512 with nogil: 

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

514 check_status(__status__) 

515  

516  

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

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

519  

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

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

522 provided to :func:`add_name_expression`. 

523  

524 Args: 

525 prog (intptr_t): CUDA Runtime Compilation program. 

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

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

528  

529 Returns: 

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

531  

532 .. seealso:: `nvrtcGetLoweredName` 

533 """ 

534 cdef const char* c_name_expression = name_expression 1 ctx1D

535 cdef const char* lowered_name = NULL 1 ctx1D

536 with nogil: 1 ctx1D

537 __status__ = nvrtcGetLoweredName(<Program>prog, c_name_expression, &lowered_name) 1 ctx1D

538 check_status(__status__) 1 ctx1D

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

540  

541  

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

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

544  

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

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

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

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

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

550 using file locking. 

551  

552 Args: 

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

554 The directory is created if it does not exist. 

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

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

557  

558 Returns: 

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

560  

561 .. seealso:: `nvrtcInstallBundledHeaders` 

562 """ 

563 cdef const char* c_install_path = install_path 

564 cdef const char* error_log = NULL 

565 with nogil: 

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

567 check_status(__status__) 

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

569  

570  

571cpdef tuple get_bundled_headers_info(): (empty)

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

573  

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

575 before calling :func:`install_bundled_headers`. 

576  

577 Returns: 

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

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

580  

581 .. seealso:: `nvrtcGetBundledHeadersInfo` 

582 """ 

583 cdef BundledHeadersInfo info = BundledHeadersInfo() 1 ctx1C

584 cdef nvrtcBundledHeadersInfo* c_info = <nvrtcBundledHeadersInfo*><intptr_t>(info._get_ptr()) 1 ctx1C

585 cdef const char* error_log = NULL 1 ctx1C

586 with nogil: 1 ctx1C

587 __status__ = nvrtcGetBundledHeadersInfo(c_info, &error_log) 1 ctx1C

588 check_status(__status__) 1 ctx1C

589 return info, (<bytes>error_log if error_log != NULL else None) 1 ctx1C

590  

591  

592cpdef remove_bundled_headers(bytes install_path): (empty)

593 """nvrtcRemoveBundledHeaders removes previously installed bundled headers. 

594  

595 Recursively removes all files and subdirectories within the installation 

596 directory to help manage disk space. 

597  

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

599 installed by NVRTC. Use with caution. 

600  

601 Args: 

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

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

604  

605 Returns: 

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

607  

608 .. seealso:: `nvrtcRemoveBundledHeaders` 

609 """ 

610 cdef const char* c_install_path = install_path 

611 cdef const char* error_log = NULL 

612 with nogil: 

613 __status__ = nvrtcRemoveBundledHeaders(c_install_path, &error_log) 

614 check_status(__status__) 

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

616  

617  

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

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

620  

621  

622 Args: 

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

624  

625 .. seealso:: `nvrtcGetErrorString` 

626 """ 

627 cdef const char *_output_cstr_ 

628 cdef bytes _output_ 

629 with nogil: 1 ctx1D

630 _output_cstr_ = nvrtcGetErrorString(<_Result>result) 1 ctx1D

631 _output_ = _output_cstr_ 1 ctx1D

632 return _output_.decode() 1 ctx1D

633  

634  

635cpdef tuple version(): (empty)

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

637  

638  

639 Returns: 

640 A 2-tuple containing: 

641 - int: CUDA Runtime Compilation major version number. 

642 - int: CUDA Runtime Compilation minor version number. 

643  

644 .. seealso:: `nvrtcVersion` 

645 """ 

646 cdef int major 

647 cdef int minor 

648 with nogil: 10 ctx1abcdefghij

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

650 check_status(__status__) 10 ctx1abcdefghij

651 return (major, minor) 10 ctx1abcdefghij

652  

653  

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

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

656  

657 see ``nvrtcGetSupportedArchs``. 

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 see ``nvrtcGetNumSupportedArchs``. 

675  

676 Returns: 

677 int: sorted array of supported architectures. 

678  

679 .. seealso:: `nvrtcGetSupportedArchs` 

680 """ 

681 cdef int numArchs 

682 with nogil: 1 ctx1E

683 __status__ = nvrtcGetNumSupportedArchs(&numArchs) 1 ctx1E

684 check_status(__status__) 1 ctx1E

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

686 cdef int *supported_archs_ptr = <int *>(_supported_archs_alloc_.data) 1 ctx1E

687 cdef object supported_archs = _supported_archs_alloc_[:numArchs] 1 ctx1E

688 if numArchs != 0: 1 ctx1E

689 with nogil: 1 ctx1E

690 __status__ = nvrtcGetSupportedArchs(supported_archs_ptr) 1 ctx1E

691 check_status(__status__) 1 ctx1E

692 return supported_archs 1 ctx1E

693  

694  

695cpdef destroy_program(intptr_t prog): (empty)

696 """nvrtcDestroyProgram destroys the given program. 

697  

698  

699 Args: 

700 prog (intptr_t): CUDA Runtime Compilation program. 

701  

702 .. seealso:: `nvrtcDestroyProgram` 

703 """ 

704 cdef Program _prog_ = <Program>prog 

705 with nogil: 

706 __status__ = nvrtcDestroyProgram(&_prog_) 

707 check_status(__status__) 

708  

709  

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

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

712  

713  

714 Args: 

715 prog (intptr_t): CUDA Runtime Compilation program. 

716  

717 Returns: 

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

719 ``NULL``). 

720  

721 .. seealso:: `nvrtcGetPTXSize` 

722 """ 

723 cdef size_t ptx_size_ret 

724 with nogil: 

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

726 check_status(__status__) 

727 return ptx_size_ret 

728  

729  

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

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

732  

733  

734 Args: 

735 prog (intptr_t): CUDA Runtime Compilation program. 

736  

737 Returns: 

738 char: Compiled result. 

739  

740 .. seealso:: `nvrtcGetPTX` 

741 """ 

742 cdef size_t ptxSizeRet 

743 with nogil: 

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

745 check_status(__status__) 

746 cdef bytes _ptx_ = bytes(ptxSizeRet) 

747 cdef char* ptx = _ptx_ 

748 if ptxSizeRet != 0: 

749 with nogil: 

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

751 check_status(__status__) 

752 return _ptx_ 

753  

754  

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

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

757  

758  

759 Args: 

760 prog (intptr_t): CUDA Runtime Compilation program. 

761  

762 Returns: 

763 size_t: Size of the generated cubin. 

764  

765 .. seealso:: `nvrtcGetCUBINSize` 

766 """ 

767 cdef size_t cubin_size_ret 

768 with nogil: 

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

770 check_status(__status__) 

771 return cubin_size_ret 

772  

773  

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

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

776  

777  

778 Args: 

779 prog (intptr_t): CUDA Runtime Compilation program. 

780  

781 Returns: 

782 char: Compiled and assembled result. 

783  

784 .. seealso:: `nvrtcGetCUBIN` 

785 """ 

786 cdef size_t cubinSizeRet 

787 with nogil: 15 ctx1bcdefghijopqrst

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

789 check_status(__status__) 15 ctx1bcdefghijopqrst

790 cdef bytes _cubin_ = bytes(cubinSizeRet) 15 ctx1bcdefghijopqrst

791 cdef char* cubin = _cubin_ 15 ctx1bcdefghijopqrst

792 if cubinSizeRet != 0: 15 ctx1bcdefghijopqrst

793 with nogil: 15 ctx1bcdefghijopqrst

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

795 check_status(__status__) 15 ctx1bcdefghijopqrst

796 return _cubin_ 15 ctx1bcdefghijopqrst

797  

798  

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

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

801  

802  

803 Args: 

804 prog (intptr_t): CUDA Runtime Compilation program. 

805  

806 Returns: 

807 size_t: Size of the generated LTO IR. 

808  

809 .. seealso:: `nvrtcGetLTOIRSize` 

810 """ 

811 cdef size_t ltoir_size_ret 

812 with nogil: 

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

814 check_status(__status__) 

815 return ltoir_size_ret 

816  

817  

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

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

820  

821  

822 Args: 

823 prog (intptr_t): CUDA Runtime Compilation program. 

824  

825 Returns: 

826 char: Compiled result. 

827  

828 .. seealso:: `nvrtcGetLTOIR` 

829 """ 

830 cdef size_t LTOIRSizeRet 

831 with nogil: 12 ctx1klmnuvwxyzAB

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

833 check_status(__status__) 12 ctx1klmnuvwxyzAB

834 cdef bytes _ltoir_ = bytes(LTOIRSizeRet) 12 ctx1klmnuvwxyzAB

835 cdef char* ltoir = _ltoir_ 12 ctx1klmnuvwxyzAB

836 if LTOIRSizeRet != 0: 12 ctx1klmnuvwxyzAB

837 with nogil: 12 ctx1klmnuvwxyzAB

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

839 check_status(__status__) 12 ctx1klmnuvwxyzAB

840 return _ltoir_ 12 ctx1klmnuvwxyzAB

841  

842  

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

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

845  

846  

847 Args: 

848 prog (intptr_t): CUDA Runtime Compilation program. 

849  

850 Returns: 

851 size_t: Size of the generated LTO IR. 

852  

853 .. seealso:: `nvrtcGetOptiXIRSize` 

854 """ 

855 cdef size_t optixir_size_ret 

856 with nogil: 

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

858 check_status(__status__) 

859 return optixir_size_ret 

860  

861  

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

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

864  

865  

866 Args: 

867 prog (intptr_t): CUDA Runtime Compilation program. 

868  

869 Returns: 

870 char: Optix IR Compiled result. 

871  

872 .. seealso:: `nvrtcGetOptiXIR` 

873 """ 

874 cdef size_t optixirSizeRet 

875 with nogil: 

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

877 check_status(__status__) 

878 cdef bytes _optixir_ = bytes(optixirSizeRet) 

879 cdef char* optixir = _optixir_ 

880 if optixirSizeRet != 0: 

881 with nogil: 

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

883 check_status(__status__) 

884 return _optixir_ 

885  

886  

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

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

889  

890 Note that compilation log may be generated with warnings and informative 

891 messages, even when the compilation of ``prog`` succeeds. 

892  

893 Args: 

894 prog (intptr_t): CUDA Runtime Compilation program. 

895  

896 Returns: 

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

898 ``NULL``). 

899  

900 .. seealso:: `nvrtcGetProgramLogSize` 

901 """ 

902 cdef size_t log_size_ret 

903 with nogil: 

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

905 check_status(__status__) 

906 return log_size_ret 

907  

908  

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

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

911  

912  

913 Args: 

914 prog (intptr_t): CUDA Runtime Compilation program. 

915  

916 Returns: 

917 char: Compilation log. 

918  

919 .. seealso:: `nvrtcGetProgramLog` 

920 """ 

921 cdef size_t logSizeRet 

922 with nogil: 9 ctx1bcdefghij

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

924 check_status(__status__) 9 ctx1bcdefghij

925 cdef bytes _log_ = bytes(logSizeRet) 9 ctx1bcdefghij

926 cdef char* log = _log_ 9 ctx1bcdefghij

927 if logSizeRet != 0: 9 ctx1bcdefghij

928 with nogil: 9 ctx1bcdefghij

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

930 check_status(__status__) 9 ctx1bcdefghij

931 return _log_ 9 ctx1bcdefghij

932  

933  

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

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

936  

937 The identical name expression string must be provided on a subsequent call to 

938 nvrtcGetLoweredName to extract the lowered name. 

939  

940 Args: 

941 prog (intptr_t): CUDA Runtime Compilation program. 

942 name_expression (str): constant expression denoting the address of 

943 a global function or device/__constant__ variable. 

944  

945 .. seealso:: `nvrtcAddNameExpression` 

946 """ 

947 if not isinstance(name_expression, str): 

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

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

950 cdef char* _name_expression_ = _temp_name_expression_ 

951 with nogil: 

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

953 check_status(__status__) 

954  

955  

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

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

958  

959  

960 Returns: 

961 size_t: pointer to location where the size of the PCH Heap will be 

962 stored. 

963  

964 .. seealso:: `nvrtcGetPCHHeapSize` 

965 """ 

966 cdef size_t ret 

967 with nogil: 

968 __status__ = nvrtcGetPCHHeapSize(&ret) 

969 check_status(__status__) 

970 return ret 

971  

972  

973cpdef set_pch_heap_size(size_t size): (empty)

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

975  

976 The requested size may be rounded up to a platform dependent alignment (e.g. 

977 page size). If the PCH Heap has already been allocated, the heap memory will be 

978 freed and a new PCH Heap will be allocated. 

979  

980 Args: 

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

982  

983 .. seealso:: `nvrtcSetPCHHeapSize` 

984 """ 

985 with nogil: 

986 __status__ = nvrtcSetPCHHeapSize(size) 

987 check_status(__status__) 

988  

989  

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

991 """returns the PCH creation status. 

992  

993 NVRTC_SUCCESS indicates that the PCH was successfully created. 

994 NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED indicates that no PCH creation was 

995 attempted, either because PCH functionality was not requested during the 

996 preceding nvrtcCompileProgram call, or automatic PCH processing was requested, 

997 and compiler chose not to create a PCH file. 

998 NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED indicates that a PCH file could 

999 potentially have been created, but the compiler ran out space in the PCH heap. 

1000 In this scenario, the :func:`get_pch_heap_size_required` can be used to query 

1001 the required heap size, the heap can be reallocated for this size with 

1002 :func:`set_pch_heap_size` and PCH creation may be reattempted again invoking 

1003 :func:`compile_program` with a new NVRTC program instance. 

1004 NVRTC_ERROR_PCH_CREATE indicates that an error condition prevented the PCH file 

1005 from being created. 

1006  

1007 Args: 

1008 prog (intptr_t): CUDA Runtime Compilation program. 

1009  

1010 .. seealso:: `nvrtcGetPCHCreateStatus` 

1011 """ 

1012 cdef int ret 

1013 with nogil: 

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

1015 return ret 

1016  

1017  

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

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

1020  

1021  

1022 Args: 

1023 prog (intptr_t): CUDA Runtime Compilation program. 

1024  

1025 Returns: 

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

1027 Heap will be stored. 

1028  

1029 .. seealso:: `nvrtcGetPCHHeapSizeRequired` 

1030 """ 

1031 cdef size_t size 

1032 with nogil: 

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

1034 check_status(__status__) 

1035 return size 

1036  

1037  

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

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

1040  

1041  

1042 Args: 

1043 prog (intptr_t): CUDA Runtime Compilation program. 

1044  

1045 Returns: 

1046 size_t: Size of the generated cuda_tile IR. 

1047  

1048 .. seealso:: `nvrtcGetTileIRSize` 

1049 """ 

1050 cdef size_t tile_ir_size_ret 

1051 with nogil: 

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

1053 check_status(__status__) 

1054 return tile_ir_size_ret 

1055  

1056  

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

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

1059  

1060  

1061 Args: 

1062 prog (intptr_t): CUDA Runtime Compilation program. 

1063  

1064 Returns: 

1065 char: Generated cuda_tile IR. 

1066  

1067 .. seealso:: `nvrtcGetTileIR` 

1068 """ 

1069 cdef size_t TileIRSizeRet 

1070 with nogil: 

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

1072 check_status(__status__) 

1073 cdef bytes _tile_ir_ = bytes(TileIRSizeRet) 

1074 cdef char* tile_ir = _tile_ir_ 

1075 if TileIRSizeRet != 0: 

1076 with nogil: 

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

1078 check_status(__status__) 

1079 return _tile_ir_ 

1080  

1081  

1082del _cyb_FastEnum (empty)