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

[FEA] Support automatic generation of link list for external code #67

Open
gmarkall opened this issue Nov 6, 2024 · 4 comments
Open
Labels
feature request New feature or request

Comments

@gmarkall
Copy link
Collaborator

gmarkall commented Nov 6, 2024

Numba-cuda extensions (e.g. nvmath-python) are frequently leaning on CUDA C++ implementations to support the core of their functionality.

One current UX limitation is that the kernel author is required to add the list of files and/or code to link with a kernel as a keyword argument to the @cuda.jit decorator, for example:

@cuda.jit(link=FFT.files)
def f(data):
    ...

from cufftdx_simple_fft_block.py

The FFT object supplies the files, and is created like:

FFT = fft(fft_type='c2c', 
          size=128,
          precision=np.float32,
          direction='forward',
          elements_per_thread=8,
          ffts_per_block=2,
          execution='Block',
          compiler='numba')

and is called inside the kernel as:

FFT(thread_data, shared_mem)

Rather than the user being required to link FFT.files, Numba should provide a mechanism to obtain and link the list of files / code (LTO-IR, PTX, CUDA C/C++ source, or binaries / objects etc.) at the point of compilation and linking from the FFT object (or any implementation of a method, property, object, etc. backed by an extension). It is expected that the implementation (of FFT, in this example) may generate code (e.g. LTO-IR) at this point just prior to returning it back to Numba.

@gmarkall gmarkall added the feature request New feature or request label Nov 6, 2024
@leofang
Copy link
Member

leofang commented Nov 8, 2024

Perhaps ask all device functions to implement a method attribute, say, __numba_cuda_link__ that returns a list of files, if they want numba-cuda to handle the linking?

@gmarkall
Copy link
Collaborator Author

gmarkall commented Nov 8, 2024

Continuing with the __numba_cuda_link__ idea, I think it might need to be a method that can accept a signature, so that it can return the appropriate files for the given signature.

@leofang
Copy link
Member

leofang commented Nov 8, 2024

How should the kernel author pass function arguments at the call site if it is a method not attribute?

@gmarkall
Copy link
Collaborator Author

gmarkall commented Nov 8, 2024

In the example above, the kernel author wrote:

FFT(thread_data, shared_mem)

assuming the Numba types of these are float32[:] and float32[::1] (for the sake of argument, they could be any Numba type really) I'd expect during compilation time that Numba would be doing the equivalent of calling

ltoir = FFT.__numba_cuda_link__(float32[:], float32[::1])

where ltoir is then an LTOIR linkable code object, i.e. an instance of

class LTOIR(LinkableCode):
"""An LTOIR file in memory"""
kind = "ltoir"
default_name = "<unnamed-ltoir>"

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
feature request New feature or request
Projects
None yet
Development

No branches or pull requests

2 participants