Command line options to use device memory for wave equation and mhp-t…

…est (#595)

benchmarks/gbench/mhp/CMakeLists.txt

@@ -37,7 +37,7 @@ target_link_libraries(mhp-bench benchmark::benchmark cxxopts DR::mpi)
 
 # cmake-format: off
 add_executable(mhp-quick-bench mhp-bench.cpp
-  ../common/dot_product.cpp
+  stencil_2d.cpp
   )
 # cmake-format: on
 target_compile_definitions(mhp-quick-bench PRIVATE BENCH_MHP)

benchmarks/gbench/mhp/wave_equation.cpp

@@ -625,6 +625,7 @@ int main(int argc, char *argv[]) {
     ("sycl", "Execute on SYCL device")
     ("l,log", "enable logging")
     ("f,fused-kernel", "Use fused kernels.", cxxopts::value<bool>()->default_value("false"))
+    ("device-memory", "Use device memory")
     ("h,help", "Print help");
   // clang-format on
 
@@ -647,7 +648,8 @@ int main(int argc, char *argv[]) {
     sycl::queue q = dr::mhp::select_queue();
     std::cout << "Run on: "
               << q.get_device().get_info<sycl::info::device::name>() << "\n";
-    dr::mhp::init(q);
+    dr::mhp::init(q, options.count("device-memory") ? sycl::usm::alloc::device
+                                                    : sycl::usm::alloc::shared);
 #else
     std::cout << "Sycl support requires icpx\n";
     exit(1);

include/dr/mhp/algorithms/reduce.hpp

@@ -38,7 +38,7 @@ inline auto dpl_reduce(rng::forward_range auto &&r, auto &&binary_op) {
       return std::reduce(dpl_policy(),
                          dr::__detail::direct_iterator(rng::begin(r) + 1),
                          dr::__detail::direct_iterator(rng::end(r)),
-                         *rng::begin(r), binary_op);
+                         sycl_get(*rng::begin(r)), binary_op);
     }
   }
 #else

test/gtest/common/reduce.cpp

@@ -14,6 +14,14 @@ TYPED_TEST_SUITE(Reduce, AllTypes);
 TYPED_TEST(Reduce, Range) {
   Ops1<TypeParam> ops(10);
 
+  auto max = [](double x, double y) { return std::max(x, y); };
+  EXPECT_EQ(std::reduce(ops.vec.begin(), ops.vec.end(), 3, max),
+            xhp::reduce(ops.dist_vec, 3, max));
+}
+
+TYPED_TEST(Reduce, Max) {
+  Ops1<TypeParam> ops(10);
+
   EXPECT_EQ(std::reduce(ops.vec.begin(), ops.vec.end(), 3, std::plus{}),
             xhp::reduce(ops.dist_vec, 3, std::plus{}));
 }

test/gtest/mhp/CMakeLists.txt

@@ -56,7 +56,7 @@ add_executable(
 
 add_executable(mhp-quick-test
   mhp-tests.cpp
-  ../common/zip.cpp
+  ../common/reduce.cpp
   )
 # cmake-format: on
 

test/gtest/mhp/halo.cpp

@@ -11,6 +11,9 @@ TYPED_TEST_SUITE(Halo, AllTypes);
 template <typename DV>
 void local_is_accessible_in_halo_region(const int halo_prev,
                                         const int halo_next) {
+  if (options.count("device-memory")) {
+    return;
+  }
   DV dv(6, dr::mhp::distribution().halo(halo_prev, halo_next));
   iota(dv, 0);
   dv.halo().exchange();

test/gtest/mhp/mdstar.cpp

@@ -138,6 +138,11 @@ TEST_F(Mdspan, GridExtents) {
 }
 
 TEST_F(Mdspan, GridLocalReference) {
+  // mdspan is not accessible for device memory
+  if (options.count("device-memory")) {
+    return;
+  }
+
   xhp::distributed_vector<T> dist(n2d, dist2d_1d);
   xhp::iota(dist, 100);
   auto dmdspan = xhp::views::mdspan(dist, extents2d);
@@ -211,6 +216,11 @@ TEST_F(Mdarray, GridExtents) {
 }
 
 TEST_F(Mdarray, GridLocalReference) {
+  // mdspan is not accessible for device memory
+  if (options.count("device-memory")) {
+    return;
+  }
+
   xhp::distributed_mdarray<T, 2> mdarray(extents2d);
   xhp::iota(mdarray, 100);
   auto grid = mdarray.grid();
@@ -225,6 +235,11 @@ TEST_F(Mdarray, GridLocalReference) {
 }
 
 TEST_F(Mdarray, Halo) {
+  // mdspan is not accessible for device memory
+  if (options.count("device-memory")) {
+    return;
+  }
+
   xhp::distributed_mdarray<T, 2> mdarray(extents2d,
                                          xhp::distribution().halo(1));
   dr::mhp::halo(mdarray);
@@ -294,6 +309,11 @@ TEST_F(Submdspan, GridExtents) {
 }
 
 TEST_F(Submdspan, GridLocalReference) {
+  // mdspan is not accessible for device memory
+  if (options.count("device-memory")) {
+    return;
+  }
+
   xhp::distributed_mdarray<T, 2> mdarray(extents2d);
   xhp::iota(mdarray, 100);
   auto sub = xhp::views::submdspan(mdarray.view(), slice_starts, slice_ends);
@@ -315,6 +335,11 @@ TEST_F(Submdspan, GridLocalReference) {
 }
 
 TEST_F(Submdspan, Segments) {
+  // mdspan is not accessible for device memory
+  if (options.count("device-memory")) {
+    return;
+  }
+
   xhp::distributed_mdarray<T, 2> mdarray(extents2d);
   xhp::iota(mdarray, 100);
   auto sub = xhp::views::submdspan(mdarray.view(), slice_starts, slice_ends);

test/gtest/mhp/mhp-tests.cpp

@@ -18,7 +18,8 @@ void dr_init() {
       fmt::print("Enable sycl device: {}\n",
                  q.get_device().get_info<sycl::info::device::name>());
     }
-    dr::mhp::init(q);
+    dr::mhp::init(q, options.count("device-memory") ? sycl::usm::alloc::device
+                                                    : sycl::usm::alloc::shared);
     return;
   }
 #endif
@@ -45,6 +46,7 @@ int main(int argc, char *argv[]) {
   options_spec.add_options()
     ("drhelp", "Print help")
     ("log", "Enable logging")
+    ("device-memory", "Use device memory")
     ("sycl", "Execute on SYCL device");
   // clang-format on
 

test/gtest/mhp/segments.cpp

@@ -34,6 +34,9 @@ template <typename T> class SegmentUtils : public testing::Test {
 TYPED_TEST_SUITE(SegmentUtils, AllTypes);
 
 TYPED_TEST(SegmentUtils, LocalSegment) {
+  if (options.count("device-memory")) {
+    return;
+  }
   Ops1<TypeParam> ops(10);
   auto segments = dr::mhp::local_segments(ops.dist_vec);
   EXPECT_EQ(dr::mhp::local_segment(ops.dist_vec), *rng::begin(segments));

test/gtest/mhp/xhp-tests.hpp

@@ -14,6 +14,7 @@
 extern MPI_Comm comm;
 extern std::size_t comm_rank;
 extern std::size_t comm_size;
+extern cxxopts::ParseResult options;
 
 namespace xhp = dr::mhp;