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.

Sign up for GitHub

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

Soft link NVRTC for cupy_backends.cuda.libs.nvrtc #7621

Merged
merged 20 commits into from
Aug 30, 2023
Merged

Soft link NVRTC for cupy_backends.cuda.libs.nvrtc #7621

Show file tree
Hide file tree
Changes from 11 commits
Commits
Show all changes
20 commits
Select commit Hold shift click to select a range
d215dcc
fix cython-lint config to align with existing rule
kmaehashi Jun 8, 2023
e21ad39
simplify SoftLink APIs
kmaehashi Jun 8, 2023
5526d8a
load all NVRTC symbols dynamically
kmaehashi Jun 8, 2023
2ec7e33
NVRTC no longer necessary for when building cupy_backends.cuda.libs.n…
kmaehashi Jun 8, 2023
897b553
remove headers
kmaehashi Jun 9, 2023
3adf560
wip: declare prototype for CreateProgram
kmaehashi Jul 7, 2023
a731cdd
raise for mandatory libraries
kmaehashi Aug 12, 2023
fe17225
ignore E211 for ctypedef in Cython
kmaehashi Aug 12, 2023
ffe14c5
soft link all APIs with typedef
kmaehashi Aug 12, 2023
3cf7831
raise from exception
kmaehashi Aug 12, 2023
c1890dc
revert rename
kmaehashi Aug 12, 2023
0613c47
remove unused import
kmaehashi Aug 12, 2023
be52b85
keep exception object to raise lazily
kmaehashi Aug 12, 2023
62d1a38
raise lazily if NVRTC is not loaded
kmaehashi Aug 12, 2023
80d13fa
make prefix public, raise when mandatory
kmaehashi Aug 12, 2023
a665de2
lazy initialize
kmaehashi Aug 12, 2023
9b8857d
inline and handle exceptions
kmaehashi Aug 29, 2023
947fc74
split initialize
kmaehashi Aug 29, 2023
97250fe
accept CUDARuntimeError when querying NVRTC version
kmaehashi Aug 29, 2023
6688747
fix compile error with CUDA Python
kmaehashi Aug 29, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion .flake8.cython
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

E211 added to avoid Whitespace before '(' error:

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

1 change: 1 addition & 0 deletions .pre-commit-config.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 36,7 @@ repos:
rev: v0.15.0
hooks:
- id: cython-lint
args: ["--max-line-length", "79"]
Copy link
Member Author

@kmaehashi kmaehashi Aug 12, 2023 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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


- repo: https://github.com/pre-commit/mirrors-mypy
rev: v1.4.1
Expand Down
1 change: 1 addition & 0 deletions cupy_backends/cuda/_softlink.pxd
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 2,7 @@ ctypedef int (*func_ptr)(...) nogil # NOQA

cdef class SoftLink:
cdef:
bint available
object _cdll
str _prefix
func_ptr get(self, str name)
14 changes: 8 additions & 6 deletions cupy_backends/cuda/_softlink.pyx
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,21 1,23 @@
import ctypes
import warnings

from libc.stdint cimport intptr_t
cimport cython


cdef class SoftLink:
def __init__(self, object libname, str prefix):
def __init__(self, object libname, str prefix, *, bint mandatory=False):
self.available = False
self._prefix = prefix
self._cdll = None
if libname is not None:
try:
self._cdll = ctypes.CDLL(libname)
self.available = True
except Exception as e:
warnings.warn(
f'Warning: CuPy failed to load "{libname}": '
f'({type(e).__name__}: {e})')
self._prefix = prefix
if mandatory:
raise RuntimeError(
f'CuPy failed to load "{libname}": '
f'{type(e).__name__}: {e}') from e

cdef func_ptr get(self, str name):
"""
Expand Down
22 changes: 0 additions & 22 deletions cupy_backends/cuda/cupy_nvrtc.h
View file
Open in desktop

This file was deleted.

97 changes: 97 additions & 0 deletions cupy_backends/cuda/libs/_cnvrtc.pxi
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 1,97 @@
import sys as _sys

from cupy_backends.cuda.api cimport runtime
from cupy_backends.cuda._softlink cimport SoftLink


cdef int _runtime_version = runtime.runtimeGetVersion()
cdef str _prefix = 'nvrtc'
cdef object _libname = None

if CUPY_CUDA_VERSION != 0:
if 11020 <= _runtime_version < 12000:
# CUDA 11.x (11.2 )
if _sys.platform == 'linux':
_libname = 'libnvrtc.so.11.2'
else:
_libname = 'nvrtc64_112_0.dll'
elif 12000 <= _runtime_version < 13000:
# CUDA 12.x
if _sys.platform == 'linux':
_libname = 'libnvrtc.so.12'
else:
_libname = 'nvrtc64_120_0.dll'
elif CUPY_HIP_VERSION != 0:
_prefix = 'hiprtc'
if _runtime_version < 5_00_00000:
# ROCm 4.x
_libname = 'libamdhip64.so.4'
elif _runtime_version < 6_00_00000:
# ROCm 5.x
_libname = 'libamdhip64.so.5'

cdef SoftLink _L = SoftLink(_libname, _prefix, mandatory=True)


ctypedef int nvrtcResult
ctypedef void* nvrtcProgram

# Aliases for compatibillity with existing CuPy codebase.
# TODO(kmaehashi): Remove these aliases.
ctypedef nvrtcProgram Program

ctypedef const char* (*F_nvrtcGetErrorString)(nvrtcResult result) nogil
cdef F_nvrtcGetErrorString nvrtcGetErrorString = <F_nvrtcGetErrorString>_L.get('GetErrorString') # NOQA

ctypedef nvrtcResult (*F_nvrtcVersion)(int *major, int *minor) nogil
cdef F_nvrtcVersion nvrtcVersion = <F_nvrtcVersion>_L.get('Version')

ctypedef nvrtcResult (*F_nvrtcCreateProgram)(
nvrtcProgram* prog, const char* src, const char* name, int numHeaders,
const char** headers, const char** includeNames) nogil
cdef F_nvrtcCreateProgram nvrtcCreateProgram = <F_nvrtcCreateProgram>_L.get('CreateProgram') # NOQA

ctypedef nvrtcResult (*F_nvrtcDestroyProgram)(nvrtcProgram *prog) nogil
cdef F_nvrtcDestroyProgram nvrtcDestroyProgram = <F_nvrtcDestroyProgram>_L.get('DestroyProgram') # NOQA

ctypedef nvrtcResult (*F_nvrtcCompileProgram)(
nvrtcProgram prog, int numOptions, const char** options) nogil
Comment on lines 26 to 27
Copy link
Member Author

@kmaehashi kmaehashi Aug 12, 2023 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

nvrtc.h:

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

cdef F_nvrtcCompileProgram nvrtcCompileProgram = <F_nvrtcCompileProgram>_L.get('CompileProgram') # NOQA

ctypedef nvrtcResult (*F_nvrtcGetPTXSize)(nvrtcProgram prog, size_t *ptxSizeRet) nogil # NOQA
cdef F_nvrtcGetPTXSize nvrtcGetPTXSize = <F_nvrtcGetPTXSize>_L.get(
'GetPTXSize' if _prefix == 'nvrtc' else 'GetCodeSize')

ctypedef nvrtcResult (*F_nvrtcGetPTX)(nvrtcProgram prog, char *ptx) nogil
cdef F_nvrtcGetPTX nvrtcGetPTX = <F_nvrtcGetPTX>_L.get(
'GetPTX' if _prefix == 'nvrtc' else 'GetCode')

ctypedef nvrtcResult (*F_nvrtcGetCUBINSize)(nvrtcProgram prog, size_t *cubinSizeRet) nogil # NOQA
cdef F_nvrtcGetCUBINSize nvrtcGetCUBINSize = <F_nvrtcGetCUBINSize>_L.get('GetCUBINSize') # NOQA

ctypedef nvrtcResult (*F_nvrtcGetCUBIN)(nvrtcProgram prog, char *cubin) nogil
cdef F_nvrtcGetCUBIN nvrtcGetCUBIN = <F_nvrtcGetCUBIN>_L.get('GetCUBIN')

ctypedef nvrtcResult (*F_nvrtcGetProgramLogSize)(nvrtcProgram prog, size_t* logSizeRet) nogil # NOQA
cdef F_nvrtcGetProgramLogSize nvrtcGetProgramLogSize = <F_nvrtcGetProgramLogSize>_L.get('GetProgramLogSize') # NOQA

ctypedef nvrtcResult (*F_nvrtcGetProgramLog)(nvrtcProgram prog, char* log) nogil # NOQA
cdef F_nvrtcGetProgramLog nvrtcGetProgramLog = <F_nvrtcGetProgramLog>_L.get('GetProgramLog') # NOQA

ctypedef nvrtcResult (*F_nvrtcAddNameExpression)(nvrtcProgram, const char*) nogil # NOQA
cdef F_nvrtcAddNameExpression nvrtcAddNameExpression = <F_nvrtcAddNameExpression>_L.get('AddNameExpression') # NOQA

ctypedef nvrtcResult (*F_nvrtcGetLoweredName)(nvrtcProgram, const char*, const char**) nogil # NOQA
cdef F_nvrtcGetLoweredName nvrtcGetLoweredName = <F_nvrtcGetLoweredName>_L.get('GetLoweredName') # NOQA

ctypedef nvrtcResult (*F_nvrtcGetNumSupportedArchs)(int* numArchs) nogil
cdef F_nvrtcGetNumSupportedArchs nvrtcGetNumSupportedArchs = <F_nvrtcGetNumSupportedArchs>_L.get('GetNumSupportedArchs') # NOQA

ctypedef nvrtcResult (*F_nvrtcGetSupportedArchs)(int* supportedArchs) nogil
cdef F_nvrtcGetSupportedArchs nvrtcGetSupportedArchs = <F_nvrtcGetSupportedArchs>_L.get('GetSupportedArchs') # NOQA

ctypedef nvrtcResult (*F_nvrtcGetNVVMSize)(nvrtcProgram prog, size_t *nvvmSizeRet) nogil # NOQA
cdef F_nvrtcGetNVVMSize nvrtcGetNVVMSize = <F_nvrtcGetNVVMSize>_L.get('GetNVVMSize') # NOQA

ctypedef nvrtcResult (*F_nvrtcGetNVVM)(nvrtcProgram prog, char *nvvm) nogil
cdef F_nvrtcGetNVVM nvrtcGetNVVM = <F_nvrtcGetNVVM>_L.get('GetNVVM')
6 changes: 0 additions & 6 deletions cupy_backends/cuda/libs/nvrtc.pxd
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,13 9,7 @@ IF CUPY_USE_CUDA_PYTHON:
from cuda.cnvrtc cimport *
# Aliases for compatibillity with existing CuPy codebase.
# TODO(kmaehashi): Remove these aliases.
ctypedef nvrtcResult Result
ctypedef nvrtcProgram Program
ELSE:
cdef extern from *:
ctypedef int Result 'nvrtcResult'
ctypedef void* Program 'nvrtcProgram'


cpdef check_status(int status)

Expand Down
47 changes: 4 additions & 43 deletions cupy_backends/cuda/libs/nvrtc.pyx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,14 11,11 @@ There are four differences compared to the original C API.
4. The resulting values are returned directly instead of references.

"""
import sys as _sys # no-cython-lint
import sys as _sys

cimport cython # NOQA
from libcpp cimport vector

from cupy_backends.cuda.api cimport runtime
from cupy_backends.cuda._softlink cimport SoftLink


###############################################################################
# Extern
Expand All @@ -27,44 24,8 @@ from cupy_backends.cuda._softlink cimport SoftLink
IF CUPY_USE_CUDA_PYTHON:
from cuda.cnvrtc cimport *
ELSE:
cdef extern from '../../cupy_rtc.h' nogil:
const char *nvrtcGetErrorString(Result result)
int nvrtcVersion(int *major, int *minor)
int nvrtcCreateProgram(
Program* prog, const char* src, const char* name, int numHeaders,
const char** headers, const char** includeNames)
int nvrtcDestroyProgram(Program *prog)
int nvrtcCompileProgram(Program prog, int numOptions,
const char** options)
int nvrtcGetPTXSize(Program prog, size_t *ptxSizeRet)
int nvrtcGetPTX(Program prog, char *ptx)
int nvrtcGetCUBINSize(Program prog, size_t *cubinSizeRet)
int nvrtcGetCUBIN(Program prog, char *cubin)
int nvrtcGetProgramLogSize(Program prog, size_t* logSizeRet)
int nvrtcGetProgramLog(Program prog, char* log)
int nvrtcAddNameExpression(Program, const char*)
int nvrtcGetLoweredName(Program, const char*, const char**)

ctypedef int (*f_type)(...) nogil # NOQA
IF 11020 <= CUPY_CUDA_VERSION < 12000:
if _sys.platform == 'linux':
_libname = 'libnvrtc.so.11.2'
else:
_libname = 'nvrtc64_112_0.dll'
ELIF 12000 <= CUPY_CUDA_VERSION < 13000:
if _sys.platform == 'linux':
_libname = 'libnvrtc.so.12'
else:
_libname = 'nvrtc64_120_0.dll'
ELSE:
_libname = None

cdef SoftLink _lib = SoftLink(_libname, 'nvrtc')
# APIs added after CUDA 11.2 .
cdef f_type nvrtcGetNumSupportedArchs = <f_type>_lib.get('GetNumSupportedArchs') # NOQA
cdef f_type nvrtcGetSupportedArchs = <f_type>_lib.get('GetSupportedArchs') # NOQA
cdef f_type nvrtcGetNVVMSize = <f_type>_lib.get('GetNVVMSize')
cdef f_type nvrtcGetNVVM = <f_type>_lib.get('GetNVVM')
include "_cnvrtc.pxi"
pass


###############################################################################
Expand All @@ -75,7 36,7 @@ class NVRTCError(RuntimeError):

def __init__(self, status):
self.status = status
cdef bytes msg = nvrtcGetErrorString(<Result>status)
cdef bytes msg = nvrtcGetErrorString(<nvrtcResult>status)
super(NVRTCError, self).__init__(
'{} ({})'.format(msg.decode(), status))

Expand Down
18 changes: 0 additions & 18 deletions cupy_backends/cupy_rtc.h
View file
Open in desktop

This file was deleted.

74 changes: 0 additions & 74 deletions cupy_backends/hip/cupy_hiprtc.h
View file
Open in desktop

This file was deleted.

2 changes: 0 additions & 2 deletions install/cupy_builder/_features.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -451,7 451,6 @@ def __init__(self, ctx: Context):
'cufft.h',
'curand.h',
'cusparse.h',
'nvrtc.h',
]
# TODO(kmaehashi): Split profiler module so that dependency to
# `cudart` can be removed when using CUDA Python.
Expand All @@ -461,7 460,6 @@ def __init__(self, ctx: Context):
'cufft',
'curand',
'cusparse',
'nvrtc',
]
)
self._version = self._UNDETERMINED
Expand Down