diff --git a/README.md b/README.md index 38703d107..8de133456 100644 --- a/README.md +++ b/README.md @@ -29,8 +29,18 @@ If you want to manage all run-time dependencies yourself, also pass the `--no-de ## Running tests +Tests must be run from the `testing` folder, which contains the pytest +configuration and code to generate binaries used during the tests. The test +binaries need to be built on the system on which the tests are run, so that +they are compiled for the appropriate compute capability. + ``` -pytest -n auto --pyargs numba.cuda.tests -v +cd testing +# Optionally, build test binaries and point to their location for the test suite +make +export NUMBA_CUDA_TEST_BIN_DIR=`pwd` +# Execute tests +pytest -n auto -v ``` diff --git a/ci/build_docs.sh b/ci/build_docs.sh index f5739b556..ee0991698 100755 --- a/ci/build_docs.sh +++ b/ci/build_docs.sh @@ -33,6 +33,10 @@ pip install nvidia-sphinx-theme rapids-print-env +# Change out of the root of the repo early so the source tree in the repo +# doesn't confuse things +pushd docs + rapids-logger "Show Numba system info" python -m numba --sysinfo @@ -41,7 +45,6 @@ trap "EXITCODE=1" ERR set +e rapids-logger "Build docs" -pushd docs make html popd diff --git a/ci/coverage_report.sh b/ci/coverage_report.sh index 38482a073..922476b83 100755 --- a/ci/coverage_report.sh +++ b/ci/coverage_report.sh @@ -16,32 +16,19 @@ python -m pip install \ pytest-cov \ coverage -GET_TEST_BINARY_DIR=" -import numba_cuda -root = numba_cuda.__file__.rstrip('__init__.py') -test_dir = root + \"numba/cuda/tests/test_binary_generation/\" -print(test_dir) -" - -rapids-logger "Build tests" - -export NUMBA_CUDA_TEST_BIN_DIR=$(python -c "$GET_TEST_BINARY_DIR") +rapids-logger "Build test binaries" +export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing pushd $NUMBA_CUDA_TEST_BIN_DIR make -popd rapids-logger "Check GPU usage" nvidia-smi -RAPIDS_TESTS_DIR=${RAPIDS_TESTS_DIR:-"${PWD}/test-results"}/ -mkdir -p "${RAPIDS_TESTS_DIR}" -pushd "${RAPIDS_TESTS_DIR}" - rapids-logger "Show Numba system info" python -m numba --sysinfo rapids-logger "Run Tests with Coverage" -python -m pytest --pyargs numba.cuda.tests -v --cov +python -m pytest -v --cov rapids-logger "Generate Markdown Coverage Report" python -m coverage report --format markdown >> $GITHUB_STEP_SUMMARY diff --git a/ci/get_test_binary_dir.py b/ci/get_test_binary_dir.py deleted file mode 100644 index fa5c7ec64..000000000 --- a/ci/get_test_binary_dir.py +++ /dev/null @@ -1,11 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: BSD-2-Clause - -import numba_cuda -import os - -root = os.path.dirname(numba_cuda.__file__) -test_dir = os.path.join( - root, "numba", "cuda", "tests", "test_binary_generation" -) -print(test_dir) diff --git a/ci/test_conda.sh b/ci/test_conda.sh index ed2d57cef..b465f966f 100755 --- a/ci/test_conda.sh +++ b/ci/test_conda.sh @@ -58,15 +58,16 @@ set -u pip install filecheck -RAPIDS_TESTS_DIR=${RAPIDS_TESTS_DIR:-"${PWD}/test-results"}/ -mkdir -p "${RAPIDS_TESTS_DIR}" -pushd "${RAPIDS_TESTS_DIR}" - rapids-print-env rapids-logger "Check GPU usage" nvidia-smi +rapids-logger "Build test binaries" +export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing +pushd $NUMBA_CUDA_TEST_BIN_DIR +make + rapids-logger "Show Numba system info" python -m numba --sysinfo @@ -77,21 +78,8 @@ set +e rapids-logger "Test importing numba.cuda" python -c "from numba import cuda" -GET_TEST_BINARY_DIR=" -import numba_cuda -root = numba_cuda.__file__.rstrip('__init__.py') -test_dir = root + \"numba/cuda/tests/test_binary_generation/\" -print(test_dir) -" - -rapids-logger "Build tests" -export NUMBA_CUDA_TEST_BIN_DIR=$(python -c "$GET_TEST_BINARY_DIR") -pushd $NUMBA_CUDA_TEST_BIN_DIR -make -popd - rapids-logger "Run Tests" -pytest --pyargs numba.cuda.tests -v +pytest -v popd diff --git a/ci/test_conda_ctypes_binding.sh b/ci/test_conda_ctypes_binding.sh index 844b35b40..644f05a7b 100755 --- a/ci/test_conda_ctypes_binding.sh +++ b/ci/test_conda_ctypes_binding.sh @@ -44,15 +44,16 @@ set -u pip install filecheck -RAPIDS_TESTS_DIR=${RAPIDS_TESTS_DIR:-"${PWD}/test-results"}/ -mkdir -p "${RAPIDS_TESTS_DIR}" -pushd "${RAPIDS_TESTS_DIR}" - rapids-print-env rapids-logger "Check GPU usage" nvidia-smi +rapids-logger "Build test binaries" +export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing +pushd $NUMBA_CUDA_TEST_BIN_DIR +make + rapids-logger "Show Numba system info" python -m numba --sysinfo @@ -60,23 +61,8 @@ EXITCODE=0 trap "EXITCODE=1" ERR set +e -rapids-logger "Build tests" - -PY_SCRIPT=" -import numba_cuda -root = numba_cuda.__file__.rstrip('__init__.py') -test_dir = root + \"numba/cuda/tests/test_binary_generation/\" -print(test_dir) -" - -NUMBA_CUDA_TEST_BIN_DIR=$(python -c "$PY_SCRIPT") -pushd $NUMBA_CUDA_TEST_BIN_DIR -make -popd - - rapids-logger "Run Tests" -NUMBA_CUDA_USE_NVIDIA_BINDING=0 NUMBA_CUDA_TEST_BIN_DIR=$NUMBA_CUDA_TEST_BIN_DIR pytest --pyargs numba.cuda.tests -v +NUMBA_CUDA_USE_NVIDIA_BINDING=0 NUMBA_CUDA_TEST_BIN_DIR=$NUMBA_CUDA_TEST_BIN_DIR pytest -v popd diff --git a/ci/test_simulator.sh b/ci/test_simulator.sh index bb85a8733..21aa7d504 100755 --- a/ci/test_simulator.sh +++ b/ci/test_simulator.sh @@ -31,12 +31,16 @@ set -u pip install filecheck -RAPIDS_TESTS_DIR=${RAPIDS_TESTS_DIR:-"${PWD}/test-results"}/ -mkdir -p "${RAPIDS_TESTS_DIR}" -pushd "${RAPIDS_TESTS_DIR}" - rapids-print-env +# The simulator doesn't actually use the test binaries, but we move into the +# test binaries folder so that we're not in the root of the repo, and therefore +# numba-cuda code from the installed package will be tested, instead of the +# code in the source repo. +rapids-logger "Move to test binaries folder" +export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing +pushd $NUMBA_CUDA_TEST_BIN_DIR + rapids-logger "Show Numba system info" python -m numba --sysinfo @@ -46,7 +50,7 @@ set +e rapids-logger "Run Tests" export NUMBA_ENABLE_CUDASIM=1 -pytest --pyargs numba.cuda.tests -v +pytest -v popd diff --git a/ci/test_wheel.ps1 b/ci/test_wheel.ps1 index 5eacf6c05..79bed36d8 100755 --- a/ci/test_wheel.ps1 +++ b/ci/test_wheel.ps1 @@ -36,13 +36,12 @@ python -m pip install "llvmlite<0.45" "numba==0.61.*" # WAR for https://github.c rapids-logger "Build tests" -$NUMBA_CUDA_TEST_BIN_DIR = (python ci\get_test_binary_dir.py) +$NUMBA_CUDA_TEST_BIN_DIR = (Join-Path -Path (Get-Location) -ChildPath "testing") echo "Test binary dir: $NUMBA_CUDA_TEST_BIN_DIR" pushd $NUMBA_CUDA_TEST_BIN_DIR Get-Location cmd.exe /c '.\build.bat' -popd rapids-logger "Check GPU usage" nvidia-smi @@ -51,4 +50,6 @@ rapids-logger "Show Numba system info" python -m numba --sysinfo rapids-logger "Run Tests" -python -m pytest --pyargs numba.cuda.tests -v +python -m pytest -v + +popd diff --git a/ci/test_wheel.sh b/ci/test_wheel.sh index a6c474fde..40589d6c2 100755 --- a/ci/test_wheel.sh +++ b/ci/test_wheel.sh @@ -23,33 +23,21 @@ fi python -m pip install "${DEPENDENCIES[@]}" -rapids-logger "Test importing numba.cuda" -python -c "from numba import cuda" - -GET_TEST_BINARY_DIR=" -import numba_cuda -root = numba_cuda.__file__.rstrip('__init__.py') -test_dir = root + \"numba/cuda/tests/test_binary_generation/\" -print(test_dir) -" - rapids-logger "Build tests" -export NUMBA_CUDA_TEST_BIN_DIR=$(python -c "$GET_TEST_BINARY_DIR") +export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing pushd $NUMBA_CUDA_TEST_BIN_DIR make -popd + +rapids-logger "Test importing numba.cuda" +python -c "from numba import cuda" rapids-logger "Check GPU usage" nvidia-smi -RAPIDS_TESTS_DIR=${RAPIDS_TESTS_DIR:-"${PWD}/test-results"}/ -mkdir -p "${RAPIDS_TESTS_DIR}" -pushd "${RAPIDS_TESTS_DIR}" - rapids-logger "Show Numba system info" python -m numba --sysinfo rapids-logger "Run Tests" -python -m pytest --pyargs numba.cuda.tests -v +python -m pytest -v popd diff --git a/ci/test_wheel_ctypes_binding.sh b/ci/test_wheel_ctypes_binding.sh index 156ade407..b386202a1 100755 --- a/ci/test_wheel_ctypes_binding.sh +++ b/ci/test_wheel_ctypes_binding.sh @@ -15,31 +15,22 @@ python -m pip install \ # FIXME: Find a way to build the tests that does not depend on the CUDA Python bindings #rapids-logger "Build tests" -#PY_SCRIPT=" -#import numba_cuda -#root = numba_cuda.__file__.rstrip('__init__.py') -#test_dir = root + \"numba/cuda/tests/test_binary_generation/\" -#print(test_dir) -#" -# -#NUMBA_CUDA_TEST_BIN_DIR=$(python -c "$PY_SCRIPT") -#pushd $NUMBA_CUDA_TEST_BIN_DIR -#NUMBA_CUDA_USE_NVIDIA_BINDING=0 make -#popd +rapids-logger "Copy and cd into test binaries dir" +export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing +pushd $NUMBA_CUDA_TEST_BIN_DIR +# make +# Prevent the testsuite trying to use the test binaries +unset NUMBA_CUDA_TEST_BIN_DIR rapids-logger "Check GPU usage" nvidia-smi -RAPIDS_TESTS_DIR=${RAPIDS_TESTS_DIR:-"${PWD}/test-results"}/ -mkdir -p "${RAPIDS_TESTS_DIR}" -pushd "${RAPIDS_TESTS_DIR}" - rapids-logger "Show Numba system info" NUMBA_CUDA_USE_NVIDIA_BINDING=0 python -m numba --sysinfo rapids-logger "Run Tests" # NUMBA_CUDA_USE_NVIDIA_BINDING=0 NUMBA_CUDA_TEST_BIN_DIR=$NUMBA_CUDA_TEST_BIN_DIR python -m pytest --pyargs numba.cuda.tests -v -NUMBA_CUDA_USE_NVIDIA_BINDING=0 python -m pytest --pyargs numba.cuda.tests -v +NUMBA_CUDA_USE_NVIDIA_BINDING=0 python -m pytest -v popd diff --git a/ci/test_wheel_deps_wheels.sh b/ci/test_wheel_deps_wheels.sh index 254745bc8..b71d64a2b 100755 --- a/ci/test_wheel_deps_wheels.sh +++ b/ci/test_wheel_deps_wheels.sh @@ -13,26 +13,15 @@ echo "Package path: ${package}" # TODO: control minor version pinning to honor TEST_MATRIX once the cuda-toolkit metapackage is up python -m pip install "${package}[cu${CUDA_VER_MAJOR},test-cu${CUDA_VER_MAJOR}]" -rapids-logger "Build tests" -PY_SCRIPT=" -import numba_cuda -root = numba_cuda.__file__.rstrip('__init__.py') -test_dir = root + \"numba/cuda/tests/test_binary_generation/\" -print(test_dir) -" -NUMBA_CUDA_TEST_BIN_DIR=$(python -c "$PY_SCRIPT") +rapids-logger "Build test binaries" + +export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing pushd $NUMBA_CUDA_TEST_BIN_DIR make -popd - rapids-logger "Check GPU usage" nvidia-smi -RAPIDS_TESTS_DIR=${RAPIDS_TESTS_DIR:-"${PWD}/test-results"}/ -mkdir -p "${RAPIDS_TESTS_DIR}" -pushd "${RAPIDS_TESTS_DIR}" - rapids-logger "Show Numba system info" python -m numba --sysinfo @@ -42,6 +31,6 @@ apt remove --purge `dpkg --get-selections | grep cuda-nvvm | awk '{print $1}'` - apt remove --purge `dpkg --get-selections | grep cuda-nvrtc | awk '{print $1}'` -y rapids-logger "Run Tests" -NUMBA_CUDA_TEST_BIN_DIR=$NUMBA_CUDA_TEST_BIN_DIR python -m pytest --pyargs numba.cuda.tests -v +NUMBA_CUDA_TEST_BIN_DIR=$NUMBA_CUDA_TEST_BIN_DIR python -m pytest -v popd diff --git a/numba_cuda/numba/cuda/cext/__init__.py b/numba_cuda/numba/cuda/cext/__init__.py index c54155bae..79b94ffc4 100644 --- a/numba_cuda/numba/cuda/cext/__init__.py +++ b/numba_cuda/numba/cuda/cext/__init__.py @@ -1,94 +1,2 @@ # SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: BSD-2-Clause - -import sys -import importlib -import importlib.util -import importlib.machinery -from pathlib import Path -from types import ModuleType -from importlib.machinery import ModuleSpec - - -def _load_ext_from_spec( - spec: ModuleSpec, fullname: str, legacy_name: str -) -> ModuleType: - assert spec.loader is not None - module = importlib.util.module_from_spec(spec) - sys.modules[fullname] = module - sys.modules[legacy_name] = ( - module # Register under legacy name for C extensions - ) - - # Ensure parent modules exist for legacy name (e.g., numba_cuda for numba_cuda._devicearray) - parts = legacy_name.split(".") - for i in range(1, len(parts)): - parent_name = ".".join(parts[:i]) - if parent_name not in sys.modules: - parent_module = ModuleType(parent_name) - sys.modules[parent_name] = parent_module - - # Set the child as an attribute of the parent - parent_module = sys.modules[parent_name] - child_name = parts[i] - if i == len(parts) - 1: # This is the final module - setattr(parent_module, child_name, module) - elif not hasattr(parent_module, child_name): - # Create intermediate module if it doesn't exist - intermediate_name = ".".join(parts[: i + 1]) - if intermediate_name not in sys.modules: - intermediate_module = ModuleType(intermediate_name) - sys.modules[intermediate_name] = intermediate_module - setattr(parent_module, child_name, intermediate_module) - - spec.loader.exec_module(module) - return module - - -def _find_in_dir( - module_name: str, directory: Path | str | None -) -> ModuleSpec | None: - if not directory: - return None - return importlib.machinery.PathFinder.find_spec( - module_name, [str(directory)] - ) - - -def _load_cext_module( - module_basename: str, required: bool = True -) -> ModuleType | None: - fullname = f"numba.cuda.cext.{module_basename}" - legacy_name = f"numba_cuda.{module_basename}" - - # 1) Try local numba_cuda directory (for development builds) - local_numba_cuda = Path(__file__).parents[ - 3 - ] # Go up from cext/ to numba_cuda/ - spec = _find_in_dir(module_basename, local_numba_cuda) - - # 2) Fallback: scan sys.path for installed numba_cuda directory - if spec is None: - for entry in sys.path: - numba_cuda_dir = Path(entry) / "numba_cuda/numba/cuda/cext" - spec = _find_in_dir(module_basename, numba_cuda_dir) - if spec is not None: - break - - if spec is None: - if required: - raise ModuleNotFoundError( - f"Could not find '{module_basename}' in numba_cuda directories" - ) - return None - - return _load_ext_from_spec(spec, fullname, legacy_name) - - -# Load known cext modules (all required) -# Load _devicearray first since _dispatcher depends on it -_devicearray = _load_cext_module("_devicearray", required=True) -_dispatcher = _load_cext_module("_dispatcher", required=True) -mviewbuf = _load_cext_module("mviewbuf", required=True) - -__all__ = ["mviewbuf", "_dispatcher", "_devicearray"] diff --git a/numba_cuda/numba/cuda/cext/_devicearray.h b/numba_cuda/numba/cuda/cext/_devicearray.h index e16726983..2e9df6cbb 100644 --- a/numba_cuda/numba/cuda/cext/_devicearray.h +++ b/numba_cuda/numba/cuda/cext/_devicearray.h @@ -8,7 +8,7 @@ extern "C" { #endif -#define NUMBA_DEVICEARRAY_IMPORT_NAME "numba_cuda._devicearray" +#define NUMBA_DEVICEARRAY_IMPORT_NAME "numba.cuda.cext._devicearray" /* These definitions should only be used by consumers of the Device Array API. * Consumers access the API through the opaque pointer stored in * _devicearray._DEVICEARRAY_API. We don't want these definitions in diff --git a/numba_cuda/numba/cuda/cext/_dispatcher.cpp b/numba_cuda/numba/cuda/cext/_dispatcher.cpp index 9b0b8f0dc..bc57be12a 100644 --- a/numba_cuda/numba/cuda/cext/_dispatcher.cpp +++ b/numba_cuda/numba/cuda/cext/_dispatcher.cpp @@ -947,14 +947,28 @@ import_devicearray(void) if (devicearray == NULL) { return -1; } - Py_DECREF(devicearray); - DeviceArray_API = (void**)PyCapsule_Import(NUMBA_DEVICEARRAY_IMPORT_NAME "._DEVICEARRAY_API", 0); - if (DeviceArray_API == NULL) { - return -1; + PyObject *d = PyModule_GetDict(devicearray); + if (d == NULL) { + Py_DECREF(devicearray); + return -1; } - return 0; + PyObject *key = PyUnicode_FromString("_DEVICEARRAY_API"); + PyObject *c_api = PyDict_GetItemWithError(d, key); + int retcode = 0; + if (PyCapsule_IsValid(c_api, NUMBA_DEVICEARRAY_IMPORT_NAME "._DEVICEARRAY_API")) { + DeviceArray_API = (void**)PyCapsule_GetPointer(c_api, NUMBA_DEVICEARRAY_IMPORT_NAME "._DEVICEARRAY_API"); + if (DeviceArray_API == NULL) { + retcode = -1; + } + } else { + retcode = -1; + } + + Py_DECREF(key); + Py_DECREF(devicearray); + return retcode; } static PyMethodDef Dispatcher_methods[] = { diff --git a/pyproject.toml b/pyproject.toml index e81e53950..6524836c3 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -82,13 +82,6 @@ include = ["numba_cuda*"] [tool.setuptools.package-data] "*" = ["*.cu", "*.h", "*.hpp", "*.ptx", "*.cuh", "VERSION", "Makefile"] -[tool.pytest.ini_options] -minversion = "8.0" -testpaths = ["numba_cuda/numba/cuda/tests"] -consider_namespace_packages = true -# loadscope ensures the grouping required by CUDATestCase -addopts = "--dist loadscope" - [tool.ruff] line-length = 80 diff --git a/numba_cuda/numba/cuda/tests/test_binary_generation/Makefile b/testing/Makefile similarity index 100% rename from numba_cuda/numba/cuda/tests/test_binary_generation/Makefile rename to testing/Makefile diff --git a/numba_cuda/numba/cuda/tests/test_binary_generation/build.bat b/testing/build.bat similarity index 100% rename from numba_cuda/numba/cuda/tests/test_binary_generation/build.bat rename to testing/build.bat diff --git a/conftest.py b/testing/conftest.py similarity index 100% rename from conftest.py rename to testing/conftest.py diff --git a/numba_cuda/numba/cuda/tests/test_binary_generation/generate_raw_ltoir.py b/testing/generate_raw_ltoir.py similarity index 100% rename from numba_cuda/numba/cuda/tests/test_binary_generation/generate_raw_ltoir.py rename to testing/generate_raw_ltoir.py diff --git a/numba_cuda/numba/cuda/tests/test_binary_generation/nrt_extern.cu b/testing/nrt_extern.cu similarity index 100% rename from numba_cuda/numba/cuda/tests/test_binary_generation/nrt_extern.cu rename to testing/nrt_extern.cu diff --git a/testing/pytest.ini b/testing/pytest.ini new file mode 100644 index 000000000..48e003878 --- /dev/null +++ b/testing/pytest.ini @@ -0,0 +1,9 @@ +# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: BSD-2-Clause + +[pytest] + +minversion = 8.0 +consider_namespace_packages = true +# loadscope ensures the grouping required by CUDATestCase +addopts = --dist loadscope --pyargs numba.cuda.tests diff --git a/numba_cuda/numba/cuda/tests/test_binary_generation/test_device_functions.cu b/testing/test_device_functions.cu similarity index 100% rename from numba_cuda/numba/cuda/tests/test_binary_generation/test_device_functions.cu rename to testing/test_device_functions.cu diff --git a/numba_cuda/numba/cuda/tests/test_binary_generation/undefined_extern.cu b/testing/undefined_extern.cu similarity index 100% rename from numba_cuda/numba/cuda/tests/test_binary_generation/undefined_extern.cu rename to testing/undefined_extern.cu