From ef95b4640574feee8a37df20a0d1f7f3974b9da1 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Thu, 6 Mar 2025 11:37:50 -0800 Subject: [PATCH 1/8] Update cuda/bindings/driver.pyx.in, nvrtc.pyx.in to return None instead of segfaulting when converting char* NULL to bytes. --- cuda_bindings/cuda/bindings/driver.pyx.in | 8 ++++---- cuda_bindings/cuda/bindings/nvrtc.pyx.in | 2 +- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/cuda_bindings/cuda/bindings/driver.pyx.in b/cuda_bindings/cuda/bindings/driver.pyx.in index 6be529571..7eec936ba 100644 --- a/cuda_bindings/cuda/bindings/driver.pyx.in +++ b/cuda_bindings/cuda/bindings/driver.pyx.in @@ -22246,7 +22246,7 @@ def cuGetErrorString(error not None : CUresult): cdef cydriver.CUresult cyerror = error.value cdef const char* pStr = NULL err = cydriver.cuGetErrorString(cyerror, &pStr) - return (CUresult(err), pStr) + return (CUresult(err), pStr if pStr != NULL else None) {{endif}} {{if 'cuGetErrorName' in found_functions}} @@ -22279,7 +22279,7 @@ def cuGetErrorName(error not None : CUresult): cdef cydriver.CUresult cyerror = error.value cdef const char* pStr = NULL err = cydriver.cuGetErrorName(cyerror, &pStr) - return (CUresult(err), pStr) + return (CUresult(err), pStr if pStr != NULL else None) {{endif}} {{if 'cuInit' in found_functions}} @@ -27132,7 +27132,7 @@ def cuKernelGetName(hfunc): cyhfunc = phfunc cdef const char* name = NULL err = cydriver.cuKernelGetName(&name, cyhfunc) - return (CUresult(err), name) + return (CUresult(err), name if name != NULL else None) {{endif}} {{if 'cuKernelGetParamInfo' in found_functions}} @@ -38744,7 +38744,7 @@ def cuFuncGetName(hfunc): cyhfunc = phfunc cdef const char* name = NULL err = cydriver.cuFuncGetName(&name, cyhfunc) - return (CUresult(err), name) + return (CUresult(err), name if name != NULL else None) {{endif}} {{if 'cuFuncGetParamInfo' in found_functions}} diff --git a/cuda_bindings/cuda/bindings/nvrtc.pyx.in b/cuda_bindings/cuda/bindings/nvrtc.pyx.in index 67808422d..6c2704c59 100644 --- a/cuda_bindings/cuda/bindings/nvrtc.pyx.in +++ b/cuda_bindings/cuda/bindings/nvrtc.pyx.in @@ -870,7 +870,7 @@ def nvrtcGetLoweredName(prog, char* name_expression): cyprog = pprog cdef const char* lowered_name = NULL err = cynvrtc.nvrtcGetLoweredName(cyprog, name_expression, &lowered_name) - return (nvrtcResult(err), lowered_name) + return (nvrtcResult(err), lowered_name if lowered_name != NULL else None) {{endif}} {{if 'nvrtcGetPCHHeapSize' in found_functions}} From 0ff0cf4678eea050306950e1cecb3eaaec1df097 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Thu, 6 Mar 2025 13:06:01 -0800 Subject: [PATCH 2/8] Add test_all_CUresult_codes Based on: https://github.com/NVIDIA/cuda-python/pull/458/commits/d3df80d6d4db1062f3dd969e2a99b3dc08837f60#diff-29c7ab322cdb6dfa72e21edffba21d51afc1a5fed8b3974206c6ba7bd4dcfd06 --- cuda_bindings/tests/test_cuda.py | 26 ++++++++++++++++++++++++++ 1 file changed, 26 insertions(+) diff --git a/cuda_bindings/tests/test_cuda.py b/cuda_bindings/tests/test_cuda.py index 7be1b0b95..e6a8ffce7 100644 --- a/cuda_bindings/tests/test_cuda.py +++ b/cuda_bindings/tests/test_cuda.py @@ -948,3 +948,29 @@ def test_conditional(): def test_CUmemDecompressParams_st(): desc = cuda.CUmemDecompressParams_st() assert int(desc.dstActBytes) == 0 + + +def test_all_CUresult_codes(max_code=1000): + num_good = 0 + for code in range(max_code + 1): + try: + error = cuda.CUresult(code) + except ValueError: + pass # cython-generated enum does not exist for this code + else: + err_name, name = cuda.cuGetErrorName(error) + if err_name == cuda.CUresult.CUDA_SUCCESS: + assert name + err_desc, desc = cuda.cuGetErrorString(error) + assert err_desc == cuda.CUresult.CUDA_SUCCESS + assert desc + num_good += 1 + else: + # cython-generated enum exists but is not known to an older driver + # (example: cuda-bindings built with CTK 12.8, driver from CTK 12.0) + assert name is None + assert err_name == cuda.CUresult.CUDA_ERROR_INVALID_VALUE + # Super-simple smoke check: Do we have at least some "good" codes? + # The number will increase over time as new enums are added and support for + # old CTKs is dropped, but it is not critical that this number is updated. + assert num_good >= 76 # CTK 11.0.3_450.51.06 From 9dd26307572107954347d390f3c750fb4ba0df3e Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Thu, 6 Mar 2025 15:00:56 -0800 Subject: [PATCH 3/8] Increase tolerance in test_timing() to avoid flaky tests. Observed failures: ``` > assert delay_seconds * 1000 <= elapsed_time_ms < delay_seconds * 1000 + 2 # tolerance 2 ms E assert 503.1598205566406 < ((0.5 * 1000) + 2) ``` ``` > assert delay_seconds * 1000 <= elapsed_time_ms < delay_seconds * 1000 + 2 # tolerance 2 ms E assert 505.0583190917969 < ((0.5 * 1000) + 2) ``` --- cuda_core/tests/test_event.py | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/cuda_core/tests/test_event.py b/cuda_core/tests/test_event.py index 384cf4586..97c6b336b 100644 --- a/cuda_core/tests/test_event.py +++ b/cuda_core/tests/test_event.py @@ -33,7 +33,12 @@ def test_timing(init_cuda, enable_timing): if enable_timing: elapsed_time_ms = e2 - e1 assert isinstance(elapsed_time_ms, float) - assert delay_seconds * 1000 <= elapsed_time_ms < delay_seconds * 1000 + 2 # tolerance 2 ms + # Using a generous tolerance, to avoid flaky tests: + # We only want to exercise the __sub__ method, this test is not meant + # to stress-test the CUDA driver or time.sleep(). + delay_ms = delay_seconds * 1000 + generous_tolerance = 20 + assert delay_ms <= elapsed_time_ms < delay_ms + generous_tolerance else: with pytest.raises(RuntimeError) as e: elapsed_time_ms = e2 - e1 From dffc603319ef2ae0e2bae0d77b44d87c920e4e2d Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Thu, 6 Mar 2025 15:11:35 -0800 Subject: [PATCH 4/8] Add to the new test_all_CUresult_codes() to also exercise cuGetErrorString() CUDA_ERROR_INVALID_VALUE --- cuda_bindings/tests/test_cuda.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/cuda_bindings/tests/test_cuda.py b/cuda_bindings/tests/test_cuda.py index e6a8ffce7..0137a4d5b 100644 --- a/cuda_bindings/tests/test_cuda.py +++ b/cuda_bindings/tests/test_cuda.py @@ -970,6 +970,9 @@ def test_all_CUresult_codes(max_code=1000): # (example: cuda-bindings built with CTK 12.8, driver from CTK 12.0) assert name is None assert err_name == cuda.CUresult.CUDA_ERROR_INVALID_VALUE + err_desc, desc = cuda.cuGetErrorString(error) + assert err_desc == cuda.CUresult.CUDA_ERROR_INVALID_VALUE + assert desc is None # Super-simple smoke check: Do we have at least some "good" codes? # The number will increase over time as new enums are added and support for # old CTKs is dropped, but it is not critical that this number is updated. From dee59ba025382ab649c8f2ea938343e69b388d13 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Thu, 6 Mar 2025 15:42:43 -0800 Subject: [PATCH 5/8] Exercise all other bindings changed in this PR. All of these would have segfaulted before. --- cuda_bindings/tests/test_cuda.py | 12 ++++++++++++ cuda_bindings/tests/test_nvrtc.py | 6 ++++++ 2 files changed, 18 insertions(+) diff --git a/cuda_bindings/tests/test_cuda.py b/cuda_bindings/tests/test_cuda.py index 0137a4d5b..c58952d14 100644 --- a/cuda_bindings/tests/test_cuda.py +++ b/cuda_bindings/tests/test_cuda.py @@ -977,3 +977,15 @@ def test_all_CUresult_codes(max_code=1000): # The number will increase over time as new enums are added and support for # old CTKs is dropped, but it is not critical that this number is updated. assert num_good >= 76 # CTK 11.0.3_450.51.06 + + +def test_cuKernelGetName_failure(): + err, name = cuda.cuKernelGetName(0) + assert err == cuda.CUresult.CUDA_ERROR_INVALID_VALUE + assert name is None + + +def test_cuFuncGetName_failure(): + err, name = cuda.cuFuncGetName(0) + assert err == cuda.CUresult.CUDA_ERROR_INVALID_VALUE + assert name is None diff --git a/cuda_bindings/tests/test_nvrtc.py b/cuda_bindings/tests/test_nvrtc.py index 077484010..27dec963a 100644 --- a/cuda_bindings/tests/test_nvrtc.py +++ b/cuda_bindings/tests/test_nvrtc.py @@ -29,3 +29,9 @@ def test_nvrtcGetSupportedArchs(): err, supportedArchs = nvrtc.nvrtcGetSupportedArchs() ASSERT_DRV(err) assert len(supportedArchs) != 0 + + +def test_nvrtcGetLoweredName_failure(): + err, name = nvrtc.nvrtcGetLoweredName(0, b"I'm an elevated name!") + assert err == nvrtc.nvrtcResult.NVRTC_ERROR_INVALID_PROGRAM + assert name is None From a39720b5c8186d966666d0ecbc05e6bab5fb8f45 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Thu, 6 Mar 2025 16:24:57 -0800 Subject: [PATCH 6/8] Remove one new test because it's generating segfaults: https://github.com/NVIDIA/cuda-python/pull/497#issuecomment-2705233606 --- cuda_bindings/tests/test_nvrtc.py | 6 ------ 1 file changed, 6 deletions(-) diff --git a/cuda_bindings/tests/test_nvrtc.py b/cuda_bindings/tests/test_nvrtc.py index 27dec963a..077484010 100644 --- a/cuda_bindings/tests/test_nvrtc.py +++ b/cuda_bindings/tests/test_nvrtc.py @@ -29,9 +29,3 @@ def test_nvrtcGetSupportedArchs(): err, supportedArchs = nvrtc.nvrtcGetSupportedArchs() ASSERT_DRV(err) assert len(supportedArchs) != 0 - - -def test_nvrtcGetLoweredName_failure(): - err, name = nvrtc.nvrtcGetLoweredName(0, b"I'm an elevated name!") - assert err == nvrtc.nvrtcResult.NVRTC_ERROR_INVALID_PROGRAM - assert name is None From 6d45bda06d4998cd8fe14ea3625edd5132edb943 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Thu, 6 Mar 2025 19:33:24 -0800 Subject: [PATCH 7/8] Revert "Increase tolerance in test_timing() to avoid flaky tests." This reverts commit 9dd26307572107954347d390f3c750fb4ba0df3e. --- cuda_core/tests/test_event.py | 7 +------ 1 file changed, 1 insertion(+), 6 deletions(-) diff --git a/cuda_core/tests/test_event.py b/cuda_core/tests/test_event.py index 97c6b336b..384cf4586 100644 --- a/cuda_core/tests/test_event.py +++ b/cuda_core/tests/test_event.py @@ -33,12 +33,7 @@ def test_timing(init_cuda, enable_timing): if enable_timing: elapsed_time_ms = e2 - e1 assert isinstance(elapsed_time_ms, float) - # Using a generous tolerance, to avoid flaky tests: - # We only want to exercise the __sub__ method, this test is not meant - # to stress-test the CUDA driver or time.sleep(). - delay_ms = delay_seconds * 1000 - generous_tolerance = 20 - assert delay_ms <= elapsed_time_ms < delay_ms + generous_tolerance + assert delay_seconds * 1000 <= elapsed_time_ms < delay_seconds * 1000 + 2 # tolerance 2 ms else: with pytest.raises(RuntimeError) as e: elapsed_time_ms = e2 - e1 From ab6db2255189ac91de50713f5cbf3846054e2471 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Fri, 7 Mar 2025 08:24:29 -0800 Subject: [PATCH 8/8] test_all_CUresult_codes(): max_code = int(max(cuda.CUresult)) as suggested by at-leofang --- cuda_bindings/tests/test_cuda.py | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/cuda_bindings/tests/test_cuda.py b/cuda_bindings/tests/test_cuda.py index c58952d14..d89420cd9 100644 --- a/cuda_bindings/tests/test_cuda.py +++ b/cuda_bindings/tests/test_cuda.py @@ -950,9 +950,12 @@ def test_CUmemDecompressParams_st(): assert int(desc.dstActBytes) == 0 -def test_all_CUresult_codes(max_code=1000): +def test_all_CUresult_codes(): + max_code = int(max(cuda.CUresult)) + # Smoke test. CUDA_ERROR_UNKNOWN = 999, but intentionally using literal value. + assert max_code >= 999 num_good = 0 - for code in range(max_code + 1): + for code in range(max_code + 2): # One past max_code try: error = cuda.CUresult(code) except ValueError: @@ -973,7 +976,7 @@ def test_all_CUresult_codes(max_code=1000): err_desc, desc = cuda.cuGetErrorString(error) assert err_desc == cuda.CUresult.CUDA_ERROR_INVALID_VALUE assert desc is None - # Super-simple smoke check: Do we have at least some "good" codes? + # Smoke test: Do we have at least some "good" codes? # The number will increase over time as new enums are added and support for # old CTKs is dropped, but it is not critical that this number is updated. assert num_good >= 76 # CTK 11.0.3_450.51.06