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 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))