gpu-direct: fix nvcc warning about host function used in device one

Author: svalat

src/mpi/buffer.hpp

@@ -31,19 +31,19 @@ class Buffer {
   }
 
   void* deviceData() {
-    return(array.data());
+    return(array.deviceView().data());
   }
 
   IdefixArray1D<real> & deviceView(void) {
-    return this->array;
+    return this->array.deviceView();
   }
 
   idefix::IdefixCommArray1D<real> & commView(void) {
     return this->array;
   }
 
   int Size() {
-    return(array.size());
+    return(array.deviceView().size());
   }
 
   void ResetPointer() {
@@ -62,7 +62,7 @@ class Buffer {
     const int kend = box[KDIR][1];
     const int offset = this->pointer;
 
-    auto arr = this->array;
+    auto arr = this->array.deviceView();
     idefix_for("LoadBuffer3D",kbeg,kend,jbeg,jend,ibeg,iend,
       KOKKOS_LAMBDA (int k, int j, int i) {
       arr(i-ibeg + (j-jbeg)*ni + (k-kbeg)*ninj + offset ) = in(k,j,i);

src/mpi/commArray.hpp

@@ -29,17 +29,26 @@ namespace idefix {
  * if we have access to WITH_MPI_GPU_DIRECT or not.
  */
 template <class T>
-class IdefixCommArrayGpuDirect : public T {
+class IdefixCommArrayGpuDirect {
  public:
-  //inherit constructors
-  using T::T;
+  template <class... Args>
+  explicit IdefixCommArrayGpuDirect(Args... args)
+    :deviceArray(args...)
+  {}
 
   /**
    * Accès it as a simple device view for usiage in kokkos kernels
    * (because need to not carry the host part).
    */
   T & deviceView(void) {
-    return *this;
+    return this->deviceArray;
+  }
+
+  /**
+   * Acces it as a simple communication view for usiage in communication routines.
+   */
+  T commView(void) {
+    return this->deviceArray;
   }
 
   /**
@@ -71,8 +80,16 @@ class IdefixCommArrayGpuDirect : public T {
    * depending on the WITH_MPI_GPU_DIRECT configuration.
    */
   void * commData(void) {
-      return this->data();
+      return this->deviceArray.data();
   }
+
+ private:
+  /**
+  * Device buffer containing the device side of the data to use for communications.
+  * This where the computation are done. Then it is transfered to/from the communication
+  * buffer if needed.
+  */
+  T deviceArray;
 };
 
 /**
@@ -81,49 +98,59 @@ class IdefixCommArrayGpuDirect : public T {
  * if we have access to WITH_MPI_GPU_DIRECT or not.
  */
 template <class T>
-class IdefixCommArrayNoGpuDirect : public T {
+class IdefixCommArrayNoGpuDirect {
  public:
-  //inherit constructors
-  using T::T;
+  template <class... Args>
+  explicit IdefixCommArrayNoGpuDirect(Args... args)
+    :deviceArray(args...)
+    ,commArray(initCommArray(this->deviceArray))
+  {}
 
   /**
-   * Accès it as a simple device view for usiage in kokkos kernels
+   * Acces it as a simple device view for usiage in kokkos kernels
    * (because need to not carry the host part).
    */
   T & deviceView(void) {
-    return *this;
+    return this->deviceArray;
+  }
+
+  /**
+   * Acces it as a simple communication view for usiage in communication routines.
+   */
+  typename T::host_mirror_type & commView(void) {
+    return this->commArray;
   }
 
   /**
    * If needed, transfers the data from the device to the host to be ready to make
    * a communication.
    */
   void syncCommData(void) {
-    Kokkos::deep_copy(this->commArray, *this);
+    Kokkos::deep_copy(this->commArray, this->deviceArray);
   }
 
   /**
    * If needed, transfers the data from the device to the host to be ready to make
    * a communication.
    */
   void syncCommDataAsync(void) {
-    Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), this->commArray, *this);
+    Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), this->commArray, this->deviceArray);
   }
 
   /**
    * If needed, transerts the data to the device after a communication to be ready
    * to use it.
    */
   void syncDeviceData(void) {
-    Kokkos::deep_copy(*this, this->commArray);
+    Kokkos::deep_copy(this->deviceArray, this->commArray);
   }
 
   /**
    * If needed, transerts the data to the device after a communication to be ready
    * to use it.
    */
   void syncDeviceDataAsync(void) {
-    Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), *this, this->commArray);
+    Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), this->deviceArray, this->commArray);
   }
 
   /**
@@ -147,11 +174,17 @@ class IdefixCommArrayNoGpuDirect : public T {
   }
 
  private:
+  /**
+   * Device buffer containing the device side of the data to use for communications.
+   * This where the computation are done. Then it is transfered to/from the communication
+   * buffer if needed.
+   */
+  T deviceArray;
   /**
    * Buffer to contain a copy of the data on host if required. If WITH_MPI_GPU_DIRECT
    * is enabled it directly points the device buffer as no copy is required.
    */
-  typename T::host_mirror_type commArray{initCommArray(*this)};
+  typename T::host_mirror_type commArray;
 };
 
 #ifdef WITH_MPI_GPU_DIRECT

src/mpi/mpi.cpp

@@ -131,11 +131,11 @@ void Mpi::CheckConfig() {
 
   // Run-time check that we can do a reduce on device arrays
   idefix::IdefixCommArray1D<int64_t> src("MPIChecksrc",1);
-  IdefixArray1D<int64_t>::host_mirror_type srcHost = Kokkos::create_mirror_view(src);
+  IdefixArray1D<int64_t>::host_mirror_type srcHost = Kokkos::create_mirror_view(src.deviceView());
 
   if(idfx::prank == 0) {
     srcHost(0) = 0;
-    Kokkos::deep_copy(src, srcHost);
+    Kokkos::deep_copy(src.deviceView(), srcHost);
   }
 
   if(idfx::psize > 1) {
@@ -160,9 +160,9 @@ void Mpi::CheckConfig() {
       } else {
         ierrRecv = idefix::MPI_Recv(src, 1, MPI_INT64_T, idfx::prank-1, 1, MPI_COMM_WORLD, &status);
         // Add our own rank to the data
-        Kokkos::deep_copy(srcHost, src);
+        Kokkos::deep_copy(srcHost, src.deviceView());
         srcHost(0) += idfx::prank;
-        Kokkos::deep_copy(src, srcHost);
+        Kokkos::deep_copy(src.deviceView(), srcHost);
         ierrSend = idefix::MPI_Send(src, 1, MPI_INT64_T, (idfx::prank+1)%idfx::psize, 1,
                                                         MPI_COMM_WORLD);
       }
@@ -199,7 +199,7 @@ void Mpi::CheckConfig() {
   }
 
   // Check that we have the proper end result
-  Kokkos::deep_copy(srcHost, src);
+  Kokkos::deep_copy(srcHost, src.deviceView());
   int64_t size = static_cast<int64_t>(idfx::psize);
   int64_t rank = static_cast<int64_t>(idfx::prank);
   int64_t result = rank == 0 ? size*(size-1)/2 : rank*(rank+1)/2;

src/mpi/mpiView.hpp

@@ -18,7 +18,6 @@
 
 #include "idefix.hpp"
 #include "commArray.hpp"
-#include "buffer.hpp"
 
 namespace idefix {