[BUG]: In cuda.cooperative linking the same algorithm twice leads to multiple definitions

[brycelelbach]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Compile-time Error

Component

cuda.cooperative (Python)

Describe the bug

If you link the same instantiation of an algorithm twice in a kernel, you'll get multiple definition errors at JIT compilation time.

This is going to be a problem when people are using multiple libraries/components built with cuda.cooperative, especially when linking is implicit in the future.

We need to do some sort of de-duplication.

How to Reproduce

import cuda.cooperative.experimental as cudax
import numba
from pynvjitlink import patch
patch.patch_numba_linker(lto=True)

load0 = cudax.block.load(numba.int32, 32, 4, algorithm='striped')
load1 = cudax.block.load(numba.int32, 32, 4, algorithm='striped')

@numba.cuda.jit(link=load0.files+load1.files)
def kernel():
  pass

kernel[1, 32]()

Expected behavior

N/A

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response