Mb/1d interpolator (#189)

* Current progress implementing 1d linear interpolator into neso (not working) * Current progress implementing 1D linear interpolator (not working) * Progress with creating 1d linear interpolator (not working) * Delete AtomicDataReader.hpp * Delete CSVAtomicDataReader.hpp * Delete One_Dimensional_Linear_Interpolator.hpp * Delete Interpolator.hpp * Added interpolators as did not push during previous commit (not working) * Remove unneeded header files * Interpolator update to use sycl_target (tested working) * Fixed bug in one_dimensional_linear_interpolator.hpp * Update to interpolator class * Update using reference to vectors instead of vector values (working) * Moved dy to constructor * Added in error handling (tested working) * Added in Class name into error handling * Fixed issues raised with pull request * Fixed bug * Checked number of x data points is equal to number of y data points * Moved division by x_data_sycl[k] - x_data_sycl[k - 1] to interpolator constructor * Fixed issues raised in pull request 189 * Update to address issues raised in pull request * Update to add comments on classes * Update to remove redundant code * Updated test so gradients change to test indexing in interpolator (tested working) * Created unequal spaces between x values in test * Moved interpolator implementation to src * Make x_data and y_data const references * Make x_input a const reference, and make use of std::numeric_limits in testing (tested working) --------- Co-authored-by: Matthew Barton <[email protected]> Co-authored-by: Matthew Barton <[email protected]>
ExCALIBUR-NEPTUNE · Jul 13, 2023 · 83d9c49 · 83d9c49
1 parent facb396
commit 83d9c49
Show file tree

Hide file tree

Showing 9 changed files with 264 additions and 6 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -94,8 +94,13 @@ endif()
 set(SRC_DIR ${CMAKE_CURRENT_SOURCE_DIR}/src)
 set(INC_DIR ${CMAKE_CURRENT_SOURCE_DIR}/include)
 set(LIB_SRC_FILES
-    ${SRC_DIR}/diagnostics.cpp ${SRC_DIR}/mesh.cpp ${SRC_DIR}/plasma.cpp
-    ${SRC_DIR}/run_info.cpp ${SRC_DIR}/simulation.cpp ${SRC_DIR}/species.cpp)
+    ${SRC_DIR}/diagnostics.cpp
+    ${SRC_DIR}/mesh.cpp
+    ${SRC_DIR}/plasma.cpp
+    ${SRC_DIR}/run_info.cpp
+    ${SRC_DIR}/simulation.cpp
+    ${SRC_DIR}/species.cpp
+    ${SRC_DIR}/linear_interpolator_1D.cpp)
 set(MAIN_SRC ${CMAKE_CURRENT_SOURCE_DIR}/src/main.cpp)
 set(HEADER_FILES
     ${INC_DIR}/custom_types.hpp
@@ -109,7 +114,11 @@ set(HEADER_FILES
     ${INC_DIR}/simulation.hpp
     ${INC_DIR}/species.hpp
     ${INC_DIR}/velocity.hpp
-    ${INC_DIR}/nektar_interface/geometry_transport_2d.hpp)
+    ${INC_DIR}/nektar_interface/geometry_transport_2d.hpp
+    ${INC_DIR}/atomic_data_reader.hpp
+    ${INC_DIR}/csv_atomic_data_reader.hpp
+    ${INC_DIR}/interpolator.hpp
+    ${INC_DIR}/linear_interpolator_1D.hpp)
 
 # Create library
 set(LIBRARY_NAME nesolib)

diff --git a/include/AtomicDataReader.hpp → include/atomic_data_reader.hpp b/include/AtomicDataReader.hpp → include/atomic_data_reader.hpp
@@ -30,7 +30,7 @@ class AtomicDataReader {
   }
 
   std::vector<double> get_temps() {
-    NP::NESOASSERT(m_rate_idx >= 0 && m_rate_idx < m_data.size(),
+    NP::NESOASSERT(m_T_idx >= 0 && m_T_idx < m_data.size(),
                    "AtomicDataReader: data temperature index isn't defined.");
     return m_data[m_T_idx];
   }

diff --git a/include/CSVAtomicDataReader.hpp → include/csv_atomic_data_reader.hpp b/include/CSVAtomicDataReader.hpp → include/csv_atomic_data_reader.hpp
@@ -1,7 +1,7 @@
 #ifndef __CSV_ATOMIC_DATA_READER_H__
 #define __CSV_ATOMIC_DATA_READER_H__
 
-#include "AtomicDataReader.hpp"
+#include "atomic_data_reader.hpp"
 
 #include <fstream>
 #include <iostream>

diff --git a/include/interpolator.hpp b/include/interpolator.hpp
@@ -0,0 +1,57 @@
+#ifndef __INTERPOLATOR_H__
+#define __INTERPOLATOR_H__
+#include <vector>
+
+#include <neso_particles.hpp>
+
+namespace NP = NESO::Particles;
+namespace NESO {
+/**
+ * This class defines the interface through which derived classes gain access to
+ * x,y data which will be interpolated
+ *
+ */
+class Interpolator {
+public:
+  /**
+   * This constructor initialises the x,y data derived classes will use
+   * checking that that the size of both vectors is the same.
+   * It also initialises the target that the sycl kernel will use
+   * in the derived classes.
+   *
+   * @param[in] x_data The x data values which the interpolator will have access
+   * to.
+   * @param[in] y_data The y data values which the interpolator will have access
+   * to.
+   * @param[in] sycl_target The target that the sycl kernels will make use of.
+   */
+  Interpolator(const std::vector<double> &x_data,
+               const std::vector<double> &y_data,
+               NP::SYCLTargetSharedPtr sycl_target)
+      : m_sycl_target(sycl_target), m_x_data(x_data), m_y_data(y_data) {
+
+    NP::NESOASSERT(m_x_data.size() == m_y_data.size(),
+                   "size of m_x_data vector doesn't equal m_y_data vector");
+  };
+  Interpolator() = delete;
+
+  /**
+   * This function returns a value of y (y_output) given a value of x (x_input)
+   * given the (x,y) data provided by the interpolator class
+   *
+   * @param[in] x_input x_input is reference to a vector of x values for which
+   * you would like the y value returned.
+   * @param[out] y_output y_output is reference to a vector of y values which
+   * the interpolator calculated, based on x_input.
+   */
+  virtual void interpolate(const std::vector<double> &x_input,
+                           std::vector<double> &y_output) = 0;
+
+protected:
+  NP::SYCLTargetSharedPtr m_sycl_target;
+  const std::vector<double> &m_x_data;
+  const std::vector<double> &m_y_data;
+  std::vector<double> dydx;
+};
+} // namespace NESO
+#endif
diff --git a/include/linear_interpolator_1D.hpp b/include/linear_interpolator_1D.hpp
@@ -0,0 +1,56 @@
+#ifndef __LINEAR_INTERPOLATOR_1D_H__
+#define __LINEAR_INTERPOLATOR_1D_H__
+
+#include "interpolator.hpp"
+#include <CL/sycl.hpp>
+#include <mpi.h>
+#include <neso_particles.hpp>
+#include <vector>
+
+using namespace NESO::Particles;
+using namespace cl;
+namespace NESO {
+
+/**
+ * Class used to output a vector of y values, given some x values, based
+ * on provided x values for input
+ */
+
+class LinearInterpolator1D : public Interpolator {
+public:
+  /**
+   * This constructor provides the class with the x_input values the
+   * interpolator will need. It creates a gradient vector based on the x,y data
+   * provided by the Interpolator class from which this class is derived. It
+   * then creates a sycl buffer for the x,y data and this gradient vector.
+   *
+   * @param[in] x_input x_input is reference to a vector of x values for which
+   * you would like the y value returned.
+   * @param[out] y_output y_output is reference to a vector of y values which
+   * the interpolator calculated based on x_input.
+   * @param[in] sycl_target The target that the sycl kernels will make use of.
+   */
+  LinearInterpolator1D(const std::vector<double> &x_data,
+                       const std::vector<double> &y_data,
+                       SYCLTargetSharedPtr sycl_target);
+
+  /**
+   * This function returns a value of y (y_output) given a value of x (x_input)
+   * given the (x,y) data provided by the interpolator class
+   *
+   * @param[in] x_input x_input is reference to a vector of x values for which
+   * you would like the y value returned.
+   * @param[out] y_output y_output is reference to a vector of y values which
+   * the interpolator calculated, based on x_input.
+   */
+  virtual void interpolate(const std::vector<double> &x_input,
+                           std::vector<double> &y_output);
+
+protected:
+  sycl::buffer<double, 1> buffer_x_data;
+  sycl::buffer<double, 1> buffer_y_data;
+  sycl::buffer<double, 1> buffer_dydx;
+};
+
+} // namespace NESO
+#endif
diff --git a/src/linear_interpolator_1D.cpp b/src/linear_interpolator_1D.cpp
@@ -0,0 +1,108 @@
+#include "linear_interpolator_1D.hpp"
+#include <CL/sycl.hpp>
+#include <mpi.h>
+#include <neso_particles.hpp>
+#include <vector>
+
+using namespace NESO::Particles;
+using namespace cl;
+namespace NESO {
+
+/**
+ * This constructor provides the class with the x_input values the
+ * interpolator will need. It creates a gradient vector based on the x,y data
+ * provided by the Interpolator class from which this class is derived. It
+ * then creates a sycl buffer for the x,y data and this gradient vector.
+ *
+ * @param[in] x_input x_input is reference to a vector of x values for which
+ * you would like the y value returned.
+ * @param[out] y_output y_output is reference to a vector of y values which
+ * the interpolator calculated based on x_input.
+ * @param[in] sycl_target The target that the sycl kernels will make use of.
+ */
+LinearInterpolator1D::LinearInterpolator1D(const std::vector<double> &x_data,
+                                           const std::vector<double> &y_data,
+                                           SYCLTargetSharedPtr sycl_target)
+    : Interpolator(x_data, y_data, sycl_target), buffer_x_data(0),
+      buffer_y_data(0), buffer_dydx(0) {
+
+  dydx.reserve(m_y_data.size());
+  for (int i = 1; i < m_y_data.size(); i++) {
+    dydx.push_back((m_y_data[i] - m_y_data[i - 1]) /
+                   ((m_x_data[i] - m_x_data[i - 1])));
+  }
+
+  // buffers for (x,y) data
+  buffer_x_data =
+      sycl::buffer<double, 1>(m_x_data.data(), sycl::range<1>{m_x_data.size()});
+  buffer_y_data =
+      sycl::buffer<double, 1>(m_y_data.data(), sycl::range<1>{m_y_data.size()});
+  buffer_dydx =
+      sycl::buffer<double, 1>(dydx.data(), sycl::range<1>{dydx.size()});
+};
+
+/**
+ * This function returns a value of y (y_output) given a value of x (x_input)
+ * given the (x,y) data provided by the interpolator class
+ *
+ * @param[in] x_input x_input is reference to a vector of x values for which
+ * you would like the y value returned.
+ * @param[out] y_output y_output is reference to a vector of y values which
+ * the interpolator calculated, based on x_input.
+ */
+void LinearInterpolator1D::interpolate(const std::vector<double> &x_input,
+                                       std::vector<double> &y_output) {
+
+  // Used in error handling of interpolator
+  ErrorPropagate ep(m_sycl_target);
+  // sycl code
+  auto k_ep = ep.device_ptr();
+  NESOASSERT(x_input.size() == y_output.size(),
+             "size of x_input vector doesn't equal y_output vector");
+  sycl::buffer<double, 1> buffer_x_input(x_input.data(),
+                                         sycl::range<1>{x_input.size()});
+  sycl::buffer<double, 1> buffer_y_output(y_output.data(),
+                                          sycl::range<1>{y_output.size()});
+  const int m_x_data_size = m_x_data.size();
+  const int x_input_size = x_input.size();
+  auto event_interpolate =
+      this->m_sycl_target->queue.submit([&](sycl::handler &cgh) {
+        auto x_data_sycl =
+            buffer_x_data.get_access<sycl::access::mode::read>(cgh);
+        auto y_data_sycl =
+            buffer_y_data.get_access<sycl::access::mode::read>(cgh);
+        auto x_input_sycl =
+            buffer_x_input.get_access<sycl::access::mode::read>(cgh);
+        auto y_output_sycl =
+            buffer_y_output.get_access<sycl::access::mode::read_write>(cgh);
+        auto dydx_sycl = buffer_dydx.get_access<sycl::access::mode::read>(cgh);
+        cgh.parallel_for<>(sycl::range<1>(x_input_size), [=](sycl::id<1> idx) {
+          int k;
+          for (int i = 0; i < m_x_data_size; i++) {
+            // error handling
+            const double x = x_input_sycl[int(idx)];
+            const double x_data_sycl_start = x_data_sycl[0];
+            const double x_data_sycl_end = x_data_sycl[m_x_data_size - 1];
+            if (x < x_data_sycl_start or x > x_data_sycl_end) {
+              // throw an error
+              NESO_KERNEL_ASSERT(false, k_ep);
+              break;
+            }
+            // set value of index to first value for which x value
+            // of input is greater than x data values
+            if (x - x_data_sycl[i] < 0.0) {
+              k = i;
+              break;
+            }
+          }
+          y_output_sycl[int(idx)] =
+              y_data_sycl[k - 1] +
+              dydx_sycl[k - 1] * (x_input_sycl[int(idx)] - x_data_sycl[k - 1]);
+        });
+      });
+  event_interpolate.wait_and_throw();
+  ep.check_and_throw("OneDimensionalLinearInterpolator: Input values are "
+                     "outside the range provided by data");
+}
+
+} // namespace NESO
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
@@ -36,6 +36,7 @@ set(UNIT_SRC_FILES
     ${UNIT_SRC}/test_plasma.cpp
     ${UNIT_SRC}/test_species.cpp
     ${UNIT_SRC}/test_atomic_data_readers.cpp
+    ${UNIT_SRC}/test_1D_linear_interpolator.cpp
     ${UNIT_SRC}/particle_utility/test_position_distribution.cpp
     ${UNIT_SRC}/particle_utility/test_particle_initialisation_line.cpp
     ${UNIT_SRC}/nektar_interface/test_particle_function_evaluation.cpp

diff --git a/test/unit/test_1D_linear_interpolator.cpp b/test/unit/test_1D_linear_interpolator.cpp
@@ -0,0 +1,27 @@
+#include "linear_interpolator_1D.hpp"
+#include <CL/sycl.hpp>
+#include <gtest/gtest.h>
+#include <mpi.h>
+#include <neso_particles.hpp>
+#include <vector>
+
+using namespace NESO;
+using namespace NESO::Particles;
+
+TEST(InterpolatorTest, 1DLinear) {
+  auto sycl_target = std::make_shared<SYCLTarget>(0, MPI_COMM_WORLD);
+  std::vector<double> test_input = {1.5, 2.5};
+  const std::vector<double> &test_input_ref = test_input;
+  std::vector<double> test_output(test_input.size());
+  std::vector<double> &test_output_ref = test_output;
+
+  std::vector<double> test_x_data = {0, 1, 2, 4};
+  const std::vector<double> &test_x_data_ref = test_x_data;
+  std::vector<double> test_y_data = {0, 1, 3, 4};
+  const std::vector<double> &test_y_data_ref = test_y_data;
+
+  LinearInterpolator1D(test_x_data_ref, test_y_data_ref, sycl_target)
+      .interpolate(test_input_ref, test_output_ref);
+  ASSERT_NEAR(test_output[0], 2.0, std::numeric_limits<double>::epsilon());
+  ASSERT_NEAR(test_output[1], 3.25, std::numeric_limits<double>::epsilon());
+}
diff --git a/test/unit/test_atomic_data_readers.cpp b/test/unit/test_atomic_data_readers.cpp
@@ -1,4 +1,4 @@
-#include "CSVAtomicDataReader.hpp"
+#include "csv_atomic_data_reader.hpp"
 #include <filesystem>
 #include <gtest/gtest.h>
 #include <string>