feat: add math nextafter

Author: kaeun97

docs/source/user/cudapysupported.rst

@@ -224,6 +224,7 @@ The following functions from the :mod:`math` module are supported:
 * :func:`math.log2`
 * :func:`math.log10`
 * :func:`math.log1p`
+* :func:`math.nextafter`
 * :func:`math.sqrt`
 * :func:`math.remainder`
 * :func:`math.pow`

numba_cuda/numba/cuda/cudamath.py

@@ -82,6 +82,7 @@ class Math_hypot(ConcreteTemplate):
 
 @infer_global(math.copysign)
 @infer_global(math.fmod)
+@infer_global(math.nextafter)
 class Math_binary(ConcreteTemplate):
     cases = [
         signature(types.float32, types.float32, types.float32),

numba_cuda/numba/cuda/mathimpl.py

@@ -64,9 +64,11 @@
 binarys += [("fmod", "fmodf", math.fmod)]
 binarys += [("hypot", "hypotf", math.hypot)]
 binarys += [("remainder", "remainderf", math.remainder)]
+binarys += [("nextafter", "nextafterf", math.nextafter)]
 
 binarys_fastmath = {}
 binarys_fastmath["powf"] = "fast_powf"
+binarys_fastmath["nextafterf"] = "fast_nextafterf"
 
 
 @lower(math.isinf, types.Integer)

numba_cuda/numba/cuda/tests/cudapy/test_fastmath.py

@@ -5,7 +5,7 @@
 from dataclasses import dataclass, field
 from numba import cuda, float32
 from numba.cuda.compiler import compile_ptx_for_current_device, compile_ptx
-from math import cos, sin, tan, exp, log, log10, log2, pow, tanh
+from math import cos, sin, tan, exp, log, log10, log2, pow, tanh, nextafter
 from operator import truediv
 import numpy as np
 from numba.cuda.testing import CUDATestCase, skip_on_cudasim, skip_unless_cc_75
@@ -179,6 +179,15 @@ def test_powf(self):
             ),
         )
 
+    def test_nextafterf(self):
+        self._test_fast_math_binary(
+            nextafter,
+            FastMathCriterion(
+                fast_expected=["lg2.approx.ftz.f32 "],  # FIX
+                prec_unexpected=["lg2.approx.ftz.f32 "],  # FIX
+            ),
+        )
+
     def test_divf(self):
         self._test_fast_math_binary(
             truediv,

numba_cuda/numba/cuda/tests/cudapy/test_math.py

@@ -138,6 +138,11 @@ def math_remainder(A, B, C):
     C[i] = math.remainder(A[i], B[i])
 
 
+def math_nextafter(A, B, C):
+    i = cuda.grid(1)
+    C[i] = math.nextafter(A[i], B[i])
+
+
 def math_sqrt(A, B):
     i = cuda.grid(1)
     B[i] = math.sqrt(A[i])
@@ -594,6 +599,15 @@ def test_0_0(r, x, y):
         test_0_0[1, 1](r, 0, 0)
         self.assertTrue(np.isnan(r[0]))
 
+    # ---------------------------------------------------------------------------
+    # test_math_nextafter
+
+    def test_math_nextafter(self):
+        self.binary_template_float32(math_nextafter, np.nextafter, start=1e-11)
+        self.binary_template_float64(math_remainder, np.remainder, start=1e-11)
+        self.binary_template_int64(math_remainder, np.remainder, start=1)
+        self.binary_template_uint64(math_remainder, np.remainder, start=1)
+
     # ---------------------------------------------------------------------------
     # test_math_sqrt