gt: link gtl on cray amd systems, more iters

Option to disable managed tests, TEST_MANAGED=OFF cmake var

CMakeLists.txt

@@ -3,11 +3,13 @@ cmake_minimum_required(VERSION 3.18 FATAL_ERROR)
 # create project
 project(mpi-daxpy-test)
 
+option(TEST_MANAGED "Test managed memory" ON)
+
 # add dependencies
 include(cmake/CPM.cmake)
 CPMFindPackage(NAME gtensor
-               GITHUB_REPOSITORY bd4/gtensor
+               GITHUB_REPOSITORY wdmapp/gtensor
-               GIT_TAG "pr/view-contiguous"
+               GIT_TAG "main"
                OPTIONS "GTENSOR_ENABLE_BLAS ON")
 
 set(MPI_CXX_SKIP_MPICXX ON)
@@ -32,7 +34,17 @@ add_executable(mpi_stencil2d_gt)
 target_sources(mpi_stencil2d_gt PRIVATE mpi_stencil2d_gt.cc)
 target_link_libraries(mpi_stencil2d_gt PRIVATE gtensor::gtensor)
 target_link_libraries(mpi_stencil2d_gt PRIVATE MPI::MPI_CXX)
-target_compile_features(mpi_stencil2d_gt PRIVATE cxx_std_17)
+#target_compile_features(mpi_stencil2d_gt PRIVATE cxx_std_17)
+if (TEST_MANAGED)
+  message(STATUS "${PROJECT_NAME}: Enabling managed memory")
+  target_compile_definitions(mpi_stencil2d_gt PRIVATE TEST_MANAGED)
+endif()
+if (GTENSOR_DEVICE STREQUAL "hip" AND DEFINED ENV{PE_MPICH_GTL_DIR_amd_gfx90a})
+  message(STATUS "${PROJECT_NAME}: Linking gtl libs for HIP backend")
+  target_link_options(mpi_stencil2d_gt PRIVATE
+                      $ENV{PE_MPICH_GTL_DIR_amd_gfx90a}
+                      $ENV{PE_MPICH_GTL_LIBS_amd_gfx90a})
+endif()
 
 if ("${GTENSOR_DEVICE}" STREQUAL "cuda") 
   set_source_files_properties(mpi_daxpy_gt.cc

mpi_stencil2d_gt.cc

@@ -27,6 +27,8 @@
 
 using namespace gt::placeholders;
 
+const double PI = 3.141592653598793;
+
 inline void check(const char* fname, int line, int mpi_rval)
 {
   if (mpi_rval != MPI_SUCCESS) {
@@ -81,7 +83,7 @@ static const gt::gtensor<double, 1> stencil5 = {1.0 / 12.0, -2.0 / 3.0, 0.0,
  */
 template <typename S>
 inline auto stencil2d_1d_5_d0(const gt::ext::gtensor2<double, 2, S>& z,
-                           const gt::gtensor<double, 1>& stencil)
+                              const gt::gtensor<double, 1>& stencil)
 {
   return stencil(0) * z.view(_s(0, -4), _all) +
          stencil(1) * z.view(_s(1, -3), _all) +
@@ -98,7 +100,7 @@ inline auto stencil2d_1d_5_d0(const gt::ext::gtensor2<double, 2, S>& z,
  */
 template <typename S>
 inline auto stencil2d_1d_5_d1(const gt::ext::gtensor2<double, 2, S>& z,
-                           const gt::gtensor<double, 1>& stencil)
+                              const gt::gtensor<double, 1>& stencil)
 {
   return stencil(0) * z.view(_all, _s(0, -4)) +
          stencil(1) * z.view(_all, _s(1, -3)) +
@@ -273,21 +275,20 @@ void boundary_exchange_y(MPI_Comm comm, int world_size, int rank,
   MPI_Request req_l[2];
   MPI_Request req_r[2];
 
-  int rank_l = (rank - 1) % world_size;
+  int rank_l = rank - 1;
-  int rank_r = (rank + 1) % world_size;
+  int rank_r = rank + 1;
 
   auto sv_l = gt::view_strided(d_z, _all, _s(n_bnd, 2 * n_bnd));
   auto sv_r = gt::view_strided(d_z, _all, _s(-2 * n_bnd, -n_bnd));
   auto rv_l = gt::view_strided(d_z, _all, _s(0, n_bnd));
   auto rv_r = gt::view_strided(d_z, _all, _s(-n_bnd, _));
+
   // start async copy of ghost points into send buffers
-  if (rank_l >= 0) {
+  if (stage_device) {
-    if (stage_device) {
+    if (rank_l >= 0) {
       sbuf_l = sv_l;
     }
-  }
+    if (rank_r <= world_size) {
-  if (rank_r <= world_size) {
-    if (stage_device) {
       sbuf_r = sv_r;
     }
   }
@@ -300,7 +301,7 @@ void boundary_exchange_y(MPI_Comm comm, int world_size, int rank,
     } else {
       rbuf_l_data = rv_l.data().get();
     }
-    CHECK(MPI_Irecv(rbuf_l_data, rbuf_l.size(), MPI_DOUBLE, rank_l, 123, comm,
+    CHECK(MPI_Irecv(rbuf_l_data, rv_l.size(), MPI_DOUBLE, rank_l, 123, comm,
                     &req_l[0]));
   }
 
@@ -311,12 +312,14 @@ void boundary_exchange_y(MPI_Comm comm, int world_size, int rank,
     } else {
       rbuf_r_data = rv_r.data().get();
     }
-    CHECK(MPI_Irecv(rbuf_r_data, rbuf_r.size(), MPI_DOUBLE, rank_r, 456, comm,
+    CHECK(MPI_Irecv(rbuf_r_data, rv_r.size(), MPI_DOUBLE, rank_r, 456, comm,
                     &req_r[0]));
   }
 
   // wait for send buffer fill
+  // if (stage_device) {
   gt::synchronize();
+  // }
 
   // initiate async sends
   if (rank_l >= 0) {
@@ -326,7 +329,7 @@ void boundary_exchange_y(MPI_Comm comm, int world_size, int rank,
     } else {
       sbuf_l_data = sv_l.data().get();
     }
-    CHECK(MPI_Isend(sbuf_l_data, sbuf_l.size(), MPI_DOUBLE, rank_l, 456, comm,
+    CHECK(MPI_Isend(sbuf_l_data, sv_l.size(), MPI_DOUBLE, rank_l, 456, comm,
                     &req_l[1]));
   }
 
@@ -337,7 +340,7 @@ void boundary_exchange_y(MPI_Comm comm, int world_size, int rank,
     } else {
       sbuf_r_data = sv_r.data().get();
     }
-    CHECK(MPI_Isend(sbuf_r_data, sbuf_r.size(), MPI_DOUBLE, rank_r, 123, comm,
+    CHECK(MPI_Isend(sbuf_r_data, sv_r.size(), MPI_DOUBLE, rank_r, 123, comm,
                     &req_r[1]));
   }
 
@@ -373,15 +376,16 @@ template <int Dim, typename S>
 void print_test_name(bool use_buffers)
 {
   if constexpr (std::is_same<S, gt::space::device>::value) {
-    printf("=== TEST dim:%d, device , buf:%d\n", Dim, use_buffers);
+    printf("TEST dim:%d, device , buf:%d", Dim, use_buffers);
   } else {
-    printf("=== TEST dim:%d, managed, buf:%d\n", Dim, use_buffers);
+    printf("TEST dim:%d, managed, buf:%d", Dim, use_buffers);
   }
 }
 
 template <typename S, int Dim>
-void test(int device_id, uint32_t vendor_id, int world_size, int world_rank,
+void test_deriv(int device_id, uint32_t vendor_id, int world_size,
-          int n_global, int n_iter, bool use_buffers, int n_warmup=5)
+                int world_rank, int n_global, int n_iter, bool use_buffers,
+                int n_warmup = 5)
 {
   // Note: domain will be n_global x n_global plus ghost points in one dimension
 
@@ -410,10 +414,6 @@ void test(int device_id, uint32_t vendor_id, int world_size, int world_rank,
     ny_local_ghost = n_local + 2 * n_bnd;
   }
 
-  if (world_rank == 0) {
-    print_test_name<Dim, S>(use_buffers);
-  }
-
   gt::shape_type<2> z_shape(nx_local_ghost, ny_local_ghost);
   gt::shape_type<2> dz_shape(nx_local, ny_local);
 
@@ -436,7 +436,7 @@ void test(int device_id, uint32_t vendor_id, int world_size, int world_rank,
   double iter_time = 0.0;
   double total_time = 0.0;
 
-  double x_start=0, y_start=0;
+  double x_start = 0, y_start = 0;
   if constexpr (Dim == 0) {
     x_start = world_rank * ln_local;
   } else {
@@ -533,8 +533,10 @@ void test(int device_id, uint32_t vendor_id, int world_size, int world_rank,
     }
     gt::synchronize();
   }
+#ifdef DEBUG
   printf("%d/%d exchange time %0.8f ms\n", world_rank, world_size,
          total_time / n_iter * 1000);
+#endif
 
   gt::copy(d_dz_numeric, h_dz_numeric);
 
@@ -552,9 +554,97 @@ void test(int device_id, uint32_t vendor_id, int world_size, int world_rank,
 
   double err_norm = std::sqrt(gt::sum_squares(h_dz_numeric - h_dz_actual));
 
+#ifdef DEBUG
   printf("%d/%d [%d:0x%08x] err_norm = %.8f\n", world_rank, world_size,
          device_id, vendor_id, err_norm);
+#endif
 
+  double time_sum;
+  MPI_Reduce(&total_time, &time_sum, 1, MPI_DOUBLE, MPI_SUM, 0, MPI_COMM_WORLD);
+
+  double err_sum;
+  MPI_Reduce(&err_norm, &err_sum, 1, MPI_DOUBLE, MPI_SUM, 0, MPI_COMM_WORLD);
+
+  if (world_rank == 0) {
+    print_test_name<Dim, S>(use_buffers);
+    printf("; %0.8f, err=%0.8f\n", time_sum, err_sum);
+  }
+}
+
+template <typename S, int Dim>
+void test_sum(int device_id, uint32_t vendor_id, int world_size, int world_rank,
+              int n_global, int n_iter, int n_warmup = 5)
+{
+  // Note: domain will be n_global x n_global plus ghost points in one dimension
+  const int n_local = n_global / world_size;
+
+  int nx_local, ny_local;
+
+  struct timespec start, end;
+  double iter_time = 0.0;
+  double total_time = 0.0;
+
+  if constexpr (Dim == 0) {
+    nx_local = n_local;
+    ny_local = n_global;
+  } else {
+    nx_local = n_global;
+    ny_local = n_local;
+  }
+
+  gt::shape_type<2> z_shape(nx_local, ny_local);
+
+  gt::ext::gtensor2<double, 2, S> d_z(z_shape, PI / world_size);
+
+  // reduction test
+  gt::shape_type<1> sum_shape;
+  if constexpr (Dim == 0) {
+    sum_shape = gt::shape(d_z.shape(0));
+  } else {
+    sum_shape = gt::shape(d_z.shape(1));
+  }
+  gt::ext::gtensor2<double, 1, S> d_sum(sum_shape);
+  gt::gtensor<double, 1> h_sum(sum_shape);
+  for (int i = 0; i < n_warmup + n_iter; i++) {
+    if constexpr (Dim == 0) {
+      gt::sum_axis_to(d_sum, d_z, 0);
+      gt::synchronize();
+
+      clock_gettime(CLOCK_MONOTONIC, &start);
+      CHECK(MPI_Allreduce(MPI_IN_PLACE, d_sum.data().get(), d_sum.size(),
+                          MPI_DOUBLE, MPI_SUM, MPI_COMM_WORLD));
+
+      clock_gettime(CLOCK_MONOTONIC, &end);
+    } else {
+      gt::sum_axis_to(d_sum, d_z, 1);
+      gt::synchronize();
+
+      clock_gettime(CLOCK_MONOTONIC, &start);
+      CHECK(MPI_Allreduce(MPI_IN_PLACE, d_sum.data().get(), d_sum.size(),
+                          MPI_DOUBLE, MPI_SUM, MPI_COMM_WORLD));
+      clock_gettime(CLOCK_MONOTONIC, &end);
+    }
+    iter_time =
+      ((end.tv_sec - start.tv_sec) + (end.tv_nsec - start.tv_nsec) * 1.0e-9);
+
+    if (i >= n_warmup) {
+      total_time += iter_time;
+    }
+  }
+#ifdef DEBUG
+  printf("%d/%d allreduce time %0.8f ms\n", world_rank, world_size,
+         total_time / n_iter * 1000);
+#endif
+
+  gt::copy(d_sum, h_sum);
+
+  double time_sum;
+  MPI_Reduce(&total_time, &time_sum, 1, MPI_DOUBLE, MPI_SUM, 0, MPI_COMM_WORLD);
+
+  if (world_rank == 0) {
+    print_test_name<Dim, S>(false);
+    printf("; allreduce=%0.8f\n", time_sum);
+  }
 }
 
 int main(int argc, char** argv)
@@ -563,8 +653,8 @@ int main(int argc, char** argv)
 
   // Note: domain will be n_global x n_global plus ghost points in one dimension
   int n_global = 8 * 1024;
-  int n_iter = 100;
+  int n_iter = 1000;
-  int n_warmup = 5;
+  int n_warmup = 10;
 
   if (argc > 1) {
     n_global = std::atoi(argv[1]) * 1024;
@@ -603,21 +693,40 @@ int main(int argc, char** argv)
 
   fflush(stdout);
 
-/*
+  test_deriv<gt::space::device, 0>(device_id, vendor_id, world_size, world_rank,
-  test<gt::space::device, 0>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, true, 5);
+                                   n_global, n_iter, true, 5);
-  test<gt::space::device, 0>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, false, 5);
+  test_deriv<gt::space::device, 0>(device_id, vendor_id, world_size, world_rank,
-  test<gt::space::managed, 0>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, true, 5);
+                                   n_global, n_iter, false, 5);
-  test<gt::space::managed, 0>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, false, 5);
+#ifdef TEST_MANAGED
+  test_deriv<gt::space::managed, 0>(device_id, vendor_id, world_size,
+                                    world_rank, n_global, n_iter, true, 5);
+  test_deriv<gt::space::managed, 0>(device_id, vendor_id, world_size,
+                                    world_rank, n_global, n_iter, false, 5);
+#endif
 
-  test<gt::space::device, 1>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, true, 5);
+  test_deriv<gt::space::device, 1>(device_id, vendor_id, world_size, world_rank,
-  test<gt::space::device, 1>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, false, 5);
+                                   n_global, n_iter, true, 5);
-  test<gt::space::managed, 1>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, true, 5);
+  test_deriv<gt::space::device, 1>(device_id, vendor_id, world_size, world_rank,
-  test<gt::space::managed, 1>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, false, 5);
+                                   n_global, n_iter, false, 5);
-  */
+#ifdef TEST_MANAGED
-  test<gt::space::managed, 0>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, true, 5);
+  test_deriv<gt::space::managed, 1>(device_id, vendor_id, world_size,
-  test<gt::space::managed, 0>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, false, 5);
+                                    world_rank, n_global, n_iter, true, 5);
-  test<gt::space::device, 0>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, true, 5);
+  test_deriv<gt::space::managed, 1>(device_id, vendor_id, world_size,
-  test<gt::space::device, 0>(device_id, vendor_id, world_size, world_rank, n_global, n_iter, false, 5);
+                                    world_rank, n_global, n_iter, false, 5);
+#endif
+
+  test_sum<gt::space::device, 0>(device_id, vendor_id, world_size, world_rank,
+                                 n_global, n_iter, 5);
+#ifdef TEST_MANAGED
+  test_sum<gt::space::managed, 0>(device_id, vendor_id, world_size, world_rank,
+                                  n_global, n_iter, 5);
+#endif
+  test_sum<gt::space::device, 1>(device_id, vendor_id, world_size, world_rank,
+                                 n_global, n_iter, 5);
+#ifdef TEST_MANAGED
+  test_sum<gt::space::managed, 1>(device_id, vendor_id, world_size, world_rank,
+                                  n_global, n_iter, 5);
+#endif
 
   MPI_Finalize();