Skip to content

Ensure all bindings have GIL released #768

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 5 commits into from
Jul 21, 2025
Merged

Ensure all bindings have GIL released #768

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
ecc7a33
ensure all functions have GIL released
leofang Jul 21, 2025
67a8837
add release note entry
leofang Jul 21, 2025
e21dee5
dst/src are not const-qualified in CUDA 12
leofang Jul 21, 2025
79d2461
ensure post-call assumption is honored
leofang Jul 21, 2025
5eb1fee
nit: avoid adding empty lines
leofang Jul 21, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1,765 changes: 1,180 additions & 585 deletions cuda_bindings/cuda/bindings/driver.pyx.in
View file
Open in desktop

Large diffs are not rendered by default.

79 changes: 53 additions & 26 deletions cuda_bindings/cuda/bindings/nvrtc.pyx.in
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,7 @@ ctypedef unsigned long long unsigned_ptr
ctypedef unsigned long long unsigned_long_long_ptr
ctypedef unsigned long long long_long_ptr
ctypedef unsigned long long size_t_ptr
ctypedef unsigned long long long_ptr
ctypedef unsigned long long float_ptr
ctypedef unsigned long long double_ptr
ctypedef unsigned long long void_ptr
Expand Down Expand Up @@ -133,7 +134,8 @@ def nvrtcGetErrorString(result not None : nvrtcResult):
Message string for the given :py:obj:`~.nvrtcResult` code.
"""
cdef cynvrtc.nvrtcResult cyresult = result.value
err = cynvrtc.nvrtcGetErrorString(cyresult)
with nogil:
err = cynvrtc.nvrtcGetErrorString(cyresult)
return (nvrtcResult.NVRTC_SUCCESS, err)
{{endif}}

Expand All @@ -155,7 +157,8 @@ def nvrtcVersion():
"""
cdef int major = 0
cdef int minor = 0
err = cynvrtc.nvrtcVersion(&major, &minor)
with nogil:
err = cynvrtc.nvrtcVersion(&major, &minor)
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None, None)
return (_dict_nvrtcResult[err], major, minor)
Expand All @@ -178,7 +181,8 @@ def nvrtcGetNumSupportedArchs():
number of supported architectures.
"""
cdef int numArchs = 0
err = cynvrtc.nvrtcGetNumSupportedArchs(&numArchs)
with nogil:
err = cynvrtc.nvrtcGetNumSupportedArchs(&numArchs)
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], numArchs)
Expand All @@ -204,7 +208,8 @@ def nvrtcGetSupportedArchs():
_, s = nvrtcGetNumSupportedArchs()
supportedArchs.resize(s)

err = cynvrtc.nvrtcGetSupportedArchs(supportedArchs.data())
with nogil:
err = cynvrtc.nvrtcGetSupportedArchs(supportedArchs.data())
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], supportedArchs)
Expand Down Expand Up @@ -261,7 +266,8 @@ def nvrtcCreateProgram(char* src, char* name, int numHeaders, headers : Optional
if numHeaders > len(includeNames): raise RuntimeError("List is too small: " + str(len(includeNames)) + " < " + str(numHeaders))
cdef vector[const char*] cyheaders = headers
cdef vector[const char*] cyincludeNames = includeNames
err = cynvrtc.nvrtcCreateProgram(<cynvrtc.nvrtcProgram*>prog._pvt_ptr, src, name, numHeaders, cyheaders.data(), cyincludeNames.data())
with nogil:
err = cynvrtc.nvrtcCreateProgram(<cynvrtc.nvrtcProgram*>prog._pvt_ptr, src, name, numHeaders, cyheaders.data(), cyincludeNames.data())
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], prog)
Expand Down Expand Up @@ -298,7 +304,8 @@ def nvrtcDestroyProgram(prog):
cyprog = <cynvrtc.nvrtcProgram*><void_ptr>prog
else:
raise TypeError("Argument 'prog' is not instance of type (expected <class 'int, nvrtc.nvrtcProgram'>, found " + str(type(prog)))
err = cynvrtc.nvrtcDestroyProgram(cyprog)
with nogil:
err = cynvrtc.nvrtcDestroyProgram(cyprog)
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down Expand Up @@ -347,7 +354,8 @@ def nvrtcCompileProgram(prog, int numOptions, options : Optional[Tuple[bytes] |
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
if numOptions > len(options): raise RuntimeError("List is too small: " + str(len(options)) + " < " + str(numOptions))
cdef vector[const char*] cyoptions = options
err = cynvrtc.nvrtcCompileProgram(cyprog, numOptions, cyoptions.data())
with nogil:
err = cynvrtc.nvrtcCompileProgram(cyprog, numOptions, cyoptions.data())
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down Expand Up @@ -384,7 +392,8 @@ def nvrtcGetPTXSize(prog):
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
cdef size_t ptxSizeRet = 0
err = cynvrtc.nvrtcGetPTXSize(cyprog, &ptxSizeRet)
with nogil:
err = cynvrtc.nvrtcGetPTXSize(cyprog, &ptxSizeRet)
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], ptxSizeRet)
Expand Down Expand Up @@ -422,7 +431,8 @@ def nvrtcGetPTX(prog, char* ptx):
else:
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
err = cynvrtc.nvrtcGetPTX(cyprog, ptx)
with nogil:
err = cynvrtc.nvrtcGetPTX(cyprog, ptx)
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down Expand Up @@ -459,7 +469,8 @@ def nvrtcGetCUBINSize(prog):
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
cdef size_t cubinSizeRet = 0
err = cynvrtc.nvrtcGetCUBINSize(cyprog, &cubinSizeRet)
with nogil:
err = cynvrtc.nvrtcGetCUBINSize(cyprog, &cubinSizeRet)
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], cubinSizeRet)
Expand Down Expand Up @@ -497,7 +508,8 @@ def nvrtcGetCUBIN(prog, char* cubin):
else:
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
err = cynvrtc.nvrtcGetCUBIN(cyprog, cubin)
with nogil:
err = cynvrtc.nvrtcGetCUBIN(cyprog, cubin)
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down Expand Up @@ -528,7 +540,8 @@ def nvrtcGetNVVMSize(prog):
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
cdef size_t nvvmSizeRet = 0
err = cynvrtc.nvrtcGetNVVMSize(cyprog, &nvvmSizeRet)
with nogil:
err = cynvrtc.nvrtcGetNVVMSize(cyprog, &nvvmSizeRet)
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], nvvmSizeRet)
Expand Down Expand Up @@ -560,7 +573,8 @@ def nvrtcGetNVVM(prog, char* nvvm):
else:
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
err = cynvrtc.nvrtcGetNVVM(cyprog, nvvm)
with nogil:
err = cynvrtc.nvrtcGetNVVM(cyprog, nvvm)
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down Expand Up @@ -597,7 +611,8 @@ def nvrtcGetLTOIRSize(prog):
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
cdef size_t LTOIRSizeRet = 0
err = cynvrtc.nvrtcGetLTOIRSize(cyprog, &LTOIRSizeRet)
with nogil:
err = cynvrtc.nvrtcGetLTOIRSize(cyprog, &LTOIRSizeRet)
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], LTOIRSizeRet)
Expand Down Expand Up @@ -635,7 +650,8 @@ def nvrtcGetLTOIR(prog, char* LTOIR):
else:
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
err = cynvrtc.nvrtcGetLTOIR(cyprog, LTOIR)
with nogil:
err = cynvrtc.nvrtcGetLTOIR(cyprog, LTOIR)
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down Expand Up @@ -672,7 +688,8 @@ def nvrtcGetOptiXIRSize(prog):
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
cdef size_t optixirSizeRet = 0
err = cynvrtc.nvrtcGetOptiXIRSize(cyprog, &optixirSizeRet)
with nogil:
err = cynvrtc.nvrtcGetOptiXIRSize(cyprog, &optixirSizeRet)
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], optixirSizeRet)
Expand Down Expand Up @@ -710,7 +727,8 @@ def nvrtcGetOptiXIR(prog, char* optixir):
else:
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
err = cynvrtc.nvrtcGetOptiXIR(cyprog, optixir)
with nogil:
err = cynvrtc.nvrtcGetOptiXIR(cyprog, optixir)
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down Expand Up @@ -750,7 +768,8 @@ def nvrtcGetProgramLogSize(prog):
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
cdef size_t logSizeRet = 0
err = cynvrtc.nvrtcGetProgramLogSize(cyprog, &logSizeRet)
with nogil:
err = cynvrtc.nvrtcGetProgramLogSize(cyprog, &logSizeRet)
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], logSizeRet)
Expand Down Expand Up @@ -788,7 +807,8 @@ def nvrtcGetProgramLog(prog, char* log):
else:
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
err = cynvrtc.nvrtcGetProgramLog(cyprog, log)
with nogil:
err = cynvrtc.nvrtcGetProgramLog(cyprog, log)
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down Expand Up @@ -829,7 +849,8 @@ def nvrtcAddNameExpression(prog, char* name_expression):
else:
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
err = cynvrtc.nvrtcAddNameExpression(cyprog, name_expression)
with nogil:
err = cynvrtc.nvrtcAddNameExpression(cyprog, name_expression)
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down Expand Up @@ -871,7 +892,8 @@ def nvrtcGetLoweredName(prog, char* name_expression):
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
cdef const char* lowered_name = NULL
err = cynvrtc.nvrtcGetLoweredName(cyprog, name_expression, &lowered_name)
with nogil:
err = cynvrtc.nvrtcGetLoweredName(cyprog, name_expression, &lowered_name)
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], <bytes>lowered_name if lowered_name != NULL else None)
Expand All @@ -892,7 +914,8 @@ def nvrtcGetPCHHeapSize():
pointer to location where the size of the PCH Heap will be stored
"""
cdef size_t ret = 0
err = cynvrtc.nvrtcGetPCHHeapSize(&ret)
with nogil:
err = cynvrtc.nvrtcGetPCHHeapSize(&ret)
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], ret)
Expand All @@ -918,7 +941,8 @@ def nvrtcSetPCHHeapSize(size_t size):
nvrtcResult
- :py:obj:`~.NVRTC_SUCCESS`
"""
err = cynvrtc.nvrtcSetPCHHeapSize(size)
with nogil:
err = cynvrtc.nvrtcSetPCHHeapSize(size)
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down Expand Up @@ -965,7 +989,8 @@ def nvrtcGetPCHCreateStatus(prog):
else:
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
err = cynvrtc.nvrtcGetPCHCreateStatus(cyprog)
with nogil:
err = cynvrtc.nvrtcGetPCHCreateStatus(cyprog)
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down Expand Up @@ -999,7 +1024,8 @@ def nvrtcGetPCHHeapSizeRequired(prog):
pprog = int(nvrtcProgram(prog))
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
cdef size_t size = 0
err = cynvrtc.nvrtcGetPCHHeapSizeRequired(cyprog, &size)
with nogil:
err = cynvrtc.nvrtcGetPCHHeapSizeRequired(cyprog, &size)
if err != cynvrtc.NVRTC_SUCCESS:
return (_dict_nvrtcResult[err], None)
return (_dict_nvrtcResult[err], size)
Expand Down Expand Up @@ -1061,7 +1087,8 @@ def nvrtcSetFlowCallback(prog, callback, payload):
cdef void* cycallback_ptr = <void*><void_ptr>cycallback.cptr
cypayload = utils.HelperInputVoidPtr(payload)
cdef void* cypayload_ptr = <void*><void_ptr>cypayload.cptr
err = cynvrtc.nvrtcSetFlowCallback(cyprog, cycallback_ptr, cypayload_ptr)
with nogil:
err = cynvrtc.nvrtcSetFlowCallback(cyprog, cycallback_ptr, cypayload_ptr)
return (_dict_nvrtcResult[err],)
{{endif}}

Expand Down
Loading
Loading