Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

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

Open
shwina opened this issue Feb 12, 2025 · 0 comments

Comments

@shwina
Copy link
Contributor

shwina commented Feb 12, 2025

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:

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). When that functionality is available, we should switch to using it.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
Status: Todo
Development

No branches or pull requests

1 participant