Skip to content

Commit b827181

Browse files
leofangleofang
authored
Ensure all bindings have GIL released (#768)
* ensure all functions have GIL released * add release note entry * dst/src are not const-qualified in CUDA 12 * ensure post-call assumption is honored * nit: avoid adding empty lines
1 parent 5d8b2d7 commit b827181

7 files changed

Lines changed: 1925 additions & 1002 deletions

File tree

‎cuda_bindings/cuda/bindings/driver.pyx.in‎

Lines changed: 1180 additions & 585 deletions
Large diffs are not rendered by default.

‎cuda_bindings/cuda/bindings/nvrtc.pyx.in‎

Lines changed: 53 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -34,6 +34,7 @@ ctypedef unsigned long long unsigned_ptr
3434
ctypedef unsigned long long unsigned_long_long_ptr
3535
ctypedef unsigned long long long_long_ptr
3636
ctypedef unsigned long long size_t_ptr
37+
ctypedef unsigned long long long_ptr
3738
ctypedef unsigned long long float_ptr
3839
ctypedef unsigned long long double_ptr
3940
ctypedef unsigned long long void_ptr
@@ -133,7 +134,8 @@ def nvrtcGetErrorString(result not None : nvrtcResult):
133134
Message string for the given :py:obj:`~.nvrtcResult` code.
134135
"""
135136
cdef cynvrtc.nvrtcResult cyresult = result.value
136-
err = cynvrtc.nvrtcGetErrorString(cyresult)
137+
with nogil:
138+
err = cynvrtc.nvrtcGetErrorString(cyresult)
137139
return (nvrtcResult.NVRTC_SUCCESS, err)
138140
{{endif}}
139141

@@ -155,7 +157,8 @@ def nvrtcVersion():
155157
"""
156158
cdef int major = 0
157159
cdef int minor = 0
158-
err = cynvrtc.nvrtcVersion(&major, &minor)
160+
with nogil:
161+
err = cynvrtc.nvrtcVersion(&major, &minor)
159162
if err != cynvrtc.NVRTC_SUCCESS:
160163
return (_dict_nvrtcResult[err], None, None)
161164
return (_dict_nvrtcResult[err], major, minor)
@@ -178,7 +181,8 @@ def nvrtcGetNumSupportedArchs():
178181
number of supported architectures.
179182
"""
180183
cdef int numArchs = 0
181-
err = cynvrtc.nvrtcGetNumSupportedArchs(&numArchs)
184+
with nogil:
185+
err = cynvrtc.nvrtcGetNumSupportedArchs(&numArchs)
182186
if err != cynvrtc.NVRTC_SUCCESS:
183187
return (_dict_nvrtcResult[err], None)
184188
return (_dict_nvrtcResult[err], numArchs)
@@ -204,7 +208,8 @@ def nvrtcGetSupportedArchs():
204208
_, s = nvrtcGetNumSupportedArchs()
205209
supportedArchs.resize(s)
206210

207-
err = cynvrtc.nvrtcGetSupportedArchs(supportedArchs.data())
211+
with nogil:
212+
err = cynvrtc.nvrtcGetSupportedArchs(supportedArchs.data())
208213
if err != cynvrtc.NVRTC_SUCCESS:
209214
return (_dict_nvrtcResult[err], None)
210215
return (_dict_nvrtcResult[err], supportedArchs)
@@ -261,7 +266,8 @@ def nvrtcCreateProgram(char* src, char* name, int numHeaders, headers : Optional
261266
if numHeaders > len(includeNames): raise RuntimeError("List is too small: " + str(len(includeNames)) + " < " + str(numHeaders))
262267
cdef vector[const char*] cyheaders = headers
263268
cdef vector[const char*] cyincludeNames = includeNames
264-
err = cynvrtc.nvrtcCreateProgram(<cynvrtc.nvrtcProgram*>prog._pvt_ptr, src, name, numHeaders, cyheaders.data(), cyincludeNames.data())
269+
with nogil:
270+
err = cynvrtc.nvrtcCreateProgram(<cynvrtc.nvrtcProgram*>prog._pvt_ptr, src, name, numHeaders, cyheaders.data(), cyincludeNames.data())
265271
if err != cynvrtc.NVRTC_SUCCESS:
266272
return (_dict_nvrtcResult[err], None)
267273
return (_dict_nvrtcResult[err], prog)
@@ -298,7 +304,8 @@ def nvrtcDestroyProgram(prog):
298304
cyprog = <cynvrtc.nvrtcProgram*><void_ptr>prog
299305
else:
300306
raise TypeError("Argument 'prog' is not instance of type (expected <class 'int, nvrtc.nvrtcProgram'>, found " + str(type(prog)))
301-
err = cynvrtc.nvrtcDestroyProgram(cyprog)
307+
with nogil:
308+
err = cynvrtc.nvrtcDestroyProgram(cyprog)
302309
return (_dict_nvrtcResult[err],)
303310
{{endif}}
304311

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

@@ -384,7 +392,8 @@ def nvrtcGetPTXSize(prog):
384392
pprog = int(nvrtcProgram(prog))
385393
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
386394
cdef size_t ptxSizeRet = 0
387-
err = cynvrtc.nvrtcGetPTXSize(cyprog, &ptxSizeRet)
395+
with nogil:
396+
err = cynvrtc.nvrtcGetPTXSize(cyprog, &ptxSizeRet)
388397
if err != cynvrtc.NVRTC_SUCCESS:
389398
return (_dict_nvrtcResult[err], None)
390399
return (_dict_nvrtcResult[err], ptxSizeRet)
@@ -422,7 +431,8 @@ def nvrtcGetPTX(prog, char* ptx):
422431
else:
423432
pprog = int(nvrtcProgram(prog))
424433
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
425-
err = cynvrtc.nvrtcGetPTX(cyprog, ptx)
434+
with nogil:
435+
err = cynvrtc.nvrtcGetPTX(cyprog, ptx)
426436
return (_dict_nvrtcResult[err],)
427437
{{endif}}
428438

@@ -459,7 +469,8 @@ def nvrtcGetCUBINSize(prog):
459469
pprog = int(nvrtcProgram(prog))
460470
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
461471
cdef size_t cubinSizeRet = 0
462-
err = cynvrtc.nvrtcGetCUBINSize(cyprog, &cubinSizeRet)
472+
with nogil:
473+
err = cynvrtc.nvrtcGetCUBINSize(cyprog, &cubinSizeRet)
463474
if err != cynvrtc.NVRTC_SUCCESS:
464475
return (_dict_nvrtcResult[err], None)
465476
return (_dict_nvrtcResult[err], cubinSizeRet)
@@ -497,7 +508,8 @@ def nvrtcGetCUBIN(prog, char* cubin):
497508
else:
498509
pprog = int(nvrtcProgram(prog))
499510
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
500-
err = cynvrtc.nvrtcGetCUBIN(cyprog, cubin)
511+
with nogil:
512+
err = cynvrtc.nvrtcGetCUBIN(cyprog, cubin)
501513
return (_dict_nvrtcResult[err],)
502514
{{endif}}
503515

@@ -528,7 +540,8 @@ def nvrtcGetNVVMSize(prog):
528540
pprog = int(nvrtcProgram(prog))
529541
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
530542
cdef size_t nvvmSizeRet = 0
531-
err = cynvrtc.nvrtcGetNVVMSize(cyprog, &nvvmSizeRet)
543+
with nogil:
544+
err = cynvrtc.nvrtcGetNVVMSize(cyprog, &nvvmSizeRet)
532545
if err != cynvrtc.NVRTC_SUCCESS:
533546
return (_dict_nvrtcResult[err], None)
534547
return (_dict_nvrtcResult[err], nvvmSizeRet)
@@ -560,7 +573,8 @@ def nvrtcGetNVVM(prog, char* nvvm):
560573
else:
561574
pprog = int(nvrtcProgram(prog))
562575
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
563-
err = cynvrtc.nvrtcGetNVVM(cyprog, nvvm)
576+
with nogil:
577+
err = cynvrtc.nvrtcGetNVVM(cyprog, nvvm)
564578
return (_dict_nvrtcResult[err],)
565579
{{endif}}
566580

@@ -597,7 +611,8 @@ def nvrtcGetLTOIRSize(prog):
597611
pprog = int(nvrtcProgram(prog))
598612
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
599613
cdef size_t LTOIRSizeRet = 0
600-
err = cynvrtc.nvrtcGetLTOIRSize(cyprog, &LTOIRSizeRet)
614+
with nogil:
615+
err = cynvrtc.nvrtcGetLTOIRSize(cyprog, &LTOIRSizeRet)
601616
if err != cynvrtc.NVRTC_SUCCESS:
602617
return (_dict_nvrtcResult[err], None)
603618
return (_dict_nvrtcResult[err], LTOIRSizeRet)
@@ -635,7 +650,8 @@ def nvrtcGetLTOIR(prog, char* LTOIR):
635650
else:
636651
pprog = int(nvrtcProgram(prog))
637652
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
638-
err = cynvrtc.nvrtcGetLTOIR(cyprog, LTOIR)
653+
with nogil:
654+
err = cynvrtc.nvrtcGetLTOIR(cyprog, LTOIR)
639655
return (_dict_nvrtcResult[err],)
640656
{{endif}}
641657

@@ -672,7 +688,8 @@ def nvrtcGetOptiXIRSize(prog):
672688
pprog = int(nvrtcProgram(prog))
673689
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
674690
cdef size_t optixirSizeRet = 0
675-
err = cynvrtc.nvrtcGetOptiXIRSize(cyprog, &optixirSizeRet)
691+
with nogil:
692+
err = cynvrtc.nvrtcGetOptiXIRSize(cyprog, &optixirSizeRet)
676693
if err != cynvrtc.NVRTC_SUCCESS:
677694
return (_dict_nvrtcResult[err], None)
678695
return (_dict_nvrtcResult[err], optixirSizeRet)
@@ -710,7 +727,8 @@ def nvrtcGetOptiXIR(prog, char* optixir):
710727
else:
711728
pprog = int(nvrtcProgram(prog))
712729
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
713-
err = cynvrtc.nvrtcGetOptiXIR(cyprog, optixir)
730+
with nogil:
731+
err = cynvrtc.nvrtcGetOptiXIR(cyprog, optixir)
714732
return (_dict_nvrtcResult[err],)
715733
{{endif}}
716734

@@ -750,7 +768,8 @@ def nvrtcGetProgramLogSize(prog):
750768
pprog = int(nvrtcProgram(prog))
751769
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
752770
cdef size_t logSizeRet = 0
753-
err = cynvrtc.nvrtcGetProgramLogSize(cyprog, &logSizeRet)
771+
with nogil:
772+
err = cynvrtc.nvrtcGetProgramLogSize(cyprog, &logSizeRet)
754773
if err != cynvrtc.NVRTC_SUCCESS:
755774
return (_dict_nvrtcResult[err], None)
756775
return (_dict_nvrtcResult[err], logSizeRet)
@@ -788,7 +807,8 @@ def nvrtcGetProgramLog(prog, char* log):
788807
else:
789808
pprog = int(nvrtcProgram(prog))
790809
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
791-
err = cynvrtc.nvrtcGetProgramLog(cyprog, log)
810+
with nogil:
811+
err = cynvrtc.nvrtcGetProgramLog(cyprog, log)
792812
return (_dict_nvrtcResult[err],)
793813
{{endif}}
794814

@@ -829,7 +849,8 @@ def nvrtcAddNameExpression(prog, char* name_expression):
829849
else:
830850
pprog = int(nvrtcProgram(prog))
831851
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
832-
err = cynvrtc.nvrtcAddNameExpression(cyprog, name_expression)
852+
with nogil:
853+
err = cynvrtc.nvrtcAddNameExpression(cyprog, name_expression)
833854
return (_dict_nvrtcResult[err],)
834855
{{endif}}
835856

@@ -871,7 +892,8 @@ def nvrtcGetLoweredName(prog, char* name_expression):
871892
pprog = int(nvrtcProgram(prog))
872893
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
873894
cdef const char* lowered_name = NULL
874-
err = cynvrtc.nvrtcGetLoweredName(cyprog, name_expression, &lowered_name)
895+
with nogil:
896+
err = cynvrtc.nvrtcGetLoweredName(cyprog, name_expression, &lowered_name)
875897
if err != cynvrtc.NVRTC_SUCCESS:
876898
return (_dict_nvrtcResult[err], None)
877899
return (_dict_nvrtcResult[err], <bytes>lowered_name if lowered_name != NULL else None)
@@ -892,7 +914,8 @@ def nvrtcGetPCHHeapSize():
892914
pointer to location where the size of the PCH Heap will be stored
893915
"""
894916
cdef size_t ret = 0
895-
err = cynvrtc.nvrtcGetPCHHeapSize(&ret)
917+
with nogil:
918+
err = cynvrtc.nvrtcGetPCHHeapSize(&ret)
896919
if err != cynvrtc.NVRTC_SUCCESS:
897920
return (_dict_nvrtcResult[err], None)
898921
return (_dict_nvrtcResult[err], ret)
@@ -918,7 +941,8 @@ def nvrtcSetPCHHeapSize(size_t size):
918941
nvrtcResult
919942
- :py:obj:`~.NVRTC_SUCCESS`
920943
"""
921-
err = cynvrtc.nvrtcSetPCHHeapSize(size)
944+
with nogil:
945+
err = cynvrtc.nvrtcSetPCHHeapSize(size)
922946
return (_dict_nvrtcResult[err],)
923947
{{endif}}
924948

@@ -965,7 +989,8 @@ def nvrtcGetPCHCreateStatus(prog):
965989
else:
966990
pprog = int(nvrtcProgram(prog))
967991
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
968-
err = cynvrtc.nvrtcGetPCHCreateStatus(cyprog)
992+
with nogil:
993+
err = cynvrtc.nvrtcGetPCHCreateStatus(cyprog)
969994
return (_dict_nvrtcResult[err],)
970995
{{endif}}
971996

@@ -999,7 +1024,8 @@ def nvrtcGetPCHHeapSizeRequired(prog):
9991024
pprog = int(nvrtcProgram(prog))
10001025
cyprog = <cynvrtc.nvrtcProgram><void_ptr>pprog
10011026
cdef size_t size = 0
1002-
err = cynvrtc.nvrtcGetPCHHeapSizeRequired(cyprog, &size)
1027+
with nogil:
1028+
err = cynvrtc.nvrtcGetPCHHeapSizeRequired(cyprog, &size)
10031029
if err != cynvrtc.NVRTC_SUCCESS:
10041030
return (_dict_nvrtcResult[err], None)
10051031
return (_dict_nvrtcResult[err], size)
@@ -1061,7 +1087,8 @@ def nvrtcSetFlowCallback(prog, callback, payload):
10611087
cdef void* cycallback_ptr = <void*><void_ptr>cycallback.cptr
10621088
cypayload = utils.HelperInputVoidPtr(payload)
10631089
cdef void* cypayload_ptr = <void*><void_ptr>cypayload.cptr
1064-
err = cynvrtc.nvrtcSetFlowCallback(cyprog, cycallback_ptr, cypayload_ptr)
1090+
with nogil:
1091+
err = cynvrtc.nvrtcSetFlowCallback(cyprog, cycallback_ptr, cypayload_ptr)
10651092
return (_dict_nvrtcResult[err],)
10661093
{{endif}}
10671094

0 commit comments

Comments
 (0)