diff --git a/cuda_bindings/build_hooks.py b/cuda_bindings/build_hooks.py index ce4745f7a0..094d8adfbf 100644 --- a/cuda_bindings/build_hooks.py +++ b/cuda_bindings/build_hooks.py @@ -94,9 +94,7 @@ def _get_cuda_path() -> str: "driver_functions.h", "cuda_profiler_api.h", ], - "nvrtc": [ - "nvrtc.h", - ], + # nvrtc: headers no longer parsed at build time (pre-generated by cybind). # During compilation, Cython will reference C headers that are not # explicitly parsed above. These are the known dependencies: # @@ -423,7 +421,8 @@ def _cleanup_dst_files(): sources_list = [ # private (["cuda/bindings/_bindings/cydriver.pyx", "cuda/bindings/_bindings/loader.cpp"], None), - (["cuda/bindings/_bindings/cynvrtc.pyx"], None), + # cynvrtc.pyx is now in cuda/bindings/ (pre-generated by cybind), + # picked up by the cuda_bindings_files glob below. (["cuda/bindings/_bindings/cyruntime.pyx"], static_runtime_libraries), (["cuda/bindings/_bindings/cyruntime_ptds.pyx"], static_runtime_libraries), # utils diff --git a/cuda_bindings/cuda/bindings/_bindings/cynvrtc.pyx.in b/cuda_bindings/cuda/bindings/_bindings/cynvrtc.pyx.in deleted file mode 100644 index 2e1c7a67cc..0000000000 --- a/cuda_bindings/cuda/bindings/_bindings/cynvrtc.pyx.in +++ /dev/null @@ -1,789 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE - -# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1422+gf4812259e.d20260318. Do not modify it directly. -{{if 'Windows' == platform.system()}} -import os -cimport cuda.bindings._lib.windll as windll -{{else}} -cimport cuda.bindings._lib.dlfcn as dlfcn -{{endif}} -from cuda.pathfinder import load_nvidia_dynamic_lib -from libc.stdint cimport intptr_t, uintptr_t -import threading - -cdef object __symbol_lock = threading.Lock() -cdef bint __cuPythonInit = False -{{if 'nvrtcGetErrorString' in found_functions}}cdef void *__nvrtcGetErrorString = NULL{{endif}} -{{if 'nvrtcVersion' in found_functions}}cdef void *__nvrtcVersion = NULL{{endif}} -{{if 'nvrtcGetNumSupportedArchs' in found_functions}}cdef void *__nvrtcGetNumSupportedArchs = NULL{{endif}} -{{if 'nvrtcGetSupportedArchs' in found_functions}}cdef void *__nvrtcGetSupportedArchs = NULL{{endif}} -{{if 'nvrtcCreateProgram' in found_functions}}cdef void *__nvrtcCreateProgram = NULL{{endif}} -{{if 'nvrtcDestroyProgram' in found_functions}}cdef void *__nvrtcDestroyProgram = NULL{{endif}} -{{if 'nvrtcCompileProgram' in found_functions}}cdef void *__nvrtcCompileProgram = NULL{{endif}} -{{if 'nvrtcGetPTXSize' in found_functions}}cdef void *__nvrtcGetPTXSize = NULL{{endif}} -{{if 'nvrtcGetPTX' in found_functions}}cdef void *__nvrtcGetPTX = NULL{{endif}} -{{if 'nvrtcGetCUBINSize' in found_functions}}cdef void *__nvrtcGetCUBINSize = NULL{{endif}} -{{if 'nvrtcGetCUBIN' in found_functions}}cdef void *__nvrtcGetCUBIN = NULL{{endif}} -{{if 'nvrtcGetLTOIRSize' in found_functions}}cdef void *__nvrtcGetLTOIRSize = NULL{{endif}} -{{if 'nvrtcGetLTOIR' in found_functions}}cdef void *__nvrtcGetLTOIR = NULL{{endif}} -{{if 'nvrtcGetOptiXIRSize' in found_functions}}cdef void *__nvrtcGetOptiXIRSize = NULL{{endif}} -{{if 'nvrtcGetOptiXIR' in found_functions}}cdef void *__nvrtcGetOptiXIR = NULL{{endif}} -{{if 'nvrtcGetProgramLogSize' in found_functions}}cdef void *__nvrtcGetProgramLogSize = NULL{{endif}} -{{if 'nvrtcGetProgramLog' in found_functions}}cdef void *__nvrtcGetProgramLog = NULL{{endif}} -{{if 'nvrtcAddNameExpression' in found_functions}}cdef void *__nvrtcAddNameExpression = NULL{{endif}} -{{if 'nvrtcGetLoweredName' in found_functions}}cdef void *__nvrtcGetLoweredName = NULL{{endif}} -{{if 'nvrtcGetPCHHeapSize' in found_functions}}cdef void *__nvrtcGetPCHHeapSize = NULL{{endif}} -{{if 'nvrtcSetPCHHeapSize' in found_functions}}cdef void *__nvrtcSetPCHHeapSize = NULL{{endif}} -{{if 'nvrtcGetPCHCreateStatus' in found_functions}}cdef void *__nvrtcGetPCHCreateStatus = NULL{{endif}} -{{if 'nvrtcGetPCHHeapSizeRequired' in found_functions}}cdef void *__nvrtcGetPCHHeapSizeRequired = NULL{{endif}} -{{if 'nvrtcSetFlowCallback' in found_functions}}cdef void *__nvrtcSetFlowCallback = NULL{{endif}} -{{if 'nvrtcGetTileIRSize' in found_functions}}cdef void *__nvrtcGetTileIRSize = NULL{{endif}} -{{if 'nvrtcGetTileIR' in found_functions}}cdef void *__nvrtcGetTileIR = NULL{{endif}} - -cdef int _cuPythonInit() except -1 nogil: - global __cuPythonInit - - # Load library - with gil, __symbol_lock: - {{if 'Windows' == platform.system()}} - handle = load_nvidia_dynamic_lib("nvrtc")._handle_uint - - # Load function - {{if 'nvrtcGetErrorString' in found_functions}} - global __nvrtcGetErrorString - __nvrtcGetErrorString = windll.GetProcAddress(handle, 'nvrtcGetErrorString') - {{endif}} - {{if 'nvrtcVersion' in found_functions}} - global __nvrtcVersion - __nvrtcVersion = windll.GetProcAddress(handle, 'nvrtcVersion') - {{endif}} - {{if 'nvrtcGetNumSupportedArchs' in found_functions}} - global __nvrtcGetNumSupportedArchs - __nvrtcGetNumSupportedArchs = windll.GetProcAddress(handle, 'nvrtcGetNumSupportedArchs') - {{endif}} - {{if 'nvrtcGetSupportedArchs' in found_functions}} - global __nvrtcGetSupportedArchs - __nvrtcGetSupportedArchs = windll.GetProcAddress(handle, 'nvrtcGetSupportedArchs') - {{endif}} - {{if 'nvrtcCreateProgram' in found_functions}} - global __nvrtcCreateProgram - __nvrtcCreateProgram = windll.GetProcAddress(handle, 'nvrtcCreateProgram') - {{endif}} - {{if 'nvrtcDestroyProgram' in found_functions}} - global __nvrtcDestroyProgram - __nvrtcDestroyProgram = windll.GetProcAddress(handle, 'nvrtcDestroyProgram') - {{endif}} - {{if 'nvrtcCompileProgram' in found_functions}} - global __nvrtcCompileProgram - __nvrtcCompileProgram = windll.GetProcAddress(handle, 'nvrtcCompileProgram') - {{endif}} - {{if 'nvrtcGetPTXSize' in found_functions}} - global __nvrtcGetPTXSize - __nvrtcGetPTXSize = windll.GetProcAddress(handle, 'nvrtcGetPTXSize') - {{endif}} - {{if 'nvrtcGetPTX' in found_functions}} - global __nvrtcGetPTX - __nvrtcGetPTX = windll.GetProcAddress(handle, 'nvrtcGetPTX') - {{endif}} - {{if 'nvrtcGetCUBINSize' in found_functions}} - global __nvrtcGetCUBINSize - __nvrtcGetCUBINSize = windll.GetProcAddress(handle, 'nvrtcGetCUBINSize') - {{endif}} - {{if 'nvrtcGetCUBIN' in found_functions}} - global __nvrtcGetCUBIN - __nvrtcGetCUBIN = windll.GetProcAddress(handle, 'nvrtcGetCUBIN') - {{endif}} - {{if 'nvrtcGetLTOIRSize' in found_functions}} - global __nvrtcGetLTOIRSize - __nvrtcGetLTOIRSize = windll.GetProcAddress(handle, 'nvrtcGetLTOIRSize') - {{endif}} - {{if 'nvrtcGetLTOIR' in found_functions}} - global __nvrtcGetLTOIR - __nvrtcGetLTOIR = windll.GetProcAddress(handle, 'nvrtcGetLTOIR') - {{endif}} - {{if 'nvrtcGetOptiXIRSize' in found_functions}} - global __nvrtcGetOptiXIRSize - __nvrtcGetOptiXIRSize = windll.GetProcAddress(handle, 'nvrtcGetOptiXIRSize') - {{endif}} - {{if 'nvrtcGetOptiXIR' in found_functions}} - global __nvrtcGetOptiXIR - __nvrtcGetOptiXIR = windll.GetProcAddress(handle, 'nvrtcGetOptiXIR') - {{endif}} - {{if 'nvrtcGetProgramLogSize' in found_functions}} - global __nvrtcGetProgramLogSize - __nvrtcGetProgramLogSize = windll.GetProcAddress(handle, 'nvrtcGetProgramLogSize') - {{endif}} - {{if 'nvrtcGetProgramLog' in found_functions}} - global __nvrtcGetProgramLog - __nvrtcGetProgramLog = windll.GetProcAddress(handle, 'nvrtcGetProgramLog') - {{endif}} - {{if 'nvrtcAddNameExpression' in found_functions}} - global __nvrtcAddNameExpression - __nvrtcAddNameExpression = windll.GetProcAddress(handle, 'nvrtcAddNameExpression') - {{endif}} - {{if 'nvrtcGetLoweredName' in found_functions}} - global __nvrtcGetLoweredName - __nvrtcGetLoweredName = windll.GetProcAddress(handle, 'nvrtcGetLoweredName') - {{endif}} - {{if 'nvrtcGetPCHHeapSize' in found_functions}} - global __nvrtcGetPCHHeapSize - __nvrtcGetPCHHeapSize = windll.GetProcAddress(handle, 'nvrtcGetPCHHeapSize') - {{endif}} - {{if 'nvrtcSetPCHHeapSize' in found_functions}} - global __nvrtcSetPCHHeapSize - __nvrtcSetPCHHeapSize = windll.GetProcAddress(handle, 'nvrtcSetPCHHeapSize') - {{endif}} - {{if 'nvrtcGetPCHCreateStatus' in found_functions}} - global __nvrtcGetPCHCreateStatus - __nvrtcGetPCHCreateStatus = windll.GetProcAddress(handle, 'nvrtcGetPCHCreateStatus') - {{endif}} - {{if 'nvrtcGetPCHHeapSizeRequired' in found_functions}} - global __nvrtcGetPCHHeapSizeRequired - __nvrtcGetPCHHeapSizeRequired = windll.GetProcAddress(handle, 'nvrtcGetPCHHeapSizeRequired') - {{endif}} - {{if 'nvrtcSetFlowCallback' in found_functions}} - global __nvrtcSetFlowCallback - __nvrtcSetFlowCallback = windll.GetProcAddress(handle, 'nvrtcSetFlowCallback') - {{endif}} - {{if 'nvrtcGetTileIRSize' in found_functions}} - global __nvrtcGetTileIRSize - __nvrtcGetTileIRSize = windll.GetProcAddress(handle, 'nvrtcGetTileIRSize') - {{endif}} - {{if 'nvrtcGetTileIR' in found_functions}} - global __nvrtcGetTileIR - __nvrtcGetTileIR = windll.GetProcAddress(handle, 'nvrtcGetTileIR') - {{endif}} - - {{else}} - handle = (load_nvidia_dynamic_lib("nvrtc")._handle_uint) - - # Load function - {{if 'nvrtcGetErrorString' in found_functions}} - global __nvrtcGetErrorString - __nvrtcGetErrorString = dlfcn.dlsym(handle, 'nvrtcGetErrorString') - {{endif}} - {{if 'nvrtcVersion' in found_functions}} - global __nvrtcVersion - __nvrtcVersion = dlfcn.dlsym(handle, 'nvrtcVersion') - {{endif}} - {{if 'nvrtcGetNumSupportedArchs' in found_functions}} - global __nvrtcGetNumSupportedArchs - __nvrtcGetNumSupportedArchs = dlfcn.dlsym(handle, 'nvrtcGetNumSupportedArchs') - {{endif}} - {{if 'nvrtcGetSupportedArchs' in found_functions}} - global __nvrtcGetSupportedArchs - __nvrtcGetSupportedArchs = dlfcn.dlsym(handle, 'nvrtcGetSupportedArchs') - {{endif}} - {{if 'nvrtcCreateProgram' in found_functions}} - global __nvrtcCreateProgram - __nvrtcCreateProgram = dlfcn.dlsym(handle, 'nvrtcCreateProgram') - {{endif}} - {{if 'nvrtcDestroyProgram' in found_functions}} - global __nvrtcDestroyProgram - __nvrtcDestroyProgram = dlfcn.dlsym(handle, 'nvrtcDestroyProgram') - {{endif}} - {{if 'nvrtcCompileProgram' in found_functions}} - global __nvrtcCompileProgram - __nvrtcCompileProgram = dlfcn.dlsym(handle, 'nvrtcCompileProgram') - {{endif}} - {{if 'nvrtcGetPTXSize' in found_functions}} - global __nvrtcGetPTXSize - __nvrtcGetPTXSize = dlfcn.dlsym(handle, 'nvrtcGetPTXSize') - {{endif}} - {{if 'nvrtcGetPTX' in found_functions}} - global __nvrtcGetPTX - __nvrtcGetPTX = dlfcn.dlsym(handle, 'nvrtcGetPTX') - {{endif}} - {{if 'nvrtcGetCUBINSize' in found_functions}} - global __nvrtcGetCUBINSize - __nvrtcGetCUBINSize = dlfcn.dlsym(handle, 'nvrtcGetCUBINSize') - {{endif}} - {{if 'nvrtcGetCUBIN' in found_functions}} - global __nvrtcGetCUBIN - __nvrtcGetCUBIN = dlfcn.dlsym(handle, 'nvrtcGetCUBIN') - {{endif}} - {{if 'nvrtcGetLTOIRSize' in found_functions}} - global __nvrtcGetLTOIRSize - __nvrtcGetLTOIRSize = dlfcn.dlsym(handle, 'nvrtcGetLTOIRSize') - {{endif}} - {{if 'nvrtcGetLTOIR' in found_functions}} - global __nvrtcGetLTOIR - __nvrtcGetLTOIR = dlfcn.dlsym(handle, 'nvrtcGetLTOIR') - {{endif}} - {{if 'nvrtcGetOptiXIRSize' in found_functions}} - global __nvrtcGetOptiXIRSize - __nvrtcGetOptiXIRSize = dlfcn.dlsym(handle, 'nvrtcGetOptiXIRSize') - {{endif}} - {{if 'nvrtcGetOptiXIR' in found_functions}} - global __nvrtcGetOptiXIR - __nvrtcGetOptiXIR = dlfcn.dlsym(handle, 'nvrtcGetOptiXIR') - {{endif}} - {{if 'nvrtcGetProgramLogSize' in found_functions}} - global __nvrtcGetProgramLogSize - __nvrtcGetProgramLogSize = dlfcn.dlsym(handle, 'nvrtcGetProgramLogSize') - {{endif}} - {{if 'nvrtcGetProgramLog' in found_functions}} - global __nvrtcGetProgramLog - __nvrtcGetProgramLog = dlfcn.dlsym(handle, 'nvrtcGetProgramLog') - {{endif}} - {{if 'nvrtcAddNameExpression' in found_functions}} - global __nvrtcAddNameExpression - __nvrtcAddNameExpression = dlfcn.dlsym(handle, 'nvrtcAddNameExpression') - {{endif}} - {{if 'nvrtcGetLoweredName' in found_functions}} - global __nvrtcGetLoweredName - __nvrtcGetLoweredName = dlfcn.dlsym(handle, 'nvrtcGetLoweredName') - {{endif}} - {{if 'nvrtcGetPCHHeapSize' in found_functions}} - global __nvrtcGetPCHHeapSize - __nvrtcGetPCHHeapSize = dlfcn.dlsym(handle, 'nvrtcGetPCHHeapSize') - {{endif}} - {{if 'nvrtcSetPCHHeapSize' in found_functions}} - global __nvrtcSetPCHHeapSize - __nvrtcSetPCHHeapSize = dlfcn.dlsym(handle, 'nvrtcSetPCHHeapSize') - {{endif}} - {{if 'nvrtcGetPCHCreateStatus' in found_functions}} - global __nvrtcGetPCHCreateStatus - __nvrtcGetPCHCreateStatus = dlfcn.dlsym(handle, 'nvrtcGetPCHCreateStatus') - {{endif}} - {{if 'nvrtcGetPCHHeapSizeRequired' in found_functions}} - global __nvrtcGetPCHHeapSizeRequired - __nvrtcGetPCHHeapSizeRequired = dlfcn.dlsym(handle, 'nvrtcGetPCHHeapSizeRequired') - {{endif}} - {{if 'nvrtcSetFlowCallback' in found_functions}} - global __nvrtcSetFlowCallback - __nvrtcSetFlowCallback = dlfcn.dlsym(handle, 'nvrtcSetFlowCallback') - {{endif}} - {{if 'nvrtcGetTileIRSize' in found_functions}} - global __nvrtcGetTileIRSize - __nvrtcGetTileIRSize = dlfcn.dlsym(handle, 'nvrtcGetTileIRSize') - {{endif}} - {{if 'nvrtcGetTileIR' in found_functions}} - global __nvrtcGetTileIR - __nvrtcGetTileIR = dlfcn.dlsym(handle, 'nvrtcGetTileIR') - {{endif}} - - {{endif}} - __cuPythonInit = True - return 0 - -# Create a very small function to check whether we are init'ed, so the C -# compiler can inline it. -cdef inline int cuPythonInit() except -1 nogil: - if __cuPythonInit: - return 0 - return _cuPythonInit() - -{{if 'nvrtcGetErrorString' in found_functions}} - -cdef const char* _nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil: - global __nvrtcGetErrorString - cuPythonInit() - if __nvrtcGetErrorString == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetErrorString" not found') - err = ( __nvrtcGetErrorString)(result) - return err -{{endif}} - -{{if 'nvrtcVersion' in found_functions}} - -cdef nvrtcResult _nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcVersion - cuPythonInit() - if __nvrtcVersion == NULL: - with gil: - raise RuntimeError('Function "nvrtcVersion" not found') - err = ( __nvrtcVersion)(major, minor) - return err -{{endif}} - -{{if 'nvrtcGetNumSupportedArchs' in found_functions}} - -cdef nvrtcResult _nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetNumSupportedArchs - cuPythonInit() - if __nvrtcGetNumSupportedArchs == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetNumSupportedArchs" not found') - err = ( __nvrtcGetNumSupportedArchs)(numArchs) - return err -{{endif}} - -{{if 'nvrtcGetSupportedArchs' in found_functions}} - -cdef nvrtcResult _nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetSupportedArchs - cuPythonInit() - if __nvrtcGetSupportedArchs == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetSupportedArchs" not found') - err = ( __nvrtcGetSupportedArchs)(supportedArchs) - return err -{{endif}} - -{{if 'nvrtcCreateProgram' in found_functions}} - -cdef nvrtcResult _nvrtcCreateProgram(nvrtcProgram* prog, const char* src, const char* name, int numHeaders, const char** headers, const char** includeNames) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcCreateProgram - cuPythonInit() - if __nvrtcCreateProgram == NULL: - with gil: - raise RuntimeError('Function "nvrtcCreateProgram" not found') - err = ( __nvrtcCreateProgram)(prog, src, name, numHeaders, headers, includeNames) - return err -{{endif}} - -{{if 'nvrtcDestroyProgram' in found_functions}} - -cdef nvrtcResult _nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcDestroyProgram - cuPythonInit() - if __nvrtcDestroyProgram == NULL: - with gil: - raise RuntimeError('Function "nvrtcDestroyProgram" not found') - err = ( __nvrtcDestroyProgram)(prog) - return err -{{endif}} - -{{if 'nvrtcCompileProgram' in found_functions}} - -cdef nvrtcResult _nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const char** options) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcCompileProgram - cuPythonInit() - if __nvrtcCompileProgram == NULL: - with gil: - raise RuntimeError('Function "nvrtcCompileProgram" not found') - err = ( __nvrtcCompileProgram)(prog, numOptions, options) - return err -{{endif}} - -{{if 'nvrtcGetPTXSize' in found_functions}} - -cdef nvrtcResult _nvrtcGetPTXSize(nvrtcProgram prog, size_t* ptxSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetPTXSize - cuPythonInit() - if __nvrtcGetPTXSize == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetPTXSize" not found') - err = ( __nvrtcGetPTXSize)(prog, ptxSizeRet) - return err -{{endif}} - -{{if 'nvrtcGetPTX' in found_functions}} - -cdef nvrtcResult _nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetPTX - cuPythonInit() - if __nvrtcGetPTX == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetPTX" not found') - err = ( __nvrtcGetPTX)(prog, ptx) - return err -{{endif}} - -{{if 'nvrtcGetCUBINSize' in found_functions}} - -cdef nvrtcResult _nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetCUBINSize - cuPythonInit() - if __nvrtcGetCUBINSize == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetCUBINSize" not found') - err = ( __nvrtcGetCUBINSize)(prog, cubinSizeRet) - return err -{{endif}} - -{{if 'nvrtcGetCUBIN' in found_functions}} - -cdef nvrtcResult _nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetCUBIN - cuPythonInit() - if __nvrtcGetCUBIN == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetCUBIN" not found') - err = ( __nvrtcGetCUBIN)(prog, cubin) - return err -{{endif}} - -{{if 'nvrtcGetLTOIRSize' in found_functions}} - -cdef nvrtcResult _nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetLTOIRSize - cuPythonInit() - if __nvrtcGetLTOIRSize == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetLTOIRSize" not found') - err = ( __nvrtcGetLTOIRSize)(prog, LTOIRSizeRet) - return err -{{endif}} - -{{if 'nvrtcGetLTOIR' in found_functions}} - -cdef nvrtcResult _nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetLTOIR - cuPythonInit() - if __nvrtcGetLTOIR == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetLTOIR" not found') - err = ( __nvrtcGetLTOIR)(prog, LTOIR) - return err -{{endif}} - -{{if 'nvrtcGetOptiXIRSize' in found_functions}} - -cdef nvrtcResult _nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetOptiXIRSize - cuPythonInit() - if __nvrtcGetOptiXIRSize == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetOptiXIRSize" not found') - err = ( __nvrtcGetOptiXIRSize)(prog, optixirSizeRet) - return err -{{endif}} - -{{if 'nvrtcGetOptiXIR' in found_functions}} - -cdef nvrtcResult _nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetOptiXIR - cuPythonInit() - if __nvrtcGetOptiXIR == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetOptiXIR" not found') - err = ( __nvrtcGetOptiXIR)(prog, optixir) - return err -{{endif}} - -{{if 'nvrtcGetProgramLogSize' in found_functions}} - -cdef nvrtcResult _nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetProgramLogSize - cuPythonInit() - if __nvrtcGetProgramLogSize == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetProgramLogSize" not found') - err = ( __nvrtcGetProgramLogSize)(prog, logSizeRet) - return err -{{endif}} - -{{if 'nvrtcGetProgramLog' in found_functions}} - -cdef nvrtcResult _nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetProgramLog - cuPythonInit() - if __nvrtcGetProgramLog == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetProgramLog" not found') - err = ( __nvrtcGetProgramLog)(prog, log) - return err -{{endif}} - -{{if 'nvrtcAddNameExpression' in found_functions}} - -cdef nvrtcResult _nvrtcAddNameExpression(nvrtcProgram prog, const char* name_expression) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcAddNameExpression - cuPythonInit() - if __nvrtcAddNameExpression == NULL: - with gil: - raise RuntimeError('Function "nvrtcAddNameExpression" not found') - err = ( __nvrtcAddNameExpression)(prog, name_expression) - return err -{{endif}} - -{{if 'nvrtcGetLoweredName' in found_functions}} - -cdef nvrtcResult _nvrtcGetLoweredName(nvrtcProgram prog, const char* name_expression, const char** lowered_name) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetLoweredName - cuPythonInit() - if __nvrtcGetLoweredName == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetLoweredName" not found') - err = ( __nvrtcGetLoweredName)(prog, name_expression, lowered_name) - return err -{{endif}} - -{{if 'nvrtcGetPCHHeapSize' in found_functions}} - -cdef nvrtcResult _nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetPCHHeapSize - cuPythonInit() - if __nvrtcGetPCHHeapSize == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetPCHHeapSize" not found') - err = ( __nvrtcGetPCHHeapSize)(ret) - return err -{{endif}} - -{{if 'nvrtcSetPCHHeapSize' in found_functions}} - -cdef nvrtcResult _nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcSetPCHHeapSize - cuPythonInit() - if __nvrtcSetPCHHeapSize == NULL: - with gil: - raise RuntimeError('Function "nvrtcSetPCHHeapSize" not found') - err = ( __nvrtcSetPCHHeapSize)(size) - return err -{{endif}} - -{{if 'nvrtcGetPCHCreateStatus' in found_functions}} - -cdef nvrtcResult _nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetPCHCreateStatus - cuPythonInit() - if __nvrtcGetPCHCreateStatus == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetPCHCreateStatus" not found') - err = ( __nvrtcGetPCHCreateStatus)(prog) - return err -{{endif}} - -{{if 'nvrtcGetPCHHeapSizeRequired' in found_functions}} - -cdef nvrtcResult _nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetPCHHeapSizeRequired - cuPythonInit() - if __nvrtcGetPCHHeapSizeRequired == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetPCHHeapSizeRequired" not found') - err = ( __nvrtcGetPCHHeapSizeRequired)(prog, size) - return err -{{endif}} - -{{if 'nvrtcSetFlowCallback' in found_functions}} - -cdef nvrtcResult _nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* payload) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcSetFlowCallback - cuPythonInit() - if __nvrtcSetFlowCallback == NULL: - with gil: - raise RuntimeError('Function "nvrtcSetFlowCallback" not found') - err = ( __nvrtcSetFlowCallback)(prog, callback, payload) - return err -{{endif}} - -{{if 'nvrtcGetTileIRSize' in found_functions}} - -cdef nvrtcResult _nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetTileIRSize - cuPythonInit() - if __nvrtcGetTileIRSize == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetTileIRSize" not found') - err = ( __nvrtcGetTileIRSize)(prog, TileIRSizeRet) - return err -{{endif}} - -{{if 'nvrtcGetTileIR' in found_functions}} - -cdef nvrtcResult _nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: - global __nvrtcGetTileIR - cuPythonInit() - if __nvrtcGetTileIR == NULL: - with gil: - raise RuntimeError('Function "nvrtcGetTileIR" not found') - err = ( __nvrtcGetTileIR)(prog, TileIR) - return err -{{endif}} - -cdef dict func_ptrs = None - -cpdef dict _inspect_function_pointers(): - global func_ptrs - if func_ptrs is not None: - return func_ptrs - - cuPythonInit() - cdef dict data = {} - - {{if 'nvrtcGetErrorString' in found_functions}} - global __nvrtcGetErrorString - data["__nvrtcGetErrorString"] = __nvrtcGetErrorString - {{else}} - data["__nvrtcGetErrorString"] = 0 - {{endif}} - - {{if 'nvrtcVersion' in found_functions}} - global __nvrtcVersion - data["__nvrtcVersion"] = __nvrtcVersion - {{else}} - data["__nvrtcVersion"] = 0 - {{endif}} - - {{if 'nvrtcGetNumSupportedArchs' in found_functions}} - global __nvrtcGetNumSupportedArchs - data["__nvrtcGetNumSupportedArchs"] = __nvrtcGetNumSupportedArchs - {{else}} - data["__nvrtcGetNumSupportedArchs"] = 0 - {{endif}} - - {{if 'nvrtcGetSupportedArchs' in found_functions}} - global __nvrtcGetSupportedArchs - data["__nvrtcGetSupportedArchs"] = __nvrtcGetSupportedArchs - {{else}} - data["__nvrtcGetSupportedArchs"] = 0 - {{endif}} - - {{if 'nvrtcCreateProgram' in found_functions}} - global __nvrtcCreateProgram - data["__nvrtcCreateProgram"] = __nvrtcCreateProgram - {{else}} - data["__nvrtcCreateProgram"] = 0 - {{endif}} - - {{if 'nvrtcDestroyProgram' in found_functions}} - global __nvrtcDestroyProgram - data["__nvrtcDestroyProgram"] = __nvrtcDestroyProgram - {{else}} - data["__nvrtcDestroyProgram"] = 0 - {{endif}} - - {{if 'nvrtcCompileProgram' in found_functions}} - global __nvrtcCompileProgram - data["__nvrtcCompileProgram"] = __nvrtcCompileProgram - {{else}} - data["__nvrtcCompileProgram"] = 0 - {{endif}} - - {{if 'nvrtcGetPTXSize' in found_functions}} - global __nvrtcGetPTXSize - data["__nvrtcGetPTXSize"] = __nvrtcGetPTXSize - {{else}} - data["__nvrtcGetPTXSize"] = 0 - {{endif}} - - {{if 'nvrtcGetPTX' in found_functions}} - global __nvrtcGetPTX - data["__nvrtcGetPTX"] = __nvrtcGetPTX - {{else}} - data["__nvrtcGetPTX"] = 0 - {{endif}} - - {{if 'nvrtcGetCUBINSize' in found_functions}} - global __nvrtcGetCUBINSize - data["__nvrtcGetCUBINSize"] = __nvrtcGetCUBINSize - {{else}} - data["__nvrtcGetCUBINSize"] = 0 - {{endif}} - - {{if 'nvrtcGetCUBIN' in found_functions}} - global __nvrtcGetCUBIN - data["__nvrtcGetCUBIN"] = __nvrtcGetCUBIN - {{else}} - data["__nvrtcGetCUBIN"] = 0 - {{endif}} - - {{if 'nvrtcGetLTOIRSize' in found_functions}} - global __nvrtcGetLTOIRSize - data["__nvrtcGetLTOIRSize"] = __nvrtcGetLTOIRSize - {{else}} - data["__nvrtcGetLTOIRSize"] = 0 - {{endif}} - - {{if 'nvrtcGetLTOIR' in found_functions}} - global __nvrtcGetLTOIR - data["__nvrtcGetLTOIR"] = __nvrtcGetLTOIR - {{else}} - data["__nvrtcGetLTOIR"] = 0 - {{endif}} - - {{if 'nvrtcGetOptiXIRSize' in found_functions}} - global __nvrtcGetOptiXIRSize - data["__nvrtcGetOptiXIRSize"] = __nvrtcGetOptiXIRSize - {{else}} - data["__nvrtcGetOptiXIRSize"] = 0 - {{endif}} - - {{if 'nvrtcGetOptiXIR' in found_functions}} - global __nvrtcGetOptiXIR - data["__nvrtcGetOptiXIR"] = __nvrtcGetOptiXIR - {{else}} - data["__nvrtcGetOptiXIR"] = 0 - {{endif}} - - {{if 'nvrtcGetProgramLogSize' in found_functions}} - global __nvrtcGetProgramLogSize - data["__nvrtcGetProgramLogSize"] = __nvrtcGetProgramLogSize - {{else}} - data["__nvrtcGetProgramLogSize"] = 0 - {{endif}} - - {{if 'nvrtcGetProgramLog' in found_functions}} - global __nvrtcGetProgramLog - data["__nvrtcGetProgramLog"] = __nvrtcGetProgramLog - {{else}} - data["__nvrtcGetProgramLog"] = 0 - {{endif}} - - {{if 'nvrtcAddNameExpression' in found_functions}} - global __nvrtcAddNameExpression - data["__nvrtcAddNameExpression"] = __nvrtcAddNameExpression - {{else}} - data["__nvrtcAddNameExpression"] = 0 - {{endif}} - - {{if 'nvrtcGetLoweredName' in found_functions}} - global __nvrtcGetLoweredName - data["__nvrtcGetLoweredName"] = __nvrtcGetLoweredName - {{else}} - data["__nvrtcGetLoweredName"] = 0 - {{endif}} - - {{if 'nvrtcGetPCHHeapSize' in found_functions}} - global __nvrtcGetPCHHeapSize - data["__nvrtcGetPCHHeapSize"] = __nvrtcGetPCHHeapSize - {{else}} - data["__nvrtcGetPCHHeapSize"] = 0 - {{endif}} - - {{if 'nvrtcSetPCHHeapSize' in found_functions}} - global __nvrtcSetPCHHeapSize - data["__nvrtcSetPCHHeapSize"] = __nvrtcSetPCHHeapSize - {{else}} - data["__nvrtcSetPCHHeapSize"] = 0 - {{endif}} - - {{if 'nvrtcGetPCHCreateStatus' in found_functions}} - global __nvrtcGetPCHCreateStatus - data["__nvrtcGetPCHCreateStatus"] = __nvrtcGetPCHCreateStatus - {{else}} - data["__nvrtcGetPCHCreateStatus"] = 0 - {{endif}} - - {{if 'nvrtcGetPCHHeapSizeRequired' in found_functions}} - global __nvrtcGetPCHHeapSizeRequired - data["__nvrtcGetPCHHeapSizeRequired"] = __nvrtcGetPCHHeapSizeRequired - {{else}} - data["__nvrtcGetPCHHeapSizeRequired"] = 0 - {{endif}} - - {{if 'nvrtcSetFlowCallback' in found_functions}} - global __nvrtcSetFlowCallback - data["__nvrtcSetFlowCallback"] = __nvrtcSetFlowCallback - {{else}} - data["__nvrtcSetFlowCallback"] = 0 - {{endif}} - - {{if 'nvrtcGetTileIRSize' in found_functions}} - global __nvrtcGetTileIRSize - data["__nvrtcGetTileIRSize"] = __nvrtcGetTileIRSize - {{else}} - data["__nvrtcGetTileIRSize"] = 0 - {{endif}} - - {{if 'nvrtcGetTileIR' in found_functions}} - global __nvrtcGetTileIR - data["__nvrtcGetTileIR"] = __nvrtcGetTileIR - {{else}} - data["__nvrtcGetTileIR"] = 0 - {{endif}} - - func_ptrs = data - return data - -cpdef _inspect_function_pointer(str name): - global func_ptrs - if func_ptrs is None: - func_ptrs = _inspect_function_pointers() - return func_ptrs[name] diff --git a/cuda_bindings/cuda/bindings/_bindings/cynvrtc.pxd.in b/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd similarity index 66% rename from cuda_bindings/cuda/bindings/_bindings/cynvrtc.pxd.in rename to cuda_bindings/cuda/bindings/_internal/nvrtc.pxd index 5a53b926d1..1964af1f16 100644 --- a/cuda_bindings/cuda/bindings/_bindings/cynvrtc.pxd.in +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd @@ -1,136 +1,63 @@ # SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE +# +# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1364+ged01d643e. Do not modify it directly. -# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1422+gf4812259e.d20260318. Do not modify it directly. -from cuda.bindings.cynvrtc cimport * +from ..cynvrtc cimport * -{{if 'nvrtcGetErrorString' in found_functions}} +############################################################################### +# Wrapper functions +############################################################################### cdef const char* _nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil -{{endif}} - -{{if 'nvrtcVersion' in found_functions}} cdef nvrtcResult _nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetNumSupportedArchs' in found_functions}} cdef nvrtcResult _nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetSupportedArchs' in found_functions}} cdef nvrtcResult _nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcCreateProgram' in found_functions}} cdef nvrtcResult _nvrtcCreateProgram(nvrtcProgram* prog, const char* src, const char* name, int numHeaders, const char** headers, const char** includeNames) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcDestroyProgram' in found_functions}} cdef nvrtcResult _nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcCompileProgram' in found_functions}} cdef nvrtcResult _nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const char** options) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetPTXSize' in found_functions}} cdef nvrtcResult _nvrtcGetPTXSize(nvrtcProgram prog, size_t* ptxSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetPTX' in found_functions}} cdef nvrtcResult _nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetCUBINSize' in found_functions}} cdef nvrtcResult _nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetCUBIN' in found_functions}} cdef nvrtcResult _nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetLTOIRSize' in found_functions}} cdef nvrtcResult _nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetLTOIR' in found_functions}} cdef nvrtcResult _nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetOptiXIRSize' in found_functions}} cdef nvrtcResult _nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetOptiXIR' in found_functions}} cdef nvrtcResult _nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetProgramLogSize' in found_functions}} cdef nvrtcResult _nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetProgramLog' in found_functions}} cdef nvrtcResult _nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcAddNameExpression' in found_functions}} cdef nvrtcResult _nvrtcAddNameExpression(nvrtcProgram prog, const char* name_expression) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetLoweredName' in found_functions}} cdef nvrtcResult _nvrtcGetLoweredName(nvrtcProgram prog, const char* name_expression, const char** lowered_name) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetPCHHeapSize' in found_functions}} cdef nvrtcResult _nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcSetPCHHeapSize' in found_functions}} cdef nvrtcResult _nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetPCHCreateStatus' in found_functions}} cdef nvrtcResult _nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetPCHHeapSizeRequired' in found_functions}} cdef nvrtcResult _nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcSetFlowCallback' in found_functions}} cdef nvrtcResult _nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* payload) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetTileIRSize' in found_functions}} cdef nvrtcResult _nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetTileIR' in found_functions}} cdef nvrtcResult _nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - diff --git a/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx b/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx new file mode 100644 index 0000000000..780042a8cd --- /dev/null +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx @@ -0,0 +1,599 @@ +# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE +# +# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1364+ged01d643e. Do not modify it directly. + +from libc.stdint cimport intptr_t, uintptr_t + +import threading +from .utils import FunctionNotFoundError, NotSupportedError + +from cuda.pathfinder import load_nvidia_dynamic_lib + +############################################################################### +# Extern +############################################################################### + +# You must 'from .utils import NotSupportedError' before using this template + +cdef extern from "" nogil: + void* dlopen(const char*, int) + char* dlerror() + void* dlsym(void*, const char*) + int dlclose(void*) + + enum: + RTLD_LAZY + RTLD_NOW + RTLD_GLOBAL + RTLD_LOCAL + + const void* RTLD_DEFAULT 'RTLD_DEFAULT' + +cdef int get_cuda_version(): + cdef void* handle = NULL + cdef int err, driver_ver = 0 + + # Load driver to check version + handle = dlopen('libcuda.so.1', RTLD_NOW | RTLD_GLOBAL) + if handle == NULL: + err_msg = dlerror() + raise NotSupportedError(f'CUDA driver is not found ({err_msg.decode()})') + cuDriverGetVersion = dlsym(handle, "cuDriverGetVersion") + if cuDriverGetVersion == NULL: + raise RuntimeError('Did not find cuDriverGetVersion symbol in libcuda.so.1') + err = (cuDriverGetVersion)(&driver_ver) + if err != 0: + raise RuntimeError(f'cuDriverGetVersion returned error code {err}') + + return driver_ver + +############################################################################### +# Wrapper init +############################################################################### + +cdef object __symbol_lock = threading.Lock() +cdef bint __py_nvrtc_init = False + +cdef void* __nvrtcGetErrorString = NULL +cdef void* __nvrtcVersion = NULL +cdef void* __nvrtcGetNumSupportedArchs = NULL +cdef void* __nvrtcGetSupportedArchs = NULL +cdef void* __nvrtcCreateProgram = NULL +cdef void* __nvrtcDestroyProgram = NULL +cdef void* __nvrtcCompileProgram = NULL +cdef void* __nvrtcGetPTXSize = NULL +cdef void* __nvrtcGetPTX = NULL +cdef void* __nvrtcGetCUBINSize = NULL +cdef void* __nvrtcGetCUBIN = NULL +cdef void* __nvrtcGetLTOIRSize = NULL +cdef void* __nvrtcGetLTOIR = NULL +cdef void* __nvrtcGetOptiXIRSize = NULL +cdef void* __nvrtcGetOptiXIR = NULL +cdef void* __nvrtcGetProgramLogSize = NULL +cdef void* __nvrtcGetProgramLog = NULL +cdef void* __nvrtcAddNameExpression = NULL +cdef void* __nvrtcGetLoweredName = NULL +cdef void* __nvrtcGetPCHHeapSize = NULL +cdef void* __nvrtcSetPCHHeapSize = NULL +cdef void* __nvrtcGetPCHCreateStatus = NULL +cdef void* __nvrtcGetPCHHeapSizeRequired = NULL +cdef void* __nvrtcSetFlowCallback = NULL +cdef void* __nvrtcGetTileIRSize = NULL +cdef void* __nvrtcGetTileIR = NULL + +cdef void* load_library() except* with gil: + cdef uintptr_t handle = load_nvidia_dynamic_lib("nvrtc")._handle_uint + return handle + +cdef int _init_nvrtc() except -1 nogil: + global __py_nvrtc_init + + cdef void* handle = NULL + + with gil, __symbol_lock: + # Recheck the flag after obtaining the locks + if __py_nvrtc_init: + return 0 + + # Load function + global __nvrtcGetErrorString + __nvrtcGetErrorString = dlsym(RTLD_DEFAULT, 'nvrtcGetErrorString') + if __nvrtcGetErrorString == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetErrorString = dlsym(handle, 'nvrtcGetErrorString') + + global __nvrtcVersion + __nvrtcVersion = dlsym(RTLD_DEFAULT, 'nvrtcVersion') + if __nvrtcVersion == NULL: + if handle == NULL: + handle = load_library() + __nvrtcVersion = dlsym(handle, 'nvrtcVersion') + + global __nvrtcGetNumSupportedArchs + __nvrtcGetNumSupportedArchs = dlsym(RTLD_DEFAULT, 'nvrtcGetNumSupportedArchs') + if __nvrtcGetNumSupportedArchs == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetNumSupportedArchs = dlsym(handle, 'nvrtcGetNumSupportedArchs') + + global __nvrtcGetSupportedArchs + __nvrtcGetSupportedArchs = dlsym(RTLD_DEFAULT, 'nvrtcGetSupportedArchs') + if __nvrtcGetSupportedArchs == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetSupportedArchs = dlsym(handle, 'nvrtcGetSupportedArchs') + + global __nvrtcCreateProgram + __nvrtcCreateProgram = dlsym(RTLD_DEFAULT, 'nvrtcCreateProgram') + if __nvrtcCreateProgram == NULL: + if handle == NULL: + handle = load_library() + __nvrtcCreateProgram = dlsym(handle, 'nvrtcCreateProgram') + + global __nvrtcDestroyProgram + __nvrtcDestroyProgram = dlsym(RTLD_DEFAULT, 'nvrtcDestroyProgram') + if __nvrtcDestroyProgram == NULL: + if handle == NULL: + handle = load_library() + __nvrtcDestroyProgram = dlsym(handle, 'nvrtcDestroyProgram') + + global __nvrtcCompileProgram + __nvrtcCompileProgram = dlsym(RTLD_DEFAULT, 'nvrtcCompileProgram') + if __nvrtcCompileProgram == NULL: + if handle == NULL: + handle = load_library() + __nvrtcCompileProgram = dlsym(handle, 'nvrtcCompileProgram') + + global __nvrtcGetPTXSize + __nvrtcGetPTXSize = dlsym(RTLD_DEFAULT, 'nvrtcGetPTXSize') + if __nvrtcGetPTXSize == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetPTXSize = dlsym(handle, 'nvrtcGetPTXSize') + + global __nvrtcGetPTX + __nvrtcGetPTX = dlsym(RTLD_DEFAULT, 'nvrtcGetPTX') + if __nvrtcGetPTX == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetPTX = dlsym(handle, 'nvrtcGetPTX') + + global __nvrtcGetCUBINSize + __nvrtcGetCUBINSize = dlsym(RTLD_DEFAULT, 'nvrtcGetCUBINSize') + if __nvrtcGetCUBINSize == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetCUBINSize = dlsym(handle, 'nvrtcGetCUBINSize') + + global __nvrtcGetCUBIN + __nvrtcGetCUBIN = dlsym(RTLD_DEFAULT, 'nvrtcGetCUBIN') + if __nvrtcGetCUBIN == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetCUBIN = dlsym(handle, 'nvrtcGetCUBIN') + + global __nvrtcGetLTOIRSize + __nvrtcGetLTOIRSize = dlsym(RTLD_DEFAULT, 'nvrtcGetLTOIRSize') + if __nvrtcGetLTOIRSize == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetLTOIRSize = dlsym(handle, 'nvrtcGetLTOIRSize') + + global __nvrtcGetLTOIR + __nvrtcGetLTOIR = dlsym(RTLD_DEFAULT, 'nvrtcGetLTOIR') + if __nvrtcGetLTOIR == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetLTOIR = dlsym(handle, 'nvrtcGetLTOIR') + + global __nvrtcGetOptiXIRSize + __nvrtcGetOptiXIRSize = dlsym(RTLD_DEFAULT, 'nvrtcGetOptiXIRSize') + if __nvrtcGetOptiXIRSize == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetOptiXIRSize = dlsym(handle, 'nvrtcGetOptiXIRSize') + + global __nvrtcGetOptiXIR + __nvrtcGetOptiXIR = dlsym(RTLD_DEFAULT, 'nvrtcGetOptiXIR') + if __nvrtcGetOptiXIR == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetOptiXIR = dlsym(handle, 'nvrtcGetOptiXIR') + + global __nvrtcGetProgramLogSize + __nvrtcGetProgramLogSize = dlsym(RTLD_DEFAULT, 'nvrtcGetProgramLogSize') + if __nvrtcGetProgramLogSize == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetProgramLogSize = dlsym(handle, 'nvrtcGetProgramLogSize') + + global __nvrtcGetProgramLog + __nvrtcGetProgramLog = dlsym(RTLD_DEFAULT, 'nvrtcGetProgramLog') + if __nvrtcGetProgramLog == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetProgramLog = dlsym(handle, 'nvrtcGetProgramLog') + + global __nvrtcAddNameExpression + __nvrtcAddNameExpression = dlsym(RTLD_DEFAULT, 'nvrtcAddNameExpression') + if __nvrtcAddNameExpression == NULL: + if handle == NULL: + handle = load_library() + __nvrtcAddNameExpression = dlsym(handle, 'nvrtcAddNameExpression') + + global __nvrtcGetLoweredName + __nvrtcGetLoweredName = dlsym(RTLD_DEFAULT, 'nvrtcGetLoweredName') + if __nvrtcGetLoweredName == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetLoweredName = dlsym(handle, 'nvrtcGetLoweredName') + + global __nvrtcGetPCHHeapSize + __nvrtcGetPCHHeapSize = dlsym(RTLD_DEFAULT, 'nvrtcGetPCHHeapSize') + if __nvrtcGetPCHHeapSize == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetPCHHeapSize = dlsym(handle, 'nvrtcGetPCHHeapSize') + + global __nvrtcSetPCHHeapSize + __nvrtcSetPCHHeapSize = dlsym(RTLD_DEFAULT, 'nvrtcSetPCHHeapSize') + if __nvrtcSetPCHHeapSize == NULL: + if handle == NULL: + handle = load_library() + __nvrtcSetPCHHeapSize = dlsym(handle, 'nvrtcSetPCHHeapSize') + + global __nvrtcGetPCHCreateStatus + __nvrtcGetPCHCreateStatus = dlsym(RTLD_DEFAULT, 'nvrtcGetPCHCreateStatus') + if __nvrtcGetPCHCreateStatus == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetPCHCreateStatus = dlsym(handle, 'nvrtcGetPCHCreateStatus') + + global __nvrtcGetPCHHeapSizeRequired + __nvrtcGetPCHHeapSizeRequired = dlsym(RTLD_DEFAULT, 'nvrtcGetPCHHeapSizeRequired') + if __nvrtcGetPCHHeapSizeRequired == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetPCHHeapSizeRequired = dlsym(handle, 'nvrtcGetPCHHeapSizeRequired') + + global __nvrtcSetFlowCallback + __nvrtcSetFlowCallback = dlsym(RTLD_DEFAULT, 'nvrtcSetFlowCallback') + if __nvrtcSetFlowCallback == NULL: + if handle == NULL: + handle = load_library() + __nvrtcSetFlowCallback = dlsym(handle, 'nvrtcSetFlowCallback') + + global __nvrtcGetTileIRSize + __nvrtcGetTileIRSize = dlsym(RTLD_DEFAULT, 'nvrtcGetTileIRSize') + if __nvrtcGetTileIRSize == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetTileIRSize = dlsym(handle, 'nvrtcGetTileIRSize') + + global __nvrtcGetTileIR + __nvrtcGetTileIR = dlsym(RTLD_DEFAULT, 'nvrtcGetTileIR') + if __nvrtcGetTileIR == NULL: + if handle == NULL: + handle = load_library() + __nvrtcGetTileIR = dlsym(handle, 'nvrtcGetTileIR') + + __py_nvrtc_init = True + return 0 + +cdef inline int _check_or_init_nvrtc() except -1 nogil: + if __py_nvrtc_init: + return 0 + + return _init_nvrtc() + +cdef dict func_ptrs = None + +cpdef dict _inspect_function_pointers(): + global func_ptrs + if func_ptrs is not None: + return func_ptrs + + _check_or_init_nvrtc() + cdef dict data = {} + + global __nvrtcGetErrorString + data["__nvrtcGetErrorString"] = __nvrtcGetErrorString + + global __nvrtcVersion + data["__nvrtcVersion"] = __nvrtcVersion + + global __nvrtcGetNumSupportedArchs + data["__nvrtcGetNumSupportedArchs"] = __nvrtcGetNumSupportedArchs + + global __nvrtcGetSupportedArchs + data["__nvrtcGetSupportedArchs"] = __nvrtcGetSupportedArchs + + global __nvrtcCreateProgram + data["__nvrtcCreateProgram"] = __nvrtcCreateProgram + + global __nvrtcDestroyProgram + data["__nvrtcDestroyProgram"] = __nvrtcDestroyProgram + + global __nvrtcCompileProgram + data["__nvrtcCompileProgram"] = __nvrtcCompileProgram + + global __nvrtcGetPTXSize + data["__nvrtcGetPTXSize"] = __nvrtcGetPTXSize + + global __nvrtcGetPTX + data["__nvrtcGetPTX"] = __nvrtcGetPTX + + global __nvrtcGetCUBINSize + data["__nvrtcGetCUBINSize"] = __nvrtcGetCUBINSize + + global __nvrtcGetCUBIN + data["__nvrtcGetCUBIN"] = __nvrtcGetCUBIN + + global __nvrtcGetLTOIRSize + data["__nvrtcGetLTOIRSize"] = __nvrtcGetLTOIRSize + + global __nvrtcGetLTOIR + data["__nvrtcGetLTOIR"] = __nvrtcGetLTOIR + + global __nvrtcGetOptiXIRSize + data["__nvrtcGetOptiXIRSize"] = __nvrtcGetOptiXIRSize + + global __nvrtcGetOptiXIR + data["__nvrtcGetOptiXIR"] = __nvrtcGetOptiXIR + + global __nvrtcGetProgramLogSize + data["__nvrtcGetProgramLogSize"] = __nvrtcGetProgramLogSize + + global __nvrtcGetProgramLog + data["__nvrtcGetProgramLog"] = __nvrtcGetProgramLog + + global __nvrtcAddNameExpression + data["__nvrtcAddNameExpression"] = __nvrtcAddNameExpression + + global __nvrtcGetLoweredName + data["__nvrtcGetLoweredName"] = __nvrtcGetLoweredName + + global __nvrtcGetPCHHeapSize + data["__nvrtcGetPCHHeapSize"] = __nvrtcGetPCHHeapSize + + global __nvrtcSetPCHHeapSize + data["__nvrtcSetPCHHeapSize"] = __nvrtcSetPCHHeapSize + + global __nvrtcGetPCHCreateStatus + data["__nvrtcGetPCHCreateStatus"] = __nvrtcGetPCHCreateStatus + + global __nvrtcGetPCHHeapSizeRequired + data["__nvrtcGetPCHHeapSizeRequired"] = __nvrtcGetPCHHeapSizeRequired + + global __nvrtcSetFlowCallback + data["__nvrtcSetFlowCallback"] = __nvrtcSetFlowCallback + + global __nvrtcGetTileIRSize + data["__nvrtcGetTileIRSize"] = __nvrtcGetTileIRSize + + global __nvrtcGetTileIR + data["__nvrtcGetTileIR"] = __nvrtcGetTileIR + + func_ptrs = data + return data + +cpdef _inspect_function_pointer(str name): + global func_ptrs + if func_ptrs is None: + func_ptrs = _inspect_function_pointers() + return func_ptrs[name] + +############################################################################### +# Wrapper functions +############################################################################### + +cdef const char* _nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil: + global __nvrtcGetErrorString + _check_or_init_nvrtc() + if __nvrtcGetErrorString == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetErrorString is not found") + return (__nvrtcGetErrorString)(result) + +cdef nvrtcResult _nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcVersion + _check_or_init_nvrtc() + if __nvrtcVersion == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcVersion is not found") + return (__nvrtcVersion)(major, minor) + +cdef nvrtcResult _nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetNumSupportedArchs + _check_or_init_nvrtc() + if __nvrtcGetNumSupportedArchs == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetNumSupportedArchs is not found") + return (__nvrtcGetNumSupportedArchs)(numArchs) + +cdef nvrtcResult _nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetSupportedArchs + _check_or_init_nvrtc() + if __nvrtcGetSupportedArchs == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetSupportedArchs is not found") + return (__nvrtcGetSupportedArchs)(supportedArchs) + +cdef nvrtcResult _nvrtcCreateProgram(nvrtcProgram* prog, const char* src, const char* name, int numHeaders, const char** headers, const char** includeNames) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcCreateProgram + _check_or_init_nvrtc() + if __nvrtcCreateProgram == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcCreateProgram is not found") + return (__nvrtcCreateProgram)(prog, src, name, numHeaders, headers, includeNames) + +cdef nvrtcResult _nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcDestroyProgram + _check_or_init_nvrtc() + if __nvrtcDestroyProgram == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcDestroyProgram is not found") + return (__nvrtcDestroyProgram)(prog) + +cdef nvrtcResult _nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const char** options) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcCompileProgram + _check_or_init_nvrtc() + if __nvrtcCompileProgram == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcCompileProgram is not found") + return (__nvrtcCompileProgram)(prog, numOptions, options) + +cdef nvrtcResult _nvrtcGetPTXSize(nvrtcProgram prog, size_t* ptxSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetPTXSize + _check_or_init_nvrtc() + if __nvrtcGetPTXSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetPTXSize is not found") + return (__nvrtcGetPTXSize)(prog, ptxSizeRet) + +cdef nvrtcResult _nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetPTX + _check_or_init_nvrtc() + if __nvrtcGetPTX == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetPTX is not found") + return (__nvrtcGetPTX)(prog, ptx) + +cdef nvrtcResult _nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetCUBINSize + _check_or_init_nvrtc() + if __nvrtcGetCUBINSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetCUBINSize is not found") + return (__nvrtcGetCUBINSize)(prog, cubinSizeRet) + +cdef nvrtcResult _nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetCUBIN + _check_or_init_nvrtc() + if __nvrtcGetCUBIN == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetCUBIN is not found") + return (__nvrtcGetCUBIN)(prog, cubin) + +cdef nvrtcResult _nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetLTOIRSize + _check_or_init_nvrtc() + if __nvrtcGetLTOIRSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetLTOIRSize is not found") + return (__nvrtcGetLTOIRSize)(prog, LTOIRSizeRet) + +cdef nvrtcResult _nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetLTOIR + _check_or_init_nvrtc() + if __nvrtcGetLTOIR == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetLTOIR is not found") + return (__nvrtcGetLTOIR)(prog, LTOIR) + +cdef nvrtcResult _nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetOptiXIRSize + _check_or_init_nvrtc() + if __nvrtcGetOptiXIRSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetOptiXIRSize is not found") + return (__nvrtcGetOptiXIRSize)(prog, optixirSizeRet) + +cdef nvrtcResult _nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetOptiXIR + _check_or_init_nvrtc() + if __nvrtcGetOptiXIR == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetOptiXIR is not found") + return (__nvrtcGetOptiXIR)(prog, optixir) + +cdef nvrtcResult _nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetProgramLogSize + _check_or_init_nvrtc() + if __nvrtcGetProgramLogSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetProgramLogSize is not found") + return (__nvrtcGetProgramLogSize)(prog, logSizeRet) + +cdef nvrtcResult _nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetProgramLog + _check_or_init_nvrtc() + if __nvrtcGetProgramLog == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetProgramLog is not found") + return (__nvrtcGetProgramLog)(prog, log) + +cdef nvrtcResult _nvrtcAddNameExpression(nvrtcProgram prog, const char* name_expression) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcAddNameExpression + _check_or_init_nvrtc() + if __nvrtcAddNameExpression == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcAddNameExpression is not found") + return (__nvrtcAddNameExpression)(prog, name_expression) + +cdef nvrtcResult _nvrtcGetLoweredName(nvrtcProgram prog, const char* name_expression, const char** lowered_name) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetLoweredName + _check_or_init_nvrtc() + if __nvrtcGetLoweredName == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetLoweredName is not found") + return (__nvrtcGetLoweredName)(prog, name_expression, lowered_name) + +cdef nvrtcResult _nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetPCHHeapSize + _check_or_init_nvrtc() + if __nvrtcGetPCHHeapSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetPCHHeapSize is not found") + return (__nvrtcGetPCHHeapSize)(ret) + +cdef nvrtcResult _nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcSetPCHHeapSize + _check_or_init_nvrtc() + if __nvrtcSetPCHHeapSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcSetPCHHeapSize is not found") + return (__nvrtcSetPCHHeapSize)(size) + +cdef nvrtcResult _nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetPCHCreateStatus + _check_or_init_nvrtc() + if __nvrtcGetPCHCreateStatus == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetPCHCreateStatus is not found") + return (__nvrtcGetPCHCreateStatus)(prog) + +cdef nvrtcResult _nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetPCHHeapSizeRequired + _check_or_init_nvrtc() + if __nvrtcGetPCHHeapSizeRequired == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetPCHHeapSizeRequired is not found") + return (__nvrtcGetPCHHeapSizeRequired)(prog, size) + +cdef nvrtcResult _nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* payload) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcSetFlowCallback + _check_or_init_nvrtc() + if __nvrtcSetFlowCallback == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcSetFlowCallback is not found") + return (__nvrtcSetFlowCallback)(prog, callback, payload) + +cdef nvrtcResult _nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetTileIRSize + _check_or_init_nvrtc() + if __nvrtcGetTileIRSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetTileIRSize is not found") + return (__nvrtcGetTileIRSize)(prog, TileIRSizeRet) + +cdef nvrtcResult _nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetTileIR + _check_or_init_nvrtc() + if __nvrtcGetTileIR == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetTileIR is not found") + return (__nvrtcGetTileIR)(prog, TileIR) diff --git a/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx b/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx new file mode 100644 index 0000000000..1fb555644e --- /dev/null +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx @@ -0,0 +1,511 @@ +# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE +# +# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1364+ged01d643e. Do not modify it directly. + +from libc.stdint cimport intptr_t + +import threading +from .utils import FunctionNotFoundError, NotSupportedError + +from cuda.pathfinder import load_nvidia_dynamic_lib + +from libc.stddef cimport wchar_t +from libc.stdint cimport uintptr_t +from cpython cimport PyUnicode_AsWideCharString, PyMem_Free + +# You must 'from .utils import NotSupportedError' before using this template + +cdef extern from "windows.h" nogil: + ctypedef void* HMODULE + ctypedef void* HANDLE + ctypedef void* FARPROC + ctypedef unsigned long DWORD + ctypedef const wchar_t *LPCWSTR + ctypedef const char *LPCSTR + + cdef DWORD LOAD_LIBRARY_SEARCH_SYSTEM32 = 0x00000800 + cdef DWORD LOAD_LIBRARY_SEARCH_DEFAULT_DIRS = 0x00001000 + cdef DWORD LOAD_LIBRARY_SEARCH_DLL_LOAD_DIR = 0x00000100 + + HMODULE _LoadLibraryExW "LoadLibraryExW"( + LPCWSTR lpLibFileName, + HANDLE hFile, + DWORD dwFlags + ) + + FARPROC _GetProcAddress "GetProcAddress"(HMODULE hModule, LPCSTR lpProcName) + +cdef inline uintptr_t LoadLibraryExW(str path, HANDLE hFile, DWORD dwFlags): + cdef uintptr_t result + cdef wchar_t* wpath = PyUnicode_AsWideCharString(path, NULL) + with nogil: + result = _LoadLibraryExW( + wpath, + hFile, + dwFlags + ) + PyMem_Free(wpath) + return result + +cdef inline void *GetProcAddress(uintptr_t hModule, const char* lpProcName) nogil: + return _GetProcAddress(hModule, lpProcName) + +cdef int get_cuda_version(): + cdef int err, driver_ver = 0 + + # Load driver to check version + handle = LoadLibraryExW("nvcuda.dll", NULL, LOAD_LIBRARY_SEARCH_SYSTEM32) + if handle == 0: + raise NotSupportedError('CUDA driver is not found') + cuDriverGetVersion = GetProcAddress(handle, 'cuDriverGetVersion') + if cuDriverGetVersion == NULL: + raise RuntimeError('Did not find cuDriverGetVersion symbol in nvcuda.dll') + err = (cuDriverGetVersion)(&driver_ver) + if err != 0: + raise RuntimeError(f'cuDriverGetVersion returned error code {err}') + + return driver_ver + +############################################################################### +# Wrapper init +############################################################################### + +cdef object __symbol_lock = threading.Lock() +cdef bint __py_nvrtc_init = False + +cdef void* __nvrtcGetErrorString = NULL +cdef void* __nvrtcVersion = NULL +cdef void* __nvrtcGetNumSupportedArchs = NULL +cdef void* __nvrtcGetSupportedArchs = NULL +cdef void* __nvrtcCreateProgram = NULL +cdef void* __nvrtcDestroyProgram = NULL +cdef void* __nvrtcCompileProgram = NULL +cdef void* __nvrtcGetPTXSize = NULL +cdef void* __nvrtcGetPTX = NULL +cdef void* __nvrtcGetCUBINSize = NULL +cdef void* __nvrtcGetCUBIN = NULL +cdef void* __nvrtcGetLTOIRSize = NULL +cdef void* __nvrtcGetLTOIR = NULL +cdef void* __nvrtcGetOptiXIRSize = NULL +cdef void* __nvrtcGetOptiXIR = NULL +cdef void* __nvrtcGetProgramLogSize = NULL +cdef void* __nvrtcGetProgramLog = NULL +cdef void* __nvrtcAddNameExpression = NULL +cdef void* __nvrtcGetLoweredName = NULL +cdef void* __nvrtcGetPCHHeapSize = NULL +cdef void* __nvrtcSetPCHHeapSize = NULL +cdef void* __nvrtcGetPCHCreateStatus = NULL +cdef void* __nvrtcGetPCHHeapSizeRequired = NULL +cdef void* __nvrtcSetFlowCallback = NULL +cdef void* __nvrtcGetTileIRSize = NULL +cdef void* __nvrtcGetTileIR = NULL + +cdef int _init_nvrtc() except -1 nogil: + global __py_nvrtc_init + + with gil, __symbol_lock: + # Recheck the flag after obtaining the locks + if __py_nvrtc_init: + return 0 + + # Load library + handle = load_nvidia_dynamic_lib("nvrtc")._handle_uint + + # Load function + global __nvrtcGetErrorString + __nvrtcGetErrorString = GetProcAddress(handle, 'nvrtcGetErrorString') + + global __nvrtcVersion + __nvrtcVersion = GetProcAddress(handle, 'nvrtcVersion') + + global __nvrtcGetNumSupportedArchs + __nvrtcGetNumSupportedArchs = GetProcAddress(handle, 'nvrtcGetNumSupportedArchs') + + global __nvrtcGetSupportedArchs + __nvrtcGetSupportedArchs = GetProcAddress(handle, 'nvrtcGetSupportedArchs') + + global __nvrtcCreateProgram + __nvrtcCreateProgram = GetProcAddress(handle, 'nvrtcCreateProgram') + + global __nvrtcDestroyProgram + __nvrtcDestroyProgram = GetProcAddress(handle, 'nvrtcDestroyProgram') + + global __nvrtcCompileProgram + __nvrtcCompileProgram = GetProcAddress(handle, 'nvrtcCompileProgram') + + global __nvrtcGetPTXSize + __nvrtcGetPTXSize = GetProcAddress(handle, 'nvrtcGetPTXSize') + + global __nvrtcGetPTX + __nvrtcGetPTX = GetProcAddress(handle, 'nvrtcGetPTX') + + global __nvrtcGetCUBINSize + __nvrtcGetCUBINSize = GetProcAddress(handle, 'nvrtcGetCUBINSize') + + global __nvrtcGetCUBIN + __nvrtcGetCUBIN = GetProcAddress(handle, 'nvrtcGetCUBIN') + + global __nvrtcGetLTOIRSize + __nvrtcGetLTOIRSize = GetProcAddress(handle, 'nvrtcGetLTOIRSize') + + global __nvrtcGetLTOIR + __nvrtcGetLTOIR = GetProcAddress(handle, 'nvrtcGetLTOIR') + + global __nvrtcGetOptiXIRSize + __nvrtcGetOptiXIRSize = GetProcAddress(handle, 'nvrtcGetOptiXIRSize') + + global __nvrtcGetOptiXIR + __nvrtcGetOptiXIR = GetProcAddress(handle, 'nvrtcGetOptiXIR') + + global __nvrtcGetProgramLogSize + __nvrtcGetProgramLogSize = GetProcAddress(handle, 'nvrtcGetProgramLogSize') + + global __nvrtcGetProgramLog + __nvrtcGetProgramLog = GetProcAddress(handle, 'nvrtcGetProgramLog') + + global __nvrtcAddNameExpression + __nvrtcAddNameExpression = GetProcAddress(handle, 'nvrtcAddNameExpression') + + global __nvrtcGetLoweredName + __nvrtcGetLoweredName = GetProcAddress(handle, 'nvrtcGetLoweredName') + + global __nvrtcGetPCHHeapSize + __nvrtcGetPCHHeapSize = GetProcAddress(handle, 'nvrtcGetPCHHeapSize') + + global __nvrtcSetPCHHeapSize + __nvrtcSetPCHHeapSize = GetProcAddress(handle, 'nvrtcSetPCHHeapSize') + + global __nvrtcGetPCHCreateStatus + __nvrtcGetPCHCreateStatus = GetProcAddress(handle, 'nvrtcGetPCHCreateStatus') + + global __nvrtcGetPCHHeapSizeRequired + __nvrtcGetPCHHeapSizeRequired = GetProcAddress(handle, 'nvrtcGetPCHHeapSizeRequired') + + global __nvrtcSetFlowCallback + __nvrtcSetFlowCallback = GetProcAddress(handle, 'nvrtcSetFlowCallback') + + global __nvrtcGetTileIRSize + __nvrtcGetTileIRSize = GetProcAddress(handle, 'nvrtcGetTileIRSize') + + global __nvrtcGetTileIR + __nvrtcGetTileIR = GetProcAddress(handle, 'nvrtcGetTileIR') + + __py_nvrtc_init = True + return 0 + +cdef inline int _check_or_init_nvrtc() except -1 nogil: + if __py_nvrtc_init: + return 0 + + return _init_nvrtc() + +cdef dict func_ptrs = None + +cpdef dict _inspect_function_pointers(): + global func_ptrs + if func_ptrs is not None: + return func_ptrs + + _check_or_init_nvrtc() + cdef dict data = {} + + global __nvrtcGetErrorString + data["__nvrtcGetErrorString"] = __nvrtcGetErrorString + + global __nvrtcVersion + data["__nvrtcVersion"] = __nvrtcVersion + + global __nvrtcGetNumSupportedArchs + data["__nvrtcGetNumSupportedArchs"] = __nvrtcGetNumSupportedArchs + + global __nvrtcGetSupportedArchs + data["__nvrtcGetSupportedArchs"] = __nvrtcGetSupportedArchs + + global __nvrtcCreateProgram + data["__nvrtcCreateProgram"] = __nvrtcCreateProgram + + global __nvrtcDestroyProgram + data["__nvrtcDestroyProgram"] = __nvrtcDestroyProgram + + global __nvrtcCompileProgram + data["__nvrtcCompileProgram"] = __nvrtcCompileProgram + + global __nvrtcGetPTXSize + data["__nvrtcGetPTXSize"] = __nvrtcGetPTXSize + + global __nvrtcGetPTX + data["__nvrtcGetPTX"] = __nvrtcGetPTX + + global __nvrtcGetCUBINSize + data["__nvrtcGetCUBINSize"] = __nvrtcGetCUBINSize + + global __nvrtcGetCUBIN + data["__nvrtcGetCUBIN"] = __nvrtcGetCUBIN + + global __nvrtcGetLTOIRSize + data["__nvrtcGetLTOIRSize"] = __nvrtcGetLTOIRSize + + global __nvrtcGetLTOIR + data["__nvrtcGetLTOIR"] = __nvrtcGetLTOIR + + global __nvrtcGetOptiXIRSize + data["__nvrtcGetOptiXIRSize"] = __nvrtcGetOptiXIRSize + + global __nvrtcGetOptiXIR + data["__nvrtcGetOptiXIR"] = __nvrtcGetOptiXIR + + global __nvrtcGetProgramLogSize + data["__nvrtcGetProgramLogSize"] = __nvrtcGetProgramLogSize + + global __nvrtcGetProgramLog + data["__nvrtcGetProgramLog"] = __nvrtcGetProgramLog + + global __nvrtcAddNameExpression + data["__nvrtcAddNameExpression"] = __nvrtcAddNameExpression + + global __nvrtcGetLoweredName + data["__nvrtcGetLoweredName"] = __nvrtcGetLoweredName + + global __nvrtcGetPCHHeapSize + data["__nvrtcGetPCHHeapSize"] = __nvrtcGetPCHHeapSize + + global __nvrtcSetPCHHeapSize + data["__nvrtcSetPCHHeapSize"] = __nvrtcSetPCHHeapSize + + global __nvrtcGetPCHCreateStatus + data["__nvrtcGetPCHCreateStatus"] = __nvrtcGetPCHCreateStatus + + global __nvrtcGetPCHHeapSizeRequired + data["__nvrtcGetPCHHeapSizeRequired"] = __nvrtcGetPCHHeapSizeRequired + + global __nvrtcSetFlowCallback + data["__nvrtcSetFlowCallback"] = __nvrtcSetFlowCallback + + global __nvrtcGetTileIRSize + data["__nvrtcGetTileIRSize"] = __nvrtcGetTileIRSize + + global __nvrtcGetTileIR + data["__nvrtcGetTileIR"] = __nvrtcGetTileIR + + func_ptrs = data + return data + +cpdef _inspect_function_pointer(str name): + global func_ptrs + if func_ptrs is None: + func_ptrs = _inspect_function_pointers() + return func_ptrs[name] + +############################################################################### +# Wrapper functions +############################################################################### + +cdef const char* _nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil: + global __nvrtcGetErrorString + _check_or_init_nvrtc() + if __nvrtcGetErrorString == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetErrorString is not found") + return (__nvrtcGetErrorString)(result) + +cdef nvrtcResult _nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcVersion + _check_or_init_nvrtc() + if __nvrtcVersion == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcVersion is not found") + return (__nvrtcVersion)(major, minor) + +cdef nvrtcResult _nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetNumSupportedArchs + _check_or_init_nvrtc() + if __nvrtcGetNumSupportedArchs == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetNumSupportedArchs is not found") + return (__nvrtcGetNumSupportedArchs)(numArchs) + +cdef nvrtcResult _nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetSupportedArchs + _check_or_init_nvrtc() + if __nvrtcGetSupportedArchs == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetSupportedArchs is not found") + return (__nvrtcGetSupportedArchs)(supportedArchs) + +cdef nvrtcResult _nvrtcCreateProgram(nvrtcProgram* prog, const char* src, const char* name, int numHeaders, const char** headers, const char** includeNames) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcCreateProgram + _check_or_init_nvrtc() + if __nvrtcCreateProgram == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcCreateProgram is not found") + return (__nvrtcCreateProgram)(prog, src, name, numHeaders, headers, includeNames) + +cdef nvrtcResult _nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcDestroyProgram + _check_or_init_nvrtc() + if __nvrtcDestroyProgram == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcDestroyProgram is not found") + return (__nvrtcDestroyProgram)(prog) + +cdef nvrtcResult _nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const char** options) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcCompileProgram + _check_or_init_nvrtc() + if __nvrtcCompileProgram == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcCompileProgram is not found") + return (__nvrtcCompileProgram)(prog, numOptions, options) + +cdef nvrtcResult _nvrtcGetPTXSize(nvrtcProgram prog, size_t* ptxSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetPTXSize + _check_or_init_nvrtc() + if __nvrtcGetPTXSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetPTXSize is not found") + return (__nvrtcGetPTXSize)(prog, ptxSizeRet) + +cdef nvrtcResult _nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetPTX + _check_or_init_nvrtc() + if __nvrtcGetPTX == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetPTX is not found") + return (__nvrtcGetPTX)(prog, ptx) + +cdef nvrtcResult _nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetCUBINSize + _check_or_init_nvrtc() + if __nvrtcGetCUBINSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetCUBINSize is not found") + return (__nvrtcGetCUBINSize)(prog, cubinSizeRet) + +cdef nvrtcResult _nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetCUBIN + _check_or_init_nvrtc() + if __nvrtcGetCUBIN == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetCUBIN is not found") + return (__nvrtcGetCUBIN)(prog, cubin) + +cdef nvrtcResult _nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetLTOIRSize + _check_or_init_nvrtc() + if __nvrtcGetLTOIRSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetLTOIRSize is not found") + return (__nvrtcGetLTOIRSize)(prog, LTOIRSizeRet) + +cdef nvrtcResult _nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetLTOIR + _check_or_init_nvrtc() + if __nvrtcGetLTOIR == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetLTOIR is not found") + return (__nvrtcGetLTOIR)(prog, LTOIR) + +cdef nvrtcResult _nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetOptiXIRSize + _check_or_init_nvrtc() + if __nvrtcGetOptiXIRSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetOptiXIRSize is not found") + return (__nvrtcGetOptiXIRSize)(prog, optixirSizeRet) + +cdef nvrtcResult _nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetOptiXIR + _check_or_init_nvrtc() + if __nvrtcGetOptiXIR == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetOptiXIR is not found") + return (__nvrtcGetOptiXIR)(prog, optixir) + +cdef nvrtcResult _nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetProgramLogSize + _check_or_init_nvrtc() + if __nvrtcGetProgramLogSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetProgramLogSize is not found") + return (__nvrtcGetProgramLogSize)(prog, logSizeRet) + +cdef nvrtcResult _nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetProgramLog + _check_or_init_nvrtc() + if __nvrtcGetProgramLog == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetProgramLog is not found") + return (__nvrtcGetProgramLog)(prog, log) + +cdef nvrtcResult _nvrtcAddNameExpression(nvrtcProgram prog, const char* name_expression) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcAddNameExpression + _check_or_init_nvrtc() + if __nvrtcAddNameExpression == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcAddNameExpression is not found") + return (__nvrtcAddNameExpression)(prog, name_expression) + +cdef nvrtcResult _nvrtcGetLoweredName(nvrtcProgram prog, const char* name_expression, const char** lowered_name) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetLoweredName + _check_or_init_nvrtc() + if __nvrtcGetLoweredName == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetLoweredName is not found") + return (__nvrtcGetLoweredName)(prog, name_expression, lowered_name) + +cdef nvrtcResult _nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetPCHHeapSize + _check_or_init_nvrtc() + if __nvrtcGetPCHHeapSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetPCHHeapSize is not found") + return (__nvrtcGetPCHHeapSize)(ret) + +cdef nvrtcResult _nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcSetPCHHeapSize + _check_or_init_nvrtc() + if __nvrtcSetPCHHeapSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcSetPCHHeapSize is not found") + return (__nvrtcSetPCHHeapSize)(size) + +cdef nvrtcResult _nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetPCHCreateStatus + _check_or_init_nvrtc() + if __nvrtcGetPCHCreateStatus == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetPCHCreateStatus is not found") + return (__nvrtcGetPCHCreateStatus)(prog) + +cdef nvrtcResult _nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetPCHHeapSizeRequired + _check_or_init_nvrtc() + if __nvrtcGetPCHHeapSizeRequired == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetPCHHeapSizeRequired is not found") + return (__nvrtcGetPCHHeapSizeRequired)(prog, size) + +cdef nvrtcResult _nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* payload) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcSetFlowCallback + _check_or_init_nvrtc() + if __nvrtcSetFlowCallback == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcSetFlowCallback is not found") + return (__nvrtcSetFlowCallback)(prog, callback, payload) + +cdef nvrtcResult _nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetTileIRSize + _check_or_init_nvrtc() + if __nvrtcGetTileIRSize == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetTileIRSize is not found") + return (__nvrtcGetTileIRSize)(prog, TileIRSizeRet) + +cdef nvrtcResult _nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: + global __nvrtcGetTileIR + _check_or_init_nvrtc() + if __nvrtcGetTileIR == NULL: + with gil: + raise FunctionNotFoundError("function nvrtcGetTileIR is not found") + return (__nvrtcGetTileIR)(prog, TileIR) diff --git a/cuda_bindings/cuda/bindings/cynvrtc.pxd.in b/cuda_bindings/cuda/bindings/cynvrtc.pxd similarity index 72% rename from cuda_bindings/cuda/bindings/cynvrtc.pxd.in rename to cuda_bindings/cuda/bindings/cynvrtc.pxd index 209b361c05..6bf1bda94e 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pxd.in +++ b/cuda_bindings/cuda/bindings/cynvrtc.pxd @@ -1,7 +1,8 @@ # SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE - -# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1422+gf4812259e.d20260318. Do not modify it directly. +# +# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1364+ged01d643e. Do not modify it directly. from libc.stdint cimport uint32_t, uint64_t @@ -31,133 +32,54 @@ cdef extern from "nvrtc.h": pass ctypedef _nvrtcProgram* nvrtcProgram -{{if 'nvrtcGetErrorString' in found_functions}} - cdef const char* nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil -{{endif}} - -{{if 'nvrtcVersion' in found_functions}} cdef nvrtcResult nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetNumSupportedArchs' in found_functions}} cdef nvrtcResult nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetSupportedArchs' in found_functions}} cdef nvrtcResult nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcCreateProgram' in found_functions}} cdef nvrtcResult nvrtcCreateProgram(nvrtcProgram* prog, const char* src, const char* name, int numHeaders, const char** headers, const char** includeNames) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcDestroyProgram' in found_functions}} cdef nvrtcResult nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcCompileProgram' in found_functions}} cdef nvrtcResult nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const char** options) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetPTXSize' in found_functions}} cdef nvrtcResult nvrtcGetPTXSize(nvrtcProgram prog, size_t* ptxSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetPTX' in found_functions}} cdef nvrtcResult nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetCUBINSize' in found_functions}} cdef nvrtcResult nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetCUBIN' in found_functions}} cdef nvrtcResult nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetLTOIRSize' in found_functions}} cdef nvrtcResult nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetLTOIR' in found_functions}} cdef nvrtcResult nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetOptiXIRSize' in found_functions}} cdef nvrtcResult nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetOptiXIR' in found_functions}} cdef nvrtcResult nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetProgramLogSize' in found_functions}} cdef nvrtcResult nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetProgramLog' in found_functions}} cdef nvrtcResult nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcAddNameExpression' in found_functions}} cdef nvrtcResult nvrtcAddNameExpression(nvrtcProgram prog, const char* name_expression) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetLoweredName' in found_functions}} cdef nvrtcResult nvrtcGetLoweredName(nvrtcProgram prog, const char* name_expression, const char** lowered_name) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetPCHHeapSize' in found_functions}} cdef nvrtcResult nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcSetPCHHeapSize' in found_functions}} cdef nvrtcResult nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetPCHCreateStatus' in found_functions}} cdef nvrtcResult nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetPCHHeapSizeRequired' in found_functions}} cdef nvrtcResult nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcSetFlowCallback' in found_functions}} cdef nvrtcResult nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* payload) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetTileIRSize' in found_functions}} cdef nvrtcResult nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - -{{if 'nvrtcGetTileIR' in found_functions}} cdef nvrtcResult nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil -{{endif}} - diff --git a/cuda_bindings/cuda/bindings/cynvrtc.pyx.in b/cuda_bindings/cuda/bindings/cynvrtc.pyx similarity index 50% rename from cuda_bindings/cuda/bindings/cynvrtc.pyx.in rename to cuda_bindings/cuda/bindings/cynvrtc.pyx index 9e64cbf9c1..2bf71611d7 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pyx.in +++ b/cuda_bindings/cuda/bindings/cynvrtc.pyx @@ -1,161 +1,85 @@ # SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE +# +# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1364+ged01d643e. Do not modify it directly. -# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1422+gf4812259e.d20260318. Do not modify it directly. -cimport cuda.bindings._bindings.cynvrtc as cynvrtc - -{{if 'nvrtcGetErrorString' in found_functions}} +from ._internal cimport nvrtc as _nvrtc cdef const char* nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil: - return cynvrtc._nvrtcGetErrorString(result) -{{endif}} - -{{if 'nvrtcVersion' in found_functions}} + return _nvrtc._nvrtcGetErrorString(result) cdef nvrtcResult nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcVersion(major, minor) -{{endif}} - -{{if 'nvrtcGetNumSupportedArchs' in found_functions}} + return _nvrtc._nvrtcVersion(major, minor) cdef nvrtcResult nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetNumSupportedArchs(numArchs) -{{endif}} - -{{if 'nvrtcGetSupportedArchs' in found_functions}} + return _nvrtc._nvrtcGetNumSupportedArchs(numArchs) cdef nvrtcResult nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetSupportedArchs(supportedArchs) -{{endif}} - -{{if 'nvrtcCreateProgram' in found_functions}} + return _nvrtc._nvrtcGetSupportedArchs(supportedArchs) cdef nvrtcResult nvrtcCreateProgram(nvrtcProgram* prog, const char* src, const char* name, int numHeaders, const char** headers, const char** includeNames) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcCreateProgram(prog, src, name, numHeaders, headers, includeNames) -{{endif}} - -{{if 'nvrtcDestroyProgram' in found_functions}} + return _nvrtc._nvrtcCreateProgram(prog, src, name, numHeaders, headers, includeNames) cdef nvrtcResult nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcDestroyProgram(prog) -{{endif}} - -{{if 'nvrtcCompileProgram' in found_functions}} + return _nvrtc._nvrtcDestroyProgram(prog) cdef nvrtcResult nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const char** options) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcCompileProgram(prog, numOptions, options) -{{endif}} - -{{if 'nvrtcGetPTXSize' in found_functions}} + return _nvrtc._nvrtcCompileProgram(prog, numOptions, options) cdef nvrtcResult nvrtcGetPTXSize(nvrtcProgram prog, size_t* ptxSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetPTXSize(prog, ptxSizeRet) -{{endif}} - -{{if 'nvrtcGetPTX' in found_functions}} + return _nvrtc._nvrtcGetPTXSize(prog, ptxSizeRet) cdef nvrtcResult nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetPTX(prog, ptx) -{{endif}} - -{{if 'nvrtcGetCUBINSize' in found_functions}} + return _nvrtc._nvrtcGetPTX(prog, ptx) cdef nvrtcResult nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetCUBINSize(prog, cubinSizeRet) -{{endif}} - -{{if 'nvrtcGetCUBIN' in found_functions}} + return _nvrtc._nvrtcGetCUBINSize(prog, cubinSizeRet) cdef nvrtcResult nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetCUBIN(prog, cubin) -{{endif}} - -{{if 'nvrtcGetLTOIRSize' in found_functions}} + return _nvrtc._nvrtcGetCUBIN(prog, cubin) cdef nvrtcResult nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetLTOIRSize(prog, LTOIRSizeRet) -{{endif}} - -{{if 'nvrtcGetLTOIR' in found_functions}} + return _nvrtc._nvrtcGetLTOIRSize(prog, LTOIRSizeRet) cdef nvrtcResult nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetLTOIR(prog, LTOIR) -{{endif}} - -{{if 'nvrtcGetOptiXIRSize' in found_functions}} + return _nvrtc._nvrtcGetLTOIR(prog, LTOIR) cdef nvrtcResult nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetOptiXIRSize(prog, optixirSizeRet) -{{endif}} - -{{if 'nvrtcGetOptiXIR' in found_functions}} + return _nvrtc._nvrtcGetOptiXIRSize(prog, optixirSizeRet) cdef nvrtcResult nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetOptiXIR(prog, optixir) -{{endif}} - -{{if 'nvrtcGetProgramLogSize' in found_functions}} + return _nvrtc._nvrtcGetOptiXIR(prog, optixir) cdef nvrtcResult nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetProgramLogSize(prog, logSizeRet) -{{endif}} - -{{if 'nvrtcGetProgramLog' in found_functions}} + return _nvrtc._nvrtcGetProgramLogSize(prog, logSizeRet) cdef nvrtcResult nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetProgramLog(prog, log) -{{endif}} - -{{if 'nvrtcAddNameExpression' in found_functions}} + return _nvrtc._nvrtcGetProgramLog(prog, log) cdef nvrtcResult nvrtcAddNameExpression(nvrtcProgram prog, const char* name_expression) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcAddNameExpression(prog, name_expression) -{{endif}} - -{{if 'nvrtcGetLoweredName' in found_functions}} + return _nvrtc._nvrtcAddNameExpression(prog, name_expression) cdef nvrtcResult nvrtcGetLoweredName(nvrtcProgram prog, const char* name_expression, const char** lowered_name) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetLoweredName(prog, name_expression, lowered_name) -{{endif}} - -{{if 'nvrtcGetPCHHeapSize' in found_functions}} + return _nvrtc._nvrtcGetLoweredName(prog, name_expression, lowered_name) cdef nvrtcResult nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetPCHHeapSize(ret) -{{endif}} - -{{if 'nvrtcSetPCHHeapSize' in found_functions}} + return _nvrtc._nvrtcGetPCHHeapSize(ret) cdef nvrtcResult nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcSetPCHHeapSize(size) -{{endif}} - -{{if 'nvrtcGetPCHCreateStatus' in found_functions}} + return _nvrtc._nvrtcSetPCHHeapSize(size) cdef nvrtcResult nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetPCHCreateStatus(prog) -{{endif}} - -{{if 'nvrtcGetPCHHeapSizeRequired' in found_functions}} + return _nvrtc._nvrtcGetPCHCreateStatus(prog) cdef nvrtcResult nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetPCHHeapSizeRequired(prog, size) -{{endif}} - -{{if 'nvrtcSetFlowCallback' in found_functions}} + return _nvrtc._nvrtcGetPCHHeapSizeRequired(prog, size) cdef nvrtcResult nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* payload) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcSetFlowCallback(prog, callback, payload) -{{endif}} - -{{if 'nvrtcGetTileIRSize' in found_functions}} + return _nvrtc._nvrtcSetFlowCallback(prog, callback, payload) cdef nvrtcResult nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetTileIRSize(prog, TileIRSizeRet) -{{endif}} - -{{if 'nvrtcGetTileIR' in found_functions}} + return _nvrtc._nvrtcGetTileIRSize(prog, TileIRSizeRet) cdef nvrtcResult nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetTileIR(prog, TileIR) -{{endif}} + return _nvrtc._nvrtcGetTileIR(prog, TileIR) diff --git a/cuda_bindings/cuda/bindings/nvrtc.pxd.in b/cuda_bindings/cuda/bindings/nvrtc.pxd similarity index 85% rename from cuda_bindings/cuda/bindings/nvrtc.pxd.in rename to cuda_bindings/cuda/bindings/nvrtc.pxd index 394bcccff0..5cb372430f 100644 --- a/cuda_bindings/cuda/bindings/nvrtc.pxd.in +++ b/cuda_bindings/cuda/bindings/nvrtc.pxd @@ -1,13 +1,11 @@ # SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE -# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1422+gf4812259e.d20260318. Do not modify it directly. +# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1364+ged01d643e. Do not modify it directly. cimport cuda.bindings.cynvrtc as cynvrtc include "_lib/utils.pxd" -{{if 'nvrtcProgram' in found_types}} - cdef class nvrtcProgram: """ nvrtcProgram is the unit of compilation, and an opaque handle for a program. @@ -21,4 +19,3 @@ cdef class nvrtcProgram: """ cdef cynvrtc.nvrtcProgram _pvt_val cdef cynvrtc.nvrtcProgram* _pvt_ptr -{{endif}} diff --git a/cuda_bindings/cuda/bindings/nvrtc.pyx.in b/cuda_bindings/cuda/bindings/nvrtc.pyx similarity index 90% rename from cuda_bindings/cuda/bindings/nvrtc.pyx.in rename to cuda_bindings/cuda/bindings/nvrtc.pyx index 2f50afa726..aca1ead365 100644 --- a/cuda_bindings/cuda/bindings/nvrtc.pyx.in +++ b/cuda_bindings/cuda/bindings/nvrtc.pyx @@ -1,7 +1,7 @@ # SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE -# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1422+gf4812259e.d20260318. Do not modify it directly. +# This code was automatically generated with version 13.2.0, generator version 0.3.1.dev1364+ged01d643e. Do not modify it directly. from typing import Any, Optional import cython import ctypes @@ -43,59 +43,50 @@ ctypedef unsigned long long float_ptr ctypedef unsigned long long double_ptr ctypedef unsigned long long void_ptr - -{{if 'nvrtcResult' in found_types}} - class nvrtcResult(_FastEnum): """ The enumerated type nvrtcResult defines API call result codes. NVRTC API functions return nvrtcResult to indicate the call result. """ - {{if 'NVRTC_SUCCESS' in found_values}} - NVRTC_SUCCESS = cynvrtc.nvrtcResult.NVRTC_SUCCESS{{endif}} - {{if 'NVRTC_ERROR_OUT_OF_MEMORY' in found_values}} - NVRTC_ERROR_OUT_OF_MEMORY = cynvrtc.nvrtcResult.NVRTC_ERROR_OUT_OF_MEMORY{{endif}} - {{if 'NVRTC_ERROR_PROGRAM_CREATION_FAILURE' in found_values}} - NVRTC_ERROR_PROGRAM_CREATION_FAILURE = cynvrtc.nvrtcResult.NVRTC_ERROR_PROGRAM_CREATION_FAILURE{{endif}} - {{if 'NVRTC_ERROR_INVALID_INPUT' in found_values}} - NVRTC_ERROR_INVALID_INPUT = cynvrtc.nvrtcResult.NVRTC_ERROR_INVALID_INPUT{{endif}} - {{if 'NVRTC_ERROR_INVALID_PROGRAM' in found_values}} - NVRTC_ERROR_INVALID_PROGRAM = cynvrtc.nvrtcResult.NVRTC_ERROR_INVALID_PROGRAM{{endif}} - {{if 'NVRTC_ERROR_INVALID_OPTION' in found_values}} - NVRTC_ERROR_INVALID_OPTION = cynvrtc.nvrtcResult.NVRTC_ERROR_INVALID_OPTION{{endif}} - {{if 'NVRTC_ERROR_COMPILATION' in found_values}} - NVRTC_ERROR_COMPILATION = cynvrtc.nvrtcResult.NVRTC_ERROR_COMPILATION{{endif}} - {{if 'NVRTC_ERROR_BUILTIN_OPERATION_FAILURE' in found_values}} - NVRTC_ERROR_BUILTIN_OPERATION_FAILURE = cynvrtc.nvrtcResult.NVRTC_ERROR_BUILTIN_OPERATION_FAILURE{{endif}} - {{if 'NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION' in found_values}} - NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = cynvrtc.nvrtcResult.NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION{{endif}} - {{if 'NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION' in found_values}} - NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = cynvrtc.nvrtcResult.NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION{{endif}} - {{if 'NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID' in found_values}} - NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID = cynvrtc.nvrtcResult.NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID{{endif}} - {{if 'NVRTC_ERROR_INTERNAL_ERROR' in found_values}} - NVRTC_ERROR_INTERNAL_ERROR = cynvrtc.nvrtcResult.NVRTC_ERROR_INTERNAL_ERROR{{endif}} - {{if 'NVRTC_ERROR_TIME_FILE_WRITE_FAILED' in found_values}} - NVRTC_ERROR_TIME_FILE_WRITE_FAILED = cynvrtc.nvrtcResult.NVRTC_ERROR_TIME_FILE_WRITE_FAILED{{endif}} - {{if 'NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED' in found_values}} - NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED = cynvrtc.nvrtcResult.NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED{{endif}} - {{if 'NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED' in found_values}} - NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED = cynvrtc.nvrtcResult.NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED{{endif}} - {{if 'NVRTC_ERROR_PCH_CREATE' in found_values}} - NVRTC_ERROR_PCH_CREATE = cynvrtc.nvrtcResult.NVRTC_ERROR_PCH_CREATE{{endif}} - {{if 'NVRTC_ERROR_CANCELLED' in found_values}} - NVRTC_ERROR_CANCELLED = cynvrtc.nvrtcResult.NVRTC_ERROR_CANCELLED{{endif}} - {{if 'NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED' in found_values}} - NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED = cynvrtc.nvrtcResult.NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED{{endif}} - -{{endif}} -cdef object _nvrtcResult = nvrtcResult -cdef object _nvrtcResult_SUCCESS = nvrtcResult.NVRTC_SUCCESS + NVRTC_SUCCESS = cynvrtc.nvrtcResult.NVRTC_SUCCESS + + NVRTC_ERROR_OUT_OF_MEMORY = cynvrtc.nvrtcResult.NVRTC_ERROR_OUT_OF_MEMORY + + NVRTC_ERROR_PROGRAM_CREATION_FAILURE = cynvrtc.nvrtcResult.NVRTC_ERROR_PROGRAM_CREATION_FAILURE + + NVRTC_ERROR_INVALID_INPUT = cynvrtc.nvrtcResult.NVRTC_ERROR_INVALID_INPUT + + NVRTC_ERROR_INVALID_PROGRAM = cynvrtc.nvrtcResult.NVRTC_ERROR_INVALID_PROGRAM + + NVRTC_ERROR_INVALID_OPTION = cynvrtc.nvrtcResult.NVRTC_ERROR_INVALID_OPTION + + NVRTC_ERROR_COMPILATION = cynvrtc.nvrtcResult.NVRTC_ERROR_COMPILATION + + NVRTC_ERROR_BUILTIN_OPERATION_FAILURE = cynvrtc.nvrtcResult.NVRTC_ERROR_BUILTIN_OPERATION_FAILURE + + NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = cynvrtc.nvrtcResult.NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION + + NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = cynvrtc.nvrtcResult.NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION + NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID = cynvrtc.nvrtcResult.NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID + NVRTC_ERROR_INTERNAL_ERROR = cynvrtc.nvrtcResult.NVRTC_ERROR_INTERNAL_ERROR -{{if 'nvrtcProgram' in found_types}} + NVRTC_ERROR_TIME_FILE_WRITE_FAILED = cynvrtc.nvrtcResult.NVRTC_ERROR_TIME_FILE_WRITE_FAILED + + NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED = cynvrtc.nvrtcResult.NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED + + NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED = cynvrtc.nvrtcResult.NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED + + NVRTC_ERROR_PCH_CREATE = cynvrtc.nvrtcResult.NVRTC_ERROR_PCH_CREATE + + NVRTC_ERROR_CANCELLED = cynvrtc.nvrtcResult.NVRTC_ERROR_CANCELLED + + NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED = cynvrtc.nvrtcResult.NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED + +cdef object _nvrtcResult = nvrtcResult +cdef object _nvrtcResult_SUCCESS = nvrtcResult.NVRTC_SUCCESS cdef class nvrtcProgram: """ nvrtcProgram is the unit of compilation, and an opaque handle for a program. @@ -130,9 +121,6 @@ cdef class nvrtcProgram: return self._pvt_ptr[0] def getPtr(self): return self._pvt_ptr -{{endif}} - -{{if 'nvrtcGetErrorString' in found_functions}} @cython.embedsignature(True) def nvrtcGetErrorString(result not None : nvrtcResult): @@ -154,9 +142,6 @@ def nvrtcGetErrorString(result not None : nvrtcResult): with nogil: err = cynvrtc.nvrtcGetErrorString(cyresult) return (nvrtcResult.NVRTC_SUCCESS, err) -{{endif}} - -{{if 'nvrtcVersion' in found_functions}} @cython.embedsignature(True) def nvrtcVersion(): @@ -179,9 +164,6 @@ def nvrtcVersion(): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None, None) return (_nvrtcResult_SUCCESS, major, minor) -{{endif}} - -{{if 'nvrtcGetNumSupportedArchs' in found_functions}} @cython.embedsignature(True) def nvrtcGetNumSupportedArchs(): @@ -203,9 +185,6 @@ def nvrtcGetNumSupportedArchs(): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, numArchs) -{{endif}} - -{{if 'nvrtcGetSupportedArchs' in found_functions}} @cython.embedsignature(True) def nvrtcGetSupportedArchs(): @@ -230,9 +209,6 @@ def nvrtcGetSupportedArchs(): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, supportedArchs) -{{endif}} - -{{if 'nvrtcCreateProgram' in found_functions}} @cython.embedsignature(True) def nvrtcCreateProgram(char* src, char* name, int numHeaders, headers : Optional[tuple[bytes] | list[bytes]], includeNames : Optional[tuple[bytes] | list[bytes]]): @@ -288,9 +264,6 @@ def nvrtcCreateProgram(char* src, char* name, int numHeaders, headers : Optional if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, prog) -{{endif}} - -{{if 'nvrtcDestroyProgram' in found_functions}} @cython.embedsignature(True) def nvrtcDestroyProgram(prog): @@ -324,9 +297,6 @@ def nvrtcDestroyProgram(prog): with nogil: err = cynvrtc.nvrtcDestroyProgram(cyprog) return (_nvrtcResult(err),) -{{endif}} - -{{if 'nvrtcCompileProgram' in found_functions}} @cython.embedsignature(True) def nvrtcCompileProgram(prog, int numOptions, options : Optional[tuple[bytes] | list[bytes]]): @@ -374,9 +344,6 @@ def nvrtcCompileProgram(prog, int numOptions, options : Optional[tuple[bytes] | with nogil: err = cynvrtc.nvrtcCompileProgram(cyprog, numOptions, cyoptions.data()) return (_nvrtcResult(err),) -{{endif}} - -{{if 'nvrtcGetPTXSize' in found_functions}} @cython.embedsignature(True) def nvrtcGetPTXSize(prog): @@ -414,9 +381,6 @@ def nvrtcGetPTXSize(prog): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, ptxSizeRet) -{{endif}} - -{{if 'nvrtcGetPTX' in found_functions}} @cython.embedsignature(True) def nvrtcGetPTX(prog, char* ptx): @@ -451,9 +415,6 @@ def nvrtcGetPTX(prog, char* ptx): with nogil: err = cynvrtc.nvrtcGetPTX(cyprog, ptx) return (_nvrtcResult(err),) -{{endif}} - -{{if 'nvrtcGetCUBINSize' in found_functions}} @cython.embedsignature(True) def nvrtcGetCUBINSize(prog): @@ -491,9 +452,6 @@ def nvrtcGetCUBINSize(prog): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, cubinSizeRet) -{{endif}} - -{{if 'nvrtcGetCUBIN' in found_functions}} @cython.embedsignature(True) def nvrtcGetCUBIN(prog, char* cubin): @@ -528,9 +486,6 @@ def nvrtcGetCUBIN(prog, char* cubin): with nogil: err = cynvrtc.nvrtcGetCUBIN(cyprog, cubin) return (_nvrtcResult(err),) -{{endif}} - -{{if 'nvrtcGetLTOIRSize' in found_functions}} @cython.embedsignature(True) def nvrtcGetLTOIRSize(prog): @@ -568,9 +523,6 @@ def nvrtcGetLTOIRSize(prog): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, LTOIRSizeRet) -{{endif}} - -{{if 'nvrtcGetLTOIR' in found_functions}} @cython.embedsignature(True) def nvrtcGetLTOIR(prog, char* LTOIR): @@ -605,9 +557,6 @@ def nvrtcGetLTOIR(prog, char* LTOIR): with nogil: err = cynvrtc.nvrtcGetLTOIR(cyprog, LTOIR) return (_nvrtcResult(err),) -{{endif}} - -{{if 'nvrtcGetOptiXIRSize' in found_functions}} @cython.embedsignature(True) def nvrtcGetOptiXIRSize(prog): @@ -645,9 +594,6 @@ def nvrtcGetOptiXIRSize(prog): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, optixirSizeRet) -{{endif}} - -{{if 'nvrtcGetOptiXIR' in found_functions}} @cython.embedsignature(True) def nvrtcGetOptiXIR(prog, char* optixir): @@ -682,9 +628,6 @@ def nvrtcGetOptiXIR(prog, char* optixir): with nogil: err = cynvrtc.nvrtcGetOptiXIR(cyprog, optixir) return (_nvrtcResult(err),) -{{endif}} - -{{if 'nvrtcGetProgramLogSize' in found_functions}} @cython.embedsignature(True) def nvrtcGetProgramLogSize(prog): @@ -725,9 +668,6 @@ def nvrtcGetProgramLogSize(prog): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, logSizeRet) -{{endif}} - -{{if 'nvrtcGetProgramLog' in found_functions}} @cython.embedsignature(True) def nvrtcGetProgramLog(prog, char* log): @@ -762,9 +702,6 @@ def nvrtcGetProgramLog(prog, char* log): with nogil: err = cynvrtc.nvrtcGetProgramLog(cyprog, log) return (_nvrtcResult(err),) -{{endif}} - -{{if 'nvrtcAddNameExpression' in found_functions}} @cython.embedsignature(True) def nvrtcAddNameExpression(prog, char* name_expression): @@ -804,9 +741,6 @@ def nvrtcAddNameExpression(prog, char* name_expression): with nogil: err = cynvrtc.nvrtcAddNameExpression(cyprog, name_expression) return (_nvrtcResult(err),) -{{endif}} - -{{if 'nvrtcGetLoweredName' in found_functions}} @cython.embedsignature(True) def nvrtcGetLoweredName(prog, char* name_expression): @@ -849,9 +783,6 @@ def nvrtcGetLoweredName(prog, char* name_expression): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, lowered_name if lowered_name != NULL else None) -{{endif}} - -{{if 'nvrtcGetPCHHeapSize' in found_functions}} @cython.embedsignature(True) def nvrtcGetPCHHeapSize(): @@ -871,9 +802,6 @@ def nvrtcGetPCHHeapSize(): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, ret) -{{endif}} - -{{if 'nvrtcSetPCHHeapSize' in found_functions}} @cython.embedsignature(True) def nvrtcSetPCHHeapSize(size_t size): @@ -896,9 +824,6 @@ def nvrtcSetPCHHeapSize(size_t size): with nogil: err = cynvrtc.nvrtcSetPCHHeapSize(size) return (_nvrtcResult(err),) -{{endif}} - -{{if 'nvrtcGetPCHCreateStatus' in found_functions}} @cython.embedsignature(True) def nvrtcGetPCHCreateStatus(prog): @@ -944,9 +869,6 @@ def nvrtcGetPCHCreateStatus(prog): with nogil: err = cynvrtc.nvrtcGetPCHCreateStatus(cyprog) return (_nvrtcResult(err),) -{{endif}} - -{{if 'nvrtcGetPCHHeapSizeRequired' in found_functions}} @cython.embedsignature(True) def nvrtcGetPCHHeapSizeRequired(prog): @@ -981,9 +903,6 @@ def nvrtcGetPCHHeapSizeRequired(prog): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, size) -{{endif}} - -{{if 'nvrtcSetFlowCallback' in found_functions}} @cython.embedsignature(True) def nvrtcSetFlowCallback(prog, callback, payload): @@ -1044,13 +963,10 @@ def nvrtcSetFlowCallback(prog, callback, payload): _helper_input_void_ptr_free(&cycallbackHelper) _helper_input_void_ptr_free(&cypayloadHelper) return (_nvrtcResult(err),) -{{endif}} - -{{if 'nvrtcGetTileIRSize' in found_functions}} @cython.embedsignature(True) def nvrtcGetTileIRSize(prog): - """ + """ Parameters ---------- @@ -1078,13 +994,10 @@ def nvrtcGetTileIRSize(prog): if err != cynvrtc.NVRTC_SUCCESS: return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, TileIRSizeRet) -{{endif}} - -{{if 'nvrtcGetTileIR' in found_functions}} @cython.embedsignature(True) def nvrtcGetTileIR(prog, char* TileIR): - """ + """ Parameters ---------- @@ -1109,7 +1022,6 @@ def nvrtcGetTileIR(prog, char* TileIR): with nogil: err = cynvrtc.nvrtcGetTileIR(cyprog, TileIR) return (_nvrtcResult(err),) -{{endif}} @cython.embedsignature(True) def sizeof(objType): @@ -1125,7 +1037,7 @@ def sizeof(objType): lowered_name : int The size of `objType` in bytes """ - {{if 'nvrtcProgram' in found_types}} + if objType == nvrtcProgram: - return sizeof(cynvrtc.nvrtcProgram){{endif}} + return sizeof(cynvrtc.nvrtcProgram) raise TypeError("Unknown type: " + str(objType))