diff --git a/docs/source/user/callconv.rst b/docs/source/user/callconv.rst
new file mode 100644
index 000000000..dbdfb5038
--- /dev/null
+++ b/docs/source/user/callconv.rst
@@ -0,0 +1,107 @@
+..
+   SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+   SPDX-License-Identifier: BSD-2-Clause
+
+
+.. _cuda-call-conventions:
+
+CUDA device call conventions
+============================
+
+Numba-CUDA supports two ABIs for device functions:
+
+- The **Numba ABI**, used internally by Numba for most compiled device code.
+- The **C ABI**, intended for interoperability with CUDA C/C++ style calls.
+
+.. important::
+
+   This is a major deviation from upstream Numba behavior: Numba-CUDA supports
+   arbitrary nesting between these ABIs. A Numba-ABI function can call a
+   C-ABI function, which can call a Numba-ABI function, and so on.
+
+
+ABI overview
+------------
+
+Numba ABI
+~~~~~~~~~
+
+The Numba ABI is described in :ref:`device-function-abi` (without the
+``extern "C"`` modifier):
+
+- The function has a **status return code**.
+- The Python return value is passed via a **pointer in the first argument**.
+- Function names are mangled using Numba's mangling rules.
+- Optional returns and exception status can be represented via the status
+  channel.
+
+C ABI
+~~~~~
+
+The C ABI behavior for compiled Python device functions is described in
+:ref:`cuda-using-the-c-abi`:
+
+- The function has a conventional C-style signature:
+  ``<return_type>(<args...>)``.
+- There is no separate status return code channel.
+- Function names are predictable (by default the Python ``__name__``), and can
+  be set explicitly with ``abi_info={"abi_name": ...}``.
+- The C ABI is supported for device functions (not kernels).
+
+
+Caller/callee matrix
+--------------------
+
+The table below summarizes what happens at each call edge:
+
+.. list-table:: Caller and callee ABI combinations
+   :header-rows: 1
+   :widths: 22 39 39
+
+   * - Caller / Callee
+     - Numba ABI callee
+     - C ABI callee
+   * - Numba ABI caller
+     - Numba-to-Numba call. Uses Numba ABI marshalling (status + return
+       pointer), and propagates lower-frame error status.
+     - Mixed call. Arguments / return are marshalled using the callee's C ABI
+       signature. No callee status channel exists to propagate Python-exception
+       status from the callee.
+   * - C ABI caller
+     - Mixed call. The call is marshalled using the callee's Numba ABI. The
+       Numba callee can still produce status, but the C ABI caller has no
+       outward status channel and does not propagate lower-frame status.
+     - C-to-C call. Conventional C-style argument / return passing with no
+       status channel.
+
+
+What arbitrary nesting means
+----------------------------
+
+Each call site is lowered using the **callee's ABI**, not by forcing one ABI
+for the whole call chain. This allows patterns like:
+
+.. code:: text
+
+   Numba ABI caller -> C ABI callee -> Numba ABI callee -> C ABI callee
+
+to compile as expected.
+
+In practice, this means mixed boundaries can appear at any depth in a call
+graph, including calls to functions declared with
+:func:`numba.cuda.declare_device` and calls to Numba-compiled device
+subroutines.
+
+
+Behavioral caveats
+------------------
+
+- The C ABI has no status channel for Python exception propagation.
+- When a C ABI caller invokes a Numba ABI callee returning ``Optional[T]``,
+  the optional is flattened to ``T`` at the C ABI boundary. A ``None`` result
+  is represented as the default-initialized value of ``T``.
+- Kernels must still use the Numba ABI entry model; compiling kernels with
+  ``abi="c"`` is unsupported.
+- For foreign CUDA C/C++ functions, use ``abi="c"`` with
+  :func:`numba.cuda.declare_device` and follow pointer-signature guidance in
+  :ref:`cuda_ffi`.
diff --git a/docs/source/user/index.rst b/docs/source/user/index.rst
index c145e0993..725b9cdc0 100644
--- a/docs/source/user/index.rst
+++ b/docs/source/user/index.rst
@@ -32,6 +32,7 @@ User guide
    bindings.rst
    cuda_ffi.rst
    cuda_compilation.rst
+   callconv.rst
    caching.rst
    minor_version_compatibility.rst
    faq.rst
diff --git a/numba_cuda/numba/cuda/compiler.py b/numba_cuda/numba/cuda/compiler.py
index d8dbb446c..1ac77a7f2 100644
--- a/numba_cuda/numba/cuda/compiler.py
+++ b/numba_cuda/numba/cuda/compiler.py
@@ -560,6 +560,8 @@ def compile_extra(
     locals,
     library=None,
     pipeline_class=CUDACompiler,
+    call_conv=None,
+    abi_info=None,
 ):
     """Compiler entry point
 
@@ -584,7 +586,15 @@ def compile_extra(
         compiler pipeline
     """
     pipeline = pipeline_class(
-        typingctx, targetctx, library, args, return_type, flags, locals
+        typingctx,
+        targetctx,
+        library,
+        args,
+        return_type,
+        flags,
+        locals,
+        call_conv,
+        abi_info,
     )
     return pipeline.compile_extra(func)
 
@@ -602,6 +612,8 @@ def compile_ir(
     is_lifted_loop=False,
     library=None,
     pipeline_class=CUDACompiler,
+    call_conv=None,
+    abi_info=None,
 ):
     """
     Compile a function with the given IR.
@@ -637,6 +649,8 @@ def compile_local(the_ir, the_flags):
                 return_type,
                 the_flags,
                 locals,
+                call_conv,
+                abi_info,
             )
             return pipeline.compile_ir(
                 func_ir=the_ir, lifted=lifted, lifted_from=lifted_from
@@ -675,13 +689,30 @@ def compile_local(the_ir, the_flags):
 
 
 def compile_internal(
-    typingctx, targetctx, library, func, args, return_type, flags, locals
+    typingctx,
+    targetctx,
+    library,
+    func,
+    args,
+    return_type,
+    flags,
+    locals,
+    call_conv=None,
+    abi_info=None,
 ):
     """
     For internal use only.
     """
     pipeline = CUDACompiler(
-        typingctx, targetctx, library, args, return_type, flags, locals
+        typingctx,
+        targetctx,
+        library,
+        args,
+        return_type,
+        flags,
+        locals,
+        call_conv,
+        abi_info,
     )
     return pipeline.compile_extra(func)
 
@@ -747,10 +778,12 @@ def compile_cuda(
     flags.lto = lto
 
     if abi == "c":
-        flags.call_conv = CUDACABICallConv(targetctx)
+        call_conv = CUDACABICallConv(targetctx)
+    else:
+        call_conv = CUDACallConv(targetctx)
 
-    if abi_info is not None:
-        flags.abi_info = abi_info
+    if abi_info is None:
+        abi_info = {}
 
     with utils.numba_target_override():
         cres = compile_extra(
@@ -762,6 +795,8 @@ def compile_cuda(
             flags=flags,
             locals={},
             pipeline_class=CUDACompiler,
+            call_conv=call_conv,
+            abi_info=abi_info,
         )
 
     library = cres.library
diff --git a/numba_cuda/numba/cuda/core/base.py b/numba_cuda/numba/cuda/core/base.py
index c518f9cb8..78b3a2028 100644
--- a/numba_cuda/numba/cuda/core/base.py
+++ b/numba_cuda/numba/cuda/core/base.py
@@ -958,26 +958,15 @@ def call_internal(self, builder, fndesc, sig, args):
         Given the function descriptor of an internally compiled function,
         emit a call to that function with the given arguments.
         """
-        status, res = self.call_internal_no_propagate(
-            builder, fndesc, sig, args
-        )
-        with cgutils.if_unlikely(builder, status.is_error):
-            fndesc.call_conv.return_status_propagate(builder, status)
-
-        res = imputils.fix_returning_optional(self, builder, sig, status, res)
-        return res
+        return self.fndesc.call_conv.call_internal(builder, fndesc, sig, args)
 
     def call_internal_no_propagate(self, builder, fndesc, sig, args):
         """Similar to `.call_internal()` but does not handle or propagate
         the return status automatically.
         """
-        # Add call to the generated function
-        llvm_mod = builder.module
-        fn = fndesc.declare_function(llvm_mod)
-        status, res = fndesc.call_conv.call_function(
-            builder, fn, sig.return_type, sig.args, args
+        return self.fndesc.call_conv.call_internal_no_propagate(
+            builder, fndesc, sig, args
         )
-        return status, res
 
     def call_unresolved(self, builder, name, sig, args):
         """
diff --git a/numba_cuda/numba/cuda/core/callconv.py b/numba_cuda/numba/cuda/core/callconv.py
index 3ada0c16b..a68e66c1d 100644
--- a/numba_cuda/numba/cuda/core/callconv.py
+++ b/numba_cuda/numba/cuda/core/callconv.py
@@ -4,6 +4,7 @@
 from numba.cuda import types
 from numba.cuda import cgutils
 from numba.cuda import itanium_mangler
+from numba.cuda.core import imputils
 from collections import namedtuple
 
 from llvmlite import ir
@@ -167,6 +168,18 @@ def mangler(self, name, argtypes, *, abi_tags=(), uid=None):
             name, argtypes, abi_tags=abi_tags, uid=uid
         )
 
+    def call_internal_no_propagate(self, builder, fndesc, sig, args):
+        """Similar to `.call_internal()` but does not handle or propagate
+        the return status automatically.
+        """
+        llvm_mod = builder.module
+        fn = fndesc.declare_function(llvm_mod)
+        # Marshal the call using the callee's ABI.
+        status, res = fndesc.call_conv.call_function(
+            builder, fn, sig.return_type, sig.args, args
+        )
+        return status, res
+
 
 class MinimalCallConv(BaseCallConv):
     """
@@ -298,6 +311,23 @@ def call_function(self, builder, callee, resty, argtys, args):
         out = self.context.get_returned_value(builder, resty, retval)
         return status, out
 
+    def call_internal(self, builder, fndesc, sig, args):
+        """
+        Given the function descriptor of an internally compiled function,
+        emit a call to that function with the given arguments.
+        """
+        status, res = self.call_internal_no_propagate(
+            builder, fndesc, sig, args
+        )
+        if status is not None:
+            with cgutils.if_unlikely(builder, status.is_error):
+                self.return_status_propagate(builder, status)
+
+        res = imputils.fix_returning_optional(
+            self.context, builder, sig, status, res
+        )
+        return res
+
 
 class _MinimalCallHelper:
     """
@@ -396,12 +426,12 @@ def return_value(self, builder, retval):
     def return_user_exc(
         self, builder, exc, exc_args=None, loc=None, func_name=None
     ):
-        msg = "Python exceptions are unsupported in the CUDA C/C++ ABI"
-        raise NotImplementedError(msg)
+        # C ABI has no status channel to propagate Python exceptions.
+        return
 
     def return_status_propagate(self, builder, status):
-        msg = "Return status is unsupported in the CUDA C/C++ ABI"
-        raise NotImplementedError(msg)
+        # C ABI has no status channel to propagate lower-frame failures.
+        return
 
     def get_function_type(self, restype, argtypes):
         """
@@ -439,6 +469,34 @@ def call_function(self, builder, callee, resty, argtys, args):
         out = self.context.get_returned_value(builder, resty, code)
         return status, out
 
+    def call_internal(self, builder, fndesc, sig, args):
+        """
+        Given the function descriptor of an internally compiled function,
+        emit a call to that function with the given arguments.
+        """
+        status, res = self.call_internal_no_propagate(
+            builder, fndesc, sig, args
+        )
+
+        # CABI intentionally ignores lower-frame error codes.
+        if not isinstance(sig.return_type, types.Optional):
+            return res
+
+        # A callee without a status channel cannot represent None.
+        if status is None:
+            return res
+
+        # Flatten Optional[T] into plain T for CABI:
+        # - if value is present, return it
+        # - if value is None, return a default-initialized T
+        value_type = sig.return_type.type
+        default_value = self.context.get_constant_null(value_type)
+
+        outptr = cgutils.alloca_once_value(builder, default_value)
+        with builder.if_then(builder.not_(status.is_none)):
+            builder.store(res, outptr)
+        return builder.load(outptr)
+
     def get_return_type(self, ty):
         return self.context.data_model_manager[ty].get_return_type()
 
@@ -456,11 +514,13 @@ def __init__(self, call_conv):
     def fp_zero_division(self, builder, exc_args=None, loc=None):
         if self.raise_on_fp_zero_division:
             self.call_conv.return_user_exc(
-                builder, ZeroDivisionError, exc_args, loc
+                builder,
+                ZeroDivisionError,
+                exc_args=exc_args,
+                loc=loc,
             )
             return True
-        else:
-            return False
+        return False
 
 
 class PythonErrorModel(ErrorModel):
diff --git a/numba_cuda/numba/cuda/core/compiler.py b/numba_cuda/numba/cuda/core/compiler.py
index c6a037e53..75a265b45 100644
--- a/numba_cuda/numba/cuda/core/compiler.py
+++ b/numba_cuda/numba/cuda/core/compiler.py
@@ -84,7 +84,16 @@ class CompilerBase:
     """
 
     def __init__(
-        self, typingctx, targetctx, library, args, return_type, flags, locals
+        self,
+        typingctx,
+        targetctx,
+        library,
+        args,
+        return_type,
+        flags,
+        locals,
+        call_conv=None,
+        abi_info=None,
     ):
         # Make sure the environment is reloaded
         config.reload_config()
@@ -116,6 +125,14 @@ def __init__(
         # hold this for e.g. with_lifting, null out on exit
         self.state.pipeline = self
 
+        if call_conv is None:
+            call_conv = CUDACallConv(self.state.targetctx)
+        if abi_info is None:
+            abi_info = {}
+
+        self.state.call_conv = call_conv
+        self.state.abi_info = abi_info
+
         self.state.status = _CompileStatus(
             can_fallback=self.state.flags.enable_pyobject
         )
diff --git a/numba_cuda/numba/cuda/core/imputils.py b/numba_cuda/numba/cuda/core/imputils.py
index 751022675..8e767282f 100644
--- a/numba_cuda/numba/cuda/core/imputils.py
+++ b/numba_cuda/numba/cuda/core/imputils.py
@@ -11,6 +11,7 @@
 
 from numba.cuda import typing, cgutils
 from numba.cuda import types
+from numba.cuda.core import callconv
 from numba.cuda.typing.templates import BaseRegistryLoader
 
 
@@ -496,8 +497,6 @@ def force_error_model(context, model_name="numpy"):
     """
     Temporarily change the context's error model.
     """
-    from numba.cuda.core import callconv
-
     old_error_model = context.error_model
     context.error_model = callconv.create_error_model(model_name, context)
     try:
diff --git a/numba_cuda/numba/cuda/core/typed_passes.py b/numba_cuda/numba/cuda/core/typed_passes.py
index a0fc80c5a..b4d887d37 100644
--- a/numba_cuda/numba/cuda/core/typed_passes.py
+++ b/numba_cuda/numba/cuda/core/typed_passes.py
@@ -333,7 +333,7 @@ def run_pass(self, state):
         metadata = state.metadata
         pre_stats = passmanagers.dump_refprune_stats()
 
-        call_conv = flags.call_conv
+        call_conv = state.call_conv
         if call_conv is None:
             call_conv = CUDACallConv(state.targetctx)
 
@@ -355,7 +355,7 @@ def run_pass(self, state):
                     noalias=flags.noalias,
                     abi_tags=[flags.get_mangle_string()],
                     call_conv=call_conv,
-                    abi_info=flags.abi_info,
+                    abi_info=state.abi_info,
                 )
             )
 
diff --git a/numba_cuda/numba/cuda/cpython/listobj.py b/numba_cuda/numba/cuda/cpython/listobj.py
index 9e2f1532f..99dc8e456 100644
--- a/numba_cuda/numba/cuda/cpython/listobj.py
+++ b/numba_cuda/numba/cuda/cpython/listobj.py
@@ -118,7 +118,9 @@ def guard_index(self, idx, msg):
         """
         with self._builder.if_then(self.is_out_of_bounds(idx), likely=False):
             self._context.fndesc.call_conv.return_user_exc(
-                self._builder, IndexError, (msg,)
+                self._builder,
+                IndexError,
+                (msg,),
             )
 
     def fix_slice(self, slice):
@@ -348,7 +350,9 @@ def allocate(cls, context, builder, list_type, nitems):
         ok, self = cls.allocate_ex(context, builder, list_type, nitems)
         with builder.if_then(builder.not_(ok), likely=False):
             context.fndesc.call_conv.return_user_exc(
-                builder, MemoryError, ("cannot allocate list",)
+                builder,
+                MemoryError,
+                ("cannot allocate list",),
             )
         return self
 
@@ -385,7 +389,9 @@ def _payload_realloc(new_allocated):
             )
             with builder.if_then(ovf, likely=False):
                 context.fndesc.call_conv.return_user_exc(
-                    builder, MemoryError, ("cannot resize list",)
+                    builder,
+                    MemoryError,
+                    ("cannot resize list",),
                 )
 
             ptr = context.nrt.meminfo_varsize_realloc_unchecked(
diff --git a/numba_cuda/numba/cuda/cpython/rangeobj.py b/numba_cuda/numba/cuda/cpython/rangeobj.py
index 99b7903d1..073de04dc 100644
--- a/numba_cuda/numba/cuda/cpython/rangeobj.py
+++ b/numba_cuda/numba/cuda/cpython/rangeobj.py
@@ -159,7 +159,9 @@ def from_range_state(cls, context, builder, state):
             with cgutils.if_unlikely(builder, zero_step):
                 # step shouldn't be zero
                 context.fndesc.call_conv.return_user_exc(
-                    builder, ValueError, ("range() arg 3 must not be zero",)
+                    builder,
+                    ValueError,
+                    ("range() arg 3 must not be zero",),
                 )
 
             with builder.if_else(sign_differs) as (then, orelse):
diff --git a/numba_cuda/numba/cuda/cpython/slicing.py b/numba_cuda/numba/cuda/cpython/slicing.py
index ed9ae163f..77f9411fc 100644
--- a/numba_cuda/numba/cuda/cpython/slicing.py
+++ b/numba_cuda/numba/cuda/cpython/slicing.py
@@ -241,13 +241,17 @@ def slice_indices(context, builder, sig, args):
 
     with builder.if_then(cgutils.is_neg_int(builder, length), likely=False):
         context.fndesc.call_conv.return_user_exc(
-            builder, ValueError, ("length should not be negative",)
+            builder,
+            ValueError,
+            ("length should not be negative",),
         )
     with builder.if_then(
         cgutils.is_scalar_zero(builder, sli.step), likely=False
     ):
         context.fndesc.call_conv.return_user_exc(
-            builder, ValueError, ("slice step cannot be zero",)
+            builder,
+            ValueError,
+            ("slice step cannot be zero",),
         )
 
     fix_slice(builder, sli, length)
diff --git a/numba_cuda/numba/cuda/flags.py b/numba_cuda/numba/cuda/flags.py
index 51d9eb658..7b9eedcb4 100644
--- a/numba_cuda/numba/cuda/flags.py
+++ b/numba_cuda/numba/cuda/flags.py
@@ -163,24 +163,6 @@ def _optional_int_type(x):
         return x
 
 
-def _call_conv_options_type(x):
-    if x is None:
-        return None
-
-    else:
-        assert isinstance(x, BaseCallConv)
-        return x
-
-
-def _abi_info_options_type(x):
-    if x is None:
-        return {}
-
-    else:
-        assert isinstance(x, dict)
-        return x
-
-
 class CUDAFlags(Flags):
     nvvm_options = Option(
         type=_nvvm_options_type,
@@ -196,7 +178,3 @@ class CUDAFlags(Flags):
         type=_optional_int_type, default=None, doc="Max registers"
     )
     lto = Option(type=bool, default=False, doc="Enable Link-time Optimization")
-
-    call_conv = Option(type=_call_conv_options_type, default=None, doc="")
-
-    abi_info = Option(type=_abi_info_options_type, default=None, doc="ABI info")
diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py
index 6184aff06..51a377a0b 100644
--- a/numba_cuda/numba/cuda/lowering.py
+++ b/numba_cuda/numba/cuda/lowering.py
@@ -15,6 +15,7 @@
     ir_utils,
     targetconfig,
     funcdesc,
+    callconv,
     config,
     generators,
     removerefctpass,
@@ -202,7 +203,7 @@ def return_dynamic_exception(self, exc_class, exc_args, nb_types, loc=None):
 
     def return_exception(self, exc_class, exc_args=None, loc=None):
         """Propagate exception to the caller."""
-        self.call_conv.return_user_exc(
+        self.fndesc.call_conv.return_user_exc(
             self.builder,
             exc_class,
             exc_args,
diff --git a/numba_cuda/numba/cuda/np/arrayobj.py b/numba_cuda/numba/cuda/np/arrayobj.py
index 8e233dd3d..299810d5e 100644
--- a/numba_cuda/numba/cuda/np/arrayobj.py
+++ b/numba_cuda/numba/cuda/np/arrayobj.py
@@ -4836,7 +4836,9 @@ def safecast_intp(context, builder, src_t, src):
         is_neg = builder.icmp_signed("<", shape, zero)
         with cgutils.if_unlikely(builder, is_neg):
             context.fndesc.call_conv.return_user_exc(
-                builder, ValueError, ("negative dimensions not allowed",)
+                builder,
+                ValueError,
+                ("negative dimensions not allowed",),
             )
 
     return shapes
@@ -6023,7 +6025,9 @@ def check_sequence_shape(context, builder, seqty, seq, shapes):
 
     def _fail():
         context.fndesc.call_conv.return_user_exc(
-            builder, ValueError, ("incompatible sequence shape",)
+            builder,
+            ValueError,
+            ("incompatible sequence shape",),
         )
 
     def check_seq_size(seqty, seq, shapes):
diff --git a/numba_cuda/numba/cuda/np/polynomial/polynomial_core.py b/numba_cuda/numba/cuda/np/polynomial/polynomial_core.py
index 1e6dc266e..38fea0743 100644
--- a/numba_cuda/numba/cuda/np/polynomial/polynomial_core.py
+++ b/numba_cuda/numba/cuda/np/polynomial/polynomial_core.py
@@ -150,12 +150,16 @@ def to_double(coef):
 
     with cgutils.if_unlikely(builder, pred1):
         context.fndesc.call_conv.return_user_exc(
-            builder, ValueError, ("Domain has wrong number of elements.",)
+            builder,
+            ValueError,
+            ("Domain has wrong number of elements.",),
         )
 
     with cgutils.if_unlikely(builder, pred2):
         context.fndesc.call_conv.return_user_exc(
-            builder, ValueError, ("Window has wrong number of elements.",)
+            builder,
+            ValueError,
+            ("Window has wrong number of elements.",),
         )
 
     polynomial.coef = coef_cast
diff --git a/numba_cuda/numba/cuda/simulator/compiler.py b/numba_cuda/numba/cuda/simulator/compiler.py
index 11f5f31d8..67c5fb8d8 100644
--- a/numba_cuda/numba/cuda/simulator/compiler.py
+++ b/numba_cuda/numba/cuda/simulator/compiler.py
@@ -30,7 +30,16 @@ def define_typed_pipeline(state, name="typed"):
 
 class CompilerBase:
     def __init__(
-        self, typingctx, targetctx, library, args, return_type, flags, locals
+        self,
+        typingctx,
+        targetctx,
+        library,
+        args,
+        return_type,
+        flags,
+        locals,
+        call_conv,
+        abi_info,
     ):
         pass
 
diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py
index 387eea09b..b28ece3e0 100644
--- a/numba_cuda/numba/cuda/target.py
+++ b/numba_cuda/numba/cuda/target.py
@@ -11,6 +11,7 @@
 
 from numba.cuda import types
 from numba.cuda import HAS_NUMBA
+from numba.cuda.core.callconv import CUDACallConv
 from numba.cuda.core.compiler_lock import global_compiler_lock
 from numba.cuda.core.errors import NumbaWarning
 from numba.cuda.core.base import BaseContext
@@ -411,6 +412,10 @@ def _compile_subroutine_no_cache(
             flags.no_cpython_wrapper = True
             flags.no_cfunc_wrapper = True
 
+            # compile_subroutine always uses CUDACallConv
+            call_conv = CUDACallConv(self)
+            abi_info = {}
+
             cres = compiler.compile_internal(
                 self.typing_context,
                 self,
@@ -420,6 +425,8 @@ def _compile_subroutine_no_cache(
                 sig.return_type,
                 flags,
                 locals=locals,
+                call_conv=call_conv,
+                abi_info=abi_info,
             )
 
             # Allow inlining the function inside callers
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_compiler.py b/numba_cuda/numba/cuda/tests/cudapy/test_compiler.py
index cf5b259c3..2a2060a78 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_compiler.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_compiler.py
@@ -4,6 +4,8 @@
 import os
 from math import sqrt
 from numba import cuda
+from numba.core.extending import intrinsic
+
 from numba.cuda import float32, int16, int32, int64, types, uint32, void
 from numba.cuda import (
     compile,
@@ -16,6 +18,8 @@
 from numba.cuda.cudadrv import nvrtc
 from numba.cuda.testing import skip_on_cudasim, unittest, CUDATestCase
 
+from numba.cuda.core.callconv import CUDACallConv
+
 TEST_BIN_DIR = os.getenv("NUMBA_CUDA_TEST_BIN_DIR")
 if TEST_BIN_DIR:
     test_device_functions_a = os.path.join(
@@ -701,6 +705,98 @@ def f(z, x, y):
             str(code_list[1].code.decode()), r"\.section\s+\.debug_info"
         )
 
+    def test_compile_jitted_subroutine(self):
+        # Reproducer from gh-781
+        # https://github.com/NVIDIA/numba-cuda/issues/781
+        def foo(x):
+            return 2 * x
+
+        # Create a wrapper that takes void* arguments
+        def create_void_ptr_wrapper():
+            """Create a wrapper that takes void* input and output pointers."""
+
+            # Make foo a device function
+            foo_device = cuda.jit(device=True)(foo)
+
+            # The inner signature: int32 -> int32
+            inner_sig = types.int32(types.int32)
+
+            # The wrapper signature: void(void*, void*) - input ptr, output ptr
+            wrapper_sig = types.void(types.voidptr, types.voidptr)
+
+            @intrinsic
+            def wrapper_impl(typingctx, arg0, arg1):
+                def codegen(context, builder, sig, args):
+                    input_ptr, output_ptr = args
+
+                    # Cast input void* to int32*, load value
+                    int32_llvm_type = context.get_value_type(types.int32)
+                    typed_input_ptr = builder.bitcast(
+                        input_ptr, int32_llvm_type.as_pointer()
+                    )
+                    input_val = builder.load(typed_input_ptr)
+
+                    # Call the inner function
+                    cres = context.compile_subroutine(
+                        builder, foo_device, inner_sig, caching=False
+                    )
+
+                    # Wrapper function is compiled with cabi, but inner function
+                    # is compiled with numba-abi. So cres should have CUDACallConv.
+                    assert isinstance(cres.fndesc.call_conv, CUDACallConv)
+
+                    result = context.call_internal(
+                        builder, cres.fndesc, inner_sig, [input_val]
+                    )
+
+                    # Cast output void* to int32*, store result
+                    typed_output_ptr = builder.bitcast(
+                        output_ptr, int32_llvm_type.as_pointer()
+                    )
+                    builder.store(result, typed_output_ptr)
+
+                    return context.get_dummy_value()
+
+                return wrapper_sig, codegen
+
+            def wrapper_func(input_ptr, output_ptr):
+                return wrapper_impl(input_ptr, output_ptr)
+
+            return wrapper_func, wrapper_sig
+
+        wrapper, wrapper_sig = create_void_ptr_wrapper()
+
+        cuda.compile(wrapper, wrapper_sig.args, output="ltoir")
+
+    def test_compile_CABI_calling_device_function_returning_optional(self):
+        # Exercise a CABI caller invoking a Numba ABI callee that can return
+        # None through Optional[int32]
+        def maybe_none(x):
+            if x > 0:
+                return x + 1
+            else:
+                return
+
+        maybe_none_device = cuda.jit(device=True)(maybe_none)
+
+        def wrapper_func(x):
+            return maybe_none_device(x)
+
+        # Compile a CABI wrapper that calls into a Numba-ABI callee returning
+        # Optional[int32]. Successful compilation exercises the ABI boundary.
+        cuda.compile(
+            wrapper_func, types.int32(types.int32), output="ltoir", abi="c"
+        )
+
+    def test_compile_complex_div_c_abi(self):
+        # Reproducer from gh-789
+        # https://github.com/NVIDIA/numba-cuda/issues/789
+        def div_by_2(x):
+            return x / 2
+
+        sig = types.complex128(types.complex128)
+        cuda.compile(div_by_2, sig, device=True, abi="c")
+
 
 @skip_on_cudasim("Compilation unsupported in the simulator")
 class TestCompileForCurrentDevice(CUDATestCase):
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ir_utils.py b/numba_cuda/numba/cuda/tests/cudapy/test_ir_utils.py
index 88c1577ae..ce941440d 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_ir_utils.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_ir_utils.py
@@ -57,6 +57,7 @@ def mk_pipeline(
                     from numba.cuda.descriptor import cuda_target
 
                     target_context = cuda_target.target_context
+
                 return cls(
                     typing_context,
                     target_context,
diff --git a/numba_cuda/numba/cuda/typing/templates.py b/numba_cuda/numba/cuda/typing/templates.py
index 42dda278c..c33483e5e 100644
--- a/numba_cuda/numba/cuda/typing/templates.py
+++ b/numba_cuda/numba/cuda/typing/templates.py
@@ -689,6 +689,7 @@ def generic(self, args, kws):
             # spoof a compiler pipline like the one that will be in use
             tyctx = fcomp.targetdescr.typing_context
             tgctx = fcomp.targetdescr.target_context
+
             compiler_inst = fcomp.pipeline_class(
                 tyctx,
                 tgctx,