Add intent-based ref handling and out-return support for bindings (#278)

In ISO C++, whether an argument is mutated or not is not directly
inferable from the signature itself. For example,

```c++
void set_as_42(int &a) { a = 42; }
const int& operator+ (const int&, const int&);
```

both accepts a reference as input, but whether the function modifies the
argument is unknown until further analysis into the body is performed.
This creates confusions for language bindings on whether the arguments
should be passed by reference or value, as other language may dictate
the argument passing semantics differently from C++.

Certain compiler provides additional annotation features to denote them.
Such as
[SAL](https://learn.microsoft.com/en-us/cpp/code-quality/best-practices-and-examples-sal?view=msvc-170).

In this PR, Numbast introduces an `argument intent` option that allows
user to configure argument passing mode on per-function, per-parameter
basis. Per argument, the following options are available:
```
- in                       (default, arguments are passed as-is by value)
- inout_ptr          (argument is both used as an input mutated, exposed in Numba as a CPointer type to itself)
- out_ptr             (argument is used to store output, thus mutated; exposed in Numba as a CPointer type to itself)
- out_return       (return arguments are pre-allocated on stack, returned in Numba bindings)
```

Take the following C++ signature as an example:
```c++
// A C++ function that processes `input`, writes result to `out`, and returns an exit code.
int mutative(int &out, int input);
```

A typical argument intent setup for Numbast looks like:
```
{"mutative": {"out": "out_return"}}
```
This indicates that argument `out` is returned as the functions return
value in the corresponding binding. And since this function already has
a return value, the binding will now return a tuple of ints, with first
corresponds to the exit code, and the second corresponds to the result.
The Python binding signature:

```python
def mutative(input: numba.int32) -> Tuple[numba.int32, numba.int32]: ...
```

Alternatively, if intent is set to: `"out": "out_ptr"`, the signature
becomes:

```python
def mutative(out: CPointer(int32), input: int32) -> int32: ...
```




<!-- This is an auto-generated comment: release notes by coderabbit.ai
-->
## Summary by CodeRabbit

* **New Features**
* Introduced a per-function argument-intent system to control parameter
semantics (in, inout, out, out-return) and pointer-passing behavior.
* Support for returning multiple values via out-parameters alongside
regular returns.
* Static binding generation and rendering now accept and propagate
per-function intent overrides.

* **Tests**
* Added end-to-end tests, fixtures and test data exercising
out-parameters, in/out-pointer semantics, out-return behavior, and
mutative device functions.

<sub>✏️ Tip: You can customize this high-level summary in your review
settings.</sub>
<!-- end of auto-generated comment: release notes by coderabbit.ai -->

---------

Co-authored-by: Michael Wang <isVoid@users.noreply.github.com>

Author: isVoid

.pre-commit-config.yaml

@@ -37,7 +37,7 @@ repos:
       - id: codespell
         additional_dependencies:
           - tomli
-        args: ["--toml", "pyproject.toml"]
+        args: ["--toml", "pyproject.toml", "--ignore-words-list", "inout"]
   - repo: https://github.com/google/yamlfmt
     rev: v0.16.0
     hooks:

numbast/src/numbast/args.py

@@ -7,7 +7,10 @@
 
 
 def prepare_ir_types(
-    context: CUDATargetContext, argtys: list[ir.Type]
+    context: CUDATargetContext,
+    argtys: list[ir.Type],
+    *,
+    pass_ptr_mask: list[bool] | None = None,
 ) -> list[ir.Type]:
     """
     Prepare IR types for passing arguments via pointers in function calls.
@@ -28,4 +31,22 @@ def prepare_ir_types(
     list[ir.Type]
         List of pointer types wrapping the value types of each argument.
     """
-    return [ir.PointerType(context.get_value_type(argty)) for argty in argtys]
+    if pass_ptr_mask is None:
+        pass_ptr_mask = [False] * len(argtys)
+
+    if len(pass_ptr_mask) != len(argtys):
+        raise ValueError(
+            f"pass_ptr_mask length ({len(pass_ptr_mask)}) must match argtys length ({len(argtys)})"
+        )
+
+    ir_types: list[ir.Type] = []
+    for argty, passthrough in zip(argtys, pass_ptr_mask):
+        vty = context.get_value_type(argty)
+        if passthrough and isinstance(vty, ir.PointerType):
+            # Pass pointer-typed values directly (e.g. C++ T& mapped to CPointer(T))
+            ir_types.append(vty)
+        else:
+            # Default ABI: pass pointer-to-value
+            ir_types.append(ir.PointerType(vty))
+
+    return ir_types

numbast/src/numbast/callconv.py

@@ -1,4 +1,10 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
 from numbast.args import prepare_ir_types
+from numbast.intent import IntentPlan
+
+# NBST:BEGIN_CALLCONV
 from numba.cuda import types, cgutils
 
 from llvmlite import ir
@@ -36,27 +42,133 @@ def __call__(self, builder, context, sig, args):
 
 
 class FunctionCallConv(BaseCallConv):
+    def __init__(
+        self,
+        itanium_mangled_name: str,
+        shim_writer: object,
+        shim_code: str,
+        *,
+        arg_is_ref: list[bool] | None = None,
+        intent_plan: IntentPlan | None = None,
+        out_return_types: list[types.Type] | None = None,
+        cxx_return_type: types.Type | None = None,
+    ):
+        super().__init__(itanium_mangled_name, shim_writer, shim_code)
+        self._arg_is_ref = list(arg_is_ref) if arg_is_ref is not None else None
+        self._intent_plan = intent_plan
+        self._out_return_types = (
+            list(out_return_types) if out_return_types is not None else None
+        )
+        self._cxx_return_type = cxx_return_type
+
     def _lower_impl(self, builder, context, sig, args):
-        return_type = sig.return_type
+        # Numba-visible return type may differ from the underlying C++ return type
+        # when out_return parameters are enabled (tuple returns, etc.).
+        cxx_return_type = (
+            self._cxx_return_type
+            if self._cxx_return_type is not None
+            else sig.return_type
+        )
         # 1. Prepare return value pointer
-        if return_type == types.void:
+        if cxx_return_type == types.void:
             # Void return type in C++ is shimmed as int& ignored
             retval_ty = ir.IntType(32)
             retval_ptr = builder.alloca(retval_ty, name="ignored")
         else:
-            retval_ty = context.get_value_type(return_type)
+            retval_ty = context.get_value_type(cxx_return_type)
             retval_ptr = builder.alloca(retval_ty, name="retval")
 
         # 2. Prepare arguments
-        arg_pointer_types = prepare_ir_types(context, sig.args)
-
-        # All arguments are passed by pointer
-        ptrs = [
-            cgutils.alloca_once(builder, context.get_value_type(argty))
-            for argty in sig.args
-        ]
-        for ptr, argty, arg in zip(ptrs, sig.args, args):
-            builder.store(arg, ptr, align=getattr(argty, "alignof_", None))
+        if self._intent_plan is None:
+            pass_ptr_mask = (
+                self._arg_is_ref
+                if self._arg_is_ref is not None
+                else [False] * len(sig.args)
+            )
+            arg_pointer_types = prepare_ir_types(
+                context, sig.args, pass_ptr_mask=pass_ptr_mask
+            )
+        else:
+            plan = self._intent_plan
+            if len(sig.args) != len(plan.visible_param_indices):
+                raise ValueError(
+                    "Signature args do not match intent plan visible params: "
+                    f"sig has {len(sig.args)} args but plan expects {len(plan.visible_param_indices)}"
+                )
+            if len(plan.pass_ptr_mask) != len(sig.args):
+                raise ValueError(
+                    "Intent plan pass_ptr_mask length does not match signature args length: "
+                    f"{len(plan.pass_ptr_mask)} != {len(sig.args)}"
+                )
+            if plan.out_return_indices:
+                if self._out_return_types is None:
+                    raise ValueError(
+                        "out_return intent plan requires out_return_types to be provided"
+                    )
+                if len(self._out_return_types) != len(plan.out_return_indices):
+                    raise ValueError(
+                        "out_return_types length does not match intent plan out_return_indices: "
+                        f"{len(self._out_return_types)} != {len(plan.out_return_indices)}"
+                    )
+            arg_pointer_types = []  # computed below alongside ptrs
+
+        # ABI:
+        # - default: pass pointer-to-value to shim (alloca + store)
+        # - for C++ reference args mapped to CPointer(T): pass pointer value directly
+        ptrs = []
+        out_return_ptrs: list[tuple[types.Type, ir.Value]] = []
+        if self._intent_plan is None:
+            for argty, arg, passthrough in zip(sig.args, args, pass_ptr_mask):
+                vty = context.get_value_type(argty)
+                if passthrough and isinstance(vty, ir.PointerType):
+                    ptrs.append(arg)
+                else:
+                    ptr = cgutils.alloca_once(builder, vty)
+                    builder.store(
+                        arg, ptr, align=getattr(argty, "alignof_", None)
+                    )
+                    ptrs.append(ptr)
+        else:
+            plan = self._intent_plan
+            n_orig = len(plan.intents)
+            # Map original parameter index -> visible signature position / out_return position
+            orig_to_vis = [None] * n_orig
+            for vis_pos, orig_idx in enumerate(plan.visible_param_indices):
+                orig_to_vis[orig_idx] = vis_pos
+            orig_to_out = [None] * n_orig
+            for out_pos, orig_idx in enumerate(plan.out_return_indices):
+                orig_to_out[orig_idx] = out_pos
+
+            for orig_idx in range(n_orig):
+                out_pos = orig_to_out[orig_idx]
+                if out_pos is not None:
+                    out_nbty = self._out_return_types[out_pos]
+                    vty = context.get_value_type(out_nbty)
+                    ptr = cgutils.alloca_once(builder, vty)
+                    ptrs.append(ptr)
+                    arg_pointer_types.append(ir.PointerType(vty))
+                    out_return_ptrs.append((out_nbty, ptr))
+                    continue
+
+                vis_pos = orig_to_vis[orig_idx]
+                if vis_pos is None:
+                    raise ValueError(
+                        f"Internal error: original param {orig_idx} is neither visible nor out_return"
+                    )
+                argty = sig.args[vis_pos]
+                arg = args[vis_pos]
+                passthrough = bool(plan.pass_ptr_mask[vis_pos])
+                vty = context.get_value_type(argty)
+                if passthrough and isinstance(vty, ir.PointerType):
+                    ptrs.append(arg)
+                    arg_pointer_types.append(vty)
+                else:
+                    ptr = cgutils.alloca_once(builder, vty)
+                    builder.store(
+                        arg, ptr, align=getattr(argty, "alignof_", None)
+                    )
+                    ptrs.append(ptr)
+                    arg_pointer_types.append(ir.PointerType(vty))
 
         # 3. Declare shim
         # Shim signature: int (retval_type*, arg0_type*, ...)
@@ -71,9 +183,39 @@ def _lower_impl(self, builder, context, sig, args):
         builder.call(fn, (retval_ptr, *ptrs))
 
         # 5. Return
-        if return_type == types.void:
-            return None
-        else:
+        if (
+            self._intent_plan is None
+            or not self._intent_plan.out_return_indices
+        ):
+            if cxx_return_type == types.void:
+                return None
             return builder.load(
-                retval_ptr, align=getattr(return_type, "alignof_", None)
+                retval_ptr, align=getattr(cxx_return_type, "alignof_", None)
+            )
+
+        # out_return enabled: return either a value or a tuple (ret, out1, out2, ...)
+        ret_vals: list[ir.Value] = []
+        if cxx_return_type != types.void:
+            ret_vals.append(
+                builder.load(
+                    retval_ptr, align=getattr(cxx_return_type, "alignof_", None)
+                )
+            )
+        for out_ty, out_ptr in out_return_ptrs:
+            ret_vals.append(
+                builder.load(out_ptr, align=getattr(out_ty, "alignof_", None))
+            )
+
+        # If Numba-visible return is a tuple, use context.make_tuple.
+        # Otherwise (void + single out), return the single out value directly.
+        if hasattr(sig.return_type, "types"):
+            return context.make_tuple(builder, sig.return_type, ret_vals)
+        if len(ret_vals) != 1:
+            raise ValueError(
+                "Non-tuple return type requires exactly one return value; "
+                f"got {len(ret_vals)}"
             )
+        return ret_vals[0]
+
+
+# NBST:END_CALLCONV