cuda.cooperative: Eliminate `link=` argument to `@cuda.jit` decorator when it becomes possible
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
Description
Currently, `cuda.cooperative` relies on passing any additional code we wish to link with the user-defined kernel via the `link=` argument to `@cuda.jit`:
```python
warp_exclusive_sum = cudax.warp.exclusive_sum(numba.int32)
@cuda.jit(link=warp_exclusive_sum.files) # <<<<<<<<---- awkward
def kernel(data):
data[cuda.threadIdx.x] = warp_exclusive_sum(data[cuda.threadIdx.x])
```
Ideally, the user wouldn't have to do this. There's work underway within Numba CUDA that would avoid this, instead letting us declare device code beforehand, and Numba would link it whenever it's used ([example](https://github.com/gmarkall/numba-cuda/blob/auto-link/link_example.py)). When that functionality is available, we should switch to using it.
Contributor guide
Research direction
Start by reviewing cuda.cooperative and the linked Numba CUDA example, link_example.py, to understand when device code can be declared and linked automatically. Confirm that the upstream functionality is available, then trace the current use of @cuda.jit(link=...) in the cooperative implementation. Done means users no longer need to pass the generated files explicitly and the relevant cooperative tests pass.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- python
- Domain
- hpc
- Issue type
- Refactor
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 35/100