[BUG] `types.void` is translated to `i8*` instead of LLVM `void` type, causing ABI and LTO mismatches

[ZzEeKkAa]

Description

When using Numba's types.void (e.g., as a return type for a device function or in signatures), Numba's datamodel translates it to an LLVM i8* (opaque pointer ptr in LLVM 15+) instead of the actual LLVM void type (ir.VoidType()).

Note: This issue originates in the core numba package itself (specifically in numba.core.datamodel where types.void is registered to OpaqueModel), but it has a severe impact on numba-cuda when linking device code.

This causes the generated LLVM IR to treat void as a pointer type. In the Numba calling convention, this results in device functions taking an i8** return pointer argument rather than simply returning void. While this might accidentally work at runtime on some architectures, it breaks strict ABI requirements. More importantly, it causes fatal errors during Link Time Optimization (LTO) when linking Numba-generated IR with C/C++ device code that correctly declares the function as returning void.

Reproducer

import numba
from numba import types, cuda
from numba.core import datamodel
from numba.cuda.descriptor import cuda_target

# 1. Numba's default void type is translated to i8* via OpaqueModel
void_type = types.void
default_model = datamodel.default_manager.lookup(void_type)
print(f"LLVM type for void: {default_model.get_value_type()}")
# Output: LLVM type for void: i8*

# 2. CUDA Context return types also reflect this mismatch
ctx = cuda_target.target_context
print(f"Context return type: {ctx.get_return_type(types.void)}")
# Output: Context return type: i8*

# 3. Demonstrate this in a compiled CUDA device function
@cuda.jit(device=True)
def foo():
    pass

@cuda.jit
def kernel():
    foo()

kernel[1, 1]()

# The device function signature takes an `i8**` return pointer instead of returning `void`
for k, v in foo.inspect_llvm().items():
    for line in v.splitlines():
        if line.startswith("define") and "foo" in line:
            print(line)
# Output: define linkonce_odr i32 @"_ZN8__main__3foo..."(i8** %".ret")

Expected Behavior

types.void should be mapped to a model that returns ir.VoidType() (i.e., void in LLVM IR), so that functions returning void have the correct LLVM IR signature without an unnecessary return pointer argument:

define void @"_ZN8__main__3foo..."()

And context return types should reflect void instead of i8*.

Actual Behavior

types.void is mapped to OpaqueModel, which returns ir.IntType(8).as_pointer() (ptr / i8*).

Because of this, the calling convention treats the return type as a pointer, generating an i8** return argument:

define linkonce_odr i32 @"_ZN8__main__3foo..."(i8** %".ret")

Impact

When exporting Numba JIT code to LLVM IR and linking it via LTO with external C/C++ device libraries, the LLVM linker detects a signature mismatch between the C++ declaration (void foo()) and the Numba definition (i32 foo(i8**)). This results in LTO linking failures and ABI correctness issues for CUDA device code.

[isVoid]

@ZzEeKkAa I realized that Numba intentionally maps a void type to i8*. In LLVM, void is strictly a return-type concept and cannot be used as an argument type. As a result, if a developer defines a function with a signature like signature(types.int32, types.NoneType), Numba lowers the NoneType argument to a pointer placeholder. Representing NoneType as a pointer simplifies much of the implementation.

In the new fix PR, I introduced a special case in the NoneTypeDataModel to distinguish between value_type and return_type. Specifically, NoneType is lowered to LLVM void only when used as a return type. For value_type, it remains i8* to preserve compatibility. This should address the problem outlined in the description as the example expects NoneType is mapped to llvm VoidType as a return type.

[gmarkall]

When exporting Numba JIT code to LLVM IR and linking it via LTO with external C/C++ device libraries, the LLVM linker detects a signature mismatch between the C++ declaration (void foo()) and the Numba definition (i32 foo(i8**)). This results in LTO linking failures and ABI correctness issues for CUDA device code.

It sounds like the C++ code is incorrect here - I think it should declare the function as i32 foo(i8**). A null pointer (or any pointer value, really) would need to be passed in the first argument. The i32 return type is the status indicator in the Numba calling convention, so the return should not be void.

Have you got a reproducer that illustrates the C++ linking problem?

[Jlisowskyy]

The main motivation stems from our use case, where we need to directly invoke cabi functions originating from C++, which are used by multiple compilers. Perhaps numba-cuda could instead offer a cleaner mechanism for calling external cabi functions and handle that edge case internally - something akin to declare_device, but for the cabi?

Edit: It looks like we can already call external cabi functions with declare_device, I missed that. Could the change perhaps be scoped to this function only, or to the cabi calling convention only? That way, it wouldn't affect the numba calling convention.

[gmarkall]

Edit: It looks like we can already call external cabi functions with declare_device, I missed that. Could the change perhaps be scoped to this function only, or to the cabi calling convention only? That way, it wouldn't affect the numba calling convention.

I think this is already covered by adding abi="c" to declare_device(): https://nvidia.github.io/numba-cuda/user/cuda_ffi.html#numba.cuda.declare_device

Is there additional flexibility needed that isn't present here?

[Jlisowskyy]

With the current assumption that types.void is always mapped to i8*, a signature mismatch can occur when declaring external cabi functions that return void.

For example, given an external cabi function with the following signature:

declare void @"foo"(i8* %"arg.arg0")

Using cuda.declare_device to declare and call it:

func = cuda.declare_device(
    "foo",
    types.void(types.CPointer(types.int8)),
    link=CODE,
    abi="c",
)

produces the following declaration in the resulting LTOIR:

declare i8* @"foo"(i8* %"arg.arg0")

This mismatch leads to linking errors. The issue is easy to overlook because cuda.declare_device works correctly for all functions that actually return a value, only void-returning functions are affected. Furthermore, the linker does not report this mismatch: it silently generates malformed code