From eba7d222d518e5a18c61f7275f46b648e59d56e4 Mon Sep 17 00:00:00 2001 From: ksimpson Date: Thu, 9 Jan 2025 16:13:03 -0800 Subject: [PATCH 1/8] add exception manager to handle options which aren't covered in all cuda CTK versions --- cuda_core/tests/test_linker.py | 21 ++++++++++++++++++++- 1 file changed, 20 insertions(+), 1 deletion(-) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 54cd8cf4..0af2d8bd 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -2,6 +2,8 @@ # # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE +from contextlib import contextmanager + import pytest from cuda.core.experimental import Linker, LinkerOptions, Program, _linker @@ -18,6 +20,8 @@ device_function_c = "__device__ int C(int a, int b) { return a + b; }" culink_backend = _linker._decide_nvjitlink_or_driver() +if not culink_backend: + from cuda.bindings import nvjitlink @pytest.fixture(scope="function") @@ -40,6 +44,19 @@ def compile_ltoir_functions(init_cuda): return object_code_a_ltoir, object_code_b_ltoir, object_code_c_ltoir +@contextmanager +def skip_version_specific_linker_options(): + if culink_backend: + return + try: + yield + except nvjitlink.nvJitLinkError as e: + if e.status == nvjitlink.Result.ERROR_UNRECOGNIZED_OPTION: + pytest.skip("current nvjitlink version does not support the option provided") + except Exception as e: + raise e + + culink_options = [ LinkerOptions(arch=ARCH, verbose=True), LinkerOptions(arch=ARCH, max_register_count=32), @@ -72,7 +89,9 @@ def compile_ltoir_functions(init_cuda): ], ) def test_linker_init(compile_ptx_functions, options): - linker = Linker(*compile_ptx_functions, options=options) + with skip_version_specific_linker_options(): + linker = Linker(*compile_ptx_functions, options=options) + object_code = linker.link("cubin") assert isinstance(object_code, ObjectCode) From 6e5114e249cafb3d45439125b9198530337900f5 Mon Sep 17 00:00:00 2001 From: Keenan Simpson Date: Wed, 15 Jan 2025 09:28:55 -0800 Subject: [PATCH 2/8] Update cuda_core/tests/test_linker.py Co-authored-by: Leo Fang --- cuda_core/tests/test_linker.py | 2 -- 1 file changed, 2 deletions(-) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 0af2d8bd..0058e349 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -53,8 +53,6 @@ def skip_version_specific_linker_options(): except nvjitlink.nvJitLinkError as e: if e.status == nvjitlink.Result.ERROR_UNRECOGNIZED_OPTION: pytest.skip("current nvjitlink version does not support the option provided") - except Exception as e: - raise e culink_options = [ From 66b03fc2b31bc455b71448ae51dcb3a5c393c92f Mon Sep 17 00:00:00 2001 From: ksimpson Date: Wed, 15 Jan 2025 10:25:33 -0800 Subject: [PATCH 3/8] address comments --- cuda_core/cuda/core/experimental/_linker.py | 1 + cuda_core/tests/test_linker.py | 34 ++++++++++++--------- 2 files changed, 20 insertions(+), 15 deletions(-) diff --git a/cuda_core/cuda/core/experimental/_linker.py b/cuda_core/cuda/core/experimental/_linker.py index 6e36a2a5..6bad7d49 100644 --- a/cuda_core/cuda/core/experimental/_linker.py +++ b/cuda_core/cuda/core/experimental/_linker.py @@ -32,6 +32,7 @@ def _decide_nvjitlink_or_driver(): _driver_ver = handle_return(cuda.cuDriverGetVersion()) _driver_ver = (_driver_ver // 1000, (_driver_ver % 1000) // 10) try: + raise ImportError from cuda.bindings import nvjitlink as _nvjitlink from cuda.bindings._internal import nvjitlink as inner_nvjitlink except ImportError: diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 0058e349..e3120137 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -2,7 +2,7 @@ # # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE -from contextlib import contextmanager +from contextlib import contextmanager, nullcontext import pytest @@ -20,9 +20,24 @@ device_function_c = "__device__ int C(int a, int b) { return a + b; }" culink_backend = _linker._decide_nvjitlink_or_driver() +skip_options = nullcontext if not culink_backend: from cuda.bindings import nvjitlink + @contextmanager + def skip_version_specific_linker_options(): + if culink_backend: + return + try: + yield + except nvjitlink.nvJitLinkError as e: + if e.status == nvjitlink.Result.ERROR_UNRECOGNIZED_OPTION: + pytest.skip("current nvjitlink version does not support the option provided") + else: + raise + + skip_options = skip_version_specific_linker_options + @pytest.fixture(scope="function") def compile_ptx_functions(init_cuda): @@ -44,17 +59,6 @@ def compile_ltoir_functions(init_cuda): return object_code_a_ltoir, object_code_b_ltoir, object_code_c_ltoir -@contextmanager -def skip_version_specific_linker_options(): - if culink_backend: - return - try: - yield - except nvjitlink.nvJitLinkError as e: - if e.status == nvjitlink.Result.ERROR_UNRECOGNIZED_OPTION: - pytest.skip("current nvjitlink version does not support the option provided") - - culink_options = [ LinkerOptions(arch=ARCH, verbose=True), LinkerOptions(arch=ARCH, max_register_count=32), @@ -87,11 +91,11 @@ def skip_version_specific_linker_options(): ], ) def test_linker_init(compile_ptx_functions, options): - with skip_version_specific_linker_options(): + with skip_options(): linker = Linker(*compile_ptx_functions, options=options) - object_code = linker.link("cubin") - assert isinstance(object_code, ObjectCode) + object_code = linker.link("cubin") + assert isinstance(object_code, ObjectCode) def test_linker_init_invalid_arch(): From df897e8e31ebf377be8fe051103913cd8e45ba3c Mon Sep 17 00:00:00 2001 From: ksimpson Date: Fri, 17 Jan 2025 10:49:40 -0800 Subject: [PATCH 4/8] switch to hardcode --- cuda_core/cuda/core/experimental/_linker.py | 1 - cuda_core/tests/test_linker.py | 76 +++++++++------------ 2 files changed, 31 insertions(+), 46 deletions(-) diff --git a/cuda_core/cuda/core/experimental/_linker.py b/cuda_core/cuda/core/experimental/_linker.py index f6f05aa5..6c2ca4d4 100644 --- a/cuda_core/cuda/core/experimental/_linker.py +++ b/cuda_core/cuda/core/experimental/_linker.py @@ -31,7 +31,6 @@ def _decide_nvjitlink_or_driver(): _driver_ver = handle_return(driver.cuDriverGetVersion()) _driver_ver = (_driver_ver // 1000, (_driver_ver % 1000) // 10) try: - raise ImportError from cuda.bindings import nvjitlink as _nvjitlink from cuda.bindings._internal import nvjitlink as inner_nvjitlink except ImportError: diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index e3120137..429e4e5f 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -2,7 +2,6 @@ # # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE -from contextlib import contextmanager, nullcontext import pytest @@ -20,23 +19,6 @@ device_function_c = "__device__ int C(int a, int b) { return a + b; }" culink_backend = _linker._decide_nvjitlink_or_driver() -skip_options = nullcontext -if not culink_backend: - from cuda.bindings import nvjitlink - - @contextmanager - def skip_version_specific_linker_options(): - if culink_backend: - return - try: - yield - except nvjitlink.nvJitLinkError as e: - if e.status == nvjitlink.Result.ERROR_UNRECOGNIZED_OPTION: - pytest.skip("current nvjitlink version does not support the option provided") - else: - raise - - skip_options = skip_version_specific_linker_options @pytest.fixture(scope="function") @@ -65,37 +47,41 @@ def compile_ltoir_functions(init_cuda): LinkerOptions(arch=ARCH, optimization_level=3), LinkerOptions(arch=ARCH, debug=True), LinkerOptions(arch=ARCH, lineinfo=True), - LinkerOptions(arch=ARCH, no_cache=True), ] +nvjitlink_base_options = [ + LinkerOptions(arch=ARCH, time=True), + LinkerOptions(arch=ARCH, ftz=True), + LinkerOptions(arch=ARCH, prec_div=True), + LinkerOptions(arch=ARCH, prec_sqrt=True), + LinkerOptions(arch=ARCH, fma=True), + LinkerOptions(arch=ARCH, kernels_used=["A"]), + LinkerOptions(arch=ARCH, kernels_used=["C", "B"]), + LinkerOptions(arch=ARCH, variables_used=["var1"]), + LinkerOptions(arch=ARCH, variables_used=["var1", "var2"]), + LinkerOptions(arch=ARCH, optimize_unused_variables=True), + LinkerOptions(arch=ARCH, xptxas=["-v"]), + LinkerOptions(arch=ARCH, split_compile=0), + LinkerOptions(arch=ARCH, split_compile_extended=1), +] -@pytest.mark.parametrize( - "options", - culink_options - if culink_backend - else culink_options - + [ - LinkerOptions(arch=ARCH, time=True), - LinkerOptions(arch=ARCH, ftz=True), - LinkerOptions(arch=ARCH, prec_div=True), - LinkerOptions(arch=ARCH, prec_sqrt=True), - LinkerOptions(arch=ARCH, fma=True), - LinkerOptions(arch=ARCH, kernels_used=["A"]), - LinkerOptions(arch=ARCH, kernels_used=["C", "B"]), - LinkerOptions(arch=ARCH, variables_used=["var1"]), - LinkerOptions(arch=ARCH, variables_used=["var1", "var2"]), - LinkerOptions(arch=ARCH, optimize_unused_variables=True), - LinkerOptions(arch=ARCH, xptxas=["-v"]), - LinkerOptions(arch=ARCH, split_compile=0), - LinkerOptions(arch=ARCH, split_compile_extended=1), - ], -) -def test_linker_init(compile_ptx_functions, options): - with skip_options(): - linker = Linker(*compile_ptx_functions, options=options) +nvjitlink_125_options = [LinkerOptions(arch=ARCH, no_cache=True)] - object_code = linker.link("cubin") - assert isinstance(object_code, ObjectCode) +options = culink_options +if not culink_backend: + from cuda.bindings import nvjitlink + + version = nvjitlink.version() + options += nvjitlink_base_options + if version >= (12, 5): + options += nvjitlink_125_options + + +@pytest.mark.parametrize("options", options) +def test_linker_init(compile_ptx_functions, options): + linker = Linker(*compile_ptx_functions, options=options) + object_code = linker.link("cubin") + assert isinstance(object_code, ObjectCode) def test_linker_init_invalid_arch(): From 94af829286fcbbbc1b7959857f70542c0fdd60ae Mon Sep 17 00:00:00 2001 From: ksimpson Date: Tue, 21 Jan 2025 14:39:53 -0800 Subject: [PATCH 5/8] construct the options class within if else logic --- cuda_core/tests/test_linker.py | 61 ++++++++++++++++------------------ 1 file changed, 29 insertions(+), 32 deletions(-) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 0b28975d..0d04a9de 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -47,42 +47,39 @@ def compile_ltoir_functions(init_cuda): return object_code_a_ltoir, object_code_b_ltoir, object_code_c_ltoir -culink_options = [ - LinkerOptions(), - LinkerOptions(arch=ARCH, verbose=True), - LinkerOptions(arch=ARCH, max_register_count=32), - LinkerOptions(arch=ARCH, optimization_level=3), - LinkerOptions(arch=ARCH, debug=True), - LinkerOptions(arch=ARCH, lineinfo=True), - LinkerOptions(arch=ARCH, no_cache=True), -] - -nvjitlink_base_options = [ - LinkerOptions(arch=ARCH, time=True), - LinkerOptions(arch=ARCH, ftz=True), - LinkerOptions(arch=ARCH, prec_div=True), - LinkerOptions(arch=ARCH, prec_sqrt=True), - LinkerOptions(arch=ARCH, fma=True), - LinkerOptions(arch=ARCH, kernels_used=["A"]), - LinkerOptions(arch=ARCH, kernels_used=["C", "B"]), - LinkerOptions(arch=ARCH, variables_used=["var1"]), - LinkerOptions(arch=ARCH, variables_used=["var1", "var2"]), - LinkerOptions(arch=ARCH, optimize_unused_variables=True), - LinkerOptions(arch=ARCH, xptxas=["-v"]), - LinkerOptions(arch=ARCH, split_compile=0), - LinkerOptions(arch=ARCH, split_compile_extended=1), -] - -nvjitlink_125_options = [LinkerOptions(arch=ARCH, no_cache=True)] - -options = culink_options -if not culink_backend: +options = [] + +if culink_backend: + options += [ + LinkerOptions(), + LinkerOptions(arch=ARCH, verbose=True), + LinkerOptions(arch=ARCH, max_register_count=32), + LinkerOptions(arch=ARCH, optimization_level=3), + LinkerOptions(arch=ARCH, debug=True), + LinkerOptions(arch=ARCH, lineinfo=True), + LinkerOptions(arch=ARCH, no_cache=True), + ] +else: from cuda.bindings import nvjitlink version = nvjitlink.version() - options += nvjitlink_base_options + options += [ + LinkerOptions(arch=ARCH, time=True), + LinkerOptions(arch=ARCH, ftz=True), + LinkerOptions(arch=ARCH, prec_div=True), + LinkerOptions(arch=ARCH, prec_sqrt=True), + LinkerOptions(arch=ARCH, fma=True), + LinkerOptions(arch=ARCH, kernels_used=["A"]), + LinkerOptions(arch=ARCH, kernels_used=["C", "B"]), + LinkerOptions(arch=ARCH, variables_used=["var1"]), + LinkerOptions(arch=ARCH, variables_used=["var1", "var2"]), + LinkerOptions(arch=ARCH, optimize_unused_variables=True), + LinkerOptions(arch=ARCH, xptxas=["-v"]), + LinkerOptions(arch=ARCH, split_compile=0), + LinkerOptions(arch=ARCH, split_compile_extended=1), + ] if version >= (12, 5): - options += nvjitlink_125_options + options += [LinkerOptions(arch=ARCH, no_cache=True)] @pytest.mark.parametrize("options", options) From 558658798941ce2edcc59902785267c31df60e96 Mon Sep 17 00:00:00 2001 From: ksimpson Date: Thu, 23 Jan 2025 09:21:29 -0800 Subject: [PATCH 6/8] fix mistake --- cuda_core/tests/test_linker.py | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 0d04a9de..fd537ae5 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -50,19 +50,17 @@ def compile_ltoir_functions(init_cuda): options = [] if culink_backend: - options += [ + options = [ LinkerOptions(), LinkerOptions(arch=ARCH, verbose=True), LinkerOptions(arch=ARCH, max_register_count=32), LinkerOptions(arch=ARCH, optimization_level=3), LinkerOptions(arch=ARCH, debug=True), LinkerOptions(arch=ARCH, lineinfo=True), - LinkerOptions(arch=ARCH, no_cache=True), ] else: from cuda.bindings import nvjitlink - version = nvjitlink.version() options += [ LinkerOptions(arch=ARCH, time=True), LinkerOptions(arch=ARCH, ftz=True), @@ -78,6 +76,7 @@ def compile_ltoir_functions(init_cuda): LinkerOptions(arch=ARCH, split_compile=0), LinkerOptions(arch=ARCH, split_compile_extended=1), ] + version = nvjitlink.version() if version >= (12, 5): options += [LinkerOptions(arch=ARCH, no_cache=True)] From b3e254df21e11821fcb6f81269a14730c090eb72 Mon Sep 17 00:00:00 2001 From: ksimpson Date: Thu, 23 Jan 2025 09:22:31 -0800 Subject: [PATCH 7/8] fix mistake --- cuda_core/tests/test_linker.py | 21 +++++++++------------ 1 file changed, 9 insertions(+), 12 deletions(-) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index fd537ae5..1ee254e4 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -47,18 +47,15 @@ def compile_ltoir_functions(init_cuda): return object_code_a_ltoir, object_code_b_ltoir, object_code_c_ltoir -options = [] - -if culink_backend: - options = [ - LinkerOptions(), - LinkerOptions(arch=ARCH, verbose=True), - LinkerOptions(arch=ARCH, max_register_count=32), - LinkerOptions(arch=ARCH, optimization_level=3), - LinkerOptions(arch=ARCH, debug=True), - LinkerOptions(arch=ARCH, lineinfo=True), - ] -else: +options = [ + LinkerOptions(), + LinkerOptions(arch=ARCH, verbose=True), + LinkerOptions(arch=ARCH, max_register_count=32), + LinkerOptions(arch=ARCH, optimization_level=3), + LinkerOptions(arch=ARCH, debug=True), + LinkerOptions(arch=ARCH, lineinfo=True), +] +if not culink_backend: from cuda.bindings import nvjitlink options += [ From 38801fe2e35db6633985a19031ea5c32569ef68e Mon Sep 17 00:00:00 2001 From: ksimpson Date: Tue, 28 Jan 2025 10:21:38 -0800 Subject: [PATCH 8/8] minor tweak --- cuda_core/tests/test_linker.py | 3 --- 1 file changed, 3 deletions(-) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 1ee254e4..288e2ce4 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -2,7 +2,6 @@ # # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE - import pytest from cuda.core.experimental import Device, Linker, LinkerOptions, Program, ProgramOptions, _linker @@ -56,8 +55,6 @@ def compile_ltoir_functions(init_cuda): LinkerOptions(arch=ARCH, lineinfo=True), ] if not culink_backend: - from cuda.bindings import nvjitlink - options += [ LinkerOptions(arch=ARCH, time=True), LinkerOptions(arch=ARCH, ftz=True),