Coverage for cuda / bindings / nvrtc.pyx: 63%
302 statements
« prev ^ index » next coverage.py v7.13.0, created at 2025-12-10 01:19 +0000
1# SPDX-FileCopyrightText: Copyright (c) 2021-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
2# SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE
4# This code was automatically generated with version 13.1.0. Do not modify it directly.
5from typing import Any, Optional
6from enum import IntEnum
7import cython
8import ctypes
9from libc.stdlib cimport calloc, malloc, free
10from libc cimport string
11from libc.stdint cimport int32_t, uint32_t, int64_t, uint64_t, uintptr_t
12from libc.stddef cimport wchar_t
13from libc.limits cimport CHAR_MIN
14from libcpp.vector cimport vector
15from cpython.buffer cimport PyObject_CheckBuffer, PyObject_GetBuffer, PyBuffer_Release, PyBUF_SIMPLE, PyBUF_ANY_CONTIGUOUS
16from cpython.bytes cimport PyBytes_FromStringAndSize
18import cuda.bindings.driver as _driver
19_driver = _driver.__dict__
20include "_lib/utils.pxi"
22ctypedef unsigned long long signed_char_ptr
23ctypedef unsigned long long unsigned_char_ptr
24ctypedef unsigned long long char_ptr
25ctypedef unsigned long long short_ptr
26ctypedef unsigned long long unsigned_short_ptr
27ctypedef unsigned long long int_ptr
28ctypedef unsigned long long long_int_ptr
29ctypedef unsigned long long long_long_int_ptr
30ctypedef unsigned long long unsigned_int_ptr
31ctypedef unsigned long long unsigned_long_int_ptr
32ctypedef unsigned long long unsigned_long_long_int_ptr
33ctypedef unsigned long long uint32_t_ptr
34ctypedef unsigned long long uint64_t_ptr
35ctypedef unsigned long long int32_t_ptr
36ctypedef unsigned long long int64_t_ptr
37ctypedef unsigned long long unsigned_ptr
38ctypedef unsigned long long unsigned_long_long_ptr
39ctypedef unsigned long long long_long_ptr
40ctypedef unsigned long long size_t_ptr
41ctypedef unsigned long long long_ptr
42ctypedef unsigned long long float_ptr
43ctypedef unsigned long long double_ptr
44ctypedef unsigned long long void_ptr
47class nvrtcResult(IntEnum):
48 """
49 The enumerated type nvrtcResult defines API call result codes.
50 NVRTC API functions return nvrtcResult to indicate the call result.
51 """
52 NVRTC_SUCCESS = cynvrtc.nvrtcResult.NVRTC_SUCCESS
53 NVRTC_ERROR_OUT_OF_MEMORY = cynvrtc.nvrtcResult.NVRTC_ERROR_OUT_OF_MEMORY
54 NVRTC_ERROR_PROGRAM_CREATION_FAILURE = cynvrtc.nvrtcResult.NVRTC_ERROR_PROGRAM_CREATION_FAILURE
55 NVRTC_ERROR_INVALID_INPUT = cynvrtc.nvrtcResult.NVRTC_ERROR_INVALID_INPUT
56 NVRTC_ERROR_INVALID_PROGRAM = cynvrtc.nvrtcResult.NVRTC_ERROR_INVALID_PROGRAM
57 NVRTC_ERROR_INVALID_OPTION = cynvrtc.nvrtcResult.NVRTC_ERROR_INVALID_OPTION
58 NVRTC_ERROR_COMPILATION = cynvrtc.nvrtcResult.NVRTC_ERROR_COMPILATION
59 NVRTC_ERROR_BUILTIN_OPERATION_FAILURE = cynvrtc.nvrtcResult.NVRTC_ERROR_BUILTIN_OPERATION_FAILURE
60 NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = cynvrtc.nvrtcResult.NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION
61 NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = cynvrtc.nvrtcResult.NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION
62 NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID = cynvrtc.nvrtcResult.NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID
63 NVRTC_ERROR_INTERNAL_ERROR = cynvrtc.nvrtcResult.NVRTC_ERROR_INTERNAL_ERROR
64 NVRTC_ERROR_TIME_FILE_WRITE_FAILED = cynvrtc.nvrtcResult.NVRTC_ERROR_TIME_FILE_WRITE_FAILED
65 NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED = cynvrtc.nvrtcResult.NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED
66 NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED = cynvrtc.nvrtcResult.NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED
67 NVRTC_ERROR_PCH_CREATE = cynvrtc.nvrtcResult.NVRTC_ERROR_PCH_CREATE
68 NVRTC_ERROR_CANCELLED = cynvrtc.nvrtcResult.NVRTC_ERROR_CANCELLED
69 NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED = cynvrtc.nvrtcResult.NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED
71_dict_nvrtcResult = dict(((int(v), v) for k, v in nvrtcResult.__members__.items()))
73cdef class nvrtcProgram:
74 """ nvrtcProgram is the unit of compilation, and an opaque handle for a program.
76 To compile a CUDA program string, an instance of nvrtcProgram must be created first with nvrtcCreateProgram, then compiled with nvrtcCompileProgram.
78 Methods
79 -------
80 getPtr()
81 Get memory address of class instance
83 """
84 def __cinit__(self, void_ptr init_value = 0, void_ptr _ptr = 0):
85 if _ptr == 0:
86 self._pvt_ptr = &self._pvt_val
87 self._pvt_ptr[0] = <cynvrtc.nvrtcProgram>init_value
88 else:
89 self._pvt_ptr = <cynvrtc.nvrtcProgram *>_ptr
90 def __init__(self, *args, **kwargs):
91 pass
92 def __repr__(self):
93 return '<nvrtcProgram ' + str(hex(self.__int__())) + '>'
94 def __index__(self):
95 return self.__int__()
96 def __eq__(self, other):
97 if not isinstance(other, nvrtcProgram):
98 return False
99 return self._pvt_ptr[0] == (<nvrtcProgram>other)._pvt_ptr[0]
100 def __hash__(self):
101 return hash(<uintptr_t><void*>(self._pvt_ptr[0]))
102 def __int__(self):
103 return <void_ptr>self._pvt_ptr[0]
104 def getPtr(self):
105 return <void_ptr>self._pvt_ptr
107@cython.embedsignature(True)
108def nvrtcGetErrorString(result not None : nvrtcResult):
109 """ 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"`.
111 Parameters
112 ----------
113 result : :py:obj:`~.nvrtcResult`
114 CUDA Runtime Compilation API result code.
116 Returns
117 -------
118 nvrtcResult.NVRTC_SUCCESS
119 nvrtcResult.NVRTC_SUCCESS
120 bytes
121 Message string for the given :py:obj:`~.nvrtcResult` code.
122 """
123 cdef cynvrtc.nvrtcResult cyresult = result.value
124 with nogil:
125 err = cynvrtc.nvrtcGetErrorString(cyresult)
126 return (nvrtcResult.NVRTC_SUCCESS, err)
128@cython.embedsignature(True)
129def nvrtcVersion():
130 """ nvrtcVersion sets the output parameters `major` and `minor` with the CUDA Runtime Compilation version number.
132 Returns
133 -------
134 nvrtcResult
135 - :py:obj:`~.NVRTC_SUCCESS`
136 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
137 major : int
138 CUDA Runtime Compilation major version number.
139 minor : int
140 CUDA Runtime Compilation minor version number.
141 """
142 cdef int major = 0
143 cdef int minor = 0
144 with nogil:
145 err = cynvrtc.nvrtcVersion(&major, &minor)
146 if err != cynvrtc.NVRTC_SUCCESS:
147 return (_dict_nvrtcResult[err], None, None)
148 return (_dict_nvrtcResult[err], major, minor)
150@cython.embedsignature(True)
151def nvrtcGetNumSupportedArchs():
152 """ nvrtcGetNumSupportedArchs sets the output parameter `numArchs` with the number of architectures supported by NVRTC. This can then be used to pass an array to :py:obj:`~.nvrtcGetSupportedArchs` to get the supported architectures.
154 see :py:obj:`~.nvrtcGetSupportedArchs`
156 Returns
157 -------
158 nvrtcResult
159 - :py:obj:`~.NVRTC_SUCCESS`
160 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
161 numArchs : int
162 number of supported architectures.
163 """
164 cdef int numArchs = 0
165 with nogil:
166 err = cynvrtc.nvrtcGetNumSupportedArchs(&numArchs)
167 if err != cynvrtc.NVRTC_SUCCESS:
168 return (_dict_nvrtcResult[err], None)
169 return (_dict_nvrtcResult[err], numArchs)
171@cython.embedsignature(True)
172def nvrtcGetSupportedArchs():
173 """ nvrtcGetSupportedArchs populates the array passed via the output parameter `supportedArchs` 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 :py:obj:`~.nvrtcGetNumSupportedArchs`.
175 see :py:obj:`~.nvrtcGetNumSupportedArchs`
177 Returns
178 -------
179 nvrtcResult
180 - :py:obj:`~.NVRTC_SUCCESS`
181 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
182 supportedArchs : list[int]
183 sorted array of supported architectures.
184 """
185 cdef vector[int] supportedArchs
186 _, s = nvrtcGetNumSupportedArchs()
187 supportedArchs.resize(s)
189 with nogil:
190 err = cynvrtc.nvrtcGetSupportedArchs(supportedArchs.data())
191 if err != cynvrtc.NVRTC_SUCCESS:
192 return (_dict_nvrtcResult[err], None)
193 return (_dict_nvrtcResult[err], supportedArchs)
195@cython.embedsignature(True)
196def nvrtcCreateProgram(char* src, char* name, int numHeaders, headers : Optional[tuple[bytes] | list[bytes]], includeNames : Optional[tuple[bytes] | list[bytes]]):
197 """ nvrtcCreateProgram creates an instance of nvrtcProgram with the given input parameters, and sets the output parameter `prog` with it.
199 Parameters
200 ----------
201 src : bytes
202 CUDA program source.
203 name : bytes
204 CUDA program name. `name` can be `NULL`; `"default_program"` is
205 used when `name` is `NULL` or "".
206 numHeaders : int
207 Number of headers used. `numHeaders` must be greater than or equal
208 to 0.
209 headers : list[bytes]
210 Sources of the headers. `headers` can be `NULL` when `numHeaders`
211 is 0.
212 includeNames : list[bytes]
213 Name of each header by which they can be included in the CUDA
214 program source. `includeNames` can be `NULL` when `numHeaders` is
215 0. These headers must be included with the exact names specified
216 here.
218 Returns
219 -------
220 nvrtcResult
221 - :py:obj:`~.NVRTC_SUCCESS`
222 - :py:obj:`~.NVRTC_ERROR_OUT_OF_MEMORY`
223 - :py:obj:`~.NVRTC_ERROR_PROGRAM_CREATION_FAILURE`
224 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
225 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
226 prog : :py:obj:`~.nvrtcProgram`
227 CUDA Runtime Compilation program.
229 See Also
230 --------
231 :py:obj:`~.nvrtcDestroyProgram`
232 """
233 includeNames = [] if includeNames is None else includeNames
234 if not all(isinstance(_x, (bytes)) for _x in includeNames):
235 raise TypeError("Argument 'includeNames' is not instance of type (expected tuple[bytes] or list[bytes]")
236 headers = [] if headers is None else headers
237 if not all(isinstance(_x, (bytes)) for _x in headers):
238 raise TypeError("Argument 'headers' is not instance of type (expected tuple[bytes] or list[bytes]")
239 cdef nvrtcProgram prog = nvrtcProgram()
240 if numHeaders > len(headers): raise RuntimeError("List is too small: " + str(len(headers)) + " < " + str(numHeaders))
241 if numHeaders > len(includeNames): raise RuntimeError("List is too small: " + str(len(includeNames)) + " < " + str(numHeaders))
242 cdef vector[const char*] cyheaders = headers
243 cdef vector[const char*] cyincludeNames = includeNames
244 with nogil:
245 err = cynvrtc.nvrtcCreateProgram(<cynvrtc.nvrtcProgram*>prog._pvt_ptr, src, name, numHeaders, cyheaders.data(), cyincludeNames.data())
246 if err != cynvrtc.NVRTC_SUCCESS:
247 return (_dict_nvrtcResult[err], None)
248 return (_dict_nvrtcResult[err], prog)
250@cython.embedsignature(True)
251def nvrtcDestroyProgram(prog):
252 """ nvrtcDestroyProgram destroys the given program.
254 Parameters
255 ----------
256 prog : :py:obj:`~.nvrtcProgram`
257 CUDA Runtime Compilation program.
259 Returns
260 -------
261 nvrtcResult
262 - :py:obj:`~.NVRTC_SUCCESS`
263 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
265 See Also
266 --------
267 :py:obj:`~.nvrtcCreateProgram`
268 """
269 cdef cynvrtc.nvrtcProgram *cyprog
270 if prog is None:
271 cyprog = <cynvrtc.nvrtcProgram*><void_ptr>NULL
272 elif isinstance(prog, (nvrtcProgram,)):
273 pprog = prog.getPtr()
274 cyprog = <cynvrtc.nvrtcProgram*><void_ptr>pprog
275 elif isinstance(prog, (int)):
276 cyprog = <cynvrtc.nvrtcProgram*><void_ptr>prog
277 else:
278 raise TypeError("Argument 'prog' is not instance of type (expected <class 'int, nvrtc.nvrtcProgram'>, found " + str(type(prog)))
279 with nogil:
280 err = cynvrtc.nvrtcDestroyProgram(cyprog)
281 return (_dict_nvrtcResult[err],)
283@cython.embedsignature(True)
284def nvrtcCompileProgram(prog, int numOptions, options : Optional[tuple[bytes] | list[bytes]]):
285 """ nvrtcCompileProgram compiles the given program.
287 It supports compile options listed in :py:obj:`~.Supported Compile
288 Options`.
290 Parameters
291 ----------
292 prog : :py:obj:`~.nvrtcProgram`
293 CUDA Runtime Compilation program.
294 numOptions : int
295 Number of compiler options passed.
296 options : list[bytes]
297 Compiler options in the form of C string array. `options` can be
298 `NULL` when `numOptions` is 0.
300 Returns
301 -------
302 nvrtcResult
303 - :py:obj:`~.NVRTC_SUCCESS`
304 - :py:obj:`~.NVRTC_ERROR_OUT_OF_MEMORY`
305 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
306 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
307 - :py:obj:`~.NVRTC_ERROR_INVALID_OPTION`
308 - :py:obj:`~.NVRTC_ERROR_COMPILATION`
309 - :py:obj:`~.NVRTC_ERROR_BUILTIN_OPERATION_FAILURE`
310 - :py:obj:`~.NVRTC_ERROR_TIME_FILE_WRITE_FAILED`
311 - :py:obj:`~.NVRTC_ERROR_CANCELLED`
312 """
313 options = [] if options is None else options
314 if not all(isinstance(_x, (bytes)) for _x in options):
315 raise TypeError("Argument 'options' is not instance of type (expected tuple[bytes] or list[bytes]")
316 cdef cynvrtc.nvrtcProgram cyprog
317 if prog is None:
318 pprog = 0
319 elif isinstance(prog, (nvrtcProgram,)):
320 pprog = int(prog)
321 else:
322 pprog = int(nvrtcProgram(prog))
323 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
324 if numOptions > len(options): raise RuntimeError("List is too small: " + str(len(options)) + " < " + str(numOptions))
325 cdef vector[const char*] cyoptions = options
326 with nogil:
327 err = cynvrtc.nvrtcCompileProgram(cyprog, numOptions, cyoptions.data())
328 return (_dict_nvrtcResult[err],)
330@cython.embedsignature(True)
331def nvrtcGetPTXSize(prog):
332 """ nvrtcGetPTXSize sets the value of `ptxSizeRet` with the size of the PTX generated by the previous compilation of `prog` (including the trailing `NULL`).
334 Parameters
335 ----------
336 prog : :py:obj:`~.nvrtcProgram`
337 CUDA Runtime Compilation program.
339 Returns
340 -------
341 nvrtcResult
342 - :py:obj:`~.NVRTC_SUCCESS`
343 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
344 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
345 ptxSizeRet : int
346 Size of the generated PTX (including the trailing `NULL`).
348 See Also
349 --------
350 :py:obj:`~.nvrtcGetPTX`
351 """
352 cdef cynvrtc.nvrtcProgram cyprog
353 if prog is None:
354 pprog = 0
355 elif isinstance(prog, (nvrtcProgram,)):
356 pprog = int(prog)
357 else:
358 pprog = int(nvrtcProgram(prog))
359 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
360 cdef size_t ptxSizeRet = 0
361 with nogil:
362 err = cynvrtc.nvrtcGetPTXSize(cyprog, &ptxSizeRet)
363 if err != cynvrtc.NVRTC_SUCCESS:
364 return (_dict_nvrtcResult[err], None)
365 return (_dict_nvrtcResult[err], ptxSizeRet)
367@cython.embedsignature(True)
368def nvrtcGetPTX(prog, char* ptx):
369 """ nvrtcGetPTX stores the PTX generated by the previous compilation of `prog` in the memory pointed by `ptx`.
371 Parameters
372 ----------
373 prog : :py:obj:`~.nvrtcProgram`
374 CUDA Runtime Compilation program.
375 ptx : bytes
376 Compiled result.
378 Returns
379 -------
380 nvrtcResult
381 - :py:obj:`~.NVRTC_SUCCESS`
382 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
383 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
385 See Also
386 --------
387 :py:obj:`~.nvrtcGetPTXSize`
388 """
389 cdef cynvrtc.nvrtcProgram cyprog
390 if prog is None:
391 pprog = 0
392 elif isinstance(prog, (nvrtcProgram,)):
393 pprog = int(prog)
394 else:
395 pprog = int(nvrtcProgram(prog))
396 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
397 with nogil:
398 err = cynvrtc.nvrtcGetPTX(cyprog, ptx)
399 return (_dict_nvrtcResult[err],)
401@cython.embedsignature(True)
402def nvrtcGetCUBINSize(prog):
403 """ nvrtcGetCUBINSize sets the value of `cubinSizeRet` with the size of the cubin generated by the previous compilation of `prog`. The value of cubinSizeRet is set to 0 if the value specified to `-arch` is a virtual architecture instead of an actual architecture.
405 Parameters
406 ----------
407 prog : :py:obj:`~.nvrtcProgram`
408 CUDA Runtime Compilation program.
410 Returns
411 -------
412 nvrtcResult
413 - :py:obj:`~.NVRTC_SUCCESS`
414 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
415 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
416 cubinSizeRet : int
417 Size of the generated cubin.
419 See Also
420 --------
421 :py:obj:`~.nvrtcGetCUBIN`
422 """
423 cdef cynvrtc.nvrtcProgram cyprog
424 if prog is None:
425 pprog = 0
426 elif isinstance(prog, (nvrtcProgram,)):
427 pprog = int(prog)
428 else:
429 pprog = int(nvrtcProgram(prog))
430 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
431 cdef size_t cubinSizeRet = 0
432 with nogil:
433 err = cynvrtc.nvrtcGetCUBINSize(cyprog, &cubinSizeRet)
434 if err != cynvrtc.NVRTC_SUCCESS:
435 return (_dict_nvrtcResult[err], None)
436 return (_dict_nvrtcResult[err], cubinSizeRet)
438@cython.embedsignature(True)
439def nvrtcGetCUBIN(prog, char* cubin):
440 """ 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.
442 Parameters
443 ----------
444 prog : :py:obj:`~.nvrtcProgram`
445 CUDA Runtime Compilation program.
446 cubin : bytes
447 Compiled and assembled result.
449 Returns
450 -------
451 nvrtcResult
452 - :py:obj:`~.NVRTC_SUCCESS`
453 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
454 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
456 See Also
457 --------
458 :py:obj:`~.nvrtcGetCUBINSize`
459 """
460 cdef cynvrtc.nvrtcProgram cyprog
461 if prog is None:
462 pprog = 0
463 elif isinstance(prog, (nvrtcProgram,)):
464 pprog = int(prog)
465 else:
466 pprog = int(nvrtcProgram(prog))
467 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
468 with nogil:
469 err = cynvrtc.nvrtcGetCUBIN(cyprog, cubin)
470 return (_dict_nvrtcResult[err],)
472@cython.embedsignature(True)
473def nvrtcGetLTOIRSize(prog):
474 """ nvrtcGetLTOIRSize sets the value of `LTOIRSizeRet` with the size of the LTO IR generated by the previous compilation of `prog`. The value of LTOIRSizeRet is set to 0 if the program was not compiled with `-dlto`.
476 Parameters
477 ----------
478 prog : :py:obj:`~.nvrtcProgram`
479 CUDA Runtime Compilation program.
481 Returns
482 -------
483 nvrtcResult
484 - :py:obj:`~.NVRTC_SUCCESS`
485 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
486 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
487 LTOIRSizeRet : int
488 Size of the generated LTO IR.
490 See Also
491 --------
492 :py:obj:`~.nvrtcGetLTOIR`
493 """
494 cdef cynvrtc.nvrtcProgram cyprog
495 if prog is None:
496 pprog = 0
497 elif isinstance(prog, (nvrtcProgram,)):
498 pprog = int(prog)
499 else:
500 pprog = int(nvrtcProgram(prog))
501 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
502 cdef size_t LTOIRSizeRet = 0
503 with nogil:
504 err = cynvrtc.nvrtcGetLTOIRSize(cyprog, <OIRSizeRet)
505 if err != cynvrtc.NVRTC_SUCCESS:
506 return (_dict_nvrtcResult[err], None)
507 return (_dict_nvrtcResult[err], LTOIRSizeRet)
509@cython.embedsignature(True)
510def nvrtcGetLTOIR(prog, char* LTOIR):
511 """ 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`.
513 Parameters
514 ----------
515 prog : :py:obj:`~.nvrtcProgram`
516 CUDA Runtime Compilation program.
517 LTOIR : bytes
518 Compiled result.
520 Returns
521 -------
522 nvrtcResult
523 - :py:obj:`~.NVRTC_SUCCESS`
524 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
525 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
527 See Also
528 --------
529 :py:obj:`~.nvrtcGetLTOIRSize`
530 """
531 cdef cynvrtc.nvrtcProgram cyprog
532 if prog is None:
533 pprog = 0
534 elif isinstance(prog, (nvrtcProgram,)):
535 pprog = int(prog)
536 else:
537 pprog = int(nvrtcProgram(prog))
538 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
539 with nogil:
540 err = cynvrtc.nvrtcGetLTOIR(cyprog, LTOIR)
541 return (_dict_nvrtcResult[err],)
543@cython.embedsignature(True)
544def nvrtcGetOptiXIRSize(prog):
545 """ nvrtcGetOptiXIRSize sets the value of `optixirSizeRet` 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.
547 Parameters
548 ----------
549 prog : :py:obj:`~.nvrtcProgram`
550 CUDA Runtime Compilation program.
552 Returns
553 -------
554 nvrtcResult
555 - :py:obj:`~.NVRTC_SUCCESS`
556 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
557 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
558 optixirSizeRet : int
559 Size of the generated LTO IR.
561 See Also
562 --------
563 :py:obj:`~.nvrtcGetOptiXIR`
564 """
565 cdef cynvrtc.nvrtcProgram cyprog
566 if prog is None:
567 pprog = 0
568 elif isinstance(prog, (nvrtcProgram,)):
569 pprog = int(prog)
570 else:
571 pprog = int(nvrtcProgram(prog))
572 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
573 cdef size_t optixirSizeRet = 0
574 with nogil:
575 err = cynvrtc.nvrtcGetOptiXIRSize(cyprog, &optixirSizeRet)
576 if err != cynvrtc.NVRTC_SUCCESS:
577 return (_dict_nvrtcResult[err], None)
578 return (_dict_nvrtcResult[err], optixirSizeRet)
580@cython.embedsignature(True)
581def nvrtcGetOptiXIR(prog, char* optixir):
582 """ 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.
584 Parameters
585 ----------
586 prog : :py:obj:`~.nvrtcProgram`
587 CUDA Runtime Compilation program.
588 optixir : bytes
589 Optix IR Compiled result.
591 Returns
592 -------
593 nvrtcResult
594 - :py:obj:`~.NVRTC_SUCCESS`
595 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
596 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
598 See Also
599 --------
600 :py:obj:`~.nvrtcGetOptiXIRSize`
601 """
602 cdef cynvrtc.nvrtcProgram cyprog
603 if prog is None:
604 pprog = 0
605 elif isinstance(prog, (nvrtcProgram,)):
606 pprog = int(prog)
607 else:
608 pprog = int(nvrtcProgram(prog))
609 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
610 with nogil:
611 err = cynvrtc.nvrtcGetOptiXIR(cyprog, optixir)
612 return (_dict_nvrtcResult[err],)
614@cython.embedsignature(True)
615def nvrtcGetProgramLogSize(prog):
616 """ nvrtcGetProgramLogSize sets `logSizeRet` with the size of the log generated by the previous compilation of `prog` (including the trailing `NULL`).
618 Note that compilation log may be generated with warnings and
619 informative messages, even when the compilation of `prog` succeeds.
621 Parameters
622 ----------
623 prog : :py:obj:`~.nvrtcProgram`
624 CUDA Runtime Compilation program.
626 Returns
627 -------
628 nvrtcResult
629 - :py:obj:`~.NVRTC_SUCCESS`
630 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
631 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
632 logSizeRet : int
633 Size of the compilation log (including the trailing `NULL`).
635 See Also
636 --------
637 :py:obj:`~.nvrtcGetProgramLog`
638 """
639 cdef cynvrtc.nvrtcProgram cyprog
640 if prog is None:
641 pprog = 0
642 elif isinstance(prog, (nvrtcProgram,)):
643 pprog = int(prog)
644 else:
645 pprog = int(nvrtcProgram(prog))
646 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
647 cdef size_t logSizeRet = 0
648 with nogil:
649 err = cynvrtc.nvrtcGetProgramLogSize(cyprog, &logSizeRet)
650 if err != cynvrtc.NVRTC_SUCCESS:
651 return (_dict_nvrtcResult[err], None)
652 return (_dict_nvrtcResult[err], logSizeRet)
654@cython.embedsignature(True)
655def nvrtcGetProgramLog(prog, char* log):
656 """ nvrtcGetProgramLog stores the log generated by the previous compilation of `prog` in the memory pointed by `log`.
658 Parameters
659 ----------
660 prog : :py:obj:`~.nvrtcProgram`
661 CUDA Runtime Compilation program.
662 log : bytes
663 Compilation log.
665 Returns
666 -------
667 nvrtcResult
668 - :py:obj:`~.NVRTC_SUCCESS`
669 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
670 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
672 See Also
673 --------
674 :py:obj:`~.nvrtcGetProgramLogSize`
675 """
676 cdef cynvrtc.nvrtcProgram cyprog
677 if prog is None:
678 pprog = 0
679 elif isinstance(prog, (nvrtcProgram,)):
680 pprog = int(prog)
681 else:
682 pprog = int(nvrtcProgram(prog))
683 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
684 with nogil:
685 err = cynvrtc.nvrtcGetProgramLog(cyprog, log)
686 return (_dict_nvrtcResult[err],)
688@cython.embedsignature(True)
689def nvrtcAddNameExpression(prog, char* name_expression):
690 """ nvrtcAddNameExpression notes the given name expression denoting the address of a global function or device/__constant__ variable.
692 The identical name expression string must be provided on a subsequent
693 call to nvrtcGetLoweredName to extract the lowered name.
695 Parameters
696 ----------
697 prog : :py:obj:`~.nvrtcProgram`
698 CUDA Runtime Compilation program.
699 name_expression : bytes
700 constant expression denoting the address of a global function or
701 device/__constant__ variable.
703 Returns
704 -------
705 nvrtcResult
706 - :py:obj:`~.NVRTC_SUCCESS`
707 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
708 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
709 - :py:obj:`~.NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION`
711 See Also
712 --------
713 :py:obj:`~.nvrtcGetLoweredName`
714 """
715 cdef cynvrtc.nvrtcProgram cyprog
716 if prog is None:
717 pprog = 0
718 elif isinstance(prog, (nvrtcProgram,)):
719 pprog = int(prog)
720 else:
721 pprog = int(nvrtcProgram(prog))
722 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
723 with nogil:
724 err = cynvrtc.nvrtcAddNameExpression(cyprog, name_expression)
725 return (_dict_nvrtcResult[err],)
727@cython.embedsignature(True)
728def nvrtcGetLoweredName(prog, char* name_expression):
729 """ nvrtcGetLoweredName extracts the lowered (mangled) name for a global function or device/__constant__ variable, and updates lowered_name to point to it. The memory containing the name is released when the NVRTC program is destroyed by nvrtcDestroyProgram. The identical name expression must have been previously provided to nvrtcAddNameExpression.
731 Parameters
732 ----------
733 prog : nvrtcProgram
734 CUDA Runtime Compilation program.
735 name_expression : bytes
736 constant expression denoting the address of a global function or
737 device/__constant__ variable.
739 Returns
740 -------
741 nvrtcResult
742 NVRTC_SUCCESS
743 NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION
744 NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID
745 lowered_name : bytes
746 initialized by the function to point to a C string containing the
747 lowered (mangled) name corresponding to the provided name
748 expression.
750 See Also
751 --------
752 nvrtcAddNameExpression
753 """
754 cdef cynvrtc.nvrtcProgram cyprog
755 if prog is None:
756 pprog = 0
757 elif isinstance(prog, (nvrtcProgram,)):
758 pprog = int(prog)
759 else:
760 pprog = int(nvrtcProgram(prog))
761 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
762 cdef const char* lowered_name = NULL
763 with nogil:
764 err = cynvrtc.nvrtcGetLoweredName(cyprog, name_expression, &lowered_name)
765 if err != cynvrtc.NVRTC_SUCCESS:
766 return (_dict_nvrtcResult[err], None)
767 return (_dict_nvrtcResult[err], <bytes>lowered_name if lowered_name != NULL else None)
769@cython.embedsignature(True)
770def nvrtcGetPCHHeapSize():
771 """ retrieve the current size of the PCH Heap.
773 Returns
774 -------
775 nvrtcResult
776 - :py:obj:`~.NVRTC_SUCCESS`
777 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
778 ret : int
779 pointer to location where the size of the PCH Heap will be stored
780 """
781 cdef size_t ret = 0
782 with nogil:
783 err = cynvrtc.nvrtcGetPCHHeapSize(&ret)
784 if err != cynvrtc.NVRTC_SUCCESS:
785 return (_dict_nvrtcResult[err], None)
786 return (_dict_nvrtcResult[err], ret)
788@cython.embedsignature(True)
789def nvrtcSetPCHHeapSize(size_t size):
790 """ set the size of the PCH Heap.
792 The requested size may be rounded up to a platform dependent alignment
793 (e.g. page size). If the PCH Heap has already been allocated, the heap
794 memory will be freed and a new PCH Heap will be allocated.
796 Parameters
797 ----------
798 size : size_t
799 requested size of the PCH Heap, in bytes
801 Returns
802 -------
803 nvrtcResult
804 - :py:obj:`~.NVRTC_SUCCESS`
805 """
806 with nogil:
807 err = cynvrtc.nvrtcSetPCHHeapSize(size)
808 return (_dict_nvrtcResult[err],)
810@cython.embedsignature(True)
811def nvrtcGetPCHCreateStatus(prog):
812 """ returns the PCH creation status.
814 NVRTC_SUCCESS indicates that the PCH was successfully created.
815 NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED indicates that no PCH creation was
816 attempted, either because PCH functionality was not requested during
817 the preceding nvrtcCompileProgram call, or automatic PCH processing was
818 requested, and compiler chose not to create a PCH file.
819 NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED indicates that a PCH file could
820 potentially have been created, but the compiler ran out space in the
821 PCH heap. In this scenario, the
822 :py:obj:`~.nvrtcGetPCHHeapSizeRequired()` can be used to query the
823 required heap size, the heap can be reallocated for this size with
824 :py:obj:`~.nvrtcSetPCHHeapSize()` and PCH creation may be reattempted
825 again invoking :py:obj:`~.nvrtcCompileProgram()` with a new NVRTC
826 program instance. NVRTC_ERROR_PCH_CREATE indicates that an error
827 condition prevented the PCH file from being created.
829 Parameters
830 ----------
831 prog : :py:obj:`~.nvrtcProgram`
832 CUDA Runtime Compilation program.
834 Returns
835 -------
836 nvrtcResult
837 - :py:obj:`~.NVRTC_SUCCESS`
838 - :py:obj:`~.NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED`
839 - :py:obj:`~.NVRTC_ERROR_PCH_CREATE`
840 - :py:obj:`~.NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED`
841 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
842 """
843 cdef cynvrtc.nvrtcProgram cyprog
844 if prog is None:
845 pprog = 0
846 elif isinstance(prog, (nvrtcProgram,)):
847 pprog = int(prog)
848 else:
849 pprog = int(nvrtcProgram(prog))
850 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
851 with nogil:
852 err = cynvrtc.nvrtcGetPCHCreateStatus(cyprog)
853 return (_dict_nvrtcResult[err],)
855@cython.embedsignature(True)
856def nvrtcGetPCHHeapSizeRequired(prog):
857 """ retrieve the required size of the PCH heap required to compile the given program.
859 Parameters
860 ----------
861 prog : :py:obj:`~.nvrtcProgram`
862 CUDA Runtime Compilation program.
864 Returns
865 -------
866 nvrtcResult
867 - :py:obj:`~.NVRTC_SUCCESS`
868 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
869 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` The size retrieved using this function is only valid if :py:obj:`~.nvrtcGetPCHCreateStatus()` returned NVRTC_SUCCESS or NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED
870 size : int
871 pointer to location where the required size of the PCH Heap will be
872 stored
873 """
874 cdef cynvrtc.nvrtcProgram cyprog
875 if prog is None:
876 pprog = 0
877 elif isinstance(prog, (nvrtcProgram,)):
878 pprog = int(prog)
879 else:
880 pprog = int(nvrtcProgram(prog))
881 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
882 cdef size_t size = 0
883 with nogil:
884 err = cynvrtc.nvrtcGetPCHHeapSizeRequired(cyprog, &size)
885 if err != cynvrtc.NVRTC_SUCCESS:
886 return (_dict_nvrtcResult[err], None)
887 return (_dict_nvrtcResult[err], size)
889@cython.embedsignature(True)
890def nvrtcSetFlowCallback(prog, callback, payload):
891 """ nvrtcSetFlowCallback registers a callback function that the compiler will invoke at different points during a call to nvrtcCompileProgram, and the callback function can decide whether to cancel compilation by returning specific values.
893 The callback function must satisfy the following constraints:
895 (1) Its signature should be:
897 **View CUDA Toolkit Documentation for a C++ code example**
899 When invoking the callback, the compiler will always pass `payload` to
900 param1 so that the callback may make decisions based on `payload` .
901 It'll always pass NULL to param2 for now which is reserved for future
902 extensions.
904 (2) It must return 1 to cancel compilation or 0 to continue. Other
905 return values are reserved for future use.
907 (3) It must return consistent values. Once it returns 1 at one point,
908 it must return 1 in all following invocations during the current
909 nvrtcCompileProgram call in progress.
911 (4) It must be thread-safe.
913 (5) It must not invoke any nvrtc/libnvvm/ptx APIs.
915 Parameters
916 ----------
917 prog : :py:obj:`~.nvrtcProgram`
918 CUDA Runtime Compilation program.
919 callback : Any
920 the callback that issues cancellation signal.
921 payload : Any
922 to be passed as a parameter when invoking the callback.
924 Returns
925 -------
926 nvrtcResult
927 - :py:obj:`~.NVRTC_SUCCESS`
928 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
929 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
930 """
931 cdef cynvrtc.nvrtcProgram cyprog
932 if prog is None:
933 pprog = 0
934 elif isinstance(prog, (nvrtcProgram,)):
935 pprog = int(prog)
936 else:
937 pprog = int(nvrtcProgram(prog))
938 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
939 cycallback = _HelperInputVoidPtr(callback)
940 cdef void* cycallback_ptr = <void*><void_ptr>cycallback.cptr
941 cypayload = _HelperInputVoidPtr(payload)
942 cdef void* cypayload_ptr = <void*><void_ptr>cypayload.cptr
943 with nogil:
944 err = cynvrtc.nvrtcSetFlowCallback(cyprog, cycallback_ptr, cypayload_ptr)
945 return (_dict_nvrtcResult[err],)
947@cython.embedsignature(True)
948def sizeof(objType):
949 """ Returns the size of provided CUDA Python structure in bytes
951 Parameters
952 ----------
953 objType : Any
954 CUDA Python object
956 Returns
957 -------
958 lowered_name : int
959 The size of `objType` in bytes
960 """
961 if objType == nvrtcProgram:
962 return sizeof(cynvrtc.nvrtcProgram)
963 raise TypeError("Unknown type: " + str(objType))