From 05303015a3fda0df91003144ffaa8e5b1a183aa8 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay Date: Tue, 20 Jan 2026 16:05:51 +0530 Subject: [PATCH 01/12] CUDA: add basic complex number operation tests\n\nAdds kernel-based tests for device-side complex64 addition, multiplication, and absolute value. Skips under cudasim. NumPy is used as the reference. Scope is intentionally limited to safe, well-defined CUDA operations. --- numba_cuda/numba/cuda/tests/test_complex.py | 85 +++++++++++++++++++++ 1 file changed, 85 insertions(+) create mode 100644 numba_cuda/numba/cuda/tests/test_complex.py diff --git a/numba_cuda/numba/cuda/tests/test_complex.py b/numba_cuda/numba/cuda/tests/test_complex.py new file mode 100644 index 000000000..f4e1f02e6 --- /dev/null +++ b/numba_cuda/numba/cuda/tests/test_complex.py @@ -0,0 +1,85 @@ +import numpy as np + +from numba import cuda +from numba.cuda.testing import CUDATestCase, skip_on_cudasim + + +@cuda.jit +def complex_add_kernel(a, b, out): + i = cuda.grid(1) + if i < out.size: + out[i] = a[i] + b[i] + + +@cuda.jit +def complex_mul_kernel(a, b, out): + i = cuda.grid(1) + if i < out.size: + out[i] = a[i] * b[i] + + +@cuda.jit +def complex_abs_kernel(inp, out): + i = cuda.grid(1) + if i < out.size: + out[i] = abs(inp[i]) + + +@skip_on_cudasim("Complex semantics differ under cudasim") +class TestCudaComplex(CUDATestCase): + + def _launch_1d(self, kernel, args, size): + threadsperblock = 128 + blockspergrid = (size + threadsperblock - 1) // threadsperblock + kernel[blockspergrid, threadsperblock](*args) + cuda.synchronize() + + + def test_complex_add(self): + a = np.array([1+2j, -3+4j, 5-6j], dtype=np.complex64) + b = np.array([2-1j, 4+3j, -5+6j], dtype=np.complex64) + out = np.zeros_like(a) + + da = cuda.to_device(a) + db = cuda.to_device(b) + dout = cuda.to_device(out) + + self._launch_1d(complex_add_kernel, (da, db, dout), a.size) + + np.testing.assert_array_equal( + dout.copy_to_host(), + a + b + ) + + + def test_complex_multiply(self): + a = np.array([1+2j, -3+4j, 5-6j], dtype=np.complex64) + b = np.array([2-1j, 4+3j, -5+6j], dtype=np.complex64) + out = np.zeros_like(a) + + da = cuda.to_device(a) + db = cuda.to_device(b) + dout = cuda.to_device(out) + + self._launch_1d(complex_mul_kernel, (da, db, dout), a.size) + + np.testing.assert_array_equal( + dout.copy_to_host(), + a * b + ) + + + def test_complex_abs(self): + inp = np.array([3+4j, 5-12j, -8+15j], dtype=np.complex64) + out = np.zeros(inp.size, dtype=np.float32) + + dinp = cuda.to_device(inp) + dout = cuda.to_device(out) + + self._launch_1d(complex_abs_kernel, (dinp, dout), inp.size) + + np.testing.assert_allclose( + dout.copy_to_host(), + np.abs(inp), + rtol=1e-6 + ) From dddd00bae96bc8635658a0cc368a270b278e4164 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 11:10:29 +0530 Subject: [PATCH 02/12] Update numba_cuda/numba/cuda/tests/test_complex.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- numba_cuda/numba/cuda/tests/test_complex.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/numba_cuda/numba/cuda/tests/test_complex.py b/numba_cuda/numba/cuda/tests/test_complex.py index f4e1f02e6..d8c065d4d 100644 --- a/numba_cuda/numba/cuda/tests/test_complex.py +++ b/numba_cuda/numba/cuda/tests/test_complex.py @@ -1,3 +1,6 @@ +# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: BSD-2-Clause + import numpy as np from numba import cuda From 3f013df541b52c5d7eefb7a170089b356f2eab64 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay Date: Wed, 21 Jan 2026 11:57:54 +0530 Subject: [PATCH 03/12] Move basic CUDA complex number operation tests to cudapy/test_complex.py\n\n- Merged device-side complex64 operation tests (add, multiply, abs) into cudapy/test_complex.py as TestCudaComplexBasicOps for organizational consistency\n- Removed standalone test_complex.py file --- .../numba/cuda/tests/cudapy/test_complex.py | 58 ++++++++++++ numba_cuda/numba/cuda/tests/test_complex.py | 88 ------------------- 2 files changed, 58 insertions(+), 88 deletions(-) delete mode 100644 numba_cuda/numba/cuda/tests/test_complex.py diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py index 8348ca96e..6b1159b55 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py @@ -1,3 +1,61 @@ +## --- Basic CUDA complex number operation tests (from test_complex.py) --- +import numpy as np +from numba import cuda +from numba.cuda.testing import CUDATestCase, skip_on_cudasim + +@cuda.jit +def complex_add_kernel(a, b, out): + i = cuda.grid(1) + if i < out.size: + out[i] = a[i] + b[i] + +@cuda.jit +def complex_mul_kernel(a, b, out): + i = cuda.grid(1) + if i < out.size: + out[i] = a[i] * b[i] + +@cuda.jit +def complex_abs_kernel(inp, out): + i = cuda.grid(1) + if i < out.size: + out[i] = abs(inp[i]) + +@skip_on_cudasim("Complex semantics differ under cudasim") +class TestCudaComplexBasicOps(CUDATestCase): + def _launch_1d(self, kernel, args, size): + threadsperblock = 128 + blockspergrid = (size + threadsperblock - 1) // threadsperblock + kernel[blockspergrid, threadsperblock](*args) + cuda.synchronize() + + def test_complex_add(self): + a = np.array([1+2j, -3+4j, 5-6j], dtype=np.complex64) + b = np.array([2-1j, 4+3j, -5+6j], dtype=np.complex64) + out = np.zeros_like(a) + da = cuda.to_device(a) + db = cuda.to_device(b) + dout = cuda.to_device(out) + self._launch_1d(complex_add_kernel, (da, db, dout), a.size) + np.testing.assert_array_equal(dout.copy_to_host(), a + b) + + def test_complex_multiply(self): + a = np.array([1+2j, -3+4j, 5-6j], dtype=np.complex64) + b = np.array([2-1j, 4+3j, -5+6j], dtype=np.complex64) + out = np.zeros_like(a) + da = cuda.to_device(a) + db = cuda.to_device(b) + dout = cuda.to_device(out) + self._launch_1d(complex_mul_kernel, (da, db, dout), a.size) + np.testing.assert_array_equal(dout.copy_to_host(), a * b) + + def test_complex_abs(self): + inp = np.array([3+4j, 5-12j, -8+15j], dtype=np.complex64) + out = np.zeros(inp.size, dtype=np.float32) + dinp = cuda.to_device(inp) + dout = cuda.to_device(out) + self._launch_1d(complex_abs_kernel, (dinp, dout), inp.size) + np.testing.assert_allclose(dout.copy_to_host(), np.abs(inp), rtol=1e-6) # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: BSD-2-Clause diff --git a/numba_cuda/numba/cuda/tests/test_complex.py b/numba_cuda/numba/cuda/tests/test_complex.py deleted file mode 100644 index d8c065d4d..000000000 --- a/numba_cuda/numba/cuda/tests/test_complex.py +++ /dev/null @@ -1,88 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: BSD-2-Clause - -import numpy as np - -from numba import cuda -from numba.cuda.testing import CUDATestCase, skip_on_cudasim - - -@cuda.jit -def complex_add_kernel(a, b, out): - i = cuda.grid(1) - if i < out.size: - out[i] = a[i] + b[i] - - -@cuda.jit -def complex_mul_kernel(a, b, out): - i = cuda.grid(1) - if i < out.size: - out[i] = a[i] * b[i] - - -@cuda.jit -def complex_abs_kernel(inp, out): - i = cuda.grid(1) - if i < out.size: - out[i] = abs(inp[i]) - - -@skip_on_cudasim("Complex semantics differ under cudasim") -class TestCudaComplex(CUDATestCase): - - def _launch_1d(self, kernel, args, size): - threadsperblock = 128 - blockspergrid = (size + threadsperblock - 1) // threadsperblock - kernel[blockspergrid, threadsperblock](*args) - cuda.synchronize() - - - def test_complex_add(self): - a = np.array([1+2j, -3+4j, 5-6j], dtype=np.complex64) - b = np.array([2-1j, 4+3j, -5+6j], dtype=np.complex64) - out = np.zeros_like(a) - - da = cuda.to_device(a) - db = cuda.to_device(b) - dout = cuda.to_device(out) - - self._launch_1d(complex_add_kernel, (da, db, dout), a.size) - - np.testing.assert_array_equal( - dout.copy_to_host(), - a + b - ) - - - def test_complex_multiply(self): - a = np.array([1+2j, -3+4j, 5-6j], dtype=np.complex64) - b = np.array([2-1j, 4+3j, -5+6j], dtype=np.complex64) - out = np.zeros_like(a) - - da = cuda.to_device(a) - db = cuda.to_device(b) - dout = cuda.to_device(out) - - self._launch_1d(complex_mul_kernel, (da, db, dout), a.size) - - np.testing.assert_array_equal( - dout.copy_to_host(), - a * b - ) - - - def test_complex_abs(self): - inp = np.array([3+4j, 5-12j, -8+15j], dtype=np.complex64) - out = np.zeros(inp.size, dtype=np.float32) - - dinp = cuda.to_device(inp) - dout = cuda.to_device(out) - - self._launch_1d(complex_abs_kernel, (dinp, dout), inp.size) - - np.testing.assert_allclose( - dout.copy_to_host(), - np.abs(inp), - rtol=1e-6 - ) From de31c05bf46bb69f1155c8de0de765818fa3cb19 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 12:00:58 +0530 Subject: [PATCH 04/12] Update numba_cuda/numba/cuda/tests/cudapy/test_complex.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- numba_cuda/numba/cuda/tests/cudapy/test_complex.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py index 6b1159b55..9f04bab70 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py @@ -1,3 +1,6 @@ +# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: BSD-2-Clause + ## --- Basic CUDA complex number operation tests (from test_complex.py) --- import numpy as np from numba import cuda From 06c51b21045a6328a0eeb84a8895d281603c7839 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 12:02:55 +0530 Subject: [PATCH 05/12] Update numba_cuda/numba/cuda/tests/cudapy/test_complex.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- numba_cuda/numba/cuda/tests/cudapy/test_complex.py | 2 -- 1 file changed, 2 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py index 9f04bab70..ddfaafc95 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py @@ -59,8 +59,6 @@ def test_complex_abs(self): dout = cuda.to_device(out) self._launch_1d(complex_abs_kernel, (dinp, dout), inp.size) np.testing.assert_allclose(dout.copy_to_host(), np.abs(inp), rtol=1e-6) -# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: BSD-2-Clause import math import itertools From e5ba15dd8333b451b5b301c5b3e3fbf06bcd0d83 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 12:04:36 +0530 Subject: [PATCH 06/12] Update numba_cuda/numba/cuda/tests/cudapy/test_complex.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- numba_cuda/numba/cuda/tests/cudapy/test_complex.py | 3 --- 1 file changed, 3 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py index ddfaafc95..d0473fe43 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py @@ -2,9 +2,6 @@ # SPDX-License-Identifier: BSD-2-Clause ## --- Basic CUDA complex number operation tests (from test_complex.py) --- -import numpy as np -from numba import cuda -from numba.cuda.testing import CUDATestCase, skip_on_cudasim @cuda.jit def complex_add_kernel(a, b, out): From 22a25204ad999d7e77dc522635caded691b98080 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 12:07:19 +0530 Subject: [PATCH 07/12] Update numba_cuda/numba/cuda/tests/cudapy/test_complex.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- .../numba/cuda/tests/cudapy/test_complex.py | 45 +++++++++++++++++++ 1 file changed, 45 insertions(+) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py index d0473fe43..eb17f2c7d 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py @@ -1,6 +1,50 @@ # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: BSD-2-Clause +## --- Basic CUDA complex number operation tests (from test_complex.py) --- + +import math +import itertools +import sys + +import numpy as np +import pytest + +from numba.cuda.testing import unittest, CUDATestCase, skip_on_cudasim +from numba.cuda import types +from numba import cuda +from numba.cuda import config +from numba.cuda.tests.cudapy.complex_usecases import ( + real_usecase, + imag_usecase, + conjugate_usecase, + phase_usecase, + polar_as_complex_usecase, + rect_usecase, + isnan_usecase, + isinf_usecase, + isfinite_usecase, + exp_usecase, + log_usecase, + log_base_usecase, + log10_usecase, + sqrt_usecase, + asin_usecase, + acos_usecase, + atan_usecase, + cos_usecase, + sin_usecase, + tan_usecase, + acosh_usecase, + asinh_usecase, + atanh_usecase, + cosh_usecase, + sinh_usecase, + tanh_usecase, +) +from numba.cuda.np import numpy_support + + ## --- Basic CUDA complex number operation tests (from test_complex.py) --- @cuda.jit @@ -57,6 +101,7 @@ def test_complex_abs(self): self._launch_1d(complex_abs_kernel, (dinp, dout), inp.size) np.testing.assert_allclose(dout.copy_to_host(), np.abs(inp), rtol=1e-6) + import math import itertools import sys From 62910e48d9f7d140e7db3562ca662c5d9fb6f41c Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 12:11:13 +0530 Subject: [PATCH 08/12] Update numba_cuda/numba/cuda/tests/cudapy/test_complex.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- .../numba/cuda/tests/cudapy/test_complex.py | 40 ------------------- 1 file changed, 40 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py index eb17f2c7d..02c90dc51 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py @@ -3,46 +3,6 @@ ## --- Basic CUDA complex number operation tests (from test_complex.py) --- -import math -import itertools -import sys - -import numpy as np -import pytest - -from numba.cuda.testing import unittest, CUDATestCase, skip_on_cudasim -from numba.cuda import types -from numba import cuda -from numba.cuda import config -from numba.cuda.tests.cudapy.complex_usecases import ( - real_usecase, - imag_usecase, - conjugate_usecase, - phase_usecase, - polar_as_complex_usecase, - rect_usecase, - isnan_usecase, - isinf_usecase, - isfinite_usecase, - exp_usecase, - log_usecase, - log_base_usecase, - log10_usecase, - sqrt_usecase, - asin_usecase, - acos_usecase, - atan_usecase, - cos_usecase, - sin_usecase, - tan_usecase, - acosh_usecase, - asinh_usecase, - atanh_usecase, - cosh_usecase, - sinh_usecase, - tanh_usecase, -) -from numba.cuda.np import numpy_support ## --- Basic CUDA complex number operation tests (from test_complex.py) --- From 7e96c5ac2f3e0734c010d08c5f14e7fe4b1084ee Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 12:13:28 +0530 Subject: [PATCH 09/12] Update numba_cuda/numba/cuda/tests/cudapy/test_complex.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- numba_cuda/numba/cuda/tests/cudapy/test_complex.py | 5 ----- 1 file changed, 5 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py index 02c90dc51..478c63db0 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py @@ -1,11 +1,6 @@ # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: BSD-2-Clause -## --- Basic CUDA complex number operation tests (from test_complex.py) --- - - - -## --- Basic CUDA complex number operation tests (from test_complex.py) --- @cuda.jit def complex_add_kernel(a, b, out): From f27eb4c98f909e4be577c795698bab131784d4b3 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay Date: Wed, 21 Jan 2026 12:15:54 +0530 Subject: [PATCH 10/12] fix: move all kernel and test class definitions after imports, remove duplicate numpy import and comment marker\n\n- Ensures no NameError at import time and code style is clean --- numba_cuda/numba/cuda/tests/cudapy/test_complex.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py index 478c63db0..cb08060f9 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py @@ -1,6 +1,5 @@ -# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: BSD-2-Clause +# --- Basic CUDA complex number operation tests (from test_complex.py) --- @cuda.jit def complex_add_kernel(a, b, out): From b4c94661557fd432de0f78ef8f98fbaab9246897 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 12:16:47 +0530 Subject: [PATCH 11/12] Update numba_cuda/numba/cuda/tests/cudapy/test_complex.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- .../numba/cuda/tests/cudapy/test_complex.py | 43 +++++++++++++++++++ 1 file changed, 43 insertions(+) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py index cb08060f9..b82febfab 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py @@ -1,6 +1,49 @@ # --- Basic CUDA complex number operation tests (from test_complex.py) --- +@cuda.jit +import math +import itertools +import sys + +import numpy as np +import pytest + +from numba.cuda.testing import unittest, CUDATestCase, skip_on_cudasim +from numba.cuda import types +from numba import cuda +from numba.cuda import config +from numba.cuda.tests.cudapy.complex_usecases import ( + real_usecase, + imag_usecase, + conjugate_usecase, + phase_usecase, + polar_as_complex_usecase, + rect_usecase, + isnan_usecase, + isinf_usecase, + isfinite_usecase, + exp_usecase, + log_usecase, + log_base_usecase, + log10_usecase, + sqrt_usecase, + asin_usecase, + acos_usecase, + atan_usecase, + cos_usecase, + sin_usecase, + tan_usecase, + acosh_usecase, + asinh_usecase, + atanh_usecase, + cosh_usecase, + sinh_usecase, + tanh_usecase, +) +from numba.cuda.np import numpy_support + + @cuda.jit def complex_add_kernel(a, b, out): i = cuda.grid(1) From 185094665c0ee642c360690e50b903ea920f9356 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay Date: Wed, 21 Jan 2026 12:19:48 +0530 Subject: [PATCH 12/12] fix: move kernels and test class after all imports\n\n- SPDX header and all imports come first\n- Removed duplicate imports\n- Kernels and TestCudaComplexBasicOps class now at end of file\n- Resolves NameError issues --- .../numba/cuda/tests/cudapy/test_complex.py | 159 +++++++----------- 1 file changed, 59 insertions(+), 100 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py index b82febfab..0773964b0 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_complex.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_complex.py @@ -1,7 +1,6 @@ +# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: BSD-2-Clause -# --- Basic CUDA complex number operation tests (from test_complex.py) --- - -@cuda.jit import math import itertools import sys @@ -44,103 +43,6 @@ from numba.cuda.np import numpy_support -@cuda.jit -def complex_add_kernel(a, b, out): - i = cuda.grid(1) - if i < out.size: - out[i] = a[i] + b[i] - -@cuda.jit -def complex_mul_kernel(a, b, out): - i = cuda.grid(1) - if i < out.size: - out[i] = a[i] * b[i] - -@cuda.jit -def complex_abs_kernel(inp, out): - i = cuda.grid(1) - if i < out.size: - out[i] = abs(inp[i]) - -@skip_on_cudasim("Complex semantics differ under cudasim") -class TestCudaComplexBasicOps(CUDATestCase): - def _launch_1d(self, kernel, args, size): - threadsperblock = 128 - blockspergrid = (size + threadsperblock - 1) // threadsperblock - kernel[blockspergrid, threadsperblock](*args) - cuda.synchronize() - - def test_complex_add(self): - a = np.array([1+2j, -3+4j, 5-6j], dtype=np.complex64) - b = np.array([2-1j, 4+3j, -5+6j], dtype=np.complex64) - out = np.zeros_like(a) - da = cuda.to_device(a) - db = cuda.to_device(b) - dout = cuda.to_device(out) - self._launch_1d(complex_add_kernel, (da, db, dout), a.size) - np.testing.assert_array_equal(dout.copy_to_host(), a + b) - - def test_complex_multiply(self): - a = np.array([1+2j, -3+4j, 5-6j], dtype=np.complex64) - b = np.array([2-1j, 4+3j, -5+6j], dtype=np.complex64) - out = np.zeros_like(a) - da = cuda.to_device(a) - db = cuda.to_device(b) - dout = cuda.to_device(out) - self._launch_1d(complex_mul_kernel, (da, db, dout), a.size) - np.testing.assert_array_equal(dout.copy_to_host(), a * b) - - def test_complex_abs(self): - inp = np.array([3+4j, 5-12j, -8+15j], dtype=np.complex64) - out = np.zeros(inp.size, dtype=np.float32) - dinp = cuda.to_device(inp) - dout = cuda.to_device(out) - self._launch_1d(complex_abs_kernel, (dinp, dout), inp.size) - np.testing.assert_allclose(dout.copy_to_host(), np.abs(inp), rtol=1e-6) - - -import math -import itertools -import sys - -import numpy as np -import pytest - -from numba.cuda.testing import unittest, CUDATestCase -from numba.cuda import types -from numba import cuda -from numba.cuda import config -from numba.cuda.tests.cudapy.complex_usecases import ( - real_usecase, - imag_usecase, - conjugate_usecase, - phase_usecase, - polar_as_complex_usecase, - rect_usecase, - isnan_usecase, - isinf_usecase, - isfinite_usecase, - exp_usecase, - log_usecase, - log_base_usecase, - log10_usecase, - sqrt_usecase, - asin_usecase, - acos_usecase, - atan_usecase, - cos_usecase, - sin_usecase, - tan_usecase, - acosh_usecase, - asinh_usecase, - atanh_usecase, - cosh_usecase, - sinh_usecase, - tanh_usecase, -) -from numba.cuda.np import numpy_support - - def compile_scalar_func(pyfunc, argtypes, restype): # First compile a scalar device function assert not any(isinstance(tp, types.Array) for tp in argtypes) @@ -477,5 +379,62 @@ def atomic_add_one_j(values): np.testing.assert_equal(arr1 + 1j, arr2) +# --- Basic CUDA complex number operation tests --- + +@cuda.jit +def complex_add_kernel(a, b, out): + i = cuda.grid(1) + if i < out.size: + out[i] = a[i] + b[i] + +@cuda.jit +def complex_mul_kernel(a, b, out): + i = cuda.grid(1) + if i < out.size: + out[i] = a[i] * b[i] + +@cuda.jit +def complex_abs_kernel(inp, out): + i = cuda.grid(1) + if i < out.size: + out[i] = abs(inp[i]) + +@skip_on_cudasim("Complex semantics differ under cudasim") +class TestCudaComplexBasicOps(CUDATestCase): + def _launch_1d(self, kernel, args, size): + threadsperblock = 128 + blockspergrid = (size + threadsperblock - 1) // threadsperblock + kernel[blockspergrid, threadsperblock](*args) + cuda.synchronize() + + def test_complex_add(self): + a = np.array([1+2j, -3+4j, 5-6j], dtype=np.complex64) + b = np.array([2-1j, 4+3j, -5+6j], dtype=np.complex64) + out = np.zeros_like(a) + da = cuda.to_device(a) + db = cuda.to_device(b) + dout = cuda.to_device(out) + self._launch_1d(complex_add_kernel, (da, db, dout), a.size) + np.testing.assert_array_equal(dout.copy_to_host(), a + b) + + def test_complex_multiply(self): + a = np.array([1+2j, -3+4j, 5-6j], dtype=np.complex64) + b = np.array([2-1j, 4+3j, -5+6j], dtype=np.complex64) + out = np.zeros_like(a) + da = cuda.to_device(a) + db = cuda.to_device(b) + dout = cuda.to_device(out) + self._launch_1d(complex_mul_kernel, (da, db, dout), a.size) + np.testing.assert_array_equal(dout.copy_to_host(), a * b) + + def test_complex_abs(self): + inp = np.array([3+4j, 5-12j, -8+15j], dtype=np.complex64) + out = np.zeros(inp.size, dtype=np.float32) + dinp = cuda.to_device(inp) + dout = cuda.to_device(out) + self._launch_1d(complex_abs_kernel, (dinp, dout), inp.size) + np.testing.assert_allclose(dout.copy_to_host(), np.abs(inp), rtol=1e-6) + + if __name__ == "__main__": unittest.main()