Fix CABI return type mismatch when calling Numba-ABI subroutines

In `user_function`, `return_status_propagate` was called with the
callee's calling convention (`fndesc.call_conv`).  When a CABI function
called a Numba-ABI subroutine (e.g. an overloaded function compiled
via `compile_subroutine`), the Numba-ABI convention emitted
`ret i32 <status>` inside a function whose declared return type was
the user type (e.g. `i8*` for NoneType), producing invalid LLVM IR
and an NVVM `ERROR_INVALID_IR` verification failure.

Use the caller's calling convention (`context.fndesc.call_conv`)
instead, which correctly delegates to `CUDACABICallConv.
return_status_propagate` (a no-op, since C ABI has no error-status
channel).

Closes #841

Made-with: Cursor

Author: isVoid

numba_cuda/numba/cuda/core/imputils.py

@@ -222,7 +222,9 @@ def imp(context, builder, sig, args, fndesc=fndesc):
 
         if status is not None:
             with cgutils.if_unlikely(builder, status.is_error):
-                fndesc.call_conv.return_status_propagate(builder, status)
+                context.fndesc.call_conv.return_status_propagate(
+                    builder, status
+                )
 
         assert sig.return_type == fndesc.restype
 

numba_cuda/numba/cuda/tests/cudapy/test_compiler.py

@@ -4,7 +4,7 @@
 import os
 from math import sqrt
 from numba import cuda
-from numba.core.extending import intrinsic
+from numba.core.extending import intrinsic, overload
 
 from numba.cuda import float32, int16, int32, int64, types, uint32, void
 from numba.cuda import (
@@ -837,6 +837,34 @@ def wrapper_func(x):
             wrapper_func, types.int32(types.int32), output="ltoir", abi="c"
         )
 
+    def test_compile_CABI_calling_overloaded_function(self):
+        # Reproducer from gh-841
+        # https://github.com/NVIDIA/numba-cuda/issues/841
+        #
+        # When a CABI function calls an overloaded function (compiled as a
+        # Numba-ABI subroutine), the error-status return path must use the
+        # caller's calling convention. Previously the callee's Numba-ABI
+        # convention was used, emitting `ret i32` (the status code) inside a
+        # function whose declared return type did not match, producing invalid
+        # LLVM IR.
+
+        def _my_func(x):
+            pass
+
+        @overload(_my_func)
+        def _ov_my_func(x):
+            if isinstance(x, types.IntegerLiteral):
+
+                def impl(x):
+                    return int32(0)
+
+                return impl
+
+        def caller():
+            _my_func(0)
+
+        cuda.compile(caller, (), abi_info={"abi_name": "caller"}, abi="c")
+
     def test_compile_complex_div_c_abi(self):
         # Reproducer from gh-789
         # https://github.com/NVIDIA/numba-cuda/issues/789