Skip to content

PERF: Inline return code checks in cuda_core #2608

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

Merged
mdboom merged 1 commit into from
Aug 13, 2026
+32 −37
Merged

PERF: Inline return code checks in cuda_core #2608

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
38 changes: 32 additions & 6 deletions cuda_core/cuda/core/_utils/cuda_utils.pxd
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,11 +18,35 @@ ctypedef fused integer_t:
cdef const cydriver.CUcontext CU_CONTEXT_INVALID = <cydriver.CUcontext>(-2)


cdef int HANDLE_RETURN(cydriver.CUresult err) except?-1 nogil
cdef int HANDLE_RETURN_NVRTC(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except?-1 nogil
cdef int HANDLE_RETURN_NVVM(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except?-1 nogil
cdef int HANDLE_RETURN_NVJITLINK(
cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except?-1 nogil
cdef inline int HANDLE_RETURN(cydriver.CUresult err) except?-1 nogil:
if err != cydriver.CUresult.CUDA_SUCCESS:
return _check_driver_error(err)
return 0


cdef inline int HANDLE_RETURN_NVRTC(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except?-1 nogil:
"""Handle NVRTC result codes, raising NVRTCError with program log on failure."""
if err == cynvrtc.nvrtcResult.NVRTC_SUCCESS:
return 0
with gil:
_raise_nvrtc_error(prog, err)


cdef inline int HANDLE_RETURN_NVVM(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except?-1 nogil:
"""Handle NVVM result codes, raising nvvmError with program log on failure."""
if err == cynvvm.nvvmResult.NVVM_SUCCESS:
return 0
with gil:
_raise_nvvm_error(prog, err)


cdef inline int HANDLE_RETURN_NVJITLINK(
cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except?-1 nogil:
"""Handle nvJitLink result codes, raising nvJitLinkError with error log on failure."""
if err == cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS:
return 0
with gil:
_raise_nvjitlink_error(handle, err)


# Helper for retrieving the current CUDA device. Raises if no active context
Expand All @@ -34,7 +58,9 @@ cdef int _get_current_device_id() except? -1
cpdef int _check_driver_error(cydriver.CUresult error) except?-1 nogil
cpdef int _check_runtime_error(error) except?-1
cpdef int _check_nvrtc_error(error) except?-1

cdef int _raise_nvrtc_error(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except -1
cdef int _raise_nvvm_error(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except -1
cdef int _raise_nvjitlink_error(cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except -1

cpdef check_or_create_options(type cls, options, str options_description=*, bint keep_none=*)

Expand Down
31 changes: 0 additions & 31 deletions cuda_core/cuda/core/_utils/cuda_utils.pyx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,12 +63,6 @@ def cast_to_3_tuple(label: str, cfg: int | tuple[int, ...]) -> tuple[int, int, i
return cfg + (1,) * (3 - len(cfg))


cdef int HANDLE_RETURN(cydriver.CUresult err) except?-1 nogil:
if err != cydriver.CUresult.CUDA_SUCCESS:
return _check_driver_error(err)
return 0


cdef int _get_current_device_id() except? -1:
"""Return the current thread's bound CUdevice ordinal."""
cdef cydriver.CUdevice dev
Expand All @@ -77,14 +71,6 @@ cdef int _get_current_device_id() except? -1:
return <int>dev


cdef int HANDLE_RETURN_NVRTC(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except?-1 nogil:
"""Handle NVRTC result codes, raising NVRTCError with program log on failure."""
if err == cynvrtc.nvrtcResult.NVRTC_SUCCESS:
return 0
with gil:
_raise_nvrtc_error(prog, err)


cdef int _raise_nvrtc_error(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except -1:
"""Build error message with program log and raise NVRTCError."""
cdef const char* err_str = cynvrtc.nvrtcGetErrorString(err)
Expand All @@ -103,14 +89,6 @@ cdef int _raise_nvrtc_error(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err)
raise NVRTCError(err_msg)


cdef int HANDLE_RETURN_NVVM(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except?-1 nogil:
"""Handle NVVM result codes, raising nvvmError with program log on failure."""
if err == cynvvm.nvvmResult.NVVM_SUCCESS:
return 0
with gil:
_raise_nvvm_error(prog, err)


cdef int _raise_nvvm_error(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except -1:
"""Raise nvvmError annotated with the program log."""
cdef size_t logsize = 0
Expand All @@ -128,15 +106,6 @@ cdef int _raise_nvvm_error(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) excep
raise exc


cdef int HANDLE_RETURN_NVJITLINK(
cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except?-1 nogil:
"""Handle nvJitLink result codes, raising nvJitLinkError with error log on failure."""
if err == cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS:
return 0
with gil:
_raise_nvjitlink_error(handle, err)


cdef int _raise_nvjitlink_error(
cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except -1:
"""Raise nvJitLinkError annotated with the error log."""
Expand Down
Loading