From 56456b19dc84903aa4c4e055cd2f78a6b7d81be8 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay Date: Tue, 20 Jan 2026 16:18:20 +0530 Subject: [PATCH 1/2] CUDA: add auto constant tests\n\nAdds kernel-based tests for device-side compile-time constants: integer, float, boolean literals, constant expressions, and constant indexing. Skips under cudasim. NumPy is used as the reference. Scope is intentionally limited to safe, well-defined CUDA constant lowering. --- .../numba/cuda/tests/test_auto_constants.py | 109 ++++++++++++++++++ 1 file changed, 109 insertions(+) create mode 100644 numba_cuda/numba/cuda/tests/test_auto_constants.py diff --git a/numba_cuda/numba/cuda/tests/test_auto_constants.py b/numba_cuda/numba/cuda/tests/test_auto_constants.py new file mode 100644 index 000000000..cb4da320d --- /dev/null +++ b/numba_cuda/numba/cuda/tests/test_auto_constants.py @@ -0,0 +1,109 @@ +import numpy as np + +from numba import cuda +from numba.cuda.testing import CUDATestCase, skip_on_cudasim + +# ----------------------------- +# Kernels using auto-constants +# ----------------------------- + +@cuda.jit +def int_constant_kernel(out): + i = cuda.grid(1) + if i < out.size: + out[i] = 42 + +@cuda.jit +def float_constant_kernel(out): + i = cuda.grid(1) + if i < out.size: + out[i] = 3.5 + +@cuda.jit +def bool_constant_kernel(out): + i = cuda.grid(1) + if i < out.size: + out[i] = True + +@cuda.jit +def arithmetic_constant_kernel(out): + i = cuda.grid(1) + if i < out.size: + out[i] = 2 + 3 * 4 # 14 + +@cuda.jit +def mixed_constant_kernel(out): + i = cuda.grid(1) + if i < out.size: + out[i] = 1.5 + 2 # 3.5 + +@cuda.jit +def constant_index_kernel(inp, out): + i = cuda.grid(1) + if i < out.size: + out[i] = inp[2] + +@skip_on_cudasim("Auto-constant lowering differs under cudasim") +class TestCudaAutoConstants(CUDATestCase): + + def _launch_1d(self, kernel, args, size): + threadsperblock = 128 + blockspergrid = (size + threadsperblock - 1) // threadsperblock + kernel[blockspergrid, threadsperblock](*args) + cuda.synchronize() + + def test_int_constant(self): + out = np.zeros(8, dtype=np.int32) + d_out = cuda.to_device(out) + self._launch_1d(int_constant_kernel, (d_out,), out.size) + np.testing.assert_array_equal( + d_out.copy_to_host(), + np.full_like(out, 42), + ) + + def test_float_constant(self): + out = np.zeros(8, dtype=np.float32) + d_out = cuda.to_device(out) + self._launch_1d(float_constant_kernel, (d_out,), out.size) + np.testing.assert_array_equal( + d_out.copy_to_host(), + np.full_like(out, 3.5), + ) + + def test_bool_constant(self): + out = np.zeros(8, dtype=np.bool_) + d_out = cuda.to_device(out) + self._launch_1d(bool_constant_kernel, (d_out,), out.size) + np.testing.assert_array_equal( + d_out.copy_to_host(), + np.ones_like(out, dtype=np.bool_), + ) + + def test_arithmetic_constant(self): + out = np.zeros(8, dtype=np.int32) + d_out = cuda.to_device(out) + self._launch_1d(arithmetic_constant_kernel, (d_out,), out.size) + np.testing.assert_array_equal( + d_out.copy_to_host(), + np.full_like(out, 14), + ) + + def test_mixed_constant(self): + out = np.zeros(8, dtype=np.float32) + d_out = cuda.to_device(out) + self._launch_1d(mixed_constant_kernel, (d_out,), out.size) + np.testing.assert_array_equal( + d_out.copy_to_host(), + np.full_like(out, 3.5), + ) + + def test_constant_indexing(self): + inp = np.arange(10, dtype=np.int32) + out = np.zeros(5, dtype=np.int32) + d_inp = cuda.to_device(inp) + d_out = cuda.to_device(out) + self._launch_1d(constant_index_kernel, (d_inp, d_out), out.size) + np.testing.assert_array_equal( + d_out.copy_to_host(), + np.full_like(out, inp[2]), + ) From 1c0a24e9ea845daeaecb1728cbfed92d67dba413 Mon Sep 17 00:00:00 2001 From: Srijan Upadhyay <104912634+CodersAcademy006@users.noreply.github.com> Date: Wed, 21 Jan 2026 11:10:11 +0530 Subject: [PATCH 2/2] Update numba_cuda/numba/cuda/tests/test_auto_constants.py Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- numba_cuda/numba/cuda/tests/test_auto_constants.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/numba_cuda/numba/cuda/tests/test_auto_constants.py b/numba_cuda/numba/cuda/tests/test_auto_constants.py index cb4da320d..3f835c93f 100644 --- a/numba_cuda/numba/cuda/tests/test_auto_constants.py +++ b/numba_cuda/numba/cuda/tests/test_auto_constants.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