BUG: Fix CUDADispatcher.compile_device ignoring abi flag

numba_cuda/numba/cuda/dispatcher.py

@@ -2053,6 +2053,7 @@ def compile_device(self, args, return_type=None):
                 forceinline = self.targetoptions.get("forceinline")
                 inline = self.targetoptions.get("inline", "never")
                 fastmath = self.targetoptions.get("fastmath")
+                abi = self.targetoptions.get("abi", "numba")
 
                 nvvm_options = {
                     "opt": 3 if self.targetoptions.get("opt") else 0,
@@ -2074,6 +2075,7 @@ def compile_device(self, args, return_type=None):
                     fastmath=fastmath,
                     nvvm_options=nvvm_options,
                     cc=cc,
+                    abi=abi,
                 )
                 self.overloads[args] = cres
 

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

@@ -0,0 +1,38 @@
+# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-2-Clause
+
+import re
+import unittest
+
+import numpy as np
+
+from numba import cuda, types
+from numba.cuda.compiler import _compile_pyfunc_with_fixup
+from numba.cuda.testing import CUDATestCase
+from numba.cuda.tests.support import override_config
+
+
+class TestCudaJitABI(CUDATestCase):
+    """
+    Tests the jit decorator with abi set to "C"
+    """
+
+    def test_abi_c(self):
+        def normalize_llvm(s):
+            return re.sub(
+                r'(; ModuleID = ".*?\$)\d+(")',
+                r"\g<1>X\g<2>",
+                s,
+                count=1,
+            )
+
+        def foo(a, b):
+            return a + b
+
+        sig = (types.int8, types.int8)
+
+        x = cuda.jit(sig, device=True, abi="c")(foo).inspect_llvm()[(sig)]
+        y = _compile_pyfunc_with_fixup(foo, sig=sig, abi="c", device=True)[
+            0
+        ].get_llvm_str()
+        self.assertTrue(normalize_llvm(x) == normalize_llvm(y))