diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index b6dd2a350..613afdb28 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) @@ -21,6 +22,12 @@ if not is_culink_backend: from cuda.bindings import nvjitlink + nvJitLinkError = nvjitlink.nvJitLinkError +else: + + class nvJitLinkError(Exception): + pass + @pytest.fixture(scope="function") def compile_ptx_functions(init_cuda): @@ -135,10 +142,23 @@ 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") - log = linker.get_error_log() - assert isinstance(log, str) + + 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:]), options=options) + try: + linker.link("cubin") + + except (nvJitLinkError, CUDAError): + log = linker.get_error_log() + assert isinstance(log, str) + # TODO when 4902246 is addressed, we can update this to cover nvjitlink as well + if is_culink_backend: + assert log.rstrip("\x00") == "error : Undefined reference to '_Z1Zv' in 'None_ptx'" def test_linker_get_info_log(compile_ptx_functions):