NVIDIA / NVIDIA/cccl

cuda.cooperative: Eliminate `link=` argument to `@cuda.jit` decorator when it becomes possible

Open
#3,786 0 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.