Regression (0.66.0): two-argument max/min fail on CUDA after #10543 refactored them into *args overloads

CUDA: two-argument max and min fail to compile on 0.66.0

Summary

On numba 0.66.0, a @cuda.jit kernel that calls max(a, b) or min(a, b) fails to compile. The compiler raises a TypeError during an untyped pass. The same kernel compiles and runs on numba 0.65.0. The CPU target (@njit) accepts the two-argument max and min on both versions. The CUDA target breaks.

Error

TypeError: Signature mismatch: 2 argument types given, but function takes 1 arguments

The error originates from the FixupArgs pass in numba/core/untyped_passes.py.

Environment

• Fails on: numba 0.66.0, llvmlite 0.48.0
• Passes on: numba 0.65.0, llvmlite 0.47.0
• Python 3.13, Linux aarch64, CUDA driver 13.0, one device
• The CPU target is unaffected on both versions.

Example script

The following code imports only numba and numpy. It runs a CPU control first,then the CUDA kernels. The CPU control passes on both versions, so the script shows on its own that the break is specific to the CUDA target.

import sys
import numpy as np
import numba
from numba import cuda, njit

@njit
def cpu_max_min(x):
    return max(0.0, x), min(0.0, x)

@cuda.jit
def cuda_max_kernel(out, x):
    out[0] = max(numba.float64(0.0), x[0])

@cuda.jit
def cuda_min_kernel(out, x):
    out[0] = min(numba.float64(0.0), x[0])

def run_cuda(kernel, x_value):
    out = cuda.to_device(np.zeros(1, dtype=np.float64))
    x = cuda.to_device(np.array([x_value], dtype=np.float64))
    kernel[1, 1](out, x)
    return float(out.copy_to_host()[0])

def main():
    print("numba", numba.__version__)
    hi, lo = cpu_max_min(-3.0)          # CPU control
    print("CPU:", hi, lo)               # 0.0 -3.0 on both versions
    if not cuda.is_available():
        return 2
    hi = run_cuda(cuda_max_kernel, -3.0)   # raises TypeError on 0.66.0
    lo = run_cuda(cuda_min_kernel, -3.0)
    print("CUDA:", hi, lo)
    return 0

if __name__ == "__main__":
    sys.exit(main())

Output on 0.66.0

numba 0.66.0
CPU: 0.0 -3.0
... TypeError: Signature mismatch: 2 argument types given, but function takes 1 arguments

Output on 0.65.0

numba 0.65.0
CPU: 0.0 -3.0
CUDA: 0.0 -3.0

Root cause

PR #10543 ("Refactored min/max functionality into overloads", commit abd1def2b) changed how numba types and lowers the max and min builtins. The PR removed the Max and Min typing templates from numba/core/typing/builtins.py. It also removed the @lower_builtin handlers from numba/cpython/builtins.py. In their place, the PR added @overload implementations that use *args:

@overload(max)
def ol_max(*x):
    for ty in x:
        if not isinstance(ty, types.Number):
            return None

    def impl(*x):          # this impl uses *args
        return max_vararg(x)
    return impl

The CUDA target cannot expand *args in a jitted function. This is the standing limitation in issue #6891 ("Cuda functions treat *args as a normal parameter"). The generated impl(*x) therefore looks like a one-parameter function to the CUDA target. A call max(a, b) supplies two arguments, so the FixupArgs pass reports a signature mismatch. The CPU target expands *args correctly, so @njit still works.

This report is not a duplicate of #6891. PR #10543 made a core builtin depend on the *args feature that CUDA does not support, turning a standing limitation into a regression for max and min.

Suggested resolution

1. Rewrite the max and min overloads so they do not use *args. A working two-argument form covers the common max(a, b) and min(a, b) case.
2. Add CUDA support for *args (issue #6891). This is the prerequisite for the current overload design to work on the CUDA target.