Skip to content
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
21 changes: 8 additions & 13 deletions numba_cuda/numba/cuda/tests/cudapy/test_boolean.py
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@
# SPDX-License-Identifier: BSD-2-Clause

import numpy as np
from numba.cuda.testing import unittest, CUDATestCase
from numba.cuda.testing import CUDATestCase
from numba import cuda


Expand All @@ -13,15 +13,10 @@ def boolean_func(A, vertial):
A[0] = 321


class TestCudaBoolean(CUDATestCase):
def test_boolean(self):
func = cuda.jit("void(float64[:], bool_)")(boolean_func)
A = np.array([0], dtype="float64")
func[1, 1](A, True)
self.assertTrue(A[0] == 123)
func[1, 1](A, False)
self.assertTrue(A[0] == 321)


if __name__ == "__main__":
unittest.main()
def test_boolean():
func = cuda.jit("void(float64[:], bool_)")(boolean_func)
A = np.array([0], dtype="float64")
func[1, 1](A, True)
assert A[0] == 123
func[1, 1](A, False)
assert A[0] == 321
138 changes: 92 additions & 46 deletions numba_cuda/numba/cuda/tests/cudapy/test_minmax.py
Original file line number Diff line number Diff line change
Expand Up @@ -2,10 +2,11 @@
# SPDX-License-Identifier: BSD-2-Clause

import numpy as np
import pytest

from numba import cuda
from numba.cuda import float64
from numba.cuda.testing import unittest, CUDATestCase, skip_on_cudasim
from numba.cuda.testing import skip_on_cudasim


def builtin_max(A, B, C):
Expand All @@ -26,53 +27,98 @@ def builtin_min(A, B, C):
C[i] = float64(min(A[i], B[i]))


@skip_on_cudasim("Tests PTX emission")
class TestCudaMinMax(CUDATestCase):
def _run(
self,
kernel,
numpy_equivalent,
ptx_instruction,
dtype_left,
dtype_right,
n=5,
):
kernel = cuda.jit(kernel)

c = np.zeros(n, dtype=np.float64)
a = np.arange(n, dtype=dtype_left) + 0.5
b = np.full(n, fill_value=2, dtype=dtype_right)

kernel[1, c.shape](a, b, c)
np.testing.assert_allclose(c, numpy_equivalent(a, b))

ptx = next(p for p in kernel.inspect_asm().values())
self.assertIn(ptx_instruction, ptx)

def test_max_f8f8(self):
self._run(builtin_max, np.maximum, "max.f64", np.float64, np.float64)
def _run(
kernel,
numpy_equivalent,
ptx_instruction,
dtype_left,
dtype_right,
n=5,
):
kernel = cuda.jit(kernel)

def test_max_f4f8(self):
self._run(builtin_max, np.maximum, "max.f64", np.float32, np.float64)
c = np.zeros(n, dtype=np.float64)
a = np.arange(n, dtype=dtype_left) + 0.5
b = np.full(n, fill_value=2, dtype=dtype_right)

def test_max_f8f4(self):
self._run(builtin_max, np.maximum, "max.f64", np.float64, np.float32)
kernel[1, c.shape](a, b, c)
np.testing.assert_allclose(c, numpy_equivalent(a, b))

def test_max_f4f4(self):
self._run(builtin_max, np.maximum, "max.f32", np.float32, np.float32)
ptx = next(p for p in kernel.inspect_asm().values())
assert ptx_instruction in ptx

def test_min_f8f8(self):
self._run(builtin_min, np.minimum, "min.f64", np.float64, np.float64)

def test_min_f4f8(self):
self._run(builtin_min, np.minimum, "min.f64", np.float32, np.float64)

def test_min_f8f4(self):
self._run(builtin_min, np.minimum, "min.f64", np.float64, np.float32)

def test_min_f4f4(self):
self._run(builtin_min, np.minimum, "min.f32", np.float32, np.float32)


if __name__ == "__main__":
unittest.main()
@skip_on_cudasim("Tests PTX emission")
@pytest.mark.parametrize(
"kernel,numpy_equivalent,ptx_instruction,dtype_left,dtype_right",
[
pytest.param(
builtin_max,
np.maximum,
"max.f64",
np.float64,
np.float64,
id="max_f8f8",
),
pytest.param(
builtin_max,
np.maximum,
"max.f64",
np.float32,
np.float64,
id="max_f4f8",
),
pytest.param(
builtin_max,
np.maximum,
"max.f64",
np.float64,
np.float32,
id="max_f8f4",
),
pytest.param(
builtin_max,
np.maximum,
"max.f32",
np.float32,
np.float32,
id="max_f4f4",
),
pytest.param(
builtin_min,
np.minimum,
"min.f64",
np.float64,
np.float64,
id="min_f8f8",
),
pytest.param(
builtin_min,
np.minimum,
"min.f64",
np.float32,
np.float64,
id="min_f4f8",
),
pytest.param(
builtin_min,
np.minimum,
"min.f64",
np.float64,
np.float32,
id="min_f8f4",
),
pytest.param(
builtin_min,
np.minimum,
"min.f32",
np.float32,
np.float32,
id="min_f4f4",
),
],
)
def test_minmax(
kernel, numpy_equivalent, ptx_instruction, dtype_left, dtype_right
):
_run(kernel, numpy_equivalent, ptx_instruction, dtype_left, dtype_right)
Loading