NVIDIA / NVIDIA/cccl

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

Open
#3,909 0 comments 0 reactions 0 assignees View on GitHub
cuda.coop
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.