From 73a6561ce72e1e99a90f642eafc35dcc252ebb1a Mon Sep 17 00:00:00 2001 From: jdolence Date: Mon, 9 Oct 2023 13:55:26 -0600 Subject: [PATCH 01/17] add generic evaluator --- singularity-eos/eos/eos_base.hpp | 8 ++++++++ singularity-eos/eos/eos_variant.hpp | 10 ++++++++++ 2 files changed, 18 insertions(+) diff --git a/singularity-eos/eos/eos_base.hpp b/singularity-eos/eos/eos_base.hpp index f206a01a58..8bf3b207dd 100644 --- a/singularity-eos/eos/eos_base.hpp +++ b/singularity-eos/eos/eos_base.hpp @@ -128,6 +128,14 @@ class EosBase { struct is_raw_pointer : std::is_same>, R *> {}; + // Generic evaluator + template + PORTABLE_INLINE_FUNCTION + void Evaluate(const Functor_t &f) const { + CRTP copy = *(static_cast(this)); + f(copy); + } + // Vector member functions template inline void diff --git a/singularity-eos/eos/eos_variant.hpp b/singularity-eos/eos/eos_variant.hpp index 04dc51a9bf..09632d1634 100644 --- a/singularity-eos/eos/eos_variant.hpp +++ b/singularity-eos/eos/eos_variant.hpp @@ -72,6 +72,16 @@ class Variant { } // Place member functions here + template + PORTABLE_INLINE_FUNCTION + void Evaluate(const Functor_t &f) const { + return mpark::visit( + [&f](const auto &eos) { + return eos.Evaluate(f); + }, + eos_); + } + PORTABLE_INLINE_FUNCTION Real TemperatureFromDensityInternalEnergy(const Real rho, const Real sie, Real *lambda = nullptr) const { From 8813cd9c82bab8155b61d78f4659b3a442a734ef Mon Sep 17 00:00:00 2001 From: jdolence Date: Mon, 9 Oct 2023 14:21:13 -0600 Subject: [PATCH 02/17] formatting --- singularity-eos/eos/eos_base.hpp | 3 +-- singularity-eos/eos/eos_variant.hpp | 9 ++------- 2 files changed, 3 insertions(+), 9 deletions(-) diff --git a/singularity-eos/eos/eos_base.hpp b/singularity-eos/eos/eos_base.hpp index 8bf3b207dd..08c292f48a 100644 --- a/singularity-eos/eos/eos_base.hpp +++ b/singularity-eos/eos/eos_base.hpp @@ -130,8 +130,7 @@ class EosBase { // Generic evaluator template - PORTABLE_INLINE_FUNCTION - void Evaluate(const Functor_t &f) const { + PORTABLE_INLINE_FUNCTION void Evaluate(const Functor_t &f) const { CRTP copy = *(static_cast(this)); f(copy); } diff --git a/singularity-eos/eos/eos_variant.hpp b/singularity-eos/eos/eos_variant.hpp index 09632d1634..9d971ff454 100644 --- a/singularity-eos/eos/eos_variant.hpp +++ b/singularity-eos/eos/eos_variant.hpp @@ -73,13 +73,8 @@ class Variant { // Place member functions here template - PORTABLE_INLINE_FUNCTION - void Evaluate(const Functor_t &f) const { - return mpark::visit( - [&f](const auto &eos) { - return eos.Evaluate(f); - }, - eos_); + PORTABLE_INLINE_FUNCTION void Evaluate(const Functor_t &f) const { + return mpark::visit([&f](const auto &eos) { return eos.Evaluate(f); }, eos_); } PORTABLE_INLINE_FUNCTION From 4b23de59eb094e41ae791a54a253b9c5d96db77c Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Mon, 9 Oct 2023 19:46:00 -0600 Subject: [PATCH 03/17] add docs and a test. make functor not have to be const --- doc/sphinx/src/using-eos.rst | 86 +++++++++++++++++++++++++++++ singularity-eos/eos/eos_base.hpp | 2 +- singularity-eos/eos/eos_variant.hpp | 2 +- test/test_eos_unit.cpp | 67 ++++++++++++++++++++++ 4 files changed, 155 insertions(+), 2 deletions(-) diff --git a/doc/sphinx/src/using-eos.rst b/doc/sphinx/src/using-eos.rst index bbb3e413a7..3405cc0cc2 100644 --- a/doc/sphinx/src/using-eos.rst +++ b/doc/sphinx/src/using-eos.rst @@ -216,6 +216,92 @@ available member function. eos.TemperatureFromDensityInternalEnergy(density.data(), energy.data(), temperature.data(), scratch.data(), density.size()); +The Evaluate Method +~~~~~~~~~~~~~~~~~~~ + +A special call related to the vector calls is the ``Evaluate`` +method. The ``Evaluate`` method requests the EOS object to evaluate +almost arbitrary code, but in a way where the type of the underlying +EOS object is resolved *before* this arbitrary code is evaluated. This +means the code required to resolve the type of the variant is only +executed *once* per ``Evaluate`` call. This can enable composite EOS +calls, non-standard vector calls, and vector calls with non-standard +loop structure. + +The ``Evaluate`` call has the signature + +.. code-block:: cpp + + template + PORTABLE_INLINE_FUNCTION + void Evaluate(Functor_t f); + +where a ``Functor_t`` is a class that *must* provide a ``void +operator() const`` method templated on EOS type. ``Evaluate`` is decorated +so that it may be evaluated on either host or device, depending on +desired use-case. + +.. note:: + + Note that for C++ versions less than C++20, the functor must be + defined as a class. Anonymous functions won't do, as they do not + support templates, and the ``operator()`` call must be + templated. + +To see the utlity of the ``Evaluate`` function, it's probably just +easiest to provide an example. The following code evaluates the EOS on +device and compares against a tabulated pressure. The total difference +is summed using the ``Kokkos::parallel_reduce`` functionality in the +``Kokkos`` performance portability library. + +.. code-block:: cpp + + // The functor we use is defined here. + // This class definition needs to be of appropriately global scope. + class CheckPofRE { + public: + CheckPofRE(Real *P, Real *rho, Real *sie, int N) : P_(P), rho_(rho), sie_(sie), N_(N) {} + template + // this is a host-only call, but if you wanted to write + // a function that you wanted to evaluate on device + // you could add the + // PORTABLE_INLINE_FUNCTION + // decorator here. + void operator()(const T &eos) const { + Real *P = P_; // gets around warnings regarding capture of "this" pointer + Real *rho = rho_; + Real *sie = sie_; + Kokkos::parallel_reduce( + "MyCheckPofRE", N_, + KOKKOS_LAMBDA(const int i, Real &diff) { + diff += std::abs(P[i] - eos.PressureFromDensityInternalEnergy(rho[i], sie[i])); + }, + tot_diff); + std::cout << "Total difference = " << tot_diff << std::endl; + } + + private: + int N_; + Real *P_; + Real *rho_; + Real *sie_; + }; + + // Here we construct our functor + // it is assumed the pointers to device memory P, rho, sie, are defined elsewhere. + CheckPofRE my_op(P, rho, sie, N); + + // Here we call the evaluate function + eos.Evaluate(my_op); + + // The above two lines could have been called "in-one" with: + // eos.Evaluate(CheckPofRE(P, rho, sie, N)); + +This is not functionality that would be available with the standard +vector calls provided by ``singularity-eos``, at least not without +chaining multiple parallel dispatch calls. Here we can do it in a +single call. + Lambdas and Optional Parameters -------------------------------- diff --git a/singularity-eos/eos/eos_base.hpp b/singularity-eos/eos/eos_base.hpp index 08c292f48a..a3f2a4e571 100644 --- a/singularity-eos/eos/eos_base.hpp +++ b/singularity-eos/eos/eos_base.hpp @@ -130,7 +130,7 @@ class EosBase { // Generic evaluator template - PORTABLE_INLINE_FUNCTION void Evaluate(const Functor_t &f) const { + PORTABLE_INLINE_FUNCTION void Evaluate(Functor_t &f) const { CRTP copy = *(static_cast(this)); f(copy); } diff --git a/singularity-eos/eos/eos_variant.hpp b/singularity-eos/eos/eos_variant.hpp index 9d971ff454..d2401636ef 100644 --- a/singularity-eos/eos/eos_variant.hpp +++ b/singularity-eos/eos/eos_variant.hpp @@ -73,7 +73,7 @@ class Variant { // Place member functions here template - PORTABLE_INLINE_FUNCTION void Evaluate(const Functor_t &f) const { + PORTABLE_INLINE_FUNCTION void Evaluate(Functor_t &f) const { return mpark::visit([&f](const auto &eos) { return eos.Evaluate(f); }, eos_); } diff --git a/test/test_eos_unit.cpp b/test/test_eos_unit.cpp index b595d4e16f..84fe7de4c6 100644 --- a/test/test_eos_unit.cpp +++ b/test/test_eos_unit.cpp @@ -509,6 +509,73 @@ SCENARIO("Ideal gas entropy", "[IdealGas][Entropy]") { } } +// A non-standard pattern where we do a reduction +class CheckPofRE { + public: + CheckPofRE(Real *P, Real *rho, Real *sie, int N) : P_(P), rho_(rho), sie_(sie), N_(N) {} + template + void operator()(const T &eos) const { + Real *P = P_; // gets around capture of this pointer + Real *rho = rho_; + Real *sie = sie_; + // must be initialized to zero because portableReduce is a simple for loop on host + int nwrong = 0; + portableReduce( + "MyCheckPofRE", 0, N_, + PORTABLE_LAMBDA(const int i, int &nw) { + nw += !(isClose(P[i], + eos.PressureFromDensityInternalEnergy(rho[i], sie[i], nullptr), + 1e-15)); + }, + nwrong); + REQUIRE(nwrong == 0); + } + + private: + int N_; + Real *P_; + Real *rho_; + Real *sie_; +}; +SCENARIO("Ideal gas vector Evaluate call", "[IdealGas][Evaluate]") { + GIVEN("An ideal gas, and some device memory") { + constexpr Real Cv = 2.0; + constexpr Real gm1 = 0.5; + constexpr int N = 100; + constexpr Real rhomin = 1; + constexpr Real rhomax = 11; + constexpr Real drho = (rhomax - rhomin) / (N - 1); + constexpr Real siemin = 1; + constexpr Real siemax = 11; + constexpr Real dsie = (siemax - siemin) / (N - 1); + + EOS eos_host = IdealGas(gm1, Cv); + auto eos_device = eos_host.GetOnDevice(); + + Real *P = (Real *)PORTABLE_MALLOC(N * sizeof(Real)); + Real *rho = (Real *)PORTABLE_MALLOC(N * sizeof(Real)); + Real *sie = (Real *)PORTABLE_MALLOC(N * sizeof(Real)); + + WHEN("We set P, rho, sie via the scalar API") { + portableFor( + "Set rho and sie", 0, N, PORTABLE_LAMBDA(const int i) { + rho[i] = rhomin + drho * i; // just some diagonal in the 2d rho/sie space + sie[i] = siemin + dsie * i; + P[i] = eos_device.PressureFromDensityInternalEnergy(rho[i], sie[i]); + }); + THEN("The vector Evaluate API can be used to compare") { + CheckPofRE myOp(P, rho, sie, N); + eos_device.Evaluate(myOp); + } + } + + eos_device.Finalize(); + PORTABLE_FREE(P); + PORTABLE_FREE(rho); + PORTABLE_FREE(sie); + } +} + SCENARIO("Vector EOS", "[VectorEOS][IdealGas]") { GIVEN("Parameters for an ideal gas") { From de6202d592269362f140d477e3aee27e514f6155 Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Mon, 9 Oct 2023 19:47:35 -0600 Subject: [PATCH 04/17] add a warning about const --- doc/sphinx/src/using-eos.rst | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/doc/sphinx/src/using-eos.rst b/doc/sphinx/src/using-eos.rst index 3405cc0cc2..bd2e93952b 100644 --- a/doc/sphinx/src/using-eos.rst +++ b/doc/sphinx/src/using-eos.rst @@ -248,6 +248,13 @@ desired use-case. support templates, and the ``operator()`` call must be templated. +.. warning:: + + It can be dangerous to use functors with side-effects. Especially + with GPUs it can produce very unintuitive behaviour. We recommend + you only make the ``operator()`` non-const if you really know what + you're doing. + To see the utlity of the ``Evaluate`` function, it's probably just easiest to provide an example. The following code evaluates the EOS on device and compares against a tabulated pressure. The total difference From 9d19dd4d57e5177f3f4da035ccd5d01eafa71a83 Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Mon, 9 Oct 2023 19:48:17 -0600 Subject: [PATCH 05/17] added a few more comments to the docs --- doc/sphinx/src/using-eos.rst | 3 +++ 1 file changed, 3 insertions(+) diff --git a/doc/sphinx/src/using-eos.rst b/doc/sphinx/src/using-eos.rst index bd2e93952b..2559cf65bc 100644 --- a/doc/sphinx/src/using-eos.rst +++ b/doc/sphinx/src/using-eos.rst @@ -278,6 +278,9 @@ is summed using the ``Kokkos::parallel_reduce`` functionality in the Real *P = P_; // gets around warnings regarding capture of "this" pointer Real *rho = rho_; Real *sie = sie_; + // reduction target + Real tot_diff; + // reduction op Kokkos::parallel_reduce( "MyCheckPofRE", N_, KOKKOS_LAMBDA(const int i, Real &diff) { From bc25dae7c912a30ade9b6075a8bbe985ea8336b1 Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Mon, 9 Oct 2023 19:53:48 -0600 Subject: [PATCH 06/17] simplify code a bit --- doc/sphinx/src/using-eos.rst | 5 +---- test/test_eos_unit.cpp | 20 +++++++++----------- 2 files changed, 10 insertions(+), 15 deletions(-) diff --git a/doc/sphinx/src/using-eos.rst b/doc/sphinx/src/using-eos.rst index 2559cf65bc..9867955870 100644 --- a/doc/sphinx/src/using-eos.rst +++ b/doc/sphinx/src/using-eos.rst @@ -275,16 +275,13 @@ is summed using the ``Kokkos::parallel_reduce`` functionality in the // PORTABLE_INLINE_FUNCTION // decorator here. void operator()(const T &eos) const { - Real *P = P_; // gets around warnings regarding capture of "this" pointer - Real *rho = rho_; - Real *sie = sie_; // reduction target Real tot_diff; // reduction op Kokkos::parallel_reduce( "MyCheckPofRE", N_, KOKKOS_LAMBDA(const int i, Real &diff) { - diff += std::abs(P[i] - eos.PressureFromDensityInternalEnergy(rho[i], sie[i])); + diff += std::abs(P_[i] - eos.PressureFromDensityInternalEnergy(rho_[i], sie_[i])); }, tot_diff); std::cout << "Total difference = " << tot_diff << std::endl; diff --git a/test/test_eos_unit.cpp b/test/test_eos_unit.cpp index 84fe7de4c6..be732ba21e 100644 --- a/test/test_eos_unit.cpp +++ b/test/test_eos_unit.cpp @@ -514,23 +514,20 @@ class CheckPofRE { public: CheckPofRE(Real *P, Real *rho, Real *sie, int N) : P_(P), rho_(rho), sie_(sie), N_(N) {} template - void operator()(const T &eos) const { - Real *P = P_; // gets around capture of this pointer - Real *rho = rho_; - Real *sie = sie_; - // must be initialized to zero because portableReduce is a simple for loop on host - int nwrong = 0; + void operator()(const T &eos) { portableReduce( "MyCheckPofRE", 0, N_, PORTABLE_LAMBDA(const int i, int &nw) { - nw += !(isClose(P[i], - eos.PressureFromDensityInternalEnergy(rho[i], sie[i], nullptr), + nw += !(isClose(P_[i], + eos.PressureFromDensityInternalEnergy(rho_[i], sie_[i], nullptr), 1e-15)); }, nwrong); - REQUIRE(nwrong == 0); } + // must be initialized to zero because portableReduce is a simple for loop on host + int nwrong = 0; + private: int N_; Real *P_; @@ -564,8 +561,9 @@ SCENARIO("Ideal gas vector Evaluate call", "[IdealGas][Evaluate]") { P[i] = eos_device.PressureFromDensityInternalEnergy(rho[i], sie[i]); }); THEN("The vector Evaluate API can be used to compare") { - CheckPofRE myOp(P, rho, sie, N); - eos_device.Evaluate(myOp); + CheckPofRE my_op(P, rho, sie, N); + eos_device.Evaluate(my_op); + REQUIRE(my_op.nwrong == 0); } } From efe6cfdb3b4c79d1321f8b9056eef0198666ffd0 Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Mon, 9 Oct 2023 19:56:55 -0600 Subject: [PATCH 07/17] formatting --- test/test_eos_unit.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/test/test_eos_unit.cpp b/test/test_eos_unit.cpp index be732ba21e..13042f3455 100644 --- a/test/test_eos_unit.cpp +++ b/test/test_eos_unit.cpp @@ -518,9 +518,9 @@ class CheckPofRE { portableReduce( "MyCheckPofRE", 0, N_, PORTABLE_LAMBDA(const int i, int &nw) { - nw += !(isClose(P_[i], - eos.PressureFromDensityInternalEnergy(rho_[i], sie_[i], nullptr), - 1e-15)); + nw += !(isClose( + P_[i], eos.PressureFromDensityInternalEnergy(rho_[i], sie_[i], nullptr), + 1e-15)); }, nwrong); } From 4d2938514f4a96800f4e63f391e3a9c80d68da2c Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Mon, 9 Oct 2023 19:59:05 -0600 Subject: [PATCH 08/17] changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index d555491177..8aef7c8d8d 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -22,6 +22,7 @@ - [[PR232]](https://github.com/lanl/singularity-eos/pull/228) Fixed uninitialized cmake path variables ### Added (new features/APIs/variables/...) +- [[PR306]](https://github.com/lanl/singularity-eos/pull/306) Added generic Evaluate method - [[PR304]](https://github.com/lanl/singularity-eos/pull/304) added a Newton-Raphson root find for use with the Helmholtz EoS - [[PR265]](https://github.com/lanl/singularity-eos/pull/265) Add missing UnitSystem modifier combinations to variant and EOSPAC - [[PR279]](https://github.com/lanl/singularity-eos/pull/279) added noble-abel EoS From a01b6b0e52251886771c04e3bf5a523c44a5ac4c Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Tue, 10 Oct 2023 18:41:53 -0600 Subject: [PATCH 09/17] Changelog and copyright --- test/test_eos_unit.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/test/test_eos_unit.cpp b/test/test_eos_unit.cpp index 13042f3455..34e86d071c 100644 --- a/test/test_eos_unit.cpp +++ b/test/test_eos_unit.cpp @@ -513,7 +513,10 @@ SCENARIO("Ideal gas entropy", "[IdealGas][Entropy]") { class CheckPofRE { public: CheckPofRE(Real *P, Real *rho, Real *sie, int N) : P_(P), rho_(rho), sie_(sie), N_(N) {} +#pragma nv_exec_check_disable +#pragma hd_warning_disable template + PORTABLE_FUNCTION void operator()(const T &eos) { portableReduce( "MyCheckPofRE", 0, N_, From 6d58fabb6f6fbd78b0943e26ff42db8831b498fa Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Tue, 10 Oct 2023 18:43:09 -0600 Subject: [PATCH 10/17] capture of this pointer --- test/test_eos_unit.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/test/test_eos_unit.cpp b/test/test_eos_unit.cpp index 34e86d071c..19b19961ea 100644 --- a/test/test_eos_unit.cpp +++ b/test/test_eos_unit.cpp @@ -518,11 +518,14 @@ class CheckPofRE { template PORTABLE_FUNCTION void operator()(const T &eos) { + int *P = P_; + int *rho = rho_; + int *sie = sie_; portableReduce( "MyCheckPofRE", 0, N_, PORTABLE_LAMBDA(const int i, int &nw) { nw += !(isClose( - P_[i], eos.PressureFromDensityInternalEnergy(rho_[i], sie_[i], nullptr), + P[i], eos.PressureFromDensityInternalEnergy(rho[i], sie[i], nullptr), 1e-15)); }, nwrong); From ec7891a188f6e40f4a02fa4d658dbc548be828c8 Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Tue, 10 Oct 2023 18:43:46 -0600 Subject: [PATCH 11/17] oops --- test/test_eos_unit.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/test/test_eos_unit.cpp b/test/test_eos_unit.cpp index 19b19961ea..96fd7dd984 100644 --- a/test/test_eos_unit.cpp +++ b/test/test_eos_unit.cpp @@ -518,9 +518,9 @@ class CheckPofRE { template PORTABLE_FUNCTION void operator()(const T &eos) { - int *P = P_; - int *rho = rho_; - int *sie = sie_; + Real *P = P_; + Real *rho = rho_; + Real *sie = sie_; portableReduce( "MyCheckPofRE", 0, N_, PORTABLE_LAMBDA(const int i, int &nw) { From a9ab2fdf0cffe706963c10e90a21ebbee71ce72d Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Tue, 10 Oct 2023 19:07:29 -0600 Subject: [PATCH 12/17] update docs to reflect auto and this pointer --- doc/sphinx/src/using-eos.rst | 41 +++++++++++++++++++++++++++--------- 1 file changed, 31 insertions(+), 10 deletions(-) diff --git a/doc/sphinx/src/using-eos.rst b/doc/sphinx/src/using-eos.rst index 9867955870..1568f6446f 100644 --- a/doc/sphinx/src/using-eos.rst +++ b/doc/sphinx/src/using-eos.rst @@ -237,23 +237,22 @@ The ``Evaluate`` call has the signature void Evaluate(Functor_t f); where a ``Functor_t`` is a class that *must* provide a ``void -operator() const`` method templated on EOS type. ``Evaluate`` is decorated -so that it may be evaluated on either host or device, depending on -desired use-case. +operator() const`` method templated on EOS type. ``Evaluate`` is +decorated so that it may be evaluated on either host or device, +depending on desired use-case. Alternatively, you may use an anonymous +function with an `auto` argument as the input, e.g., -.. note:: +.. code-block:: - Note that for C++ versions less than C++20, the functor must be - defined as a class. Anonymous functions won't do, as they do not - support templates, and the ``operator()`` call must be - templated. + eos.Evaluate([=](auto eos) { /* my code snippet */ }); .. warning:: It can be dangerous to use functors with side-effects. Especially with GPUs it can produce very unintuitive behaviour. We recommend you only make the ``operator()`` non-const if you really know what - you're doing. + you're doing. And in the anonymous function case, we recommend you + capture by value, not reference. To see the utlity of the ``Evaluate`` function, it's probably just easiest to provide an example. The following code evaluates the EOS on @@ -275,13 +274,19 @@ is summed using the ``Kokkos::parallel_reduce`` functionality in the // PORTABLE_INLINE_FUNCTION // decorator here. void operator()(const T &eos) const { + // Capturing member functions of a class in a lambda typically causes problems + // when launching a GPU kernel. + // Better to pull out new variables to capture before launching a kernel. + Real *P = P_; + Real *rho = rho_; + Real *sie = sie_; // reduction target Real tot_diff; // reduction op Kokkos::parallel_reduce( "MyCheckPofRE", N_, KOKKOS_LAMBDA(const int i, Real &diff) { - diff += std::abs(P_[i] - eos.PressureFromDensityInternalEnergy(rho_[i], sie_[i])); + diff += std::abs(P[i] - eos.PressureFromDensityInternalEnergy(rho[i], sie[i])); }, tot_diff); std::cout << "Total difference = " << tot_diff << std::endl; @@ -304,6 +309,22 @@ is summed using the ``Kokkos::parallel_reduce`` functionality in the // The above two lines could have been called "in-one" with: // eos.Evaluate(CheckPofRE(P, rho, sie, N)); +Alternatively, you could eliminate the functor and use an anonymous +function with: + +.. code-block:: cpp + + eos.Evaluate([=](auto eos) { + Real tot_diff; + Kokkos::parallel_reduce( + "MyCheckPofRE", N_, + KOKKOS_LAMBDA(const int i, Real &diff) { + diff += std::abs(P[i] - eos.PressureFromDensityInternalEnergy(rho[i], sie[i])); + }, + tot_diff); + std::cout << "Total difference = " << tot_diff << std::endl; + }); + This is not functionality that would be available with the standard vector calls provided by ``singularity-eos``, at least not without chaining multiple parallel dispatch calls. Here we can do it in a From 08d3c25592c5b0cec1396d9d465a29e5f82f8b30 Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Tue, 10 Oct 2023 19:07:49 -0600 Subject: [PATCH 13/17] formatting --- test/test_eos_unit.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/test/test_eos_unit.cpp b/test/test_eos_unit.cpp index 96fd7dd984..f934b0fbed 100644 --- a/test/test_eos_unit.cpp +++ b/test/test_eos_unit.cpp @@ -516,17 +516,16 @@ class CheckPofRE { #pragma nv_exec_check_disable #pragma hd_warning_disable template - PORTABLE_FUNCTION - void operator()(const T &eos) { + PORTABLE_FUNCTION void operator()(const T &eos) { Real *P = P_; Real *rho = rho_; Real *sie = sie_; portableReduce( "MyCheckPofRE", 0, N_, PORTABLE_LAMBDA(const int i, int &nw) { - nw += !(isClose( - P[i], eos.PressureFromDensityInternalEnergy(rho[i], sie[i], nullptr), - 1e-15)); + nw += !(isClose(P[i], + eos.PressureFromDensityInternalEnergy(rho[i], sie[i], nullptr), + 1e-15)); }, nwrong); } From 2fe37236549002e94797cea0d848a5b08ebde13c Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Tue, 10 Oct 2023 19:23:23 -0600 Subject: [PATCH 14/17] slight tweak --- doc/sphinx/src/using-eos.rst | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/doc/sphinx/src/using-eos.rst b/doc/sphinx/src/using-eos.rst index 1568f6446f..e86aac815d 100644 --- a/doc/sphinx/src/using-eos.rst +++ b/doc/sphinx/src/using-eos.rst @@ -244,7 +244,8 @@ function with an `auto` argument as the input, e.g., .. code-block:: - eos.Evaluate([=](auto eos) { /* my code snippet */ }); + // equivalent to [=], but with device markings + eos.Evaluate(PORTABLE_LAMBDA(auto eos) { /* my code snippet */ }); .. warning:: From 621c3ae8d03a81ffbfb5325542fccf6cd0d44fb1 Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Tue, 10 Oct 2023 19:27:28 -0600 Subject: [PATCH 15/17] add blurm in parallelism model as suggested by jhp --- doc/sphinx/src/using-eos.rst | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/doc/sphinx/src/using-eos.rst b/doc/sphinx/src/using-eos.rst index e86aac815d..402f340e96 100644 --- a/doc/sphinx/src/using-eos.rst +++ b/doc/sphinx/src/using-eos.rst @@ -35,6 +35,10 @@ conditions. The type of parallelism used depends on how ``singularity-eos`` is compiled. If the ``Kokkos`` backend is used, any parallel dispatch supported by ``Kokkos`` is supported. +A more generic version of the vector calls exists in the ``Evaluate`` +method, which allows the user to specify arbitrary parallel dispatch +models by writing their own loops. See the relevant section below. + .. _variant section: Variants From bcfb597983735c7180ef17056c855f670384f2f4 Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Tue, 10 Oct 2023 19:35:42 -0600 Subject: [PATCH 16/17] Add cuda checks. Disable annoying warnings in the base class, rather than the test code --- singularity-eos/eos/eos_base.hpp | 13 +++++++++++++ test/test_eos_unit.cpp | 4 +--- 2 files changed, 14 insertions(+), 3 deletions(-) diff --git a/singularity-eos/eos/eos_base.hpp b/singularity-eos/eos/eos_base.hpp index a3f2a4e571..5e6abecec5 100644 --- a/singularity-eos/eos/eos_base.hpp +++ b/singularity-eos/eos/eos_base.hpp @@ -129,6 +129,19 @@ class EosBase { : std::is_same>, R *> {}; // Generic evaluator + // JMM: These macros disable checks that the functor called by evaluate + // is __host__ __device__ + // This is desirable, because you may want to call device-side batched + // evaluations (or inner loop dispatch with hierarchical parallelism) + // or you may want to call Evaluate from the host to, e.g., launch + // a reduction operation or somesuch. + // Note that this is a little dangerous as the compiler won't tell you + // if you messed up and ACTUALLY called a host-side function from + // device. The code will just segfault. +#if defined(__CUDACC__) +#pragma nv_exec_check_disable +#pragma hd_warning_disable +#endif // __CUDACC__ template PORTABLE_INLINE_FUNCTION void Evaluate(Functor_t &f) const { CRTP copy = *(static_cast(this)); diff --git a/test/test_eos_unit.cpp b/test/test_eos_unit.cpp index f934b0fbed..04f5ff009c 100644 --- a/test/test_eos_unit.cpp +++ b/test/test_eos_unit.cpp @@ -513,10 +513,8 @@ SCENARIO("Ideal gas entropy", "[IdealGas][Entropy]") { class CheckPofRE { public: CheckPofRE(Real *P, Real *rho, Real *sie, int N) : P_(P), rho_(rho), sie_(sie), N_(N) {} -#pragma nv_exec_check_disable -#pragma hd_warning_disable template - PORTABLE_FUNCTION void operator()(const T &eos) { + void operator()(const T &eos) { Real *P = P_; Real *rho = rho_; Real *sie = sie_; From da54c64ac692a39b9c302980257f5e058179b123 Mon Sep 17 00:00:00 2001 From: Jonah Maxwell Miller Date: Wed, 11 Oct 2023 18:39:37 -0600 Subject: [PATCH 17/17] switch Evaluate to be constexpr, as suggested by danl --- singularity-eos/eos/eos_base.hpp | 15 +-------------- singularity-eos/eos/eos_variant.hpp | 2 +- 2 files changed, 2 insertions(+), 15 deletions(-) diff --git a/singularity-eos/eos/eos_base.hpp b/singularity-eos/eos/eos_base.hpp index 5e6abecec5..9588a57fa9 100644 --- a/singularity-eos/eos/eos_base.hpp +++ b/singularity-eos/eos/eos_base.hpp @@ -129,21 +129,8 @@ class EosBase { : std::is_same>, R *> {}; // Generic evaluator - // JMM: These macros disable checks that the functor called by evaluate - // is __host__ __device__ - // This is desirable, because you may want to call device-side batched - // evaluations (or inner loop dispatch with hierarchical parallelism) - // or you may want to call Evaluate from the host to, e.g., launch - // a reduction operation or somesuch. - // Note that this is a little dangerous as the compiler won't tell you - // if you messed up and ACTUALLY called a host-side function from - // device. The code will just segfault. -#if defined(__CUDACC__) -#pragma nv_exec_check_disable -#pragma hd_warning_disable -#endif // __CUDACC__ template - PORTABLE_INLINE_FUNCTION void Evaluate(Functor_t &f) const { + constexpr void Evaluate(Functor_t &f) const { CRTP copy = *(static_cast(this)); f(copy); } diff --git a/singularity-eos/eos/eos_variant.hpp b/singularity-eos/eos/eos_variant.hpp index d2401636ef..3e72d87662 100644 --- a/singularity-eos/eos/eos_variant.hpp +++ b/singularity-eos/eos/eos_variant.hpp @@ -73,7 +73,7 @@ class Variant { // Place member functions here template - PORTABLE_INLINE_FUNCTION void Evaluate(Functor_t &f) const { + constexpr void Evaluate(Functor_t &f) const { return mpark::visit([&f](const auto &eos) { return eos.Evaluate(f); }, eos_); }