From c31c4945e92a236043890fee6ef1f30526c9aac6 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay Date: Tue, 20 Jan 2026 16:26:13 +0530 Subject: [PATCH 1/5] CUDA: add closure tests\n\nAdds kernel-based tests for device-side closure capture: constants, scalars, kernel factories, and read-only captures. Skips under cudasim. NumPy is used as the reference. Scope is intentionally limited to safe, well-defined CUDA closure semantics. --- numba_cuda/numba/cuda/tests/test_closure.py | 108 ++++++++++++++++++++ 1 file changed, 108 insertions(+) create mode 100644 numba_cuda/numba/cuda/tests/test_closure.py diff --git a/numba_cuda/numba/cuda/tests/test_closure.py b/numba_cuda/numba/cuda/tests/test_closure.py new file mode 100644 index 000000000..a9dbd62af --- /dev/null +++ b/numba_cuda/numba/cuda/tests/test_closure.py @@ -0,0 +1,108 @@ +import numpy as np + +from numba import cuda +from numba.cuda.testing import CUDATestCase, skip_on_cudasim + +# ----------------------------- +# Kernel factories with closures +# ----------------------------- + +def make_add_kernel(const): + @cuda.jit + def kernel(inp, out): + i = cuda.grid(1) + if i < inp.size: + out[i] = inp[i] + const + return kernel + +def make_mul_kernel(factor): + @cuda.jit + def kernel(inp, out): + i = cuda.grid(1) + if i < inp.size: + out[i] = inp[i] * factor + return kernel + +def make_bool_closure_kernel(flag): + @cuda.jit + def kernel(out): + i = cuda.grid(1) + if i < out.size: + if flag: + out[i] = 1 + else: + out[i] = 0 + return kernel + +def make_nested_constant_kernel(): + x = 3 + y = 4 + @cuda.jit + def kernel(out): + i = cuda.grid(1) + if i < out.size: + out[i] = x + y + return kernel + +@skip_on_cudasim("Closure capture semantics differ under cudasim") +class TestCudaClosure(CUDATestCase): + + def _launch_1d(self, kernel, args, size): + threadsperblock = 128 + blockspergrid = (size + threadsperblock - 1) // threadsperblock + kernel[blockspergrid, threadsperblock](*args) + cuda.synchronize() + + def test_closure_add_constant(self): + add5 = make_add_kernel(5) + inp = np.arange(10, dtype=np.int32) + out = np.zeros_like(inp) + d_inp = cuda.to_device(inp) + d_out = cuda.to_device(out) + self._launch_1d(add5, (d_inp, d_out), inp.size) + np.testing.assert_array_equal( + d_out.copy_to_host(), + inp + 5, + ) + + def test_closure_multiply_constant(self): + mul3 = make_mul_kernel(3) + inp = np.array([1, 2, 3, 4], dtype=np.int32) + out = np.zeros_like(inp) + d_inp = cuda.to_device(inp) + d_out = cuda.to_device(out) + self._launch_1d(mul3, (d_inp, d_out), inp.size) + np.testing.assert_array_equal( + d_out.copy_to_host(), + inp * 3, + ) + + def test_bool_closure_true(self): + kernel = make_bool_closure_kernel(True) + out = np.zeros(6, dtype=np.int32) + d_out = cuda.to_device(out) + self._launch_1d(kernel, (d_out,), out.size) + np.testing.assert_array_equal( + d_out.copy_to_host(), + np.ones_like(out), + ) + + def test_bool_closure_false(self): + kernel = make_bool_closure_kernel(False) + out = np.zeros(6, dtype=np.int32) + d_out = cuda.to_device(out) + self._launch_1d(kernel, (d_out,), out.size) + np.testing.assert_array_equal( + d_out.copy_to_host(), + np.zeros_like(out), + ) + + def test_multiple_constant_capture(self): + kernel = make_nested_constant_kernel() + out = np.zeros(8, dtype=np.int32) + d_out = cuda.to_device(out) + self._launch_1d(kernel, (d_out,), out.size) + np.testing.assert_array_equal( + d_out.copy_to_host(), + np.full_like(out, 7), + ) From 7dc1fabb4335dc1511d03911886476bd9c057327 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 11:09:48 +0530 Subject: [PATCH 2/5] Update numba_cuda/numba/cuda/tests/test_closure.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- numba_cuda/numba/cuda/tests/test_closure.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/numba_cuda/numba/cuda/tests/test_closure.py b/numba_cuda/numba/cuda/tests/test_closure.py index a9dbd62af..36aed2051 100644 --- a/numba_cuda/numba/cuda/tests/test_closure.py +++ b/numba_cuda/numba/cuda/tests/test_closure.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 34d5222782f1aba38bac2b1e18017e204071a1d2 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 12:07:59 +0530 Subject: [PATCH 3/5] Update numba_cuda/numba/cuda/tests/test_closure.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- numba_cuda/numba/cuda/tests/test_closure.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/test_closure.py b/numba_cuda/numba/cuda/tests/test_closure.py index 36aed2051..d23bddff6 100644 --- a/numba_cuda/numba/cuda/tests/test_closure.py +++ b/numba_cuda/numba/cuda/tests/test_closure.py @@ -91,8 +91,7 @@ def test_bool_closure_true(self): ) def test_bool_closure_false(self): - kernel = make_bool_closure_kernel(False) - out = np.zeros(6, dtype=np.int32) + out = np.ones(6, dtype=np.int32) d_out = cuda.to_device(out) self._launch_1d(kernel, (d_out,), out.size) np.testing.assert_array_equal( From 20f47b5d8128b241d60b6d4e8ccd901157daf671 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 12:10:59 +0530 Subject: [PATCH 4/5] Update numba_cuda/numba/cuda/tests/test_closure.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- numba_cuda/numba/cuda/tests/test_closure.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/numba_cuda/numba/cuda/tests/test_closure.py b/numba_cuda/numba/cuda/tests/test_closure.py index d23bddff6..4863555b5 100644 --- a/numba_cuda/numba/cuda/tests/test_closure.py +++ b/numba_cuda/numba/cuda/tests/test_closure.py @@ -93,7 +93,8 @@ def test_bool_closure_true(self): def test_bool_closure_false(self): out = np.ones(6, dtype=np.int32) d_out = cuda.to_device(out) - self._launch_1d(kernel, (d_out,), out.size) + kernel = make_bool_closure_kernel(False) + d_out = cuda.to_device(out) np.testing.assert_array_equal( d_out.copy_to_host(), np.zeros_like(out), From 905efbf2511498fa8a44f7837658dcca686b9ab6 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 12:13:16 +0530 Subject: [PATCH 5/5] Update numba_cuda/numba/cuda/tests/test_closure.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- numba_cuda/numba/cuda/tests/test_closure.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/test_closure.py b/numba_cuda/numba/cuda/tests/test_closure.py index 4863555b5..c5612e239 100644 --- a/numba_cuda/numba/cuda/tests/test_closure.py +++ b/numba_cuda/numba/cuda/tests/test_closure.py @@ -91,10 +91,10 @@ def test_bool_closure_true(self): ) def test_bool_closure_false(self): - out = np.ones(6, dtype=np.int32) - d_out = cuda.to_device(out) kernel = make_bool_closure_kernel(False) + out = np.ones(6, dtype=np.int32) d_out = cuda.to_device(out) + self._launch_1d(kernel, (d_out,), out.size) np.testing.assert_array_equal( d_out.copy_to_host(), np.zeros_like(out),