Fix tests and add runner script

libNeonDomain/include/Neon/domain/details/bGrid/bField_imp.h

@@ -100,12 +100,22 @@ auto bField<T, C, SBlock>::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;

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 <typename Grid>
 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 <typename Grid>
 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 <typename Grid>
@@ -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 <typename Grid, typename Type, int Card>

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;

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"

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 = "<DDDData_view: addr=%ld, sizeof %ld>" % (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
 

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)

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

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

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
 

py/tests/axpy.py → 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()

py/tests/jacobi.py → 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()