Skip to content

[BUG] Linker error when lto and debug are both set to True #194

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
isVoid opened this issue Apr 9, 2025 · 0 comments
Open

[BUG] Linker error when lto and debug are both set to True #194

isVoid opened this issue Apr 9, 2025 · 0 comments
Labels
bug Something isn't working

Comments

@isVoid
Copy link
Collaborator

isVoid commented Apr 9, 2025

Describe the bug
When kernel invoked with jit(lto=True, debug=True, there's a linker error.

Steps/Code to reproduce bug

from numba import cuda, config

config.CUDA_ENABLE_PYNVJITLINK=True

add = cuda.declare_device(
    "add",
    "int32(int32, int32)",
    link=["./add.cu"]
)

@cuda.jit(lto=True, debug=True)
def kernel(a, b, c):
    i = cuda.grid(1)
    c[i] = add(a[i], b[i])
    assert c[i] > 0

a = cuda.to_device([1, 2, 3, 4, 5])
b = cuda.to_device([6, 7, 8, 9, 10])
c = cuda.device_array_like(a)
kernel[1, 5](a, b, c)
print(c.copy_to_host())
extern "C"
int __device__ add(int *ret, int *left, int *right)
{
    *ret = *left + *right;
    return 0;
}
Linker Error
Traceback (most recent call last):
  File "/opt/venv/lib/python3.12/site-packages/pynvjitlink/api.py", line 84, in get_linked_cubin
    _nvjitlinklib.complete(self.handle)
RuntimeError: NVJITLINK_ERROR_NVVM_COMPILE error when calling nvJitLinkComplete

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/workspace/numba-cuda/numba_cuda/numba/cuda/cudadrv/driver.py", line 3195, in complete
    return self._linker.get_linked_cubin()
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/pynvjitlink/api.py", line 89, in get_linked_cubin
    raise NvJitLinkError(f"{e}\n{self.error_log}")
pynvjitlink.api.NvJitLinkError: NVJITLINK_ERROR_NVVM_COMPILE error when calling nvJitLinkComplete
add.ltoir: link error: error: linking module flags 'Debug Info Version': IDs have conflicting behaviors

ERROR 9 in nvvmCompileProgram callback


The above exception was the direct cause of the following exception:

Traceback (most recent call last):
  File "/workspace/numba-cuda-tests/lto_debug_fault/test.py", line 19, in <module>
    kernel[1, 5](a, b, c)
  File "/workspace/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 674, in __call__
    return self.dispatcher.call(args, self.griddim, self.blockdim,
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/workspace/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 816, in call
    kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/workspace/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 824, in _compile_for_args
    return self.compile(tuple(argtypes))
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/workspace/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 1075, in compile
    kernel.bind()
  File "/workspace/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 317, in bind
    cufunc = self._codelibrary.get_cufunc()
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/workspace/numba-cuda/numba_cuda/numba/cuda/codegen.py", line 255, in get_cufunc
    cubin = self.get_cubin(cc=device.compute_capability)
            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/workspace/numba-cuda/numba_cuda/numba/cuda/codegen.py", line 235, in get_cubin
    cubin = linker.complete()
            ^^^^^^^^^^^^^^^^^
  File "/workspace/numba-cuda/numba_cuda/numba/cuda/cudadrv/driver.py", line 3197, in complete
    raise LinkerError from e
numba.cuda.cudadrv.driver.LinkerError

Expected behavior
No error should surface.

Additional context
I ran into this error when trying out assert statements inside kernel, to learned that it is only activated when debug is set to true.

@isVoid isVoid added the bug Something isn't working label Apr 9, 2025
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working
Projects
None yet
Development

No branches or pull requests

1 participant