# Copyright 2021-2024 NVIDIA Corporation. All rights reserved. # # Please refer to the NVIDIA end user license agreement (EULA) associated # with this source code for terms and conditions that govern your use of # this software. Any use, reproduction, disclosure, or distribution of # this software and related documentation outside the terms of the EULA # is strictly prohibited. from typing import List, Tuple, Any, Optional from enum import IntEnum import cython import ctypes from libc.stdlib cimport calloc, free from libc cimport string from libc.stdint cimport int32_t, uint32_t, int64_t, uint64_t from libc.stddef cimport wchar_t from libc.limits cimport CHAR_MIN from libcpp.vector cimport vector from cpython.buffer cimport PyObject_CheckBuffer, PyObject_GetBuffer, PyBuffer_Release, PyBUF_SIMPLE, PyBUF_ANY_CONTIGUOUS from cpython.bytes cimport PyBytes_FromStringAndSize ctypedef unsigned long long signed_char_ptr ctypedef unsigned long long unsigned_char_ptr ctypedef unsigned long long char_ptr ctypedef unsigned long long short_ptr ctypedef unsigned long long unsigned_short_ptr ctypedef unsigned long long int_ptr ctypedef unsigned long long long_int_ptr ctypedef unsigned long long long_long_int_ptr ctypedef unsigned long long unsigned_int_ptr ctypedef unsigned long long unsigned_long_int_ptr ctypedef unsigned long long unsigned_long_long_int_ptr ctypedef unsigned long long uint32_t_ptr ctypedef unsigned long long uint64_t_ptr ctypedef unsigned long long int32_t_ptr ctypedef unsigned long long int64_t_ptr ctypedef unsigned long long unsigned_ptr ctypedef unsigned long long unsigned_long_long_ptr ctypedef unsigned long long long_long_ptr ctypedef unsigned long long size_t_ptr ctypedef unsigned long long float_ptr ctypedef unsigned long long double_ptr ctypedef unsigned long long void_ptr class nvrtcResult(IntEnum): """ The enumerated type nvrtcResult defines API call result codes. NVRTC API functions return nvrtcResult to indicate the call result. """ NVRTC_SUCCESS = cnvrtc.nvrtcResult.NVRTC_SUCCESS NVRTC_ERROR_OUT_OF_MEMORY = cnvrtc.nvrtcResult.NVRTC_ERROR_OUT_OF_MEMORY NVRTC_ERROR_PROGRAM_CREATION_FAILURE = cnvrtc.nvrtcResult.NVRTC_ERROR_PROGRAM_CREATION_FAILURE NVRTC_ERROR_INVALID_INPUT = cnvrtc.nvrtcResult.NVRTC_ERROR_INVALID_INPUT NVRTC_ERROR_INVALID_PROGRAM = cnvrtc.nvrtcResult.NVRTC_ERROR_INVALID_PROGRAM NVRTC_ERROR_INVALID_OPTION = cnvrtc.nvrtcResult.NVRTC_ERROR_INVALID_OPTION NVRTC_ERROR_COMPILATION = cnvrtc.nvrtcResult.NVRTC_ERROR_COMPILATION NVRTC_ERROR_BUILTIN_OPERATION_FAILURE = cnvrtc.nvrtcResult.NVRTC_ERROR_BUILTIN_OPERATION_FAILURE NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = cnvrtc.nvrtcResult.NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = cnvrtc.nvrtcResult.NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID = cnvrtc.nvrtcResult.NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID NVRTC_ERROR_INTERNAL_ERROR = cnvrtc.nvrtcResult.NVRTC_ERROR_INTERNAL_ERROR NVRTC_ERROR_TIME_FILE_WRITE_FAILED = cnvrtc.nvrtcResult.NVRTC_ERROR_TIME_FILE_WRITE_FAILED cdef class nvrtcProgram: """ nvrtcProgram is the unit of compilation, and an opaque handle for a program. To compile a CUDA program string, an instance of nvrtcProgram must be created first with nvrtcCreateProgram, then compiled with nvrtcCompileProgram. Methods ------- getPtr() Get memory address of class instance """ def __cinit__(self, void_ptr init_value = 0, void_ptr _ptr = 0): if _ptr == 0: self._ptr = &self.__val self._ptr[0] = init_value else: self._ptr = _ptr def __init__(self, *args, **kwargs): pass def __repr__(self): return '' def __index__(self): return self.__int__() def __int__(self): return self._ptr[0] def getPtr(self): return self._ptr @cython.embedsignature(True) def nvrtcGetErrorString(result not None : nvrtcResult): """ 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"`. Parameters ---------- result : :py:obj:`~.nvrtcResult` CUDA Runtime Compilation API result code. Returns ------- nvrtcResult.NVRTC_SUCCESS nvrtcResult.NVRTC_SUCCESS bytes Message string for the given :py:obj:`~.nvrtcResult` code. """ cdef cnvrtc.nvrtcResult cresult = result.value err = cnvrtc.nvrtcGetErrorString(cresult) return (nvrtcResult.NVRTC_SUCCESS, err) @cython.embedsignature(True) def nvrtcVersion(): """ nvrtcVersion sets the output parameters `major` and `minor` with the CUDA Runtime Compilation version number. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` major : int CUDA Runtime Compilation major version number. minor : int CUDA Runtime Compilation minor version number. """ cdef int major = 0 cdef int minor = 0 err = cnvrtc.nvrtcVersion(&major, &minor) return (nvrtcResult(err), major, minor) @cython.embedsignature(True) def nvrtcGetNumSupportedArchs(): """ 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. see :py:obj:`~.nvrtcGetSupportedArchs` Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` numArchs : int number of supported architectures. """ cdef int numArchs = 0 err = cnvrtc.nvrtcGetNumSupportedArchs(&numArchs) return (nvrtcResult(err), numArchs) @cython.embedsignature(True) def nvrtcGetSupportedArchs(): """ 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`. see :py:obj:`~.nvrtcGetNumSupportedArchs` Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` supportedArchs : List[int] sorted array of supported architectures. """ cdef vector[int] supportedArchs _, s = nvrtcGetNumSupportedArchs() supportedArchs.resize(s) err = cnvrtc.nvrtcGetSupportedArchs(supportedArchs.data()) return (nvrtcResult(err), supportedArchs) @cython.embedsignature(True) def nvrtcCreateProgram(char* src, char* name, int numHeaders, headers : Optional[Tuple[bytes] | List[bytes]], includeNames : Optional[Tuple[bytes] | List[bytes]]): """ nvrtcCreateProgram creates an instance of nvrtcProgram with the given input parameters, and sets the output parameter `prog` with it. Parameters ---------- src : bytes CUDA program source. name : bytes CUDA program name. `name` can be `NULL`; `"default_program"` is used when `name` is `NULL` or "". numHeaders : int Number of headers used. `numHeaders` must be greater than or equal to 0. headers : List[bytes] Sources of the headers. `headers` can be `NULL` when `numHeaders` is 0. includeNames : List[bytes] Name of each header by which they can be included in the CUDA program source. `includeNames` can be `NULL` when `numHeaders` is 0. These headers must be included with the exact names specified here. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_OUT_OF_MEMORY` - :py:obj:`~.NVRTC_ERROR_PROGRAM_CREATION_FAILURE` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. See Also -------- :py:obj:`~.nvrtcDestroyProgram` """ includeNames = [] if includeNames is None else includeNames if not all(isinstance(_x, (bytes)) for _x in includeNames): raise TypeError("Argument 'includeNames' is not instance of type (expected Tuple[bytes] or List[bytes]") headers = [] if headers is None else headers if not all(isinstance(_x, (bytes)) for _x in headers): raise TypeError("Argument 'headers' is not instance of type (expected Tuple[bytes] or List[bytes]") cdef nvrtcProgram prog = nvrtcProgram() if numHeaders > len(headers): raise RuntimeError("List is too small: " + str(len(headers)) + " < " + str(numHeaders)) if numHeaders > len(includeNames): raise RuntimeError("List is too small: " + str(len(includeNames)) + " < " + str(numHeaders)) cdef vector[const char*] cheaders = headers cdef vector[const char*] cincludeNames = includeNames err = cnvrtc.nvrtcCreateProgram(prog._ptr, src, name, numHeaders, cheaders.data(), cincludeNames.data()) return (nvrtcResult(err), prog) @cython.embedsignature(True) def nvrtcDestroyProgram(prog): """ nvrtcDestroyProgram destroys the given program. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` See Also -------- :py:obj:`~.nvrtcCreateProgram` """ cdef cnvrtc.nvrtcProgram *cprog if prog is None: cprog = NULL elif isinstance(prog, (nvrtcProgram,)): pprog = prog.getPtr() cprog = pprog elif isinstance(prog, (int)): cprog = prog else: raise TypeError("Argument 'prog' is not instance of type (expected , found " + str(type(prog))) err = cnvrtc.nvrtcDestroyProgram(cprog) return (nvrtcResult(err),) @cython.embedsignature(True) def nvrtcCompileProgram(prog, int numOptions, options : Optional[Tuple[bytes] | List[bytes]]): """ nvrtcCompileProgram compiles the given program. It supports compile options listed in :py:obj:`~.Supported Compile Options`. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. numOptions : int Number of compiler options passed. options : List[bytes] Compiler options in the form of C string array. `options` can be `NULL` when `numOptions` is 0. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_OUT_OF_MEMORY` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` - :py:obj:`~.NVRTC_ERROR_INVALID_OPTION` - :py:obj:`~.NVRTC_ERROR_COMPILATION` - :py:obj:`~.NVRTC_ERROR_BUILTIN_OPERATION_FAILURE` - :py:obj:`~.NVRTC_ERROR_TIME_FILE_WRITE_FAILED` """ options = [] if options is None else options if not all(isinstance(_x, (bytes)) for _x in options): raise TypeError("Argument 'options' is not instance of type (expected Tuple[bytes] or List[bytes]") cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog if numOptions > len(options): raise RuntimeError("List is too small: " + str(len(options)) + " < " + str(numOptions)) cdef vector[const char*] coptions = options err = cnvrtc.nvrtcCompileProgram(cprog, numOptions, coptions.data()) return (nvrtcResult(err),) @cython.embedsignature(True) def nvrtcGetPTXSize(prog): """ nvrtcGetPTXSize sets the value of `ptxSizeRet` with the size of the PTX generated by the previous compilation of `prog` (including the trailing `NULL`). Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` ptxSizeRet : int Size of the generated PTX (including the trailing `NULL`). See Also -------- :py:obj:`~.nvrtcGetPTX` """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog cdef size_t ptxSizeRet = 0 err = cnvrtc.nvrtcGetPTXSize(cprog, &ptxSizeRet) return (nvrtcResult(err), ptxSizeRet) @cython.embedsignature(True) def nvrtcGetPTX(prog, char* ptx): """ nvrtcGetPTX stores the PTX generated by the previous compilation of `prog` in the memory pointed by `ptx`. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. ptx : bytes Compiled result. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` See Also -------- :py:obj:`~.nvrtcGetPTXSize` """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog err = cnvrtc.nvrtcGetPTX(cprog, ptx) return (nvrtcResult(err),) @cython.embedsignature(True) def nvrtcGetCUBINSize(prog): """ 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. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` cubinSizeRet : int Size of the generated cubin. See Also -------- :py:obj:`~.nvrtcGetCUBIN` """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog cdef size_t cubinSizeRet = 0 err = cnvrtc.nvrtcGetCUBINSize(cprog, &cubinSizeRet) return (nvrtcResult(err), cubinSizeRet) @cython.embedsignature(True) def nvrtcGetCUBIN(prog, char* cubin): """ 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. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. cubin : bytes Compiled and assembled result. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` See Also -------- :py:obj:`~.nvrtcGetCUBINSize` """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog err = cnvrtc.nvrtcGetCUBIN(cprog, cubin) return (nvrtcResult(err),) @cython.embedsignature(True) def nvrtcGetNVVMSize(prog): """ DEPRECATION NOTICE: This function will be removed in a future release. Please use nvrtcGetLTOIRSize (and nvrtcGetLTOIR) instead. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` None Returns ------- nvrtcResult nvvmSizeRet : int None """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog cdef size_t nvvmSizeRet = 0 err = cnvrtc.nvrtcGetNVVMSize(cprog, &nvvmSizeRet) return (nvrtcResult(err), nvvmSizeRet) @cython.embedsignature(True) def nvrtcGetNVVM(prog, char* nvvm): """ DEPRECATION NOTICE: This function will be removed in a future release. Please use nvrtcGetLTOIR (and nvrtcGetLTOIRSize) instead. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` None nvvm : bytes None Returns ------- nvrtcResult """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog err = cnvrtc.nvrtcGetNVVM(cprog, nvvm) return (nvrtcResult(err),) @cython.embedsignature(True) def nvrtcGetLTOIRSize(prog): """ 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`. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` LTOIRSizeRet : int Size of the generated LTO IR. See Also -------- :py:obj:`~.nvrtcGetLTOIR` """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog cdef size_t LTOIRSizeRet = 0 err = cnvrtc.nvrtcGetLTOIRSize(cprog, <OIRSizeRet) return (nvrtcResult(err), LTOIRSizeRet) @cython.embedsignature(True) def nvrtcGetLTOIR(prog, char* LTOIR): """ 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`. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. LTOIR : bytes Compiled result. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` See Also -------- :py:obj:`~.nvrtcGetLTOIRSize` """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog err = cnvrtc.nvrtcGetLTOIR(cprog, LTOIR) return (nvrtcResult(err),) @cython.embedsignature(True) def nvrtcGetOptiXIRSize(prog): """ 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. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` optixirSizeRet : int Size of the generated LTO IR. See Also -------- :py:obj:`~.nvrtcGetOptiXIR` """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog cdef size_t optixirSizeRet = 0 err = cnvrtc.nvrtcGetOptiXIRSize(cprog, &optixirSizeRet) return (nvrtcResult(err), optixirSizeRet) @cython.embedsignature(True) def nvrtcGetOptiXIR(prog, char* optixir): """ 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. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. Optix : bytes IR Compiled result. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` See Also -------- :py:obj:`~.nvrtcGetOptiXIRSize` """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog err = cnvrtc.nvrtcGetOptiXIR(cprog, optixir) return (nvrtcResult(err),) @cython.embedsignature(True) def nvrtcGetProgramLogSize(prog): """ nvrtcGetProgramLogSize sets `logSizeRet` with the size of the log generated by the previous compilation of `prog` (including the trailing `NULL`). Note that compilation log may be generated with warnings and informative messages, even when the compilation of `prog` succeeds. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` logSizeRet : int Size of the compilation log (including the trailing `NULL`). See Also -------- :py:obj:`~.nvrtcGetProgramLog` """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog cdef size_t logSizeRet = 0 err = cnvrtc.nvrtcGetProgramLogSize(cprog, &logSizeRet) return (nvrtcResult(err), logSizeRet) @cython.embedsignature(True) def nvrtcGetProgramLog(prog, char* log): """ nvrtcGetProgramLog stores the log generated by the previous compilation of `prog` in the memory pointed by `log`. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. log : bytes Compilation log. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` See Also -------- :py:obj:`~.nvrtcGetProgramLogSize` """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog err = cnvrtc.nvrtcGetProgramLog(cprog, log) return (nvrtcResult(err),) @cython.embedsignature(True) def nvrtcAddNameExpression(prog, char* name_expression): """ nvrtcAddNameExpression notes the given name expression denoting the address of a global function or device/__constant__ variable. The identical name expression string must be provided on a subsequent call to nvrtcGetLoweredName to extract the lowered name. Parameters ---------- prog : :py:obj:`~.nvrtcProgram` CUDA Runtime Compilation program. name_expression : bytes constant expression denoting the address of a global function or device/__constant__ variable. Returns ------- nvrtcResult - :py:obj:`~.NVRTC_SUCCESS` - :py:obj:`~.NVRTC_ERROR_INVALID_PROGRAM` - :py:obj:`~.NVRTC_ERROR_INVALID_INPUT` - :py:obj:`~.NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION` See Also -------- :py:obj:`~.nvrtcGetLoweredName` """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog err = cnvrtc.nvrtcAddNameExpression(cprog, name_expression) return (nvrtcResult(err),) @cython.embedsignature(True) def nvrtcGetLoweredName(prog, char* name_expression): """ 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. Parameters ---------- prog : nvrtcProgram CUDA Runtime Compilation program. name_expression : bytes constant expression denoting the address of a global function or device/__constant__ variable. Returns ------- nvrtcResult NVRTC_SUCCESS NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID lowered_name : bytes initialized by the function to point to a C string containing the lowered (mangled) name corresponding to the provided name expression. See Also -------- nvrtcAddNameExpression """ cdef cnvrtc.nvrtcProgram cprog if prog is None: cprog = 0 elif isinstance(prog, (nvrtcProgram,)): pprog = int(prog) cprog = pprog else: pprog = int(nvrtcProgram(prog)) cprog = pprog cdef const char* lowered_name = NULL err = cnvrtc.nvrtcGetLoweredName(cprog, name_expression, &lowered_name) return (nvrtcResult(err), lowered_name) @cython.embedsignature(True) def sizeof(objType): """ Returns the size of provided CUDA Python structure in bytes Parameters ---------- objType : Any CUDA Python object Returns ------- lowered_name : int The size of `objType` in bytes """ if objType == nvrtcProgram: return sizeof(cnvrtc.nvrtcProgram) raise TypeError("Unknown type: " + str(objType))