PERF: Inline return code checks in cuda_core (#2608)

Author: mdboom

cuda_core/cuda/core/_utils/cuda_utils.pxd

@@ -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
@@ -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=*)
 

cuda_core/cuda/core/_utils/cuda_utils.pyx

@@ -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
@@ -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)
@@ -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
@@ -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."""