When I use cuda.jit(..., device=True, abi="c").inspect_llvm(...), it appears the resulting function doesn't follow the C ABI calling convention (note, pointer return):
from numba import cuda, types
def foo(a, b):
return a + b
sig = (types.int8, types.int8)
print(cuda.jit(sig, device=True, abi="c")(foo).inspect_llvm()[(sig)])output:
; ModuleID = "foo$7"
target triple = "nvptx64-nvidia-cuda"
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
@"_ZN08NumbaEnv8__main__3fooB2v7B100cw51cXTLSUwv1sDUaKthoaNgqamjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdSBL3UrBWEdBHRauQC7YhbWaAA_3d_3dEaa" = common global i8* null
define i32 @"_ZN8__main__3fooB2v7B100cw51cXTLSUwv1sDUaKthoaNgqamjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdSBL3UrBWEdBHRauQC7YhbWaAA_3d_3dEaa"(i64* %".ret", i8 %"a", i8 %"b")
{
entry:
br label %"B0"
B0:
%".5" = sext i8 %"a" to i64
%".6" = sext i8 %"b" to i64
%".7" = add nsw i64 %".5", %".6"
store i64 %".7", i64* %".ret"
ret i32 0
}
!nvvmir.version = !{ !0 }
!0 = !{ i32 2, i32 0, i32 3, i32 2 }
If I reach for numba internals to get the LLVM IR instead, it works correctly (note, return by value)
from numba.cuda import types
from numba.cuda.compiler import _compile_pyfunc_with_fixup
def foo(a, b):
return a + b
sig = (types.int8, types.int8)
print(_compile_pyfunc_with_fixup(foo, sig=sig, abi="c", device=True)[0].get_llvm_str())prints:
; ModuleID = "foo$2"
target triple = "nvptx64-nvidia-cuda"
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
define i64 @"foo"(i8 %"arg.a", i8 %"arg.b")
{
entry:
br label %"B0"
B0:
%".4" = sext i8 %"arg.a" to i64
%".5" = sext i8 %"arg.b" to i64
%".6" = add nsw i64 %".4", %".5"
ret i64 %".6"
}
!nvvmir.version = !{ !0 }
!0 = !{ i32 2, i32 0, i32 3, i32 2 }
When I use
cuda.jit(..., device=True, abi="c").inspect_llvm(...), it appears the resulting function doesn't follow the C ABI calling convention (note, pointer return):output:
If I reach for numba internals to get the LLVM IR instead, it works correctly (note, return by value)
prints: