Soft link NVRTC for cupy_backends.cuda.libs.nvrtc

[kmaehashi]

Part of #7620.

This PR changes the cupy_backends.cuda.libs.nvrtc module to load NVRTC symbols dynamically at runtime, instead of build-time linking.

Pros:

• Make the module buildable & importable without NVRTC installed. (note: Jitify module still links to NVRTC, so this PR is not enough to run import cupy without NVRTC)
• No longer need to maintain headers in cupy_backend.

Cons:

• When calling NVRTC APIs, maintainers are responsible for passing the correct number/type of arguments, as the compiler can no longer detect such errors. Resolved by adding ctypedef for each API.

[emcastillo]

This seems to be affecting rocm build somehow?

[emcastillo]

/test mini

[kmaehashi]

Will resume after #7647

[leofang]

Q: This might change the design significantly, but I wonder if it's possible to import the function signatures as is, and reinterpret the fetched symbol using the function signature? That is, instead of

     cdef F_t nvrtcCreateProgram = <F_t>_L.get('CreateProgram')

we do

    ctypedef int (*createProgramFn)(  # naming convention TBD
             Program* prog, const char* src, const char* name, int numHeaders,
             const char** headers, const char** includeNames)
    cdef createProgramFn nvrtcCreateProgram = < createProgramFn>_L.get('CreateProgram')

Basically, we'd be trading this pro

No longer need to maintain headers in cupy_backend.

with this con

When calling NVRTC APIs, maintainers are responsible for passing the correct number/type of arguments, as the compiler can no longer detect such errors.

The main blocker for this would be the maintainability of the Cython layer. We still have to keep the function signatures...

[kmaehashi]

@leofang This is the part I was still wondering... "maintainers are responsible for passing the correct number/type of arguments" might be a bit dangerous as this immediately causes SegV on failure.

I'm now thinking that maintaining ctypedefs could be easier if we stop the convention of renaming typedefs (ctypedef nvrtcProgram Program). In that case, we can (almost) copy-paste the arguments from NVRTC header:

nvrtcResult nvrtcCreateProgram(nvrtcProgram *prog,
                               const char *src,
                               const char *name,
                               int numHeaders,
                               const char * const *headers,
                               const char * const *includeNames);

[leofang]

Sorry for some reason I missed your reply, Kenichi 😅

if we stop the convention of renaming typedefs

Yeah, sounds reasonable to me. I think at some point my hope was that all these are automated (say, #4395/#4427/#4606/#5876), and we can pick whatever style we like, but if it is no less work than this route, allowing easier copy/paste sounds like the right thing to do indeed.

[kmaehashi]

E211 added to avoid Whitespace before '(' error:

ctypedef nvrtcResult (*F_nvrtcVersion)(int *major, int *minor) nogil
                    ^

[kmaehashi]

This is the same limitation as current coding standard (.flake8.cython)

[kmaehashi]

ctypedefs can be generated by copy-pasting arguments from the header file.

nvrtc.h:

nvrtcResult nvrtcCompileProgram(nvrtcProgram prog,
                                int numOptions, const char * const *options);

[kmaehashi]

/test mini

[kmaehashi]

/test cuda-python

[leofang]

LGTM!

[kmaehashi]

Test passed!

[emcastillo]

Nice!!

[emcastillo]

LGTM!!!!