From b8183a8539891cea7702885201d74b139d330871 Mon Sep 17 00:00:00 2001 From: Leo Fang Date: Mon, 13 Apr 2026 21:04:32 +0000 Subject: [PATCH 1/6] [DRAFT] Replace nvrtc Tempita templates with pre-generated Cython files MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replace .pyx.in / .pxd.in Tempita templates with plain .pyx / .pxd files generated by the refactored cython-gen (cybind leof/cybind\!369). Changes: - Remove 6 Tempita template files (.pyx.in / .pxd.in) for nvrtc - Add 7 pre-generated Cython files with zero Tempita - Move internal layer from _bindings/ to _internal/ following cybind convention - Split platform-specific code into nvrtc_linux.pyx / nvrtc_windows.pyx - Internal layer now uses cybind-style loading (RTLD_DEFAULT + fallback, FunctionNotFoundError, _init_nvrtc/_check_or_init_nvrtc naming) This is a companion PR to cybind leof/cybind\!369. Build system changes (build_hooks.py, pyproject.toml) to remove pyclibrary are not yet included — this PR shows the generated output for review. Part of https://gitlab-master.nvidia.com/leof/cybind/-/issues/173 Co-Authored-By: Claude Opus 4.6 (1M context) --- .../cuda/bindings/_bindings/cynvrtc.pyx.in | 789 ------------------ .../cynvrtc.pxd.in => _internal/nvrtc.pxd} | 106 +-- .../cuda/bindings/_internal/nvrtc_linux.pyx | 631 ++++++++++++++ .../cuda/bindings/_internal/nvrtc_windows.pyx | 551 ++++++++++++ .../bindings/{cynvrtc.pxd.in => cynvrtc.pxd} | 106 +-- .../bindings/{cynvrtc.pyx.in => cynvrtc.pyx} | 106 +-- .../cuda/bindings/{nvrtc.pxd.in => nvrtc.pxd} | 6 +- .../cuda/bindings/{nvrtc.pyx.in => nvrtc.pyx} | 192 ++--- 8 files changed, 1440 insertions(+), 1047 deletions(-) delete mode 100644 cuda_bindings/cuda/bindings/_bindings/cynvrtc.pyx.in rename cuda_bindings/cuda/bindings/{_bindings/cynvrtc.pxd.in => _internal/nvrtc.pxd} (67%) create mode 100644 cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx create mode 100644 cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx rename cuda_bindings/cuda/bindings/{cynvrtc.pxd.in => cynvrtc.pxd} (72%) rename cuda_bindings/cuda/bindings/{cynvrtc.pyx.in => cynvrtc.pyx} (75%) rename cuda_bindings/cuda/bindings/{nvrtc.pxd.in => nvrtc.pxd} (85%) rename cuda_bindings/cuda/bindings/{nvrtc.pyx.in => nvrtc.pyx} (90%) 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 67% rename from cuda_bindings/cuda/bindings/_bindings/cynvrtc.pxd.in rename to cuda_bindings/cuda/bindings/_internal/nvrtc.pxd index 5a53b926d1..3e8b1f61cc 100644 --- a/cuda_bindings/cuda/bindings/_bindings/cynvrtc.pxd.in +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd @@ -1,136 +1,136 @@ # 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 cuda.bindings.cynvrtc cimport * -{{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/_internal/nvrtc_linux.pyx b/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx new file mode 100644 index 0000000000..6e8371aea2 --- /dev/null +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx @@ -0,0 +1,631 @@ +# 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..4518d3b4fd --- /dev/null +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx @@ -0,0 +1,551 @@ +# 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 + +from libc.stddef cimport wchar_t +from cpython cimport PyUnicode_AsWideCharString, PyMem_Free + +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 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 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..36e3dab00c 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pxd.in +++ b/cuda_bindings/cuda/bindings/cynvrtc.pxd @@ -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 libc.stdint cimport uint32_t, uint64_t @@ -31,133 +31,133 @@ 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 75% rename from cuda_bindings/cuda/bindings/cynvrtc.pyx.in rename to cuda_bindings/cuda/bindings/cynvrtc.pyx index 9e64cbf9c1..5b6a57c3c3 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pyx.in +++ b/cuda_bindings/cuda/bindings/cynvrtc.pyx @@ -1,161 +1,161 @@ # 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._bindings.cynvrtc as cynvrtc -{{if 'nvrtcGetErrorString' in found_functions}} + cdef const char* nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil: return cynvrtc._nvrtcGetErrorString(result) -{{endif}} -{{if 'nvrtcVersion' in found_functions}} + + cdef nvrtcResult nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcVersion(major, minor) -{{endif}} -{{if 'nvrtcGetNumSupportedArchs' in found_functions}} + + cdef nvrtcResult nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetNumSupportedArchs(numArchs) -{{endif}} -{{if 'nvrtcGetSupportedArchs' in found_functions}} + + cdef nvrtcResult nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetSupportedArchs(supportedArchs) -{{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: return cynvrtc._nvrtcCreateProgram(prog, src, name, numHeaders, headers, includeNames) -{{endif}} -{{if 'nvrtcDestroyProgram' in found_functions}} + + cdef nvrtcResult nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcDestroyProgram(prog) -{{endif}} -{{if 'nvrtcCompileProgram' in found_functions}} + + 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}} + + 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}} + + cdef nvrtcResult nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetPTX(prog, ptx) -{{endif}} -{{if 'nvrtcGetCUBINSize' in found_functions}} + + 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}} + + cdef nvrtcResult nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetCUBIN(prog, cubin) -{{endif}} -{{if 'nvrtcGetLTOIRSize' in found_functions}} + + 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}} + + cdef nvrtcResult nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetLTOIR(prog, LTOIR) -{{endif}} -{{if 'nvrtcGetOptiXIRSize' in found_functions}} + + 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}} + + cdef nvrtcResult nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetOptiXIR(prog, optixir) -{{endif}} -{{if 'nvrtcGetProgramLogSize' in found_functions}} + + 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}} + + cdef nvrtcResult nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetProgramLog(prog, log) -{{endif}} -{{if 'nvrtcAddNameExpression' in found_functions}} + + 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}} + + 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}} + + cdef nvrtcResult nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetPCHHeapSize(ret) -{{endif}} -{{if 'nvrtcSetPCHHeapSize' in found_functions}} + + cdef nvrtcResult nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcSetPCHHeapSize(size) -{{endif}} -{{if 'nvrtcGetPCHCreateStatus' in found_functions}} + + cdef nvrtcResult nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetPCHCreateStatus(prog) -{{endif}} -{{if 'nvrtcGetPCHHeapSizeRequired' in found_functions}} + + 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}} + + 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}} + + 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}} + + cdef nvrtcResult nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetTileIR(prog, TileIR) -{{endif}} + 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..96be715ea4 100644 --- a/cuda_bindings/cuda/bindings/nvrtc.pxd.in +++ b/cuda_bindings/cuda/bindings/nvrtc.pxd @@ -1,12 +1,12 @@ # 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 +21,4 @@ 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..2cef2dd799 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 @@ -44,58 +44,58 @@ 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}} + + 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 + + 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 -{{if 'nvrtcProgram' in found_types}} + cdef class nvrtcProgram: """ nvrtcProgram is the unit of compilation, and an opaque handle for a program. @@ -130,9 +130,9 @@ 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 +154,9 @@ 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 +179,9 @@ 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 +203,9 @@ 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 +230,9 @@ 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 +288,9 @@ 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 +324,9 @@ 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 +374,9 @@ 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 +414,9 @@ 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 +451,9 @@ 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 +491,9 @@ 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 +528,9 @@ 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 +568,9 @@ 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 +605,9 @@ 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 +645,9 @@ 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 +682,9 @@ 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 +725,9 @@ 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 +762,9 @@ 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 +804,9 @@ 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 +849,9 @@ 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 +871,9 @@ 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 +896,9 @@ 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 +944,9 @@ 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 +981,9 @@ 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,9 +1044,9 @@ 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): @@ -1078,9 +1078,9 @@ 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): @@ -1109,7 +1109,7 @@ def nvrtcGetTileIR(prog, char* TileIR): with nogil: err = cynvrtc.nvrtcGetTileIR(cyprog, TileIR) return (_nvrtcResult(err),) -{{endif}} + @cython.embedsignature(True) def sizeof(objType): @@ -1125,7 +1125,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)) From 7082b2da06a61e811ae0cda40c01fdaf7989f248 Mon Sep 17 00:00:00 2001 From: Leo Fang Date: Mon, 13 Apr 2026 21:17:42 +0000 Subject: [PATCH 2/6] Regenerate nvrtc files with template-based generation Update generated artifacts to use cybind-style templates: - Fix excessive blank lines from Tempita stripping - Internal layer files now generated from template files under cybind/assets/templates/ via string.Template.substitute() - Proper license headers from templates Co-Authored-By: Claude Opus 4.6 (1M context) --- .../cuda/bindings/_internal/nvrtc.pxd | 86 ++---------------- .../cuda/bindings/_internal/nvrtc_linux.pyx | 89 ++++++------------- .../cuda/bindings/_internal/nvrtc_windows.pyx | 47 +--------- cuda_bindings/cuda/bindings/cynvrtc.pxd | 78 ---------------- cuda_bindings/cuda/bindings/cynvrtc.pyx | 77 ---------------- cuda_bindings/cuda/bindings/nvrtc.pxd | 2 - cuda_bindings/cuda/bindings/nvrtc.pyx | 88 ------------------ 7 files changed, 40 insertions(+), 427 deletions(-) diff --git a/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd b/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd index 3e8b1f61cc..ad1b9059f2 100644 --- a/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd @@ -1,136 +1,64 @@ # 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. 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 cuda.bindings.cynvrtc cimport * - +from ..cynvrtc cimport * +############################################################################### +# Wrapper functions +############################################################################### cdef const char* _nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil - - - cdef nvrtcResult _nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - 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 - - - cdef nvrtcResult _nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const char** options) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetPTXSize(nvrtcProgram prog, size_t* ptxSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcAddNameExpression(nvrtcProgram prog, const char* name_expression) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetLoweredName(nvrtcProgram prog, const char* name_expression, const char** lowered_name) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* payload) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult _nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil - diff --git a/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx b/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx index 6e8371aea2..9e138ce37f 100644 --- a/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx @@ -1,7 +1,9 @@ # 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. 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 intptr_t, uintptr_t import threading @@ -9,7 +11,6 @@ from .utils import FunctionNotFoundError, NotSupportedError from cuda.pathfinder import load_nvidia_dynamic_lib - ############################################################################### # Extern ############################################################################### @@ -48,7 +49,6 @@ cdef int get_cuda_version(): return driver_ver - ############################################################################### # Wrapper init ############################################################################### @@ -83,12 +83,10 @@ 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 @@ -105,187 +103,186 @@ cdef int _init_nvrtc() except -1 nogil: if __nvrtcGetErrorString == NULL: if handle == NULL: handle = load_library() - __nvrtcGetErrorString = dlsym(handle, 'nvrtcGetErrorString') + __nvrtcGetErrorString = dlsym(handle, 'nvrtcGetErrorString') global __nvrtcVersion __nvrtcVersion = dlsym(RTLD_DEFAULT, 'nvrtcVersion') if __nvrtcVersion == NULL: if handle == NULL: handle = load_library() - __nvrtcVersion = dlsym(handle, 'nvrtcVersion') + __nvrtcVersion = dlsym(handle, 'nvrtcVersion') global __nvrtcGetNumSupportedArchs __nvrtcGetNumSupportedArchs = dlsym(RTLD_DEFAULT, 'nvrtcGetNumSupportedArchs') if __nvrtcGetNumSupportedArchs == NULL: if handle == NULL: handle = load_library() - __nvrtcGetNumSupportedArchs = dlsym(handle, 'nvrtcGetNumSupportedArchs') + __nvrtcGetNumSupportedArchs = dlsym(handle, 'nvrtcGetNumSupportedArchs') global __nvrtcGetSupportedArchs __nvrtcGetSupportedArchs = dlsym(RTLD_DEFAULT, 'nvrtcGetSupportedArchs') if __nvrtcGetSupportedArchs == NULL: if handle == NULL: handle = load_library() - __nvrtcGetSupportedArchs = dlsym(handle, 'nvrtcGetSupportedArchs') + __nvrtcGetSupportedArchs = dlsym(handle, 'nvrtcGetSupportedArchs') global __nvrtcCreateProgram __nvrtcCreateProgram = dlsym(RTLD_DEFAULT, 'nvrtcCreateProgram') if __nvrtcCreateProgram == NULL: if handle == NULL: handle = load_library() - __nvrtcCreateProgram = dlsym(handle, 'nvrtcCreateProgram') + __nvrtcCreateProgram = dlsym(handle, 'nvrtcCreateProgram') global __nvrtcDestroyProgram __nvrtcDestroyProgram = dlsym(RTLD_DEFAULT, 'nvrtcDestroyProgram') if __nvrtcDestroyProgram == NULL: if handle == NULL: handle = load_library() - __nvrtcDestroyProgram = dlsym(handle, 'nvrtcDestroyProgram') + __nvrtcDestroyProgram = dlsym(handle, 'nvrtcDestroyProgram') global __nvrtcCompileProgram __nvrtcCompileProgram = dlsym(RTLD_DEFAULT, 'nvrtcCompileProgram') if __nvrtcCompileProgram == NULL: if handle == NULL: handle = load_library() - __nvrtcCompileProgram = dlsym(handle, 'nvrtcCompileProgram') + __nvrtcCompileProgram = dlsym(handle, 'nvrtcCompileProgram') global __nvrtcGetPTXSize __nvrtcGetPTXSize = dlsym(RTLD_DEFAULT, 'nvrtcGetPTXSize') if __nvrtcGetPTXSize == NULL: if handle == NULL: handle = load_library() - __nvrtcGetPTXSize = dlsym(handle, 'nvrtcGetPTXSize') + __nvrtcGetPTXSize = dlsym(handle, 'nvrtcGetPTXSize') global __nvrtcGetPTX __nvrtcGetPTX = dlsym(RTLD_DEFAULT, 'nvrtcGetPTX') if __nvrtcGetPTX == NULL: if handle == NULL: handle = load_library() - __nvrtcGetPTX = dlsym(handle, 'nvrtcGetPTX') + __nvrtcGetPTX = dlsym(handle, 'nvrtcGetPTX') global __nvrtcGetCUBINSize __nvrtcGetCUBINSize = dlsym(RTLD_DEFAULT, 'nvrtcGetCUBINSize') if __nvrtcGetCUBINSize == NULL: if handle == NULL: handle = load_library() - __nvrtcGetCUBINSize = dlsym(handle, 'nvrtcGetCUBINSize') + __nvrtcGetCUBINSize = dlsym(handle, 'nvrtcGetCUBINSize') global __nvrtcGetCUBIN __nvrtcGetCUBIN = dlsym(RTLD_DEFAULT, 'nvrtcGetCUBIN') if __nvrtcGetCUBIN == NULL: if handle == NULL: handle = load_library() - __nvrtcGetCUBIN = dlsym(handle, 'nvrtcGetCUBIN') + __nvrtcGetCUBIN = dlsym(handle, 'nvrtcGetCUBIN') global __nvrtcGetLTOIRSize __nvrtcGetLTOIRSize = dlsym(RTLD_DEFAULT, 'nvrtcGetLTOIRSize') if __nvrtcGetLTOIRSize == NULL: if handle == NULL: handle = load_library() - __nvrtcGetLTOIRSize = dlsym(handle, 'nvrtcGetLTOIRSize') + __nvrtcGetLTOIRSize = dlsym(handle, 'nvrtcGetLTOIRSize') global __nvrtcGetLTOIR __nvrtcGetLTOIR = dlsym(RTLD_DEFAULT, 'nvrtcGetLTOIR') if __nvrtcGetLTOIR == NULL: if handle == NULL: handle = load_library() - __nvrtcGetLTOIR = dlsym(handle, 'nvrtcGetLTOIR') + __nvrtcGetLTOIR = dlsym(handle, 'nvrtcGetLTOIR') global __nvrtcGetOptiXIRSize __nvrtcGetOptiXIRSize = dlsym(RTLD_DEFAULT, 'nvrtcGetOptiXIRSize') if __nvrtcGetOptiXIRSize == NULL: if handle == NULL: handle = load_library() - __nvrtcGetOptiXIRSize = dlsym(handle, 'nvrtcGetOptiXIRSize') + __nvrtcGetOptiXIRSize = dlsym(handle, 'nvrtcGetOptiXIRSize') global __nvrtcGetOptiXIR __nvrtcGetOptiXIR = dlsym(RTLD_DEFAULT, 'nvrtcGetOptiXIR') if __nvrtcGetOptiXIR == NULL: if handle == NULL: handle = load_library() - __nvrtcGetOptiXIR = dlsym(handle, 'nvrtcGetOptiXIR') + __nvrtcGetOptiXIR = dlsym(handle, 'nvrtcGetOptiXIR') global __nvrtcGetProgramLogSize __nvrtcGetProgramLogSize = dlsym(RTLD_DEFAULT, 'nvrtcGetProgramLogSize') if __nvrtcGetProgramLogSize == NULL: if handle == NULL: handle = load_library() - __nvrtcGetProgramLogSize = dlsym(handle, 'nvrtcGetProgramLogSize') + __nvrtcGetProgramLogSize = dlsym(handle, 'nvrtcGetProgramLogSize') global __nvrtcGetProgramLog __nvrtcGetProgramLog = dlsym(RTLD_DEFAULT, 'nvrtcGetProgramLog') if __nvrtcGetProgramLog == NULL: if handle == NULL: handle = load_library() - __nvrtcGetProgramLog = dlsym(handle, 'nvrtcGetProgramLog') + __nvrtcGetProgramLog = dlsym(handle, 'nvrtcGetProgramLog') global __nvrtcAddNameExpression __nvrtcAddNameExpression = dlsym(RTLD_DEFAULT, 'nvrtcAddNameExpression') if __nvrtcAddNameExpression == NULL: if handle == NULL: handle = load_library() - __nvrtcAddNameExpression = dlsym(handle, 'nvrtcAddNameExpression') + __nvrtcAddNameExpression = dlsym(handle, 'nvrtcAddNameExpression') global __nvrtcGetLoweredName __nvrtcGetLoweredName = dlsym(RTLD_DEFAULT, 'nvrtcGetLoweredName') if __nvrtcGetLoweredName == NULL: if handle == NULL: handle = load_library() - __nvrtcGetLoweredName = dlsym(handle, 'nvrtcGetLoweredName') + __nvrtcGetLoweredName = dlsym(handle, 'nvrtcGetLoweredName') global __nvrtcGetPCHHeapSize __nvrtcGetPCHHeapSize = dlsym(RTLD_DEFAULT, 'nvrtcGetPCHHeapSize') if __nvrtcGetPCHHeapSize == NULL: if handle == NULL: handle = load_library() - __nvrtcGetPCHHeapSize = dlsym(handle, 'nvrtcGetPCHHeapSize') + __nvrtcGetPCHHeapSize = dlsym(handle, 'nvrtcGetPCHHeapSize') global __nvrtcSetPCHHeapSize __nvrtcSetPCHHeapSize = dlsym(RTLD_DEFAULT, 'nvrtcSetPCHHeapSize') if __nvrtcSetPCHHeapSize == NULL: if handle == NULL: handle = load_library() - __nvrtcSetPCHHeapSize = dlsym(handle, 'nvrtcSetPCHHeapSize') + __nvrtcSetPCHHeapSize = dlsym(handle, 'nvrtcSetPCHHeapSize') global __nvrtcGetPCHCreateStatus __nvrtcGetPCHCreateStatus = dlsym(RTLD_DEFAULT, 'nvrtcGetPCHCreateStatus') if __nvrtcGetPCHCreateStatus == NULL: if handle == NULL: handle = load_library() - __nvrtcGetPCHCreateStatus = dlsym(handle, 'nvrtcGetPCHCreateStatus') + __nvrtcGetPCHCreateStatus = dlsym(handle, 'nvrtcGetPCHCreateStatus') global __nvrtcGetPCHHeapSizeRequired __nvrtcGetPCHHeapSizeRequired = dlsym(RTLD_DEFAULT, 'nvrtcGetPCHHeapSizeRequired') if __nvrtcGetPCHHeapSizeRequired == NULL: if handle == NULL: handle = load_library() - __nvrtcGetPCHHeapSizeRequired = dlsym(handle, 'nvrtcGetPCHHeapSizeRequired') + __nvrtcGetPCHHeapSizeRequired = dlsym(handle, 'nvrtcGetPCHHeapSizeRequired') global __nvrtcSetFlowCallback __nvrtcSetFlowCallback = dlsym(RTLD_DEFAULT, 'nvrtcSetFlowCallback') if __nvrtcSetFlowCallback == NULL: if handle == NULL: handle = load_library() - __nvrtcSetFlowCallback = dlsym(handle, 'nvrtcSetFlowCallback') + __nvrtcSetFlowCallback = dlsym(handle, 'nvrtcSetFlowCallback') global __nvrtcGetTileIRSize __nvrtcGetTileIRSize = dlsym(RTLD_DEFAULT, 'nvrtcGetTileIRSize') if __nvrtcGetTileIRSize == NULL: if handle == NULL: handle = load_library() - __nvrtcGetTileIRSize = dlsym(handle, 'nvrtcGetTileIRSize') + __nvrtcGetTileIRSize = dlsym(handle, 'nvrtcGetTileIRSize') global __nvrtcGetTileIR __nvrtcGetTileIR = dlsym(RTLD_DEFAULT, 'nvrtcGetTileIR') if __nvrtcGetTileIR == NULL: if handle == NULL: handle = load_library() - __nvrtcGetTileIR = dlsym(handle, 'nvrtcGetTileIR') + __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 @@ -294,7 +291,6 @@ cdef inline int _check_or_init_nvrtc() except -1 nogil: cdef dict func_ptrs = None - cpdef dict _inspect_function_pointers(): global func_ptrs if func_ptrs is not None: @@ -384,14 +380,12 @@ cpdef dict _inspect_function_pointers(): 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 ############################################################################### @@ -404,7 +398,6 @@ cdef const char* _nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil: 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() @@ -413,7 +406,6 @@ cdef nvrtcResult _nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVAL 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() @@ -422,7 +414,6 @@ cdef nvrtcResult _nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_I 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() @@ -431,7 +422,6 @@ cdef nvrtcResult _nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERRO 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() @@ -440,7 +430,6 @@ cdef nvrtcResult _nvrtcCreateProgram(nvrtcProgram* prog, const char* src, const 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() @@ -449,7 +438,6 @@ cdef nvrtcResult _nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_IN 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() @@ -458,7 +446,6 @@ cdef nvrtcResult _nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const c 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() @@ -467,7 +454,6 @@ cdef nvrtcResult _nvrtcGetPTXSize(nvrtcProgram prog, size_t* ptxSizeRet) except 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() @@ -476,7 +462,6 @@ cdef nvrtcResult _nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_ 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() @@ -485,7 +470,6 @@ cdef nvrtcResult _nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) exc 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() @@ -494,7 +478,6 @@ cdef nvrtcResult _nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ER 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() @@ -503,7 +486,6 @@ cdef nvrtcResult _nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) exc 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() @@ -512,7 +494,6 @@ cdef nvrtcResult _nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ER 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() @@ -521,7 +502,6 @@ cdef nvrtcResult _nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) 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() @@ -530,7 +510,6 @@ cdef nvrtcResult _nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRT 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() @@ -539,7 +518,6 @@ cdef nvrtcResult _nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) 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() @@ -548,7 +526,6 @@ cdef nvrtcResult _nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC 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() @@ -557,7 +534,6 @@ cdef nvrtcResult _nvrtcAddNameExpression(nvrtcProgram prog, const char* name_exp 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() @@ -566,7 +542,6 @@ cdef nvrtcResult _nvrtcGetLoweredName(nvrtcProgram prog, const char* name_expres 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() @@ -575,7 +550,6 @@ cdef nvrtcResult _nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_I 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() @@ -584,7 +558,6 @@ cdef nvrtcResult _nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_I 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() @@ -593,7 +566,6 @@ cdef nvrtcResult _nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR 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() @@ -602,7 +574,6 @@ cdef nvrtcResult _nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) e 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() @@ -611,7 +582,6 @@ cdef nvrtcResult _nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* 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() @@ -620,7 +590,6 @@ cdef nvrtcResult _nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) e 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() diff --git a/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx b/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx index 4518d3b4fd..a40b48fc34 100644 --- a/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx @@ -1,17 +1,16 @@ # 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. 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 intptr_t, uintptr_t +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 cpython cimport PyUnicode_AsWideCharString, PyMem_Free - from libc.stddef cimport wchar_t from libc.stdint cimport uintptr_t from cpython cimport PyUnicode_AsWideCharString, PyMem_Free @@ -69,7 +68,6 @@ cdef int get_cuda_version(): return driver_ver - ############################################################################### # Wrapper init ############################################################################### @@ -104,17 +102,9 @@ 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: @@ -205,7 +195,6 @@ cdef int _init_nvrtc() except -1 nogil: __py_nvrtc_init = True return 0 - cdef inline int _check_or_init_nvrtc() except -1 nogil: if __py_nvrtc_init: return 0 @@ -214,7 +203,6 @@ cdef inline int _check_or_init_nvrtc() except -1 nogil: cdef dict func_ptrs = None - cpdef dict _inspect_function_pointers(): global func_ptrs if func_ptrs is not None: @@ -304,14 +292,12 @@ cpdef dict _inspect_function_pointers(): 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 ############################################################################### @@ -324,7 +310,6 @@ cdef const char* _nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil: 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() @@ -333,7 +318,6 @@ cdef nvrtcResult _nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVAL 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() @@ -342,7 +326,6 @@ cdef nvrtcResult _nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_I 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() @@ -351,7 +334,6 @@ cdef nvrtcResult _nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERRO 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() @@ -360,7 +342,6 @@ cdef nvrtcResult _nvrtcCreateProgram(nvrtcProgram* prog, const char* src, const 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() @@ -369,7 +350,6 @@ cdef nvrtcResult _nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_IN 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() @@ -378,7 +358,6 @@ cdef nvrtcResult _nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const c 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() @@ -387,7 +366,6 @@ cdef nvrtcResult _nvrtcGetPTXSize(nvrtcProgram prog, size_t* ptxSizeRet) except 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() @@ -396,7 +374,6 @@ cdef nvrtcResult _nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_ 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() @@ -405,7 +382,6 @@ cdef nvrtcResult _nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) exc 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() @@ -414,7 +390,6 @@ cdef nvrtcResult _nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ER 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() @@ -423,7 +398,6 @@ cdef nvrtcResult _nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) exc 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() @@ -432,7 +406,6 @@ cdef nvrtcResult _nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ER 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() @@ -441,7 +414,6 @@ cdef nvrtcResult _nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) 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() @@ -450,7 +422,6 @@ cdef nvrtcResult _nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRT 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() @@ -459,7 +430,6 @@ cdef nvrtcResult _nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) 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() @@ -468,7 +438,6 @@ cdef nvrtcResult _nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC 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() @@ -477,7 +446,6 @@ cdef nvrtcResult _nvrtcAddNameExpression(nvrtcProgram prog, const char* name_exp 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() @@ -486,7 +454,6 @@ cdef nvrtcResult _nvrtcGetLoweredName(nvrtcProgram prog, const char* name_expres 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() @@ -495,7 +462,6 @@ cdef nvrtcResult _nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_I 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() @@ -504,7 +470,6 @@ cdef nvrtcResult _nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_I 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() @@ -513,7 +478,6 @@ cdef nvrtcResult _nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR 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() @@ -522,7 +486,6 @@ cdef nvrtcResult _nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) e 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() @@ -531,7 +494,6 @@ cdef nvrtcResult _nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* 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() @@ -540,7 +502,6 @@ cdef nvrtcResult _nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) e 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() diff --git a/cuda_bindings/cuda/bindings/cynvrtc.pxd b/cuda_bindings/cuda/bindings/cynvrtc.pxd index 36e3dab00c..71ea3442f0 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pxd +++ b/cuda_bindings/cuda/bindings/cynvrtc.pxd @@ -31,133 +31,55 @@ cdef extern from "nvrtc.h": pass ctypedef _nvrtcProgram* nvrtcProgram - - cdef const char* nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil - - - cdef nvrtcResult nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - 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 - - - cdef nvrtcResult nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const char** options) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetPTXSize(nvrtcProgram prog, size_t* ptxSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcAddNameExpression(nvrtcProgram prog, const char* name_expression) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetLoweredName(nvrtcProgram prog, const char* name_expression, const char** lowered_name) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* payload) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil - - - cdef nvrtcResult nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil - diff --git a/cuda_bindings/cuda/bindings/cynvrtc.pyx b/cuda_bindings/cuda/bindings/cynvrtc.pyx index 5b6a57c3c3..d93cf98ae7 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pyx +++ b/cuda_bindings/cuda/bindings/cynvrtc.pyx @@ -4,158 +4,81 @@ # 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._bindings.cynvrtc as cynvrtc - - cdef const char* nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil: return cynvrtc._nvrtcGetErrorString(result) - - - cdef nvrtcResult nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcVersion(major, minor) - - - cdef nvrtcResult nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetNumSupportedArchs(numArchs) - - - cdef nvrtcResult nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._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) - - - cdef nvrtcResult nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcDestroyProgram(prog) - - - cdef nvrtcResult nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const char** options) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcCompileProgram(prog, numOptions, options) - - - cdef nvrtcResult nvrtcGetPTXSize(nvrtcProgram prog, size_t* ptxSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetPTXSize(prog, ptxSizeRet) - - - cdef nvrtcResult nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetPTX(prog, ptx) - - - cdef nvrtcResult nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetCUBINSize(prog, cubinSizeRet) - - - cdef nvrtcResult nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetCUBIN(prog, cubin) - - - cdef nvrtcResult nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetLTOIRSize(prog, LTOIRSizeRet) - - - cdef nvrtcResult nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetLTOIR(prog, LTOIR) - - - cdef nvrtcResult nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetOptiXIRSize(prog, optixirSizeRet) - - - cdef nvrtcResult nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetOptiXIR(prog, optixir) - - - cdef nvrtcResult nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetProgramLogSize(prog, logSizeRet) - - - cdef nvrtcResult nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetProgramLog(prog, log) - - - cdef nvrtcResult nvrtcAddNameExpression(nvrtcProgram prog, const char* name_expression) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._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) - - - cdef nvrtcResult nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetPCHHeapSize(ret) - - - cdef nvrtcResult nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcSetPCHHeapSize(size) - - - cdef nvrtcResult nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetPCHCreateStatus(prog) - - - cdef nvrtcResult nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetPCHHeapSizeRequired(prog, size) - - - cdef nvrtcResult nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* payload) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcSetFlowCallback(prog, callback, payload) - - - cdef nvrtcResult nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetTileIRSize(prog, TileIRSizeRet) - - - cdef nvrtcResult nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: return cynvrtc._nvrtcGetTileIR(prog, TileIR) diff --git a/cuda_bindings/cuda/bindings/nvrtc.pxd b/cuda_bindings/cuda/bindings/nvrtc.pxd index 96be715ea4..fd49d8bc39 100644 --- a/cuda_bindings/cuda/bindings/nvrtc.pxd +++ b/cuda_bindings/cuda/bindings/nvrtc.pxd @@ -6,8 +6,6 @@ cimport cuda.bindings.cynvrtc as cynvrtc include "_lib/utils.pxd" - - cdef class nvrtcProgram: """ nvrtcProgram is the unit of compilation, and an opaque handle for a program. diff --git a/cuda_bindings/cuda/bindings/nvrtc.pyx b/cuda_bindings/cuda/bindings/nvrtc.pyx index 2cef2dd799..24c6e1557e 100644 --- a/cuda_bindings/cuda/bindings/nvrtc.pyx +++ b/cuda_bindings/cuda/bindings/nvrtc.pyx @@ -43,9 +43,6 @@ ctypedef unsigned long long float_ptr ctypedef unsigned long long double_ptr ctypedef unsigned long long void_ptr - - - class nvrtcResult(_FastEnum): """ The enumerated type nvrtcResult defines API call result codes. @@ -88,15 +85,9 @@ class nvrtcResult(_FastEnum): 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. @@ -131,9 +122,6 @@ cdef class nvrtcProgram: def getPtr(self): return self._pvt_ptr - - - @cython.embedsignature(True) def nvrtcGetErrorString(result not None : nvrtcResult): """ nvrtcGetErrorString is a helper function that returns a string describing the given nvrtcResult code, e.g., NVRTC_SUCCESS to `"NVRTC_SUCCESS"`. For unrecognized enumeration values, it returns `"NVRTC_ERROR unknown"`. @@ -155,9 +143,6 @@ def nvrtcGetErrorString(result not None : nvrtcResult): err = cynvrtc.nvrtcGetErrorString(cyresult) return (nvrtcResult.NVRTC_SUCCESS, err) - - - @cython.embedsignature(True) def nvrtcVersion(): """ nvrtcVersion sets the output parameters `major` and `minor` with the CUDA Runtime Compilation version number. @@ -180,9 +165,6 @@ def nvrtcVersion(): return (_nvrtcResult(err), None, None) return (_nvrtcResult_SUCCESS, major, minor) - - - @cython.embedsignature(True) def nvrtcGetNumSupportedArchs(): """ nvrtcGetNumSupportedArchs sets the output parameter `numArchs` with the number of architectures supported by NVRTC. This can then be used to pass an array to :py:obj:`~.nvrtcGetSupportedArchs` to get the supported architectures. @@ -204,9 +186,6 @@ def nvrtcGetNumSupportedArchs(): return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, numArchs) - - - @cython.embedsignature(True) def nvrtcGetSupportedArchs(): """ nvrtcGetSupportedArchs populates the array passed via the output parameter `supportedArchs` with the architectures supported by NVRTC. The array is sorted in the ascending order. The size of the array to be passed can be determined using :py:obj:`~.nvrtcGetNumSupportedArchs`. @@ -231,9 +210,6 @@ def nvrtcGetSupportedArchs(): return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, supportedArchs) - - - @cython.embedsignature(True) def nvrtcCreateProgram(char* src, char* name, int numHeaders, headers : Optional[tuple[bytes] | list[bytes]], includeNames : Optional[tuple[bytes] | list[bytes]]): """ nvrtcCreateProgram creates an instance of nvrtcProgram with the given input parameters, and sets the output parameter `prog` with it. @@ -289,9 +265,6 @@ def nvrtcCreateProgram(char* src, char* name, int numHeaders, headers : Optional return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, prog) - - - @cython.embedsignature(True) def nvrtcDestroyProgram(prog): """ nvrtcDestroyProgram destroys the given program. @@ -325,9 +298,6 @@ def nvrtcDestroyProgram(prog): err = cynvrtc.nvrtcDestroyProgram(cyprog) return (_nvrtcResult(err),) - - - @cython.embedsignature(True) def nvrtcCompileProgram(prog, int numOptions, options : Optional[tuple[bytes] | list[bytes]]): """ nvrtcCompileProgram compiles the given program. @@ -375,9 +345,6 @@ def nvrtcCompileProgram(prog, int numOptions, options : Optional[tuple[bytes] | err = cynvrtc.nvrtcCompileProgram(cyprog, numOptions, cyoptions.data()) return (_nvrtcResult(err),) - - - @cython.embedsignature(True) def nvrtcGetPTXSize(prog): """ nvrtcGetPTXSize sets the value of `ptxSizeRet` with the size of the PTX generated by the previous compilation of `prog` (including the trailing `NULL`). @@ -415,9 +382,6 @@ def nvrtcGetPTXSize(prog): return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, ptxSizeRet) - - - @cython.embedsignature(True) def nvrtcGetPTX(prog, char* ptx): """ nvrtcGetPTX stores the PTX generated by the previous compilation of `prog` in the memory pointed by `ptx`. @@ -452,9 +416,6 @@ def nvrtcGetPTX(prog, char* ptx): err = cynvrtc.nvrtcGetPTX(cyprog, ptx) return (_nvrtcResult(err),) - - - @cython.embedsignature(True) def nvrtcGetCUBINSize(prog): """ nvrtcGetCUBINSize sets the value of `cubinSizeRet` with the size of the cubin generated by the previous compilation of `prog`. The value of cubinSizeRet is set to 0 if the value specified to `-arch` is a virtual architecture instead of an actual architecture. @@ -492,9 +453,6 @@ def nvrtcGetCUBINSize(prog): return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, cubinSizeRet) - - - @cython.embedsignature(True) def nvrtcGetCUBIN(prog, char* cubin): """ nvrtcGetCUBIN stores the cubin generated by the previous compilation of `prog` in the memory pointed by `cubin`. No cubin is available if the value specified to `-arch` is a virtual architecture instead of an actual architecture. @@ -529,9 +487,6 @@ def nvrtcGetCUBIN(prog, char* cubin): err = cynvrtc.nvrtcGetCUBIN(cyprog, cubin) return (_nvrtcResult(err),) - - - @cython.embedsignature(True) def nvrtcGetLTOIRSize(prog): """ nvrtcGetLTOIRSize sets the value of `LTOIRSizeRet` with the size of the LTO IR generated by the previous compilation of `prog`. The value of LTOIRSizeRet is set to 0 if the program was not compiled with `-dlto`. @@ -569,9 +524,6 @@ def nvrtcGetLTOIRSize(prog): return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, LTOIRSizeRet) - - - @cython.embedsignature(True) def nvrtcGetLTOIR(prog, char* LTOIR): """ nvrtcGetLTOIR stores the LTO IR generated by the previous compilation of `prog` in the memory pointed by `LTOIR`. No LTO IR is available if the program was compiled without `-dlto`. @@ -606,9 +558,6 @@ def nvrtcGetLTOIR(prog, char* LTOIR): err = cynvrtc.nvrtcGetLTOIR(cyprog, LTOIR) return (_nvrtcResult(err),) - - - @cython.embedsignature(True) def nvrtcGetOptiXIRSize(prog): """ nvrtcGetOptiXIRSize sets the value of `optixirSizeRet` with the size of the OptiX IR generated by the previous compilation of `prog`. The value of nvrtcGetOptiXIRSize is set to 0 if the program was compiled with options incompatible with OptiX IR generation. @@ -646,9 +595,6 @@ def nvrtcGetOptiXIRSize(prog): return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, optixirSizeRet) - - - @cython.embedsignature(True) def nvrtcGetOptiXIR(prog, char* optixir): """ nvrtcGetOptiXIR stores the OptiX IR generated by the previous compilation of `prog` in the memory pointed by `optixir`. No OptiX IR is available if the program was compiled with options incompatible with OptiX IR generation. @@ -683,9 +629,6 @@ def nvrtcGetOptiXIR(prog, char* optixir): err = cynvrtc.nvrtcGetOptiXIR(cyprog, optixir) return (_nvrtcResult(err),) - - - @cython.embedsignature(True) def nvrtcGetProgramLogSize(prog): """ nvrtcGetProgramLogSize sets `logSizeRet` with the size of the log generated by the previous compilation of `prog` (including the trailing `NULL`). @@ -726,9 +669,6 @@ def nvrtcGetProgramLogSize(prog): return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, logSizeRet) - - - @cython.embedsignature(True) def nvrtcGetProgramLog(prog, char* log): """ nvrtcGetProgramLog stores the log generated by the previous compilation of `prog` in the memory pointed by `log`. @@ -763,9 +703,6 @@ def nvrtcGetProgramLog(prog, char* log): err = cynvrtc.nvrtcGetProgramLog(cyprog, log) return (_nvrtcResult(err),) - - - @cython.embedsignature(True) def nvrtcAddNameExpression(prog, char* name_expression): """ nvrtcAddNameExpression notes the given name expression denoting the address of a global function or device/__constant__ variable. @@ -805,9 +742,6 @@ def nvrtcAddNameExpression(prog, char* name_expression): err = cynvrtc.nvrtcAddNameExpression(cyprog, name_expression) return (_nvrtcResult(err),) - - - @cython.embedsignature(True) def nvrtcGetLoweredName(prog, char* name_expression): """ nvrtcGetLoweredName extracts the lowered (mangled) name for a global function or device/__constant__ variable, and updates lowered_name to point to it. The memory containing the name is released when the NVRTC program is destroyed by nvrtcDestroyProgram. The identical name expression must have been previously provided to nvrtcAddNameExpression. @@ -850,9 +784,6 @@ def nvrtcGetLoweredName(prog, char* name_expression): return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, lowered_name if lowered_name != NULL else None) - - - @cython.embedsignature(True) def nvrtcGetPCHHeapSize(): """ retrieve the current size of the PCH Heap. @@ -872,9 +803,6 @@ def nvrtcGetPCHHeapSize(): return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, ret) - - - @cython.embedsignature(True) def nvrtcSetPCHHeapSize(size_t size): """ set the size of the PCH Heap. @@ -897,9 +825,6 @@ def nvrtcSetPCHHeapSize(size_t size): err = cynvrtc.nvrtcSetPCHHeapSize(size) return (_nvrtcResult(err),) - - - @cython.embedsignature(True) def nvrtcGetPCHCreateStatus(prog): """ returns the PCH creation status. @@ -945,9 +870,6 @@ def nvrtcGetPCHCreateStatus(prog): err = cynvrtc.nvrtcGetPCHCreateStatus(cyprog) return (_nvrtcResult(err),) - - - @cython.embedsignature(True) def nvrtcGetPCHHeapSizeRequired(prog): """ retrieve the required size of the PCH heap required to compile the given program. @@ -982,9 +904,6 @@ def nvrtcGetPCHHeapSizeRequired(prog): return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, size) - - - @cython.embedsignature(True) def nvrtcSetFlowCallback(prog, callback, payload): """ nvrtcSetFlowCallback registers a callback function that the compiler will invoke at different points during a call to nvrtcCompileProgram, and the callback function can decide whether to cancel compilation by returning specific values. @@ -1045,9 +964,6 @@ def nvrtcSetFlowCallback(prog, callback, payload): _helper_input_void_ptr_free(&cypayloadHelper) return (_nvrtcResult(err),) - - - @cython.embedsignature(True) def nvrtcGetTileIRSize(prog): """ @@ -1079,9 +995,6 @@ def nvrtcGetTileIRSize(prog): return (_nvrtcResult(err), None) return (_nvrtcResult_SUCCESS, TileIRSizeRet) - - - @cython.embedsignature(True) def nvrtcGetTileIR(prog, char* TileIR): """ @@ -1110,7 +1023,6 @@ def nvrtcGetTileIR(prog, char* TileIR): err = cynvrtc.nvrtcGetTileIR(cyprog, TileIR) return (_nvrtcResult(err),) - @cython.embedsignature(True) def sizeof(objType): """ Returns the size of provided CUDA Python structure in bytes From 37a929fe6ebe8fa6c777510b84984f0a974e6e2b Mon Sep 17 00:00:00 2001 From: Leo Fang Date: Mon, 13 Apr 2026 21:34:33 +0000 Subject: [PATCH 3/6] Update cynvrtc.pyx to use _internal cimport (cybind convention) - cynvrtc.pyx now uses "from ._internal cimport nvrtc as _nvrtc" instead of the legacy "cimport _bindings.cynvrtc as cynvrtc" - Function wrappers reference _nvrtc._funcName() matching cybind style - cynvrtc.pxd gets proper template header Co-Authored-By: Claude Opus 4.6 (1M context) --- cuda_bindings/cuda/bindings/cynvrtc.pxd | 5 ++- cuda_bindings/cuda/bindings/cynvrtc.pyx | 58 +++++++++++++------------ 2 files changed, 33 insertions(+), 30 deletions(-) diff --git a/cuda_bindings/cuda/bindings/cynvrtc.pxd b/cuda_bindings/cuda/bindings/cynvrtc.pxd index 71ea3442f0..2f092faa0d 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pxd +++ 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.dev1364+ged01d643e. Do not modify it directly. +# +# This code was automatically generated with version 13.2.0. Do not modify it directly. from libc.stdint cimport uint32_t, uint64_t diff --git a/cuda_bindings/cuda/bindings/cynvrtc.pyx b/cuda_bindings/cuda/bindings/cynvrtc.pyx index d93cf98ae7..976dccf500 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pyx +++ b/cuda_bindings/cuda/bindings/cynvrtc.pyx @@ -1,84 +1,86 @@ # 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. 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._bindings.cynvrtc as cynvrtc +from ._internal cimport nvrtc as _nvrtc cdef const char* nvrtcGetErrorString(nvrtcResult result) except ?NULL nogil: - return cynvrtc._nvrtcGetErrorString(result) + return _nvrtc._nvrtcGetErrorString(result) cdef nvrtcResult nvrtcVersion(int* major, int* minor) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcVersion(major, minor) + return _nvrtc._nvrtcVersion(major, minor) cdef nvrtcResult nvrtcGetNumSupportedArchs(int* numArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetNumSupportedArchs(numArchs) + return _nvrtc._nvrtcGetNumSupportedArchs(numArchs) cdef nvrtcResult nvrtcGetSupportedArchs(int* supportedArchs) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetSupportedArchs(supportedArchs) + 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) + return _nvrtc._nvrtcCreateProgram(prog, src, name, numHeaders, headers, includeNames) cdef nvrtcResult nvrtcDestroyProgram(nvrtcProgram* prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcDestroyProgram(prog) + 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) + 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) + return _nvrtc._nvrtcGetPTXSize(prog, ptxSizeRet) cdef nvrtcResult nvrtcGetPTX(nvrtcProgram prog, char* ptx) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetPTX(prog, ptx) + return _nvrtc._nvrtcGetPTX(prog, ptx) cdef nvrtcResult nvrtcGetCUBINSize(nvrtcProgram prog, size_t* cubinSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetCUBINSize(prog, cubinSizeRet) + return _nvrtc._nvrtcGetCUBINSize(prog, cubinSizeRet) cdef nvrtcResult nvrtcGetCUBIN(nvrtcProgram prog, char* cubin) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetCUBIN(prog, cubin) + return _nvrtc._nvrtcGetCUBIN(prog, cubin) cdef nvrtcResult nvrtcGetLTOIRSize(nvrtcProgram prog, size_t* LTOIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetLTOIRSize(prog, LTOIRSizeRet) + return _nvrtc._nvrtcGetLTOIRSize(prog, LTOIRSizeRet) cdef nvrtcResult nvrtcGetLTOIR(nvrtcProgram prog, char* LTOIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetLTOIR(prog, LTOIR) + return _nvrtc._nvrtcGetLTOIR(prog, LTOIR) cdef nvrtcResult nvrtcGetOptiXIRSize(nvrtcProgram prog, size_t* optixirSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetOptiXIRSize(prog, optixirSizeRet) + return _nvrtc._nvrtcGetOptiXIRSize(prog, optixirSizeRet) cdef nvrtcResult nvrtcGetOptiXIR(nvrtcProgram prog, char* optixir) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetOptiXIR(prog, optixir) + return _nvrtc._nvrtcGetOptiXIR(prog, optixir) cdef nvrtcResult nvrtcGetProgramLogSize(nvrtcProgram prog, size_t* logSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetProgramLogSize(prog, logSizeRet) + return _nvrtc._nvrtcGetProgramLogSize(prog, logSizeRet) cdef nvrtcResult nvrtcGetProgramLog(nvrtcProgram prog, char* log) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetProgramLog(prog, log) + 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) + 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) + return _nvrtc._nvrtcGetLoweredName(prog, name_expression, lowered_name) cdef nvrtcResult nvrtcGetPCHHeapSize(size_t* ret) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetPCHHeapSize(ret) + return _nvrtc._nvrtcGetPCHHeapSize(ret) cdef nvrtcResult nvrtcSetPCHHeapSize(size_t size) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcSetPCHHeapSize(size) + return _nvrtc._nvrtcSetPCHHeapSize(size) cdef nvrtcResult nvrtcGetPCHCreateStatus(nvrtcProgram prog) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetPCHCreateStatus(prog) + return _nvrtc._nvrtcGetPCHCreateStatus(prog) cdef nvrtcResult nvrtcGetPCHHeapSizeRequired(nvrtcProgram prog, size_t* size) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetPCHHeapSizeRequired(prog, size) + 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) + 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) + return _nvrtc._nvrtcGetTileIRSize(prog, TileIRSizeRet) cdef nvrtcResult nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: - return cynvrtc._nvrtcGetTileIR(prog, TileIR) + return _nvrtc._nvrtcGetTileIR(prog, TileIR) From 47712e09037d5173bac550a4ccd7775e2e332252 Mon Sep 17 00:00:00 2001 From: Leo Fang Date: Tue, 14 Apr 2026 01:10:15 +0000 Subject: [PATCH 4/6] Update build system: skip pyclibrary for nvrtc, fix cynvrtc source path - Remove nvrtc from _REQUIRED_HEADERS (no pyclibrary parsing needed; nvrtc .pyx/.pxd files are now pre-generated by cybind) - Update sources_list: cynvrtc.pyx moved from _bindings/ to bindings/ - _rename_architecture_specific_files already handles _internal/nvrtc_linux.pyx -> _internal/nvrtc.pyx Co-Authored-By: Claude Opus 4.6 (1M context) --- cuda_bindings/build_hooks.py | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/cuda_bindings/build_hooks.py b/cuda_bindings/build_hooks.py index ce4745f7a0..957a98c19e 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,7 @@ def _cleanup_dst_files(): sources_list = [ # private (["cuda/bindings/_bindings/cydriver.pyx", "cuda/bindings/_bindings/loader.cpp"], None), - (["cuda/bindings/_bindings/cynvrtc.pyx"], None), + (["cuda/bindings/cynvrtc.pyx"], None), (["cuda/bindings/_bindings/cyruntime.pyx"], static_runtime_libraries), (["cuda/bindings/_bindings/cyruntime_ptds.pyx"], static_runtime_libraries), # utils From 1476959e065ac5d502f9fb212536137eecb3ed89 Mon Sep 17 00:00:00 2001 From: Leo Fang Date: Tue, 14 Apr 2026 01:17:43 +0000 Subject: [PATCH 5/6] Restore generator version in headers, remove duplicate cynvrtc source entry - Regenerate files with generator version in header comments - Remove explicit cynvrtc.pyx from sources_list (already covered by cuda_bindings_files glob) Co-Authored-By: Claude Opus 4.6 (1M context) --- cuda_bindings/build_hooks.py | 3 ++- cuda_bindings/cuda/bindings/_internal/nvrtc.pxd | 2 +- cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx | 2 +- cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx | 2 +- cuda_bindings/cuda/bindings/cynvrtc.pxd | 2 +- cuda_bindings/cuda/bindings/cynvrtc.pyx | 2 +- 6 files changed, 7 insertions(+), 6 deletions(-) diff --git a/cuda_bindings/build_hooks.py b/cuda_bindings/build_hooks.py index 957a98c19e..094d8adfbf 100644 --- a/cuda_bindings/build_hooks.py +++ b/cuda_bindings/build_hooks.py @@ -421,7 +421,8 @@ def _cleanup_dst_files(): sources_list = [ # private (["cuda/bindings/_bindings/cydriver.pyx", "cuda/bindings/_bindings/loader.cpp"], None), - (["cuda/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/_internal/nvrtc.pxd b/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd index ad1b9059f2..f56cacab92 100644 --- a/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd @@ -2,7 +2,7 @@ # # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE # -# This code was automatically generated with version 13.2.0. 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 ..cynvrtc cimport * diff --git a/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx b/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx index 9e138ce37f..8e6062a423 100644 --- a/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx @@ -2,7 +2,7 @@ # # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE # -# This code was automatically generated with version 13.2.0. 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 intptr_t, uintptr_t diff --git a/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx b/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx index a40b48fc34..619bb5bdd5 100644 --- a/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx @@ -2,7 +2,7 @@ # # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE # -# This code was automatically generated with version 13.2.0. 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 intptr_t diff --git a/cuda_bindings/cuda/bindings/cynvrtc.pxd b/cuda_bindings/cuda/bindings/cynvrtc.pxd index 2f092faa0d..06c6c90587 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pxd +++ b/cuda_bindings/cuda/bindings/cynvrtc.pxd @@ -2,7 +2,7 @@ # # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE # -# This code was automatically generated with version 13.2.0. 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 diff --git a/cuda_bindings/cuda/bindings/cynvrtc.pyx b/cuda_bindings/cuda/bindings/cynvrtc.pyx index 976dccf500..8818ee2131 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pyx +++ b/cuda_bindings/cuda/bindings/cynvrtc.pyx @@ -2,7 +2,7 @@ # # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE # -# This code was automatically generated with version 13.2.0. 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 ._internal cimport nvrtc as _nvrtc From 3317324b2a6d5f9dad4a57f27a73709737de38ab Mon Sep 17 00:00:00 2001 From: Leo Fang Date: Tue, 14 Apr 2026 01:22:21 +0000 Subject: [PATCH 6/6] Fix pre-commit: strip trailing whitespace and ensure final newline Regenerated from latest cybind which now cleans up trailing whitespace and ensures files end with exactly one newline. Co-Authored-By: Claude Opus 4.6 (1M context) --- .../cuda/bindings/_internal/nvrtc.pxd | 1 - .../cuda/bindings/_internal/nvrtc_linux.pyx | 1 - .../cuda/bindings/_internal/nvrtc_windows.pyx | 1 - cuda_bindings/cuda/bindings/cynvrtc.pxd | 1 - cuda_bindings/cuda/bindings/cynvrtc.pyx | 1 - cuda_bindings/cuda/bindings/nvrtc.pxd | 1 - cuda_bindings/cuda/bindings/nvrtc.pyx | 42 +++++++++---------- 7 files changed, 21 insertions(+), 27 deletions(-) diff --git a/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd b/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd index f56cacab92..1964af1f16 100644 --- a/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc.pxd @@ -61,4 +61,3 @@ cdef nvrtcResult _nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* cdef nvrtcResult _nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil cdef nvrtcResult _nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil - diff --git a/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx b/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx index 8e6062a423..780042a8cd 100644 --- a/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc_linux.pyx @@ -597,4 +597,3 @@ cdef nvrtcResult _nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ 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 index 619bb5bdd5..1fb555644e 100644 --- a/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx +++ b/cuda_bindings/cuda/bindings/_internal/nvrtc_windows.pyx @@ -509,4 +509,3 @@ cdef nvrtcResult _nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ with gil: raise FunctionNotFoundError("function nvrtcGetTileIR is not found") return (__nvrtcGetTileIR)(prog, TileIR) - diff --git a/cuda_bindings/cuda/bindings/cynvrtc.pxd b/cuda_bindings/cuda/bindings/cynvrtc.pxd index 06c6c90587..6bf1bda94e 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pxd +++ b/cuda_bindings/cuda/bindings/cynvrtc.pxd @@ -83,4 +83,3 @@ cdef nvrtcResult nvrtcSetFlowCallback(nvrtcProgram prog, void* callback, void* p cdef nvrtcResult nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) except ?NVRTC_ERROR_INVALID_INPUT nogil cdef nvrtcResult nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil - diff --git a/cuda_bindings/cuda/bindings/cynvrtc.pyx b/cuda_bindings/cuda/bindings/cynvrtc.pyx index 8818ee2131..2bf71611d7 100644 --- a/cuda_bindings/cuda/bindings/cynvrtc.pyx +++ b/cuda_bindings/cuda/bindings/cynvrtc.pyx @@ -83,4 +83,3 @@ cdef nvrtcResult nvrtcGetTileIRSize(nvrtcProgram prog, size_t* TileIRSizeRet) ex cdef nvrtcResult nvrtcGetTileIR(nvrtcProgram prog, char* TileIR) except ?NVRTC_ERROR_INVALID_INPUT nogil: return _nvrtc._nvrtcGetTileIR(prog, TileIR) - diff --git a/cuda_bindings/cuda/bindings/nvrtc.pxd b/cuda_bindings/cuda/bindings/nvrtc.pxd index fd49d8bc39..5cb372430f 100644 --- a/cuda_bindings/cuda/bindings/nvrtc.pxd +++ b/cuda_bindings/cuda/bindings/nvrtc.pxd @@ -19,4 +19,3 @@ cdef class nvrtcProgram: """ cdef cynvrtc.nvrtcProgram _pvt_val cdef cynvrtc.nvrtcProgram* _pvt_ptr - diff --git a/cuda_bindings/cuda/bindings/nvrtc.pyx b/cuda_bindings/cuda/bindings/nvrtc.pyx index 24c6e1557e..aca1ead365 100644 --- a/cuda_bindings/cuda/bindings/nvrtc.pyx +++ b/cuda_bindings/cuda/bindings/nvrtc.pyx @@ -48,41 +48,41 @@ class nvrtcResult(_FastEnum): The enumerated type nvrtcResult defines API call result codes. NVRTC API functions return nvrtcResult to indicate the call result. """ - + 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 - + 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 @@ -966,7 +966,7 @@ def nvrtcSetFlowCallback(prog, callback, payload): @cython.embedsignature(True) def nvrtcGetTileIRSize(prog): - """ + """ Parameters ---------- @@ -997,7 +997,7 @@ def nvrtcGetTileIRSize(prog): @cython.embedsignature(True) def nvrtcGetTileIR(prog, char* TileIR): - """ + """ Parameters ---------- @@ -1037,7 +1037,7 @@ def sizeof(objType): lowered_name : int The size of `objType` in bytes """ - + if objType == nvrtcProgram: return sizeof(cynvrtc.nvrtcProgram) raise TypeError("Unknown type: " + str(objType))