From a957f731cbd6d789c58cc1a9be555dc79adef992 Mon Sep 17 00:00:00 2001 From: Massimiliano Meneghin Date: Thu, 20 Oct 2022 21:41:44 -0400 Subject: [PATCH] Cleaning up. --- .../Neon/domain/internal/eGrid/eGrid.h | 7 +- .../src/domain/internal/eGrid/eGrid.cpp | 16 +- .../src/RunHelper.h | 12 +- .../src/containers.cu | 1 + .../src/containers.h | 7 +- .../src/staggeredGrid.cpp | 91 ++++++---- tutorials/staggered-grids/src/containers.cu | 90 +++++---- tutorials/staggered-grids/src/containers.h | 6 +- .../staggered-grids/src/staggeredGrid.cpp | 171 ++++++++++-------- 9 files changed, 235 insertions(+), 166 deletions(-) diff --git a/libNeonDomain/include/Neon/domain/internal/eGrid/eGrid.h b/libNeonDomain/include/Neon/domain/internal/eGrid/eGrid.h index c173787e..ad62c4fb 100644 --- a/libNeonDomain/include/Neon/domain/internal/eGrid/eGrid.h +++ b/libNeonDomain/include/Neon/domain/internal/eGrid/eGrid.h @@ -130,8 +130,11 @@ class eGrid : public Neon::domain::interface::GridBaseTemplate auto isInsideDomain(const Neon::index_3d& idx) const -> bool final; - auto getProperties(const Neon::index_3d& idx) const - -> GridBaseTemplate::CellProperties final; + auto getProperties(const Neon::index_3d& idx) + const -> GridBaseTemplate::CellProperties final; + + auto setReduceEngine(Neon::sys::patterns::Engine eng) + -> void; private: using GridBaseTemplate = Neon::domain::interface::GridBaseTemplate; diff --git a/libNeonDomain/src/domain/internal/eGrid/eGrid.cpp b/libNeonDomain/src/domain/internal/eGrid/eGrid.cpp index 01988334..4a77518e 100644 --- a/libNeonDomain/src/domain/internal/eGrid/eGrid.cpp +++ b/libNeonDomain/src/domain/internal/eGrid/eGrid.cpp @@ -169,7 +169,7 @@ auto eGrid::isInsideDomain(const index_3d& idx) const -> bool isInsideBox = (idx.x >= 0) && isInsideBox; isInsideBox = (idx.y >= 0) && isInsideBox; isInsideBox = (idx.z >= 0) && isInsideBox; - if(!isInsideBox){ + if (!isInsideBox) { return false; } const auto& GtoL = frame()->globalToLocal(); @@ -197,9 +197,11 @@ auto eGrid::getProperties(const index_3d& idx) const -> GridBaseTemplate::CellPr return cellProperties; } -auto eGrid::getLaunchParameters(Neon::DataView dataView, - const index_3d& blockDim, - const size_t& shareMem) const -> Neon::set::LaunchParameters +auto eGrid:: + getLaunchParameters(Neon::DataView dataView, + const index_3d& blockDim, + const size_t& shareMem) + const -> Neon::set::LaunchParameters { if (blockDim.y != 1 || blockDim.z != 1) { NeonException exc("eGrid"); @@ -218,4 +220,10 @@ auto eGrid::getLaunchParameters(Neon::DataView dataView, return newLaunchParameters; } +auto eGrid::setReduceEngine(Neon::sys::patterns::Engine ) + -> void +{ + return; +} + } // namespace Neon::domain::internal::eGrid diff --git a/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/RunHelper.h b/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/RunHelper.h index 3a255731..e354a164 100644 --- a/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/RunHelper.h +++ b/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/RunHelper.h @@ -33,11 +33,12 @@ void runAllTestConfiguration(const std::string& , } // std::vector nGpuTest{2,4,6,8}; std::vector cardinalityTest{1}; - nGpuTest = std::vector(1,1); - //std::vector dimTest{{60, 10, 250}}; - //std::vector runtimeE{Neon::Runtime::openmp, Neon::Runtime::stream}; - std::vector dimTest{{3}}; - std::vector runtimeE{Neon::Runtime::stream}; + //nGpuTest = std::vector(1,1); + + std::vector dimTest{{60, 10, 250}, {10, 50, 40}}; + std::vector runtimeE{Neon::Runtime::openmp, Neon::Runtime::stream}; + //std::vector dimTest{{3}}; + //std::vector runtimeE{Neon::Runtime::stream}; std::vector geos; @@ -50,7 +51,6 @@ void runAllTestConfiguration(const std::string& , Geometry::FullDomain, Geometry::Sphere, Geometry::HollowSphere, - }; } diff --git a/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/containers.cu b/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/containers.cu index 0e073274..b272c18a 100644 --- a/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/containers.cu +++ b/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/containers.cu @@ -166,3 +166,4 @@ auto Containers::sumVoxelsOnNodes(Self::NodeField_3 }); } template struct Containers, TEST_TYPE>; +//template struct Containers, TEST_TYPE>; diff --git a/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/containers.h b/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/containers.h index acd419df..45a31bfc 100644 --- a/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/containers.h +++ b/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/containers.h @@ -1,5 +1,7 @@ #include "Neon/domain/dGrid.h" -#include "Neon/domain/internal/experimental/staggeredGrid/StaggeredGrid.h" +#include "Neon/domain/eGrid.h" + +#include "Neon/domain/StaggeredGrid.h" #define TEST_TYPE float @@ -37,4 +39,5 @@ struct Containers }; -extern template struct Containers, TEST_TYPE>; +extern template struct Containers, TEST_TYPE>; +//extern template struct Containers, TEST_TYPE>; diff --git a/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/staggeredGrid.cpp b/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/staggeredGrid.cpp index 58747517..3ad73302 100644 --- a/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/staggeredGrid.cpp +++ b/libNeonDomain/tests/unit/domain-unit-test-staggered-grid/src/staggeredGrid.cpp @@ -19,7 +19,7 @@ void StaggeredGrid_Map(TestData& data) { // Neon::int32_3d voxDim = [&] { - auto dim = data.getGrid().getDimension(); + auto dim = data.getGrid().getDimension(); return dim; }(); @@ -36,12 +36,7 @@ void StaggeredGrid_Map(TestData& data) temperature.forEachActiveCell([](const Neon::index_3d& idx, const int& /*cardinality*/, TEST_TYPE& value) { - // if (((idx.x + idx.y + idx.z) % 2) == 0) { - // value = 10; - // } else { - // value = 0; - // } - value = TEST_TYPE(idx.x); + value = TEST_TYPE(idx.rMul()); }); @@ -57,24 +52,46 @@ void StaggeredGrid_Map(TestData& data) density.updateCompute(Neon::Backend::mainStreamIdx); const std::string appName(testFilePrefix + "_Map_" + data.getGrid().getImplementationName()); - if (EXECUTE_IO_TO_VTK == 1) { - temperature.ioToVtk(appName + "-temperature_0000", "temperature"); + if constexpr (EXECUTE_IO_TO_VTK == 1) { temperature.ioToVtk(appName + "-temperature_asVoxels", "temperature", false, Neon::IoFileType::ASCII, false); density.ioToVtk(appName + "-density_0000", "density"); } - Containers::addConstOnNodes(temperature, 50).run(Neon::Backend::mainStreamIdx); + const TEST_TYPE valueToAddOnNodes = 50; + const TEST_TYPE valueToAddOnVoxels = 33; + + Containers::addConstOnNodes(temperature, valueToAddOnNodes).run(Neon::Backend::mainStreamIdx); + Containers::addConstOnVoxels(density, valueToAddOnVoxels).run(Neon::Backend::mainStreamIdx); + temperature.updateIO(Neon::Backend::mainStreamIdx); density.updateIO(Neon::Backend::mainStreamIdx); data.getBackend().sync(Neon::Backend::mainStreamIdx); - temperature.ioToVtk(appName + "-temperature_0001", "temperature"); - // temperature.ioToVtk(appName + "-temperature_asVoxels", "temperature", false, Neon::IoFileType::ASCII, false); + if constexpr (EXECUTE_IO_TO_VTK == 1) { + temperature.ioToVtk(appName + "-temperature_0001", "temperature", false, Neon::IoFileType::ASCII); + density.ioToVtk(appName + "-density_0001", "density"); + } + + bool errorDetected = false; + + temperature.forEachActiveCell([valueToAddOnNodes, &errorDetected](const Neon::index_3d& idx, + const int& /*cardinality*/, + TEST_TYPE& value) { + auto target = TEST_TYPE(idx.rMul()) + valueToAddOnNodes; + if (target != value) { + errorDetected = true; + } + }); - density.ioToVtk(appName + "-density_0001", "density"); - // const Type alpha = 11; - // NEON_INFO(temperature.toString()); - // data.resetValuesToLinear(1, 100); + density.forEachActiveCell([valueToAddOnVoxels, &errorDetected](const Neon::index_3d& idx, + const int& /*cardinality*/, + TEST_TYPE& value) { + auto target = TEST_TYPE(idx.rSum()) + valueToAddOnVoxels; + if (target != value) { + errorDetected = true; + } + }); + ASSERT_TRUE(!errorDetected); } template @@ -111,11 +128,6 @@ void StaggeredGrid_VoxToNodes(TestData& data) nodeIDX.updateCompute(Neon::Backend::mainStreamIdx); voxelIDX.updateCompute(Neon::Backend::mainStreamIdx); - // const std::string appName(testFilePrefix + "_VoxToNodes_" + grid.getImplementationName()); - // - // nodeIDX.ioToVtk(appName + "-nodeIDX_0000", "density"); - // voxelIDX.ioToVtk(appName + "-voxelIDX_0000", "density"); - Containers::sumNodesOnVoxels(voxelIDX, nodeIDX, @@ -126,13 +138,13 @@ void StaggeredGrid_VoxToNodes(TestData& data) bool errorDetected = false; - nodeIDX.forEachActiveCell([&](const Neon::index_3d& idx, + nodeIDX.forEachActiveCell([&](const Neon::index_3d& /*idx*/, const int& /*cardinality*/, TEST_TYPE& value) { if (value != 0) { if (value == Containers::errorCode) { errorDetected = true; - std::cout << "Error detected at " << idx << std::endl; + /// std::cout << "Error detected at " << idx << std::endl; } } }); @@ -191,13 +203,12 @@ void StaggeredGrid_NodeToVoxels(TestData& data) bool errorDetected = false; - nodeIDX.forEachActiveCell([&](const Neon::index_3d& idx, + nodeIDX.forEachActiveCell([&](const Neon::index_3d& /*idx*/, const int& /*cardinality*/, TEST_TYPE& value) { if (value != 0) { if (value == Containers::errorCode) { errorDetected = true; - std::cout << "Error detected at " << idx << std::endl; } } }); @@ -225,7 +236,7 @@ TEST(Map, dGrid) Neon::init(); int nGpus = getNGpus(); using Grid = Neon::domain::dGrid; - using Type = int32_t; + using Type = TEST_TYPE; runAllTestConfiguration("staggeredGrid", StaggeredGrid_Map, nGpus, 1); } @@ -234,7 +245,7 @@ TEST(VoxToNodes, dGrid) Neon::init(); int nGpus = getNGpus(); using Grid = Neon::domain::dGrid; - using Type = int32_t; + using Type = TEST_TYPE; runAllTestConfiguration("staggeredGrid", StaggeredGrid_VoxToNodes, nGpus, 1); } @@ -243,17 +254,35 @@ TEST(NodeToVoxels, dGrid) Neon::init(); int nGpus = getNGpus(); using Grid = Neon::domain::dGrid; - using Type = int32_t; + using Type = TEST_TYPE; runAllTestConfiguration("staggeredGrid", StaggeredGrid_NodeToVoxels, nGpus, 1); } + +//TEST(Map, eGrid) +//{ +// Neon::init(); +// int nGpus = getNGpus(); +// using Grid = Neon::domain::eGrid; +// using Type = TEST_TYPE; +// runAllTestConfiguration("staggereeGrid", StaggereeGrid_Map, nGpus, 1); +//} +// +//TEST(VoxToNodes, eGrid) +//{ +// Neon::init(); +// int nGpus = getNGpus(); +// using Grid = Neon::domain::eGrid; +// using Type = TEST_TYPE; +// runAllTestConfiguration("staggereeGrid", StaggereeGrid_VoxToNodes, nGpus, 1); +//} // -// TEST(Swap, eGrid) +//TEST(NodeToVoxels, eGrid) //{ // Neon::init(); // int nGpus = getNGpus(); // using Grid = Neon::domain::eGrid; -// using Type = int32_t; -// runAllTestConfiguration("sGrid", staggeredGrid, nGpus, 1); +// using Type = TEST_TYPE; +// runAllTestConfiguration("staggereeGrid", StaggereeGrid_NodeToVoxels, nGpus, 1); //} int main(int argc, char** argv) diff --git a/tutorials/staggered-grids/src/containers.cu b/tutorials/staggered-grids/src/containers.cu index a8b0b2a5..d686da3e 100644 --- a/tutorials/staggered-grids/src/containers.cu +++ b/tutorials/staggered-grids/src/containers.cu @@ -1,7 +1,7 @@ #include "Neon/domain/tools/Geometries.h" #include "containers.h" - +namespace tools { template auto Containers::resetValue(Self::NodeField field, const Self::Type alpha) -> Neon::set::Container @@ -42,78 +42,76 @@ auto Containers::sumNodesOnVoxels(Self::VoxelField& densi const Self::NodeField& temperatureField) -> Neon::set::Container { - using Type = typename Self::NodeField::Type; - return densityField.getGrid().getContainerOnVoxels( "sumNodesOnVoxels", + // Neon Loading Lambda [&](Neon::set::Loader& loader) { auto& density = loader.load(densityField); const auto& temperature = loader.load(temperatureField, Neon::Compute::STENCIL); + // Neon Compute Lambda return [=] NEON_CUDA_HOST_DEVICE(const typename Self::VoxelField::Voxel& voxHandle) mutable { - Type sum = 0; - -#define CHECK_DIRECTION(X, Y, Z) \ - { \ - Type nghNodeValue = temperature.template getNghNodeValue(voxHandle, 0); \ - \ - sum += nghNodeValue; \ - } - - CHECK_DIRECTION(1, 1, 1); - CHECK_DIRECTION(1, 1, -1); - CHECK_DIRECTION(1, -1, 1); - CHECK_DIRECTION(1, -1, -1); - CHECK_DIRECTION(-1, 1, 1); - CHECK_DIRECTION(-1, 1, -1); - CHECK_DIRECTION(-1, -1, 1); - CHECK_DIRECTION(-1, -1, -1); - + Type sum = 0; + constexpr int componentId = 0; + + // We visit all the neighbour nodes around voxHandle. + // Relative discrete offers are used to identify the neighbour node. + // As by construction all nodes of a voxel are active, we don't have to do any extra check. + sum += temperature.template getNghNodeValue<1, 1, 1>(voxHandle, componentId); + sum += temperature.template getNghNodeValue<1, 1, -1>(voxHandle, componentId); + sum += temperature.template getNghNodeValue<1, -1, 1>(voxHandle, componentId); + sum += temperature.template getNghNodeValue<1, -1, -1>(voxHandle, componentId); + sum += temperature.template getNghNodeValue<-1, 1, 1>(voxHandle, componentId); + sum += temperature.template getNghNodeValue<-1, 1, -1>(voxHandle, componentId); + sum += temperature.template getNghNodeValue<-1, -1, 1>(voxHandle, componentId); + sum += temperature.template getNghNodeValue<-1, -1, -1>(voxHandle, componentId); + + // Storing the final result in the target voxel. density(voxHandle, 0) = sum; -#undef CHECK_DIRECTION }; }); } template -auto Containers::sumVoxelsOnNodes(Self::NodeField& temperatureField, - const Self::VoxelField& densityField) -> Neon::set::Container +auto Containers::sumVoxelsOnNodesAndDivideBy8(Self::NodeField& temperatureField, + const Self::VoxelField& densityField) -> Neon::set::Container { - - using Type = typename Self::NodeField::Type; - return temperatureField.getGrid().getContainerOnNodes( - "sumVoxelsOnNodes", + "sumVoxelsOnNodesAndDivideBy8", + // Neon Loading Lambda [&](Neon::set::Loader& loader) { const auto& density = loader.load(densityField, Neon::Compute::STENCIL); auto& temperature = loader.load(temperatureField); auto nodeSpaceDim = temperatureField.getGrid().getDimension(); + // Neon Compute Lambda return [=] NEON_CUDA_HOST_DEVICE(const typename Self::NodeField::Node& nodeHandle) mutable { Type sum = 0; -#define CHECK_DIRECTION(X, Y, Z) \ - { \ - Type nghDensity = density.template getNghVoxelValue(nodeHandle, 0, 0).value; \ - sum += nghDensity; \ - } - - CHECK_DIRECTION(1, 1, 1); - CHECK_DIRECTION(1, 1, -1); - CHECK_DIRECTION(1, -1, 1); - CHECK_DIRECTION(1, -1, -1); - CHECK_DIRECTION(-1, 1, 1); - CHECK_DIRECTION(-1, 1, -1); - CHECK_DIRECTION(-1, -1, 1); - CHECK_DIRECTION(-1, -1, -1); - - temperature(nodeHandle, 0) = sum; - ; -#undef CHECK_DIRECTION + constexpr int componentId = 0; + constexpr int returnValueIfVoxelIsNotActive = 0; + + // We visit all the neighbouring voxels around nodeHandle. + // Relative discrete offers are used to identify the neighbour node. + // Note that some neighbouring nodes may be not active. + // Rather than explicitly checking we ask Neon to return 0 if the node is not active. + sum += density.template getNghVoxelValue<1, 1, 1>(nodeHandle, componentId, returnValueIfVoxelIsNotActive).value; + sum += density.template getNghVoxelValue<1, 1, -1>(nodeHandle, componentId, returnValueIfVoxelIsNotActive).value; + sum += density.template getNghVoxelValue<1, -1, 1>(nodeHandle, componentId, returnValueIfVoxelIsNotActive).value; + sum += density.template getNghVoxelValue<1, -1, -1>(nodeHandle, componentId, returnValueIfVoxelIsNotActive).value; + sum += density.template getNghVoxelValue<-1, 1, 1>(nodeHandle, componentId, returnValueIfVoxelIsNotActive).value; + sum += density.template getNghVoxelValue<-1, 1, -1>(nodeHandle, componentId, returnValueIfVoxelIsNotActive).value; + sum += density.template getNghVoxelValue<-1, -1, 1>(nodeHandle, componentId, returnValueIfVoxelIsNotActive).value; + sum += density.template getNghVoxelValue<-1, -1, -1>(nodeHandle, componentId, returnValueIfVoxelIsNotActive).value; + + // Storing the final result in the target node. + temperature(nodeHandle, 0) = sum / 8; }; }); } + template struct Containers, double>; template struct Containers, float>; +} // namespace tools \ No newline at end of file diff --git a/tutorials/staggered-grids/src/containers.h b/tutorials/staggered-grids/src/containers.h index b8522c06..e6833776 100644 --- a/tutorials/staggered-grids/src/containers.h +++ b/tutorials/staggered-grids/src/containers.h @@ -1,6 +1,8 @@ #include "Neon/domain/dGrid.h" #include "Neon/domain/internal/experimental/staggeredGrid/StaggeredGrid.h" +namespace tools { + template struct Containers { @@ -25,7 +27,7 @@ struct Containers const Self::NodeField& fieldNode) -> Neon::set::Container; - static auto sumVoxelsOnNodes(Self::NodeField& fieldNode, + static auto sumVoxelsOnNodesAndDivideBy8(Self::NodeField& fieldNode, const Self::VoxelField& fieldVox) -> Neon::set::Container; }; @@ -33,3 +35,5 @@ struct Containers extern template struct Containers, double>; extern template struct Containers, float>; + +} // namespace tools \ No newline at end of file diff --git a/tutorials/staggered-grids/src/staggeredGrid.cpp b/tutorials/staggered-grids/src/staggeredGrid.cpp index f1b92224..68c1e036 100644 --- a/tutorials/staggered-grids/src/staggeredGrid.cpp +++ b/tutorials/staggered-grids/src/staggeredGrid.cpp @@ -1,26 +1,25 @@ #include "gtest/gtest.h" #include "Neon/Neon.h" + #include "Neon/domain/dGrid.h" -#include "Neon/domain/internal/experimental/staggeredGrid/StaggeredGrid.h" +#include "Neon/domain/StaggeredGrid.h" #include "containers.h" +/** + * A simple tutorial demonstrating the use of staggered grid in Neon. + */ int main() { - // Step 0 -> initialize Neon runtime - Neon::init(); - - // Step 1 -> Neon backend: choosing the hardware for the computation + // Selecting the hardware for the computation Neon::Backend backend = [] { - // auto runtime = Neon::Runtime::openmp; + Neon::init(); + // Our XPU will be a CPU device. auto runtime = Neon::Runtime::openmp; - // We are overbooking GPU 0 three times - std::vector gpu_ids{0,0}; + // We are overbooking XPU 0 two times + std::vector gpu_ids{0, 0}; Neon::Backend backend(gpu_ids, runtime); - - // Printing some information - NEON_INFO(backend.toString()); return backend; }(); @@ -28,13 +27,15 @@ int main() using UniformGrid = Neon::domain::dGrid; using StaggeredGrid = Neon::domain::internal::experimental::staggeredGrid::StaggeredGrid; using FP = double; + using UserContainers = tools::Containers; + // Initializing a staggered grid based on dGrid StaggeredGrid grid = [&] { // Setting the dimension for the voxel grid - Neon::int32_3d voxDims{2, 3, 9}; + Neon::int32_3d voxDims{2, 3, 9}; - // For our example we don't need to add new stencil. - // By default the staggered grid will add the stencil + // For our tutorial, we don't need to add new stencil. + // By default, the staggered grid will add the stencil // to move from voxel to node and vice versa std::vector noExtraStencils; @@ -45,80 +46,102 @@ int main() return true; }, noExtraStencils); + return newGrid; }(); - // Defining a voxel field to represent the density of each voxel - auto density = grid.template newVoxelField("Density", 1, 0); - density.forEachActiveCell([](const Neon::index_3d& /*idx*/, - const int& /*cardinality*/, - FP& value) { - value = 1; - }); - - // Defining a node field to represent the temperature at each node - auto temperature = grid.template newNodeField("Temperature", 1, 0); - temperature.forEachActiveCell([](const Neon::index_3d& /*idx*/, - const int& /*cardinality*/, - FP& value) { - value = 0; - }); - - // Exporting the initial values of the fields - // For the temperature field we choose to use binary format for vti file. - const std::string appName("staggered-grid"); - temperature.ioToVtk(appName + "-temperature_binary_0000", "temperature", false, Neon::IoFileType::BINARY); - density.ioToVtk(appName + "-density_0000", "density"); - - { // Moving memory from CPU to the GPUs - temperature.updateCompute(Neon::Backend::mainStreamIdx); + auto density = [&grid] { + // Defining a voxel field to represent the density of each voxel + // We then set its initial values to zero in the host (CPU), + // and finally we transfer the data to the XPU + auto density = grid.template newVoxelField("Density", 1, 0); + density.forEachActiveCell([](const Neon::index_3d& /*idx*/, + const int& /*cardinality*/, + FP& value) { + value = 0; + }); density.updateCompute(Neon::Backend::mainStreamIdx); - } + return density; + }(); + auto temperature = [&grid] { + // Defining a node field to represent the temperature at each node + // We then set its initial values to one, but to make it more interesting, + // we do the initialization directly on the device calling a Containers. + // Note that in this case we don't need to transfer data from host to XPUs + auto temperature = grid.template newNodeField("Temperature", 1, 0); + temperature.forEachActiveCell([](const Neon::index_3d& /*idx*/, + const int& /*cardinality*/, + FP& value) { + value = 0; + }); + UserContainers::resetValue(temperature, FP(1.0)).run(Neon::Backend::mainStreamIdx); + return temperature; + }(); - { // Accessing voxels from nodes - { - // We use halo update as we are working only at the domain level. - // When leveraging the skeleton abstractions, hlo update are automatically - // handled by Neon. - Neon::set::HuOptions huOptions(Neon::set::TransferMode::get, - true, - Neon::Backend::mainStreamIdx); - density.haloUpdate(huOptions); + // We define a simple function to export to vtk both + // temperature and density field during each step of the tutorial. + auto exportingToVti = [&](const std::string& tutorialStepId) { + { // Moving memory from XPUs to CPU + temperature.updateIO(Neon::Backend::mainStreamIdx); + density.updateIO(Neon::Backend::mainStreamIdx); + backend.sync(Neon::Backend::mainStreamIdx); } - - // For each nodes we loop around active voxel and we sum their values. - Containers::sumVoxelsOnNodes(temperature, density).run(Neon::Backend::mainStreamIdx); - - temperature.updateIO(Neon::Backend::mainStreamIdx); - density.updateIO(Neon::Backend::mainStreamIdx); - backend.sync(Neon::Backend::mainStreamIdx); - - // Exporting the results - temperature.ioToVtk(appName + "-temperature_binary_0001", "temperature", false, Neon::IoFileType::BINARY); - density.ioToVtk(appName + "-density_0001", "density"); - } - - { // Accessing nodes from voxels + { // Exporting the results + const std::string appName("staggered-grid"); + temperature.ioToVtk(appName + "-temperature-" + tutorialStepId, "temperature", false, Neon::IoFileType::BINARY); + density.ioToVtk(appName + "-density-" + tutorialStepId, "density"); + } + }; + + // We are exporting to vtk the values of the fields after the initialization. + // We expect all temperature nodes to be set to one, + // and all density voxels to be set to zero. + exportingToVti("0000"); + + { // As the previous container changes the values of the voxel field, + // a halo update must be called before the next Container that uses + // the temperature as input for a stencil operation. + // + // Because we wrote this tutorial only using the Domain level, we have to + // execute the halo update manually. + // + // However, when using the Neon skeleton level halo update are + // automatically managed. Neon::set::HuOptions huOptions(Neon::set::TransferMode::get, true, Neon::Backend::mainStreamIdx); - - // We reset all node value to 1. - Containers::resetValue(temperature, FP(1.0)).run(Neon::Backend::mainStreamIdx); temperature.haloUpdate(huOptions); + } - // For each voxel we loop around all nodes and we sum their values. - // Note that all nodes of a voxel are always active. - Containers::sumNodesOnVoxels(density, temperature).run(Neon::Backend::mainStreamIdx); + { // Accessing voxels from nodes + // For each node we loop around active voxel, and we sum their values. + // At the end of this operation we expect all voxel to store a value of 8, + // as there are 8 nodes for each voxel, each one set to one. + UserContainers::sumNodesOnVoxels(density, temperature) + .run(Neon::Backend::mainStreamIdx); - temperature.updateIO(Neon::Backend::mainStreamIdx); - density.updateIO(Neon::Backend::mainStreamIdx); - backend.sync(Neon::Backend::mainStreamIdx); + exportingToVti("0001"); + } - // Exporting the results - temperature.ioToVtk(appName + "-temperature_binary_0002", "temperature", false, Neon::IoFileType::BINARY); - density.ioToVtk(appName + "-density_0002", "density"); + { // The previous comment on halo updates is applicable here. + Neon::set::HuOptions huOptions(Neon::set::TransferMode::get, + true, + Neon::Backend::mainStreamIdx); + density.haloUpdate(huOptions); + } + { // Accessing nodes from voxels + // For each voxel we loop around all nodes, we sum their values and divide the result by 8. + // At the end of the container we expect to have the following values: + // -- 1 for any corner node + // -- 2 for any node on an edge of our domain + // -- 3 for any node on a face + // -- 4 for internal nodes + UserContainers::sumVoxelsOnNodesAndDivideBy8(temperature, + density) + .run(Neon::Backend::mainStreamIdx); + + exportingToVti("0002"); } return 0; }