diff --git a/libNeonDomain/include/Neon/domain/details/bGrid/bField_imp.h b/libNeonDomain/include/Neon/domain/details/bGrid/bField_imp.h index c9c7dae7..9b9b0b7a 100644 --- a/libNeonDomain/include/Neon/domain/details/bGrid/bField_imp.h +++ b/libNeonDomain/include/Neon/domain/details/bGrid/bField_imp.h @@ -100,12 +100,22 @@ auto bField::getReference(const Neon::index_3d& cartesianIdx, NEON_THROW(exp); } auto [setIdx, bIdx] = grid.helpGetSetIdxAndGridIdx(uniformCartesianIdx); + if (setIdx.idx() == -1) { + NeonException exp("bField::getReference"); + exp << "Cannot get a writable reference to an inactive cell. Index = " << cartesianIdx; + NEON_THROW(exp); + } auto& partition = getPartition(Neon::Execution::host, setIdx, Neon::DataView::STANDARD); auto& result = partition(bIdx, cardinality); return result; } else { auto& grid = this->getGrid(); auto [setIdx, bIdx] = grid.helpGetSetIdxAndGridIdx(cartesianIdx); + if (setIdx.idx() == -1) { + NeonException exp("bField::getReference"); + exp << "Cannot get a writable reference to an inactive cell. Index = " << cartesianIdx; + NEON_THROW(exp); + } auto& partition = getPartition(Neon::Execution::host, setIdx, Neon::DataView::STANDARD); auto& result = partition(bIdx, cardinality); return result; diff --git a/libNeonPy/src/Neon/py/Container.cpp b/libNeonPy/src/Neon/py/Container.cpp index 89227269..111f471c 100644 --- a/libNeonPy/src/Neon/py/Container.cpp +++ b/libNeonPy/src/Neon/py/Container.cpp @@ -384,6 +384,7 @@ extern "C" auto warp_mGrid_container_new( using dGrid = Neon::domain::details::dGrid::dGrid; using mGrid = Neon::domain::mGrid; +using bGrid = Neon::bGrid; template auto warp_container_delete( @@ -409,6 +410,7 @@ auto warp_container_delete( DO_EXPORT(dGrid, 1, warp_container_delete, int, void**, handle); DO_EXPORT(mGrid, 1, warp_container_delete, int, void**, handle); +DO_EXPORT(bGrid, 1, warp_container_delete, int, void**, handle); template auto warp_container_parse( @@ -428,6 +430,7 @@ auto warp_container_parse( DO_EXPORT(dGrid, 1, warp_container_parse, int, void*, handle); DO_EXPORT(mGrid, 1, warp_container_parse, int, void*, handle); +DO_EXPORT(bGrid, 1, warp_container_parse, int, void*, handle); template @@ -451,6 +454,7 @@ auto warp_container_run( DO_EXPORT(dGrid, 3, warp_container_run, int, void*, handle, int, streamIdx, Neon::DataView, dataView); DO_EXPORT(mGrid, 3, warp_container_run, int, void*, handle, int, streamIdx, Neon::DataView, dataView); +DO_EXPORT(bGrid, 3, warp_container_run, int, void*, handle, int, streamIdx, Neon::DataView, dataView); template diff --git a/libNeonPy/src/Neon/py/mGrid.cpp b/libNeonPy/src/Neon/py/mGrid.cpp index e6d47bad..e8f1cccc 100644 --- a/libNeonPy/src/Neon/py/mGrid.cpp +++ b/libNeonPy/src/Neon/py/mGrid.cpp @@ -469,7 +469,12 @@ auto mGrid_mField_write( return -1; } - fieldPtr->getReference(*idx, cardinality, resolution_level) = newValue; + try { + fieldPtr->getReference(*idx, cardinality, resolution_level) = newValue; + } catch (const std::exception& e) { + NEON_ERROR("mGrid Python bindings: mField_write failed (writing to an inactive/invalid cell?): {}", e.what()); + return -1; + } // NEON_PY_DBG_COUT << "mGrid_mField_write end" << std::endl; return 0; diff --git a/py/neon/block/bGrid.py b/py/neon/block/bGrid.py index faa298d6..5d72cec2 100644 --- a/py/neon/block/bGrid.py +++ b/py/neon/block/bGrid.py @@ -177,5 +177,10 @@ def get_backend(self): def get_handle(self): return self.handle + @property + def name(self) -> str: + """Grid type name identifier.""" + return "bGrid" + def get_name(self): return "bGrid" \ No newline at end of file diff --git a/py/neon/dataView.py b/py/neon/dataView.py index e8a7aaff..6136b97e 100644 --- a/py/neon/dataView.py +++ b/py/neon/dataView.py @@ -41,7 +41,7 @@ class DataView(ctypes.Structure): Use the static factory methods to create instances. """ - _fields_ = [("data_view", ctypes.c_char)] + _fields_ = [("data_view", ctypes.c_int8)] class Values(Enum): """ @@ -68,23 +68,23 @@ def __init__(self, data_view: 'DataView.Values'): or ``DataView.boundary()`` factory methods. """ if data_view == DataView.Values.standard: - self.data_view = ctypes.c_char(b'\x00') + self.data_view = 0 elif data_view == DataView.Values.internal: - self.data_view = ctypes.c_char(b'\x01') + self.data_view = 1 elif data_view == DataView.Values.boundary: - self.data_view = ctypes.c_char(b'\x02') + self.data_view = 2 def __int__(self) -> int: """Return the integer value of this data view.""" - return int.from_bytes(self.data_view, byteorder='little') + return int(self.data_view) def __str__(self): str_repr = "" % (ctypes.addressof(self), ctypes.sizeof(self)) - if self.data_view == ctypes.c_char(b'\x00'): + if self.data_view == 0: str_repr += f"\n\tdataView: {'standard'}" - elif self.data_view == ctypes.c_char(b'\x01'): + elif self.data_view == 1: str_repr += f"\n\tdataView: {'internal'}" - elif self.data_view == ctypes.c_char(b'\x02'): + elif self.data_view == 2: str_repr += f"\n\tdataView: {'boundary'}" return str_repr diff --git a/py/tests/neon_test_utils.py b/py/tests/neon_test_utils.py index 9829217e..69e269fa 100644 --- a/py/tests/neon_test_utils.py +++ b/py/tests/neon_test_utils.py @@ -65,14 +65,6 @@ def gpu_available(): return gpu_count() > 0 -def wpne_available(): - try: - import wpne # noqa: F401 - return True - except ImportError: - return False - - MRES_FP16_REQUIRED_SYMBOLS = ( "mGrid_mField_new_float16", "mGrid_mField_delete_float16", @@ -110,16 +102,6 @@ def mres_fp16_container_available() -> bool: return False -def setup_wpne_build(script_dir): - import wpne - - wp.build.set_cpp_standard("c++17") - wp.build.add_include_directory(script_dir) - wp.build.add_preprocessor_macro_definition("NEON_WARP_COMPILATION") - wp.build.clear_kernel_cache() - wpne.init() - - def make_shell_mask(dim, shell_depth): mask = np.zeros((dim, dim, dim), dtype=np.int32) for i in range(dim): @@ -173,7 +155,3 @@ def export_vti_if_requested(field, filename, **kwargs): def require_gpu(test_item): return unittest.skipUnless(gpu_available(), "CUDA GPU not available")(test_item) - - -def require_wpne(test_item): - return unittest.skipUnless(wpne_available(), "wpne module not available")(test_item) diff --git a/py/tests/run_tests.sh b/py/tests/run_tests.sh new file mode 100755 index 00000000..7006cdda --- /dev/null +++ b/py/tests/run_tests.sh @@ -0,0 +1,54 @@ +#!/usr/bin/env bash +# Run the Neon Python test suite. +# +# Each test_*.py is run in its OWN process. Warp + Neon initialize global GPU +# state (CUDA contexts via cuCtxCreate_v4, kernel cache, etc.) that does not +# survive being re-initialized many times in a single interpreter, so running +# everything in one `unittest discover` process segfaults. Per-process +# isolation is the reliable way to run them as a suite. +# +# Usage (from an environment with the `neon` deps and CUDA on PATH): +# ./run_tests.sh +# Honors $PYTHON (default: python3) so it works under `conda run -n neon` too. + +set -u + +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]:-$0}")" && pwd)" +REPO_ROOT="$(cd "${SCRIPT_DIR}/../../.." && pwd)" +PYTHON="${PYTHON:-python3}" + +# Set PYTHONPATH / LD_LIBRARY_PATH for the local Neon build. +# env.sh derives these from $PWD, so it must be sourced from the repo root. +# (It also references LD_LIBRARY_PATH unguarded, so seed it for `set -u`.) +export LD_LIBRARY_PATH="${LD_LIBRARY_PATH:-}" +cd "${REPO_ROOT}" +# shellcheck disable=SC1091 +source "./env.sh" export >/dev/null + +cd "${SCRIPT_DIR}" + +pass=0 +fail=0 +failed_modules=() + +for f in test_*.py; do + mod="${f%.py}" + printf '=== %-32s ' "${mod}" + # shellcheck disable=SC2086 # $PYTHON may be a multi-word command (e.g. "conda run -n neon python3") + if ${PYTHON} -m unittest "${mod}" >"/tmp/neon_suite_${mod}.log" 2>&1; then + echo "PASS" + pass=$((pass + 1)) + else + echo "FAIL (see /tmp/neon_suite_${mod}.log)" + fail=$((fail + 1)) + failed_modules+=("${mod}") + fi +done + +echo +echo "===================================================" +echo "Suite result: ${pass} passed, ${fail} failed" +if [ "${fail}" -ne 0 ]; then + printf ' failed: %s\n' "${failed_modules[*]}" + exit 1 +fi diff --git a/py/tests/test_02_data_view.py b/py/tests/test_02_data_view.py index 99caab7b..165c1138 100644 --- a/py/tests/test_02_data_view.py +++ b/py/tests/test_02_data_view.py @@ -1,4 +1,3 @@ -import os import unittest from env_setup import update_pythonpath @@ -7,25 +6,15 @@ import warp as wp -from neon_test_utils import gpu_available, init_warp_neon, require_gpu, require_wpne, setup_wpne_build - -try: - import wpne - HAS_WPNE = True -except ImportError: - HAS_WPNE = False - -if HAS_WPNE: - from neon import DataView +from neon import DataView +from neon_test_utils import init_warp_neon, require_gpu @require_gpu -@require_wpne class TestDataView(unittest.TestCase): @classmethod def setUpClass(cls): init_warp_neon(verbose=False) - setup_wpne_build(os.path.dirname(os.path.abspath(__file__))) def test_print_data_views(self): @wp.kernel diff --git a/py/tests/test_04_closure.py b/py/tests/test_04_closure.py index 3191617c..cab23b70 100644 --- a/py/tests/test_04_closure.py +++ b/py/tests/test_04_closure.py @@ -1,4 +1,3 @@ -import os import unittest from env_setup import update_pythonpath @@ -11,27 +10,19 @@ from neon.dense import dSpan from neon.dense.dPartition import dPartition_int32 -from neon_test_utils import gpu_available, init_warp_neon, require_gpu, require_wpne, setup_wpne_build - -try: - import wpne - HAS_WPNE = True -except ImportError: - HAS_WPNE = False +from neon_test_utils import init_warp_neon, require_gpu @require_gpu -@require_wpne class TestClosure(unittest.TestCase): @classmethod def setUpClass(cls): init_warp_neon(verbose=False) - setup_wpne_build(os.path.dirname(os.path.abspath(__file__))) def test_kernel_without_closure(self): @wp.kernel def kernel(): - wp.neon_print(wp.NeonDenseIdx_create(11, 22, 33)) + wp.neon_print(wp.neon_idx_3d(11, 22, 33)) with wp.ScopedDevice("cuda:0"): wp.launch(kernel, dim=1, inputs=[]) @@ -56,8 +47,8 @@ def make_kernel(idx, data_view, span, partition): def kernel(): wp.neon_print(idx) wp.NeonDataView_print(data_view) - wp.NeonDenseSpan_print(span) - wp.neon_print(partition) + wp.neon_print(span) + wp.neon_print_dbg(partition) return kernel diff --git a/py/tests/axpy.py b/py/tests/test_axpy.py similarity index 97% rename from py/tests/axpy.py rename to py/tests/test_axpy.py index 0dc6892f..d9b5da75 100644 --- a/py/tests/axpy.py +++ b/py/tests/test_axpy.py @@ -226,7 +226,16 @@ def float_apxpy(dimx, neon_ngpus: int = 1): # execution(nun_devs=1, num_card=1, dim=neon.Index_3d(10, 10, 10), dtype=ctypes.c_float, # container_runtime=neon.Container.ContainerRuntime.neon) + +import unittest +from neon_test_utils import require_gpu + + +@require_gpu +class TestAxpy(unittest.TestCase): + def test_run(self): + int_apxpy(100, 1) + + if __name__ == "__main__": - # gpu1_int() - # gpu1_int() - int_apxpy(100, 1) + unittest.main() diff --git a/py/tests/block_grid.py b/py/tests/test_block_grid.py similarity index 100% rename from py/tests/block_grid.py rename to py/tests/test_block_grid.py diff --git a/py/tests/jacobi.py b/py/tests/test_jacobi.py similarity index 92% rename from py/tests/jacobi.py rename to py/tests/test_jacobi.py index dbdfbc0a..c033ad15 100644 --- a/py/tests/jacobi.py +++ b/py/tests/test_jacobi.py @@ -5,7 +5,6 @@ import os import warp as wp -import wpne import neon as ne from neon import Index_3d from neon import Ngh_idx @@ -38,9 +37,9 @@ def warp_jacobi( y[k, j, i] = tmp / wp.float(6) -@wpne.Container.factory +@ne.Container.factory() def get_jacobi(f_X, f_Y): - def axpy(loader: wpne.Loader): + def axpy(loader: ne.Loader): loader.set_grid(f_Y.get_grid()) f_x = loader.get_read_handle(f_X) @@ -63,7 +62,7 @@ def foo(idx: typing.Any): wp.int8(dj), wp.int8(dk)) unused_is_valid = wp.bool(False) - tmp = tmp + wp.neon_ngh_data(f_x, + tmp = tmp + wp.neon_read_ngh(f_x, idx, ngh, wp.int32(0), @@ -79,7 +78,7 @@ def foo(idx: typing.Any): def execution(nun_devs: int, dim: ne.Index_3d, dtype, - container_runtime: wpne.Container.ContainerRuntime): + container_runtime: ne.Container.ContainerRuntime): num_card = 1 # Get the path of the current script script_path = __file__ @@ -101,7 +100,7 @@ def execution(nun_devs: int, wp.build.clear_kernel_cache() # !!! DO THIS BEFORE DEFINING/USING ANY KERNELS WITH CUSTOM TYPES - wpne.init() + ne.init() dev_idx_list = list(range(nun_devs)) bk = ne.Backend(runtime=ne.Backend.Runtime.stream, @@ -218,19 +217,28 @@ def golden_jacobi_input(idx: ne.Index_3d): def gpu1_int(dimx, neon_ngpus: int = 1): execution(nun_devs=neon_ngpus, dim=ne.Index_3d(dimx, dimx, dimx), dtype=int, - container_runtime=wpne.Container.ContainerRuntime.neon) + container_runtime=ne.Container.ContainerRuntime.neon) def gpu1_float(dimx, neon_ngpus: int = 1): execution(nun_devs=neon_ngpus, dim=ne.Index_3d(dimx, dimx, dimx), dtype=wp.float32, - container_runtime=wpne.Container.ContainerRuntime.neon) + container_runtime=ne.Container.ContainerRuntime.neon) # def gpu1_float(): # execution(nun_devs=1, num_card=1, dim=ne.Index_3d(10, 10, 10), dtype=ctypes.c_float, -# container_runtime=wpne.Container.ContainerRuntime.neon) +# container_runtime=ne.Container.ContainerRuntime.neon) + + +import unittest +from neon_test_utils import require_gpu + + +@require_gpu +class TestJacobi(unittest.TestCase): + def test_run(self): + gpu1_float(10, 1) + if __name__ == "__main__": - # gpu1_int() - # gpu1_int() - gpu1_float(10, 1) + unittest.main() diff --git a/py/tests/mres_finer_ngh.py b/py/tests/test_mres_finer_ngh.py similarity index 92% rename from py/tests/mres_finer_ngh.py rename to py/tests/test_mres_finer_ngh.py index 31597db7..9ec2c29c 100644 --- a/py/tests/mres_finer_ngh.py +++ b/py/tests/test_mres_finer_ngh.py @@ -133,8 +133,8 @@ def peel(dim, idx, peel_level, outwards): ], stencil=[[0, 0, 0], [1, 0, 0]], ) print(grid) - A = grid.new_field(cardinality=1, dtype=wp.int32) - # B = grid.new_field(cardinality=1, dtype=wp.int32) + A = grid.new_field(cardinality=1, dtype=wp.int32, memory_type=neon.MemoryType.host_device()) + # B = grid.new_field(cardinality=1, dtype=wp.int32, memory_type=neon.MemoryType.host_device()) print("Field created") @@ -149,5 +149,15 @@ def peel(dim, idx, peel_level, outwards): A.export_vti("mres_finer_ngh","test") +import unittest +from neon_test_utils import require_gpu + + +@require_gpu +class TestMresFinerNgh(unittest.TestCase): + def test_run(self): + block_grid_try() + + if __name__ == "__main__": - block_grid_try() + unittest.main() diff --git a/py/tests/mres_global_idx.py b/py/tests/test_mres_global_idx.py similarity index 94% rename from py/tests/mres_global_idx.py rename to py/tests/test_mres_global_idx.py index b41145eb..b92913e5 100644 --- a/py/tests/mres_global_idx.py +++ b/py/tests/test_mres_global_idx.py @@ -100,5 +100,15 @@ def main(): export_vti_if_requested(field, "mres_global_idx", field_name="test") +import unittest +from neon_test_utils import require_gpu + + +@require_gpu +class TestMresGlobalIdx(unittest.TestCase): + def test_run(self): + main() + + if __name__ == "__main__": - main() + unittest.main() diff --git a/py/tests/mres_grid.py b/py/tests/test_mres_grid.py similarity index 100% rename from py/tests/mres_grid.py rename to py/tests/test_mres_grid.py diff --git a/py/tests/mres_grid_export.py b/py/tests/test_mres_grid_export.py similarity index 94% rename from py/tests/mres_grid_export.py rename to py/tests/test_mres_grid_export.py index 6600229a..76708b2e 100644 --- a/py/tests/mres_grid_export.py +++ b/py/tests/test_mres_grid_export.py @@ -166,8 +166,8 @@ def peel(dim, idx, peel_level, outwards): ], stencil=[[0, 0, 0], [1, 0, 0]], ) print(grid) - field = grid.new_field(cardinality=1, dtype=wp.int32) - c3_field = grid.new_field(cardinality=3, dtype=wp.float32) + field = grid.new_field(cardinality=1, dtype=wp.int32, memory_type=neon.MemoryType.host_device()) + c3_field = grid.new_field(cardinality=3, dtype=wp.float32, memory_type=neon.MemoryType.host_device()) print("Field created") field.export_vti("export_test","ut") @@ -212,5 +212,15 @@ def peel(dim, idx, peel_level, outwards): field.export_vti("export_test_after_has_parent_operator","has_parent_operator") +import unittest +from neon_test_utils import require_gpu + + +@require_gpu +class TestMresGridExport(unittest.TestCase): + def test_run(self): + block_grid_try() + + if __name__ == "__main__": - block_grid_try() + unittest.main() diff --git a/py/tests/mres_grid_single_level.py b/py/tests/test_mres_grid_single_level.py similarity index 90% rename from py/tests/mres_grid_single_level.py rename to py/tests/test_mres_grid_single_level.py index f099b85a..c109f6ac 100644 --- a/py/tests/mres_grid_single_level.py +++ b/py/tests/test_mres_grid_single_level.py @@ -97,6 +97,8 @@ def main(): for z in range(GRID_DIM): for y in range(GRID_DIM): for x in range(GRID_DIM): + if level_zero_mask[x, y, z] == 0: + continue idx = neon.Index_3d(x, y, z) field.write(level=0, idx=idx, cardinality=0, newValue=coord_sum(idx)) @@ -104,6 +106,8 @@ def main(): for z in range(coarse): for y in range(coarse): for x in range(coarse): + if level_one_mask[x, y, z] == 0: + continue idx = neon.Index_3d(x * 2, y * 2, z * 2) field.write(level=1, idx=idx, cardinality=0, newValue=coord_sum(idx)) @@ -114,5 +118,15 @@ def main(): export_vti_if_requested(field, "out.vti") +import unittest +from neon_test_utils import require_gpu + + +@require_gpu +class TestMresGridSingleLevel(unittest.TestCase): + def test_run(self): + main() + + if __name__ == "__main__": - main() + unittest.main() diff --git a/py/tests/mres_skeleton.py b/py/tests/test_mres_skeleton.py similarity index 94% rename from py/tests/mres_skeleton.py rename to py/tests/test_mres_skeleton.py index cd20a500..c5041fb9 100644 --- a/py/tests/mres_skeleton.py +++ b/py/tests/test_mres_skeleton.py @@ -174,8 +174,8 @@ def peel(dim, idx, peel_level, outwards): stencil=[[0, 0, 0], [1, 0, 0]], ) print(grid) - field_a = grid.new_field(cardinality=3, dtype=wp.int32) - field_b = grid.new_field(cardinality=3, dtype=wp.int32) + field_a = grid.new_field(cardinality=3, dtype=wp.int32, memory_type=neon.MemoryType.host_device()) + field_b = grid.new_field(cardinality=3, dtype=wp.int32, memory_type=neon.MemoryType.host_device()) print("Field created") @@ -203,5 +203,15 @@ def peel(dim, idx, peel_level, outwards): field_a.export_vti("mres_skeleton_field_a", "test") +import unittest +from neon_test_utils import require_gpu + + +@require_gpu +class TestMresSkeleton(unittest.TestCase): + def test_run(self): + block_grid_try() + + if __name__ == "__main__": - block_grid_try() + unittest.main() diff --git a/py/tests/mres_skeleton_stencil_up.py b/py/tests/test_mres_skeleton_stencil_up.py similarity index 93% rename from py/tests/mres_skeleton_stencil_up.py rename to py/tests/test_mres_skeleton_stencil_up.py index 9539ee0d..6ea15895 100644 --- a/py/tests/mres_skeleton_stencil_up.py +++ b/py/tests/test_mres_skeleton_stencil_up.py @@ -16,7 +16,7 @@ def set(field, level): def kernel(loader: neon.Loader): loader.set_mres_grid(field.get_grid(), level=level) - if level + 1 < field.get_grid().get_num_levels(): + if level + 1 < field.get_grid().num_levels: f = loader.get_mres_write_handle(field, operation=neon.Loader.Operation.stencil_up) else: f = loader.get_mres_write_handle(field, operation=neon.Loader.Operation.map ) @@ -177,8 +177,8 @@ def peel(dim, idx, peel_level, outwards): stencil=[[0, 0, 0], [1, 0, 0]], ) print(grid) - field_a = grid.new_field(cardinality=3, dtype=wp.int32) - field_b = grid.new_field(cardinality=3, dtype=wp.int32) + field_a = grid.new_field(cardinality=3, dtype=wp.int32, memory_type=neon.MemoryType.host_device()) + field_b = grid.new_field(cardinality=3, dtype=wp.int32, memory_type=neon.MemoryType.host_device()) print("Field created") @@ -206,5 +206,15 @@ def peel(dim, idx, peel_level, outwards): field_a.export_vti("mres_skeleton_field_a", "test") +import unittest +from neon_test_utils import require_gpu + + +@require_gpu +class TestMresSkeletonStencilUp(unittest.TestCase): + def test_run(self): + block_grid_try() + + if __name__ == "__main__": - block_grid_try() + unittest.main()