From abd36f730dee16b98e8e7ac7a239b44398b63e15 Mon Sep 17 00:00:00 2001
From: Andres Guzman-Ballen <acosmicflamingo@gmail.com>
Date: Thu, 28 May 2026 20:34:13 -0500
Subject: [PATCH 1/2] BUG: Fix CUDADispatcher.compile_device ignoring abi flag

---
 numba_cuda/numba/cuda/dispatcher.py | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/numba_cuda/numba/cuda/dispatcher.py b/numba_cuda/numba/cuda/dispatcher.py
index 3b3b30a0b..e7a56b7bf 100644
--- a/numba_cuda/numba/cuda/dispatcher.py
+++ b/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
 

From 0b2f14403b503bb21205b8f64693d86b720db32f Mon Sep 17 00:00:00 2001
From: Andres Guzman-Ballen <acosmicflamingo@gmail.com>
Date: Wed, 3 Jun 2026 17:16:29 -0500
Subject: [PATCH 2/2] Add test asserting cuda.jit(abi='c') output

---
 .../cuda/tests/cudapy/test_cuda_jit_abi.py    | 38 +++++++++++++++++++
 1 file changed, 38 insertions(+)
 create mode 100644 numba_cuda/numba/cuda/tests/cudapy/test_cuda_jit_abi.py

diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_cuda_jit_abi.py b/numba_cuda/numba/cuda/tests/cudapy/test_cuda_jit_abi.py
new file mode 100644
index 000000000..a30a4b130
--- /dev/null
+++ b/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))