[BUG]: In cuda.cooperative linking the same algorithm twice leads to multiple definitions
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
### Is this a duplicate?
- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### 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_
Contributor guide
Assessment
This issue has not been assessed yet.