Coverage for cuda / bindings / nvrtc.pyx: 54.46%
303 statements
« prev ^ index » next coverage.py v7.13.4, created at 2026-03-08 01:07 +0000
1# SPDX-FileCopyrightText: Copyright (c) 2021-2026 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, generator version 49a8141. Do not modify it directly.
5from typing import Any, Optional
6import cython
7import ctypes
8from libc.stdlib cimport calloc, malloc, free
9from libc cimport string
10from libc.stdint cimport int32_t, uint32_t, int64_t, uint64_t, uintptr_t
11from libc.stddef cimport wchar_t
12from libc.limits cimport CHAR_MIN
13from libcpp.vector cimport vector
14from cpython.buffer cimport PyObject_CheckBuffer, PyObject_GetBuffer, PyBuffer_Release, PyBUF_SIMPLE, PyBUF_ANY_CONTIGUOUS
15from cpython.bytes cimport PyBytes_FromStringAndSize
16from ._internal._fast_enum import FastEnum as _FastEnum
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(_FastEnum):
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
71cdef object _nvrtcResult = nvrtcResult
72cdef object _nvrtcResult_SUCCESS = nvrtcResult.NVRTC_SUCCESS
76cdef class nvrtcProgram:
77 """ nvrtcProgram is the unit of compilation, and an opaque handle for a program.
79 To compile a CUDA program string, an instance of nvrtcProgram must be created first with nvrtcCreateProgram, then compiled with nvrtcCompileProgram.
81 Methods
82 -------
83 getPtr()
84 Get memory address of class instance
86 """
87 def __cinit__(self, void_ptr init_value = 0, void_ptr _ptr = 0):
88 if _ptr == 0: 2a b c d e f g h j k l m n o p q r s t u v w x qbrb
89 self._pvt_ptr = &self._pvt_val 2a b c d e f g h j k l m n o p q r s t u v w x qbrb
90 self._pvt_ptr[0] = <cynvrtc.nvrtcProgram>init_value 2a b c d e f g h j k l m n o p q r s t u v w x qbrb
91 else:
92 self._pvt_ptr = <cynvrtc.nvrtcProgram *>_ptr
93 def __init__(self, *args, **kwargs):
94 pass 2a b c d e f g h j k l m n o p q r s t u v w x qbrb
95 def __repr__(self):
96 return '<nvrtcProgram ' + str(hex(self.__int__())) + '>'
97 def __index__(self):
98 return self.__int__()
99 def __eq__(self, other):
100 if not isinstance(other, nvrtcProgram):
101 return False
102 return self._pvt_ptr[0] == (<nvrtcProgram>other)._pvt_ptr[0]
103 def __hash__(self):
104 return hash(<uintptr_t><void*>(self._pvt_ptr[0]))
105 def __int__(self):
106 return <void_ptr>self._pvt_ptr[0] 1abcdefghjklmnopqrstuvwx
107 def getPtr(self):
108 return <void_ptr>self._pvt_ptr 1jklmnopqrstuvw
110@cython.embedsignature(True)
111def nvrtcGetErrorString(result not None : nvrtcResult):
112 """ 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"`.
114 Parameters
115 ----------
116 result : :py:obj:`~.nvrtcResult`
117 CUDA Runtime Compilation API result code.
119 Returns
120 -------
121 nvrtcResult.NVRTC_SUCCESS
122 nvrtcResult.NVRTC_SUCCESS
123 bytes
124 Message string for the given :py:obj:`~.nvrtcResult` code.
125 """
126 cdef cynvrtc.nvrtcResult cyresult = int(result)
127 with nogil:
128 err = cynvrtc.nvrtcGetErrorString(cyresult)
129 return (nvrtcResult.NVRTC_SUCCESS, err)
131@cython.embedsignature(True)
132def nvrtcVersion():
133 """ nvrtcVersion sets the output parameters `major` and `minor` with the CUDA Runtime Compilation version number.
135 Returns
136 -------
137 nvrtcResult
138 - :py:obj:`~.NVRTC_SUCCESS`
139 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
140 major : int
141 CUDA Runtime Compilation major version number.
142 minor : int
143 CUDA Runtime Compilation minor version number.
144 """
145 cdef int major = 0 2i a b c d e f g h z A B C D E F G H I J K L M N O P Q R S T U V W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbgbhbibjbkblbmbnbobpb
146 cdef int minor = 0 2i a b c d e f g h z A B C D E F G H I J K L M N O P Q R S T U V W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbgbhbibjbkblbmbnbobpb
147 with nogil: 2i a b c d e f g h z A B C D E F G H I J K L M N O P Q R S T U V W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbgbhbibjbkblbmbnbobpb
148 err = cynvrtc.nvrtcVersion(&major, &minor) 2i a b c d e f g h z A B C D E F G H I J K L M N O P Q R S T U V W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbgbhbibjbkblbmbnbobpb
149 if err != cynvrtc.NVRTC_SUCCESS: 2i a b c d e f g h z A B C D E F G H I J K L M N O P Q R S T U V W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbgbhbibjbkblbmbnbobpb
150 return (_nvrtcResult(err), None, None)
151 return (_nvrtcResult_SUCCESS, major, minor) 2i a b c d e f g h z A B C D E F G H I J K L M N O P Q R S T U V W X Y Z 0 1 2 3 4 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbgbhbibjbkblbmbnbobpb
153@cython.embedsignature(True)
154def nvrtcGetNumSupportedArchs():
155 """ 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.
157 see :py:obj:`~.nvrtcGetSupportedArchs`
159 Returns
160 -------
161 nvrtcResult
162 - :py:obj:`~.NVRTC_SUCCESS`
163 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
164 numArchs : int
165 number of supported architectures.
166 """
167 cdef int numArchs = 0 1y
168 with nogil: 1y
169 err = cynvrtc.nvrtcGetNumSupportedArchs(&numArchs) 1y
170 if err != cynvrtc.NVRTC_SUCCESS: 1y
171 return (_nvrtcResult(err), None)
172 return (_nvrtcResult_SUCCESS, numArchs) 1y
174@cython.embedsignature(True)
175def nvrtcGetSupportedArchs():
176 """ 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`.
178 see :py:obj:`~.nvrtcGetNumSupportedArchs`
180 Returns
181 -------
182 nvrtcResult
183 - :py:obj:`~.NVRTC_SUCCESS`
184 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
185 supportedArchs : list[int]
186 sorted array of supported architectures.
187 """
188 cdef vector[int] supportedArchs
189 _, s = nvrtcGetNumSupportedArchs() 1y
190 supportedArchs.resize(s) 1y
192 with nogil: 1y
193 err = cynvrtc.nvrtcGetSupportedArchs(supportedArchs.data()) 1y
194 if err != cynvrtc.NVRTC_SUCCESS: 1y
195 return (_nvrtcResult(err), None)
196 return (_nvrtcResult_SUCCESS, supportedArchs) 1y
198@cython.embedsignature(True)
199def nvrtcCreateProgram(char* src, char* name, int numHeaders, headers : Optional[tuple[bytes] | list[bytes]], includeNames : Optional[tuple[bytes] | list[bytes]]):
200 """ nvrtcCreateProgram creates an instance of nvrtcProgram with the given input parameters, and sets the output parameter `prog` with it.
202 Parameters
203 ----------
204 src : bytes
205 CUDA program source.
206 name : bytes
207 CUDA program name. `name` can be `NULL`; `"default_program"` is
208 used when `name` is `NULL` or "".
209 numHeaders : int
210 Number of headers used. `numHeaders` must be greater than or equal
211 to 0.
212 headers : list[bytes]
213 Sources of the headers. `headers` can be `NULL` when `numHeaders`
214 is 0.
215 includeNames : list[bytes]
216 Name of each header by which they can be included in the CUDA
217 program source. `includeNames` can be `NULL` when `numHeaders` is
218 0. These headers must be included with the exact names specified
219 here.
221 Returns
222 -------
223 nvrtcResult
224 - :py:obj:`~.NVRTC_SUCCESS`
225 - :py:obj:`~.NVRTC_ERROR_OUT_OF_MEMORY`
226 - :py:obj:`~.NVRTC_ERROR_PROGRAM_CREATION_FAILURE`
227 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
228 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
229 prog : :py:obj:`~.nvrtcProgram`
230 CUDA Runtime Compilation program.
232 See Also
233 --------
234 :py:obj:`~.nvrtcDestroyProgram`
235 """
236 includeNames = [] if includeNames is None else includeNames 1abcdefghjklmnopqrstuvw
237 if not all(isinstance(_x, (bytes)) for _x in includeNames): 1abcdefghjklmnopqrstuvw
238 raise TypeError("Argument 'includeNames' is not instance of type (expected tuple[bytes] or list[bytes]")
239 headers = [] if headers is None else headers 1abcdefghjklmnopqrstuvw
240 if not all(isinstance(_x, (bytes)) for _x in headers): 1abcdefghjklmnopqrstuvw
241 raise TypeError("Argument 'headers' is not instance of type (expected tuple[bytes] or list[bytes]")
242 cdef nvrtcProgram prog = nvrtcProgram() 1abcdefghjklmnopqrstuvw
243 if numHeaders > len(headers): raise RuntimeError("List is too small: " + str(len(headers)) + " < " + str(numHeaders)) 1abcdefghjklmnopqrstuvw
244 if numHeaders > len(includeNames): raise RuntimeError("List is too small: " + str(len(includeNames)) + " < " + str(numHeaders)) 1abcdefghjklmnopqrstuvw
245 cdef vector[const char*] cyheaders = headers 1abcdefghjklmnopqrstuvw
246 cdef vector[const char*] cyincludeNames = includeNames 1abcdefghjklmnopqrstuvw
247 with nogil: 1abcdefghjklmnopqrstuvw
248 err = cynvrtc.nvrtcCreateProgram(<cynvrtc.nvrtcProgram*>prog._pvt_ptr, src, name, numHeaders, cyheaders.data(), cyincludeNames.data()) 1abcdefghjklmnopqrstuvw
249 if err != cynvrtc.NVRTC_SUCCESS: 1abcdefghjklmnopqrstuvw
250 return (_nvrtcResult(err), None)
251 return (_nvrtcResult_SUCCESS, prog) 1abcdefghjklmnopqrstuvw
253@cython.embedsignature(True)
254def nvrtcDestroyProgram(prog):
255 """ nvrtcDestroyProgram destroys the given program.
257 Parameters
258 ----------
259 prog : :py:obj:`~.nvrtcProgram`
260 CUDA Runtime Compilation program.
262 Returns
263 -------
264 nvrtcResult
265 - :py:obj:`~.NVRTC_SUCCESS`
266 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
268 See Also
269 --------
270 :py:obj:`~.nvrtcCreateProgram`
271 """
272 cdef cynvrtc.nvrtcProgram *cyprog
273 if prog is None: 1jklmnopqrstuvw
274 cyprog = <cynvrtc.nvrtcProgram*><void_ptr>NULL
275 elif isinstance(prog, (nvrtcProgram,)): 1jklmnopqrstuvw
276 pprog = prog.getPtr() 1jklmnopqrstuvw
277 cyprog = <cynvrtc.nvrtcProgram*><void_ptr>pprog 1jklmnopqrstuvw
278 elif isinstance(prog, (int)):
279 cyprog = <cynvrtc.nvrtcProgram*><void_ptr>prog
280 else:
281 raise TypeError("Argument 'prog' is not instance of type (expected <class 'int, nvrtc.nvrtcProgram'>, found " + str(type(prog)))
282 with nogil: 1jklmnopqrstuvw
283 err = cynvrtc.nvrtcDestroyProgram(cyprog) 1jklmnopqrstuvw
284 return (_nvrtcResult(err),) 1jklmnopqrstuvw
286@cython.embedsignature(True)
287def nvrtcCompileProgram(prog, int numOptions, options : Optional[tuple[bytes] | list[bytes]]):
288 """ nvrtcCompileProgram compiles the given program.
290 It supports compile options listed in :py:obj:`~.Supported Compile
291 Options`.
293 Parameters
294 ----------
295 prog : :py:obj:`~.nvrtcProgram`
296 CUDA Runtime Compilation program.
297 numOptions : int
298 Number of compiler options passed.
299 options : list[bytes]
300 Compiler options in the form of C string array. `options` can be
301 `NULL` when `numOptions` is 0.
303 Returns
304 -------
305 nvrtcResult
306 - :py:obj:`~.NVRTC_SUCCESS`
307 - :py:obj:`~.NVRTC_ERROR_OUT_OF_MEMORY`
308 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
309 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
310 - :py:obj:`~.NVRTC_ERROR_INVALID_OPTION`
311 - :py:obj:`~.NVRTC_ERROR_COMPILATION`
312 - :py:obj:`~.NVRTC_ERROR_BUILTIN_OPERATION_FAILURE`
313 - :py:obj:`~.NVRTC_ERROR_TIME_FILE_WRITE_FAILED`
314 - :py:obj:`~.NVRTC_ERROR_CANCELLED`
315 """
316 options = [] if options is None else options 1abcdefghjklmnopqrstuvw
317 if not all(isinstance(_x, (bytes)) for _x in options): 1abcdefghjklmnopqrstuvw
318 raise TypeError("Argument 'options' is not instance of type (expected tuple[bytes] or list[bytes]")
319 cdef cynvrtc.nvrtcProgram cyprog
320 if prog is None: 1abcdefghjklmnopqrstuvw
321 pprog = 0
322 elif isinstance(prog, (nvrtcProgram,)): 1abcdefghjklmnopqrstuvw
323 pprog = int(prog) 1abcdefghjklmnopqrstuvw
324 else:
325 pprog = int(nvrtcProgram(prog))
326 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog 1abcdefghjklmnopqrstuvw
327 if numOptions > len(options): raise RuntimeError("List is too small: " + str(len(options)) + " < " + str(numOptions)) 1abcdefghjklmnopqrstuvw
328 cdef vector[const char*] cyoptions = options 1abcdefghjklmnopqrstuvw
329 with nogil: 1abcdefghjklmnopqrstuvw
330 err = cynvrtc.nvrtcCompileProgram(cyprog, numOptions, cyoptions.data()) 1abcdefghjklmnopqrstuvw
331 return (_nvrtcResult(err),) 1abcdefghjklmnopqrstuvw
333@cython.embedsignature(True)
334def nvrtcGetPTXSize(prog):
335 """ nvrtcGetPTXSize sets the value of `ptxSizeRet` with the size of the PTX generated by the previous compilation of `prog` (including the trailing `NULL`).
337 Parameters
338 ----------
339 prog : :py:obj:`~.nvrtcProgram`
340 CUDA Runtime Compilation program.
342 Returns
343 -------
344 nvrtcResult
345 - :py:obj:`~.NVRTC_SUCCESS`
346 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
347 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
348 ptxSizeRet : int
349 Size of the generated PTX (including the trailing `NULL`).
351 See Also
352 --------
353 :py:obj:`~.nvrtcGetPTX`
354 """
355 cdef cynvrtc.nvrtcProgram cyprog
356 if prog is None:
357 pprog = 0
358 elif isinstance(prog, (nvrtcProgram,)):
359 pprog = int(prog)
360 else:
361 pprog = int(nvrtcProgram(prog))
362 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
363 cdef size_t ptxSizeRet = 0
364 with nogil:
365 err = cynvrtc.nvrtcGetPTXSize(cyprog, &ptxSizeRet)
366 if err != cynvrtc.NVRTC_SUCCESS:
367 return (_nvrtcResult(err), None)
368 return (_nvrtcResult_SUCCESS, ptxSizeRet)
370@cython.embedsignature(True)
371def nvrtcGetPTX(prog, char* ptx):
372 """ nvrtcGetPTX stores the PTX generated by the previous compilation of `prog` in the memory pointed by `ptx`.
374 Parameters
375 ----------
376 prog : :py:obj:`~.nvrtcProgram`
377 CUDA Runtime Compilation program.
378 ptx : bytes
379 Compiled result.
381 Returns
382 -------
383 nvrtcResult
384 - :py:obj:`~.NVRTC_SUCCESS`
385 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
386 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
388 See Also
389 --------
390 :py:obj:`~.nvrtcGetPTXSize`
391 """
392 cdef cynvrtc.nvrtcProgram cyprog
393 if prog is None:
394 pprog = 0
395 elif isinstance(prog, (nvrtcProgram,)):
396 pprog = int(prog)
397 else:
398 pprog = int(nvrtcProgram(prog))
399 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
400 with nogil:
401 err = cynvrtc.nvrtcGetPTX(cyprog, ptx)
402 return (_nvrtcResult(err),)
404@cython.embedsignature(True)
405def nvrtcGetCUBINSize(prog):
406 """ 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.
408 Parameters
409 ----------
410 prog : :py:obj:`~.nvrtcProgram`
411 CUDA Runtime Compilation program.
413 Returns
414 -------
415 nvrtcResult
416 - :py:obj:`~.NVRTC_SUCCESS`
417 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
418 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
419 cubinSizeRet : int
420 Size of the generated cubin.
422 See Also
423 --------
424 :py:obj:`~.nvrtcGetCUBIN`
425 """
426 cdef cynvrtc.nvrtcProgram cyprog
427 if prog is None: 1abcdefghnopqrs
428 pprog = 0
429 elif isinstance(prog, (nvrtcProgram,)): 1abcdefghnopqrs
430 pprog = int(prog) 1abcdefghnopqrs
431 else:
432 pprog = int(nvrtcProgram(prog))
433 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog 1abcdefghnopqrs
434 cdef size_t cubinSizeRet = 0 1abcdefghnopqrs
435 with nogil: 1abcdefghnopqrs
436 err = cynvrtc.nvrtcGetCUBINSize(cyprog, &cubinSizeRet) 1abcdefghnopqrs
437 if err != cynvrtc.NVRTC_SUCCESS: 1abcdefghnopqrs
438 return (_nvrtcResult(err), None)
439 return (_nvrtcResult_SUCCESS, cubinSizeRet) 1abcdefghnopqrs
441@cython.embedsignature(True)
442def nvrtcGetCUBIN(prog, char* cubin):
443 """ 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.
445 Parameters
446 ----------
447 prog : :py:obj:`~.nvrtcProgram`
448 CUDA Runtime Compilation program.
449 cubin : bytes
450 Compiled and assembled result.
452 Returns
453 -------
454 nvrtcResult
455 - :py:obj:`~.NVRTC_SUCCESS`
456 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
457 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
459 See Also
460 --------
461 :py:obj:`~.nvrtcGetCUBINSize`
462 """
463 cdef cynvrtc.nvrtcProgram cyprog
464 if prog is None: 1abcdefghnopqrs
465 pprog = 0
466 elif isinstance(prog, (nvrtcProgram,)): 1abcdefghnopqrs
467 pprog = int(prog) 1abcdefghnopqrs
468 else:
469 pprog = int(nvrtcProgram(prog))
470 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog 1abcdefghnopqrs
471 with nogil: 1abcdefghnopqrs
472 err = cynvrtc.nvrtcGetCUBIN(cyprog, cubin) 1abcdefghnopqrs
473 return (_nvrtcResult(err),) 1abcdefghnopqrs
475@cython.embedsignature(True)
476def nvrtcGetLTOIRSize(prog):
477 """ 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`.
479 Parameters
480 ----------
481 prog : :py:obj:`~.nvrtcProgram`
482 CUDA Runtime Compilation program.
484 Returns
485 -------
486 nvrtcResult
487 - :py:obj:`~.NVRTC_SUCCESS`
488 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
489 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
490 LTOIRSizeRet : int
491 Size of the generated LTO IR.
493 See Also
494 --------
495 :py:obj:`~.nvrtcGetLTOIR`
496 """
497 cdef cynvrtc.nvrtcProgram cyprog
498 if prog is None: 1jklmtuvw
499 pprog = 0
500 elif isinstance(prog, (nvrtcProgram,)): 1jklmtuvw
501 pprog = int(prog) 1jklmtuvw
502 else:
503 pprog = int(nvrtcProgram(prog))
504 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog 1jklmtuvw
505 cdef size_t LTOIRSizeRet = 0 1jklmtuvw
506 with nogil: 1jklmtuvw
507 err = cynvrtc.nvrtcGetLTOIRSize(cyprog, <OIRSizeRet) 1jklmtuvw
508 if err != cynvrtc.NVRTC_SUCCESS: 1jklmtuvw
509 return (_nvrtcResult(err), None)
510 return (_nvrtcResult_SUCCESS, LTOIRSizeRet) 1jklmtuvw
512@cython.embedsignature(True)
513def nvrtcGetLTOIR(prog, char* LTOIR):
514 """ 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`.
516 Parameters
517 ----------
518 prog : :py:obj:`~.nvrtcProgram`
519 CUDA Runtime Compilation program.
520 LTOIR : bytes
521 Compiled result.
523 Returns
524 -------
525 nvrtcResult
526 - :py:obj:`~.NVRTC_SUCCESS`
527 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
528 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
530 See Also
531 --------
532 :py:obj:`~.nvrtcGetLTOIRSize`
533 """
534 cdef cynvrtc.nvrtcProgram cyprog
535 if prog is None: 1jklmtuvw
536 pprog = 0
537 elif isinstance(prog, (nvrtcProgram,)): 1jklmtuvw
538 pprog = int(prog) 1jklmtuvw
539 else:
540 pprog = int(nvrtcProgram(prog))
541 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog 1jklmtuvw
542 with nogil: 1jklmtuvw
543 err = cynvrtc.nvrtcGetLTOIR(cyprog, LTOIR) 1jklmtuvw
544 return (_nvrtcResult(err),) 1jklmtuvw
546@cython.embedsignature(True)
547def nvrtcGetOptiXIRSize(prog):
548 """ 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.
550 Parameters
551 ----------
552 prog : :py:obj:`~.nvrtcProgram`
553 CUDA Runtime Compilation program.
555 Returns
556 -------
557 nvrtcResult
558 - :py:obj:`~.NVRTC_SUCCESS`
559 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
560 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
561 optixirSizeRet : int
562 Size of the generated LTO IR.
564 See Also
565 --------
566 :py:obj:`~.nvrtcGetOptiXIR`
567 """
568 cdef cynvrtc.nvrtcProgram cyprog
569 if prog is None:
570 pprog = 0
571 elif isinstance(prog, (nvrtcProgram,)):
572 pprog = int(prog)
573 else:
574 pprog = int(nvrtcProgram(prog))
575 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
576 cdef size_t optixirSizeRet = 0
577 with nogil:
578 err = cynvrtc.nvrtcGetOptiXIRSize(cyprog, &optixirSizeRet)
579 if err != cynvrtc.NVRTC_SUCCESS:
580 return (_nvrtcResult(err), None)
581 return (_nvrtcResult_SUCCESS, optixirSizeRet)
583@cython.embedsignature(True)
584def nvrtcGetOptiXIR(prog, char* optixir):
585 """ 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.
587 Parameters
588 ----------
589 prog : :py:obj:`~.nvrtcProgram`
590 CUDA Runtime Compilation program.
591 optixir : bytes
592 Optix IR Compiled result.
594 Returns
595 -------
596 nvrtcResult
597 - :py:obj:`~.NVRTC_SUCCESS`
598 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
599 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
601 See Also
602 --------
603 :py:obj:`~.nvrtcGetOptiXIRSize`
604 """
605 cdef cynvrtc.nvrtcProgram cyprog
606 if prog is None:
607 pprog = 0
608 elif isinstance(prog, (nvrtcProgram,)):
609 pprog = int(prog)
610 else:
611 pprog = int(nvrtcProgram(prog))
612 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
613 with nogil:
614 err = cynvrtc.nvrtcGetOptiXIR(cyprog, optixir)
615 return (_nvrtcResult(err),)
617@cython.embedsignature(True)
618def nvrtcGetProgramLogSize(prog):
619 """ nvrtcGetProgramLogSize sets `logSizeRet` with the size of the log generated by the previous compilation of `prog` (including the trailing `NULL`).
621 Note that compilation log may be generated with warnings and
622 informative messages, even when the compilation of `prog` succeeds.
624 Parameters
625 ----------
626 prog : :py:obj:`~.nvrtcProgram`
627 CUDA Runtime Compilation program.
629 Returns
630 -------
631 nvrtcResult
632 - :py:obj:`~.NVRTC_SUCCESS`
633 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
634 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
635 logSizeRet : int
636 Size of the compilation log (including the trailing `NULL`).
638 See Also
639 --------
640 :py:obj:`~.nvrtcGetProgramLog`
641 """
642 cdef cynvrtc.nvrtcProgram cyprog
643 if prog is None: 1abcdefgh
644 pprog = 0
645 elif isinstance(prog, (nvrtcProgram,)): 1abcdefgh
646 pprog = int(prog) 1abcdefgh
647 else:
648 pprog = int(nvrtcProgram(prog))
649 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog 1abcdefgh
650 cdef size_t logSizeRet = 0 1abcdefgh
651 with nogil: 1abcdefgh
652 err = cynvrtc.nvrtcGetProgramLogSize(cyprog, &logSizeRet) 1abcdefgh
653 if err != cynvrtc.NVRTC_SUCCESS: 1abcdefgh
654 return (_nvrtcResult(err), None)
655 return (_nvrtcResult_SUCCESS, logSizeRet) 1abcdefgh
657@cython.embedsignature(True)
658def nvrtcGetProgramLog(prog, char* log):
659 """ nvrtcGetProgramLog stores the log generated by the previous compilation of `prog` in the memory pointed by `log`.
661 Parameters
662 ----------
663 prog : :py:obj:`~.nvrtcProgram`
664 CUDA Runtime Compilation program.
665 log : bytes
666 Compilation log.
668 Returns
669 -------
670 nvrtcResult
671 - :py:obj:`~.NVRTC_SUCCESS`
672 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
673 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
675 See Also
676 --------
677 :py:obj:`~.nvrtcGetProgramLogSize`
678 """
679 cdef cynvrtc.nvrtcProgram cyprog
680 if prog is None: 1abcdefgh
681 pprog = 0
682 elif isinstance(prog, (nvrtcProgram,)): 1abcdefgh
683 pprog = int(prog) 1abcdefgh
684 else:
685 pprog = int(nvrtcProgram(prog))
686 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog 1abcdefgh
687 with nogil: 1abcdefgh
688 err = cynvrtc.nvrtcGetProgramLog(cyprog, log) 1abcdefgh
689 return (_nvrtcResult(err),) 1abcdefgh
691@cython.embedsignature(True)
692def nvrtcAddNameExpression(prog, char* name_expression):
693 """ nvrtcAddNameExpression notes the given name expression denoting the address of a global function or device/__constant__ variable.
695 The identical name expression string must be provided on a subsequent
696 call to nvrtcGetLoweredName to extract the lowered name.
698 Parameters
699 ----------
700 prog : :py:obj:`~.nvrtcProgram`
701 CUDA Runtime Compilation program.
702 name_expression : bytes
703 constant expression denoting the address of a global function or
704 device/__constant__ variable.
706 Returns
707 -------
708 nvrtcResult
709 - :py:obj:`~.NVRTC_SUCCESS`
710 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
711 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
712 - :py:obj:`~.NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION`
714 See Also
715 --------
716 :py:obj:`~.nvrtcGetLoweredName`
717 """
718 cdef cynvrtc.nvrtcProgram cyprog
719 if prog is None:
720 pprog = 0
721 elif isinstance(prog, (nvrtcProgram,)):
722 pprog = int(prog)
723 else:
724 pprog = int(nvrtcProgram(prog))
725 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
726 with nogil:
727 err = cynvrtc.nvrtcAddNameExpression(cyprog, name_expression)
728 return (_nvrtcResult(err),)
730@cython.embedsignature(True)
731def nvrtcGetLoweredName(prog, char* name_expression):
732 """ 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.
734 Parameters
735 ----------
736 prog : nvrtcProgram
737 CUDA Runtime Compilation program.
738 name_expression : bytes
739 constant expression denoting the address of a global function or
740 device/__constant__ variable.
742 Returns
743 -------
744 nvrtcResult
745 NVRTC_SUCCESS
746 NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION
747 NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID
748 lowered_name : bytes
749 initialized by the function to point to a C string containing the
750 lowered (mangled) name corresponding to the provided name
751 expression.
753 See Also
754 --------
755 nvrtcAddNameExpression
756 """
757 cdef cynvrtc.nvrtcProgram cyprog
758 if prog is None: 1x
759 pprog = 0 1x
760 elif isinstance(prog, (nvrtcProgram,)): 1x
761 pprog = int(prog)
762 else:
763 pprog = int(nvrtcProgram(prog)) 1x
764 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog 1x
765 cdef const char* lowered_name = NULL 1x
766 with nogil: 1x
767 err = cynvrtc.nvrtcGetLoweredName(cyprog, name_expression, &lowered_name) 1x
768 if err != cynvrtc.NVRTC_SUCCESS: 1x
769 return (_nvrtcResult(err), None) 1x
770 return (_nvrtcResult_SUCCESS, <bytes>lowered_name if lowered_name != NULL else None)
772@cython.embedsignature(True)
773def nvrtcGetPCHHeapSize():
774 """ retrieve the current size of the PCH Heap.
776 Returns
777 -------
778 nvrtcResult
779 - :py:obj:`~.NVRTC_SUCCESS`
780 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
781 ret : int
782 pointer to location where the size of the PCH Heap will be stored
783 """
784 cdef size_t ret = 0
785 with nogil:
786 err = cynvrtc.nvrtcGetPCHHeapSize(&ret)
787 if err != cynvrtc.NVRTC_SUCCESS:
788 return (_nvrtcResult(err), None)
789 return (_nvrtcResult_SUCCESS, ret)
791@cython.embedsignature(True)
792def nvrtcSetPCHHeapSize(size_t size):
793 """ set the size of the PCH Heap.
795 The requested size may be rounded up to a platform dependent alignment
796 (e.g. page size). If the PCH Heap has already been allocated, the heap
797 memory will be freed and a new PCH Heap will be allocated.
799 Parameters
800 ----------
801 size : size_t
802 requested size of the PCH Heap, in bytes
804 Returns
805 -------
806 nvrtcResult
807 - :py:obj:`~.NVRTC_SUCCESS`
808 """
809 with nogil:
810 err = cynvrtc.nvrtcSetPCHHeapSize(size)
811 return (_nvrtcResult(err),)
813@cython.embedsignature(True)
814def nvrtcGetPCHCreateStatus(prog):
815 """ returns the PCH creation status.
817 NVRTC_SUCCESS indicates that the PCH was successfully created.
818 NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED indicates that no PCH creation was
819 attempted, either because PCH functionality was not requested during
820 the preceding nvrtcCompileProgram call, or automatic PCH processing was
821 requested, and compiler chose not to create a PCH file.
822 NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED indicates that a PCH file could
823 potentially have been created, but the compiler ran out space in the
824 PCH heap. In this scenario, the
825 :py:obj:`~.nvrtcGetPCHHeapSizeRequired()` can be used to query the
826 required heap size, the heap can be reallocated for this size with
827 :py:obj:`~.nvrtcSetPCHHeapSize()` and PCH creation may be reattempted
828 again invoking :py:obj:`~.nvrtcCompileProgram()` with a new NVRTC
829 program instance. NVRTC_ERROR_PCH_CREATE indicates that an error
830 condition prevented the PCH file from being created.
832 Parameters
833 ----------
834 prog : :py:obj:`~.nvrtcProgram`
835 CUDA Runtime Compilation program.
837 Returns
838 -------
839 nvrtcResult
840 - :py:obj:`~.NVRTC_SUCCESS`
841 - :py:obj:`~.NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED`
842 - :py:obj:`~.NVRTC_ERROR_PCH_CREATE`
843 - :py:obj:`~.NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED`
844 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
845 """
846 cdef cynvrtc.nvrtcProgram cyprog
847 if prog is None:
848 pprog = 0
849 elif isinstance(prog, (nvrtcProgram,)):
850 pprog = int(prog)
851 else:
852 pprog = int(nvrtcProgram(prog))
853 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
854 with nogil:
855 err = cynvrtc.nvrtcGetPCHCreateStatus(cyprog)
856 return (_nvrtcResult(err),)
858@cython.embedsignature(True)
859def nvrtcGetPCHHeapSizeRequired(prog):
860 """ retrieve the required size of the PCH heap required to compile the given program.
862 Parameters
863 ----------
864 prog : :py:obj:`~.nvrtcProgram`
865 CUDA Runtime Compilation program.
867 Returns
868 -------
869 nvrtcResult
870 - :py:obj:`~.NVRTC_SUCCESS`
871 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
872 - :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
873 size : int
874 pointer to location where the required size of the PCH Heap will be
875 stored
876 """
877 cdef cynvrtc.nvrtcProgram cyprog
878 if prog is None:
879 pprog = 0
880 elif isinstance(prog, (nvrtcProgram,)):
881 pprog = int(prog)
882 else:
883 pprog = int(nvrtcProgram(prog))
884 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
885 cdef size_t size = 0
886 with nogil:
887 err = cynvrtc.nvrtcGetPCHHeapSizeRequired(cyprog, &size)
888 if err != cynvrtc.NVRTC_SUCCESS:
889 return (_nvrtcResult(err), None)
890 return (_nvrtcResult_SUCCESS, size)
892@cython.embedsignature(True)
893def nvrtcSetFlowCallback(prog, callback, payload):
894 """ 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.
896 The callback function must satisfy the following constraints:
898 (1) Its signature should be:
900 **View CUDA Toolkit Documentation for a C++ code example**
902 When invoking the callback, the compiler will always pass `payload` to
903 param1 so that the callback may make decisions based on `payload` .
904 It'll always pass NULL to param2 for now which is reserved for future
905 extensions.
907 (2) It must return 1 to cancel compilation or 0 to continue. Other
908 return values are reserved for future use.
910 (3) It must return consistent values. Once it returns 1 at one point,
911 it must return 1 in all following invocations during the current
912 nvrtcCompileProgram call in progress.
914 (4) It must be thread-safe.
916 (5) It must not invoke any nvrtc/libnvvm/ptx APIs.
918 Parameters
919 ----------
920 prog : :py:obj:`~.nvrtcProgram`
921 CUDA Runtime Compilation program.
922 callback : Any
923 the callback that issues cancellation signal.
924 payload : Any
925 to be passed as a parameter when invoking the callback.
927 Returns
928 -------
929 nvrtcResult
930 - :py:obj:`~.NVRTC_SUCCESS`
931 - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM`
932 - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT`
933 """
934 cdef cynvrtc.nvrtcProgram cyprog
935 if prog is None:
936 pprog = 0
937 elif isinstance(prog, (nvrtcProgram,)):
938 pprog = int(prog)
939 else:
940 pprog = int(nvrtcProgram(prog))
941 cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
942 cdef _HelperInputVoidPtrStruct cycallbackHelper
943 cdef void* cycallback = _helper_input_void_ptr(callback, &cycallbackHelper)
944 cdef _HelperInputVoidPtrStruct cypayloadHelper
945 cdef void* cypayload = _helper_input_void_ptr(payload, &cypayloadHelper)
946 with nogil:
947 err = cynvrtc.nvrtcSetFlowCallback(cyprog, cycallback, cypayload)
948 _helper_input_void_ptr_free(&cycallbackHelper)
949 _helper_input_void_ptr_free(&cypayloadHelper)
950 return (_nvrtcResult(err),)
952@cython.embedsignature(True)
953def sizeof(objType):
954 """ Returns the size of provided CUDA Python structure in bytes
956 Parameters
957 ----------
958 objType : Any
959 CUDA Python object
961 Returns
962 -------
963 lowered_name : int
964 The size of `objType` in bytes
965 """
966 if objType == nvrtcProgram:
967 return sizeof(cynvrtc.nvrtcProgram)
968 raise TypeError("Unknown type: " + str(objType))