Soft link NVRTC for cupy_backends.cuda.libs.nvrtc
#7621
Conversation
|
This seems to be affecting rocm build somehow? |
|
/test mini |
|
Will resume after #7647 |
|
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
with this con
The main blocker for this would be the maintainability of the Cython layer. We still have to keep the function signatures... |
|
@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 nvrtcResult nvrtcCreateProgram(nvrtcProgram *prog,
const char *src,
const char *name,
int numHeaders,
const char * const *headers,
const char * const *includeNames); |
|
Sorry for some reason I missed your reply, Kenichi 😅
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
force-pushed
the
softlink-nvrtc
branch
from
5c9cd11 to
fe17225
Compare
| @@ -1,4 1,4 @@ | |||
| [flake8] | |||
| filename = *.pyx, *.pxd, *.pxi | |||
| exclude = .git, .eggs, *.py | |||
| ignore = W503,W504,E225,E226,E227,E275,E402,E999 | |||
| ignore = W503,W504,E211,E225,E226,E227,E275,E402,E999 | |||
There was a problem hiding this comment.
E211 added to avoid Whitespace before '(' error:
ctypedef nvrtcResult (*F_nvrtcVersion)(int *major, int *minor) nogil
^
| @@ -36,6 36,7 @@ repos: | |||
| rev: v0.15.0 | |||
| hooks: | |||
| - id: cython-lint | |||
| args: ["--max-line-length", "79"] | |||
There was a problem hiding this comment.
This is the same limitation as current coding standard (.flake8.cython)
| ctypedef nvrtcResult (*F_nvrtcCompileProgram)( | ||
| nvrtcProgram prog, int numOptions, const char** options) nogil |
There was a problem hiding this comment.
ctypedefs can be generated by copy-pasting arguments from the header file.
nvrtcResult nvrtcCompileProgram(nvrtcProgram prog,
int numOptions, const char * const *options);
kmaehashi
added
the
blocking
|
/test mini |
|
/test cuda-python |
|
LGTM! |
|
Test passed! |
| if libname is None: | ||
| # Stub build or CUDA/HIP only library. | ||
| self.error = RuntimeError( | ||
| 'The library is unavailable in the current platform.') |
Part of #7620.
This PR changes the
cupy_backends.cuda.libs.nvrtcmodule to load NVRTC symbols dynamically at runtime, instead of build-time linking.Pros:
import cupywithout NVRTC)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 addingctypedeffor each API.