From 586551ac55f0eda52fe211fb969f314a95e8f91e Mon Sep 17 00:00:00 2001 From: ksimpson Date: Tue, 28 Jan 2025 15:08:02 -0800 Subject: [PATCH 1/7] skip nvjitlink backend for now --- cuda_core/cuda/core/experimental/_linker.py | 2 ++ cuda_core/tests/test_linker.py | 17 +++++++++++++++-- 2 files changed, 17 insertions(+), 2 deletions(-) diff --git a/cuda_core/cuda/core/experimental/_linker.py b/cuda_core/cuda/core/experimental/_linker.py index b5a6b675c..33a72f879 100644 --- a/cuda_core/cuda/core/experimental/_linker.py +++ b/cuda_core/cuda/core/experimental/_linker.py @@ -436,7 +436,9 @@ def link(self, target_type) -> ObjectCode: raise ValueError(f"Unsupported target type: {target_type}") with _exception_manager(self): if _nvjitlink: + # print('completing') _nvjitlink.complete(self._mnff.handle) + # print('done completing') if target_type == "cubin": get_size = _nvjitlink.get_linked_cubin_size get_code = _nvjitlink.get_linked_cubin diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index b81c1654a..75427927b 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -6,6 +6,7 @@ from cuda.core.experimental import Device, Linker, LinkerOptions, Program, ProgramOptions, _linker from cuda.core.experimental._module import ObjectCode +from cuda.core.experimental._utils import CUDAError ARCH = "sm_" + "".join(f"{i}" for i in Device().compute_capability) @@ -124,9 +125,21 @@ def test_linker_link_invalid_target_type(compile_ptx_functions): def test_linker_get_error_log(compile_ptx_functions): options = LinkerOptions(arch=ARCH) - linker = Linker(*compile_ptx_functions, options=options) - linker.link("cubin") + + replacement_kernel = """ +extern __device__ int Z(); +extern __device__ int C(int a, int b); +__global__ void A() { int result = C(Z(), 1);} +""" + dummy_program = Program(replacement_kernel, "c++", ProgramOptions(relocatable_device_code=True)).compile("ptx") + linker = Linker(dummy_program, compile_ptx_functions[1], compile_ptx_functions[2], options=options) + with pytest.raises((CUDAError, nvjitlink.nvJitLinkError)): + linker.link("cubin") + log = linker.get_error_log() + # TODO when 4902246 is addressed, we can update this to cover nvjitlink as well + if culink_backend: + assert log.rstrip("\x00") == "error : Undefined reference to '_Z1Zv' in 'None_ptx'" assert isinstance(log, str) From aa45a6eaf13db36ef5af9ea7172873cca1c1a811 Mon Sep 17 00:00:00 2001 From: ksimpson Date: Tue, 28 Jan 2025 15:10:12 -0800 Subject: [PATCH 2/7] remove prints --- cuda_core/cuda/core/experimental/_linker.py | 2 -- cuda_core/tests/test_linker.py | 1 + 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/cuda_core/cuda/core/experimental/_linker.py b/cuda_core/cuda/core/experimental/_linker.py index 33a72f879..b5a6b675c 100644 --- a/cuda_core/cuda/core/experimental/_linker.py +++ b/cuda_core/cuda/core/experimental/_linker.py @@ -436,9 +436,7 @@ def link(self, target_type) -> ObjectCode: raise ValueError(f"Unsupported target type: {target_type}") with _exception_manager(self): if _nvjitlink: - # print('completing') _nvjitlink.complete(self._mnff.handle) - # print('done completing') if target_type == "cubin": get_size = _nvjitlink.get_linked_cubin_size get_code = _nvjitlink.get_linked_cubin diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 75427927b..1f85e8cf5 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -148,4 +148,5 @@ def test_linker_get_info_log(compile_ptx_functions): linker = Linker(*compile_ptx_functions, options=options) linker.link("cubin") log = linker.get_info_log() + print(log) assert isinstance(log, str) From 066804b40410da61cd805bfab5ecb3b41eb9791f Mon Sep 17 00:00:00 2001 From: ksimpson Date: Wed, 5 Feb 2025 16:19:27 -0800 Subject: [PATCH 3/7] fix test --- cuda_core/tests/test_linker.py | 13 ++++++------- 1 file changed, 6 insertions(+), 7 deletions(-) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 1f85e8cf5..4af0ce65e 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -6,7 +6,6 @@ from cuda.core.experimental import Device, Linker, LinkerOptions, Program, ProgramOptions, _linker from cuda.core.experimental._module import ObjectCode -from cuda.core.experimental._utils import CUDAError ARCH = "sm_" + "".join(f"{i}" for i in Device().compute_capability) @@ -133,14 +132,14 @@ def test_linker_get_error_log(compile_ptx_functions): """ dummy_program = Program(replacement_kernel, "c++", ProgramOptions(relocatable_device_code=True)).compile("ptx") linker = Linker(dummy_program, compile_ptx_functions[1], compile_ptx_functions[2], options=options) - with pytest.raises((CUDAError, nvjitlink.nvJitLinkError)): + try: linker.link("cubin") - log = linker.get_error_log() - # TODO when 4902246 is addressed, we can update this to cover nvjitlink as well - if culink_backend: - assert log.rstrip("\x00") == "error : Undefined reference to '_Z1Zv' in 'None_ptx'" - assert isinstance(log, str) + except: # noqa E722 + log = linker.get_error_log() + # TODO when 4902246 is addressed, we can update this to cover nvjitlink as well + if culink_backend: + assert log.rstrip("\x00") == "error : Undefined reference to '_Z1Zv' in 'None_ptx'" def test_linker_get_info_log(compile_ptx_functions): From 651b37740b248beb9f6dee802669044fbe9eb739 Mon Sep 17 00:00:00 2001 From: Keenan Simpson Date: Wed, 19 Feb 2025 10:05:38 -0800 Subject: [PATCH 4/7] Update cuda_core/tests/test_linker.py Co-authored-by: Leo Fang --- cuda_core/tests/test_linker.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index a42e74c68..030c6f3a8 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -129,7 +129,7 @@ def test_linker_get_error_log(compile_ptx_functions): __global__ void A() { int result = C(Z(), 1);} """ dummy_program = Program(replacement_kernel, "c++", ProgramOptions(relocatable_device_code=True)).compile("ptx") - linker = Linker(dummy_program, compile_ptx_functions[1], compile_ptx_functions[2], options=options) + linker = Linker(dummy_program, *(compile_ptx_functions[1:]), options=options) try: linker.link("cubin") From 642539ac5110f09c7fe290a4e14fb7c652cc1c64 Mon Sep 17 00:00:00 2001 From: Keenan Simpson Date: Wed, 19 Feb 2025 10:05:45 -0800 Subject: [PATCH 5/7] Update cuda_core/tests/test_linker.py Co-authored-by: Leo Fang --- cuda_core/tests/test_linker.py | 1 - 1 file changed, 1 deletion(-) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 030c6f3a8..485f92457 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -145,5 +145,4 @@ def test_linker_get_info_log(compile_ptx_functions): linker = Linker(*compile_ptx_functions, options=options) linker.link("cubin") log = linker.get_info_log() - print(log) assert isinstance(log, str) From 949897f2dbf65d06623bae38f16e877560b0875b Mon Sep 17 00:00:00 2001 From: Keenan Simpson Date: Wed, 19 Feb 2025 10:05:58 -0800 Subject: [PATCH 6/7] Update cuda_core/tests/test_linker.py Co-authored-by: Leo Fang --- cuda_core/tests/test_linker.py | 1 + 1 file changed, 1 insertion(+) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 485f92457..a8633b7ed 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -135,6 +135,7 @@ def test_linker_get_error_log(compile_ptx_functions): except: # noqa E722 log = linker.get_error_log() + assert isinstance(log, str) # TODO when 4902246 is addressed, we can update this to cover nvjitlink as well if culink_backend: assert log.rstrip("\x00") == "error : Undefined reference to '_Z1Zv' in 'None_ptx'" From 8e45f16dbf38177dc68705a42eee37a9d14a542f Mon Sep 17 00:00:00 2001 From: ksimpson Date: Wed, 19 Feb 2025 10:14:00 -0800 Subject: [PATCH 7/7] remove the noqa --- cuda_core/tests/test_linker.py | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index a42e74c68..4898b4f57 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -6,6 +6,7 @@ from cuda.core.experimental import Device, Linker, LinkerOptions, Program, ProgramOptions, _linker from cuda.core.experimental._module import ObjectCode +from cuda.core.experimental._utils import CUDAError ARCH = "sm_" + "".join(f"{i}" for i in Device().compute_capability) @@ -18,9 +19,17 @@ device_function_c = "__device__ int C(int a, int b) { return a + b; }" culink_backend = _linker._decide_nvjitlink_or_driver() + + +class nvJitLinkError_(Exception): + pass + + if not culink_backend: from cuda.bindings import nvjitlink + nvJitLinkError_ = nvjitlink.nvJitLinkError + @pytest.fixture(scope="function") def compile_ptx_functions(init_cuda): @@ -133,7 +142,7 @@ def test_linker_get_error_log(compile_ptx_functions): try: linker.link("cubin") - except: # noqa E722 + except (nvJitLinkError_, CUDAError): log = linker.get_error_log() # TODO when 4902246 is addressed, we can update this to cover nvjitlink as well if culink_backend: