diff --git a/CMakeLists.txt b/CMakeLists.txt
index fad0058ba..84974d3b4 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -16,6 +16,7 @@ option(Idefix_DEBUG "Enable Idefix debug features (makes the code very slow)" OF
 option(Idefix_RUNTIME_CHECKS "Enable runtime sanity checks" OFF)
 option(Idefix_WERROR "Treat compiler warnings as errors" OFF)
 option(Idefix_PYTHON "Enable python bindings (requires pybind11)" OFF)
+option(Idefix_MPI_GPU_DIRECT "Enable usage of GPU direct to avoid CPU-GPU copies for MPI messages." ON)
 set(Idefix_PROBLEM_DIR "${CMAKE_BINARY_DIR}" CACHE STRING "Problem directory to build for.")
 set(Idefix_CXX_FLAGS "" CACHE STRING "Additional compiler/linker flag")
 set(Idefix_DEFS "definitions.hpp" CACHE FILEPATH "Problem definition header file")
@@ -104,6 +105,10 @@ if(Idefix_MPI)
   add_subdirectory(src/mpi)
 endif()
 
+if (Idefix_MPI_GPU_DIRECT)
+  add_compile_definitions("WITH_MPI_GPU_DIRECT")
+endif()
+
 if(Idefix_HDF5)
   add_compile_definitions("WITH_HDF5")
   if(Idefix_MPI)
@@ -264,6 +269,7 @@ endif()
 message(STATUS "    MPI:  ${Idefix_MPI}")
 message(STATUS "    HDF5: ${Idefix_HDF5}")
 message(STATUS "    Python: ${Idefix_PYTHON}")
+message(STATUS "    GPU direct: ${Idefix_MPI_GPU_DIRECT}")
 message(STATUS "    Reconstruction: ${Idefix_RECONSTRUCTION}")
 message(STATUS "    Precision: ${Idefix_PRECISION}")
 message(STATUS "    Version: ${Idefix_VERSION}")
diff --git a/src/fluid/boundary/axis.cpp b/src/fluid/boundary/axis.cpp
index 45ea49ea5..11bf593c1 100644
--- a/src/fluid/boundary/axis.cpp
+++ b/src/fluid/boundary/axis.cpp
@@ -25,7 +25,7 @@ void Axis::ShowConfig() {
 void Axis::SymmetrizeEx1Side(int jref, IdefixArray3D<real> Ex1) {
 #if DIMENSIONS == 3
 
-  IdefixArray1D<real> Ex1Avg = this->Ex1Avg;
+  auto Ex1Avg = this->Ex1Avg.deviceView();
 
   idefix_for("Ex1_ini",0,data->np_tot[IDIR],
       KOKKOS_LAMBDA(int i) {
@@ -40,7 +40,7 @@ void Axis::SymmetrizeEx1Side(int jref, IdefixArray3D<real> Ex1) {
     #ifdef WITH_MPI
       Kokkos::fence();
       // sum along all of the processes on the same r
-      MPI_Allreduce(MPI_IN_PLACE, Ex1Avg.data(), data->np_tot[IDIR], realMPI,
+      idefix::MPI_Allreduce(MPI_IN_PLACE, this->Ex1Avg, data->np_tot[IDIR], realMPI,
                     MPI_SUM, data->mygrid->AxisComm);
     #endif
   }
@@ -88,7 +88,8 @@ void Axis::RegularizeCurrentSide(int side) {
       jc = data->end[JDIR]-1;
       sign = -1;
     }
-    IdefixArray1D<real> BAvg = this->Ex1Avg;
+    auto BAvg = this->Ex1Avg.deviceView();
+    auto BAvgComm = this->Ex1Avg;
     IdefixArray1D<real> x1 = data->x[IDIR];
     IdefixArray1D<real> dx3 = data->dx[KDIR];
     IdefixArray1D<real> dx2 = data->dx[JDIR];
@@ -107,7 +108,7 @@ void Axis::RegularizeCurrentSide(int side) {
       #ifdef WITH_MPI
         Kokkos::fence();
         // sum along all of the processes on the same r
-        MPI_Allreduce(MPI_IN_PLACE, BAvg.data(), data->np_tot[IDIR], realMPI,
+        MPI_Allreduce(MPI_IN_PLACE, BAvgComm, data->np_tot[IDIR], realMPI,
                       MPI_SUM, data->mygrid->AxisComm);
       #endif
     }
@@ -169,7 +170,7 @@ void Axis::FixBx2sAxis(int side) {
   // Compute the values of Bx and By that are consistent with BX2 along the axis
   #if DIMENSIONS == 3
     IdefixArray4D<real> Vs = this->Vs;
-    IdefixArray2D<real> BAvg = this->BAvg;
+    auto BAvg = this->BAvg.deviceView();
     IdefixArray1D<real> phi = data->x[KDIR];
 
     int jin = 0;
@@ -208,7 +209,7 @@ void Axis::FixBx2sAxis(int side) {
       Kokkos::fence();
       #ifdef WITH_MPI
         // sum along all of the processes on the same r
-        MPI_Allreduce(MPI_IN_PLACE, BAvg.data(), 2*data->np_tot[IDIR], realMPI,
+        idefix::MPI_Allreduce(MPI_IN_PLACE, this->BAvg, 2*data->np_tot[IDIR], realMPI,
                       MPI_SUM, data->mygrid->AxisComm);
       #endif
     }
@@ -392,7 +393,7 @@ void Axis::ExchangeMPI(int side) {
   idfx::pushRegion("Axis::ExchangeMPI");
   #ifdef WITH_MPI
   // Load  the buffers with data
-  int ibeg,iend,jbeg,jend,kbeg,kend,offset;
+  int offset;
   int ny;
   Buffer bufferSend = this->bufferSend;
   IdefixArray1D<int> map = this->mapVars;
@@ -408,7 +409,7 @@ void Axis::ExchangeMPI(int side) {
   MPI_Status recvStatus;
 
   double tStart = MPI_Wtime();
-  MPI_SAFE_CALL(MPI_Start(&recvRequest));
+  MPI_SAFE_CALL(idefix::MPI_Start(&recvRequest));
   idfx::mpiCallsTimer += MPI_Wtime() - tStart;
 
   // Coordinates of the ghost region which needs to be transfered
@@ -470,8 +471,8 @@ void Axis::ExchangeMPI(int side) {
   Kokkos::fence();
 
   tStart = MPI_Wtime();
-  MPI_SAFE_CALL(MPI_Start(&sendRequest));
-  MPI_Wait(&recvRequest,&recvStatus);
+  MPI_SAFE_CALL(idefix::MPI_Start(&sendRequest));
+  idefix::MPI_Wait(&recvRequest,&recvStatus);
   idfx::mpiCallsTimer += MPI_Wtime() - tStart;
 
   // Unpack
@@ -525,8 +526,7 @@ void Axis::ExchangeMPI(int side) {
     } // MHD
   }
 
-  MPI_Wait(&sendRequest, &sendStatus);
-
+  idefix::MPI_Wait(&sendRequest, &sendStatus);
   idfx::mpiCallsTimer += MPI_Wtime() - tStart;
 
 
@@ -593,11 +593,11 @@ void Axis::InitMPI() {
   MPI_SAFE_CALL(MPI_Cart_shift(data->mygrid->AxisComm,0,data->mygrid->nproc[KDIR]/2,
                                &procRecv,&procSend ));
 
-  MPI_SAFE_CALL(MPI_Send_init(bufferSend.data(), bufferSend.Size(), realMPI, procSend,
-                650, data->mygrid->AxisComm, &sendRequest));
+  MPI_SAFE_CALL(idefix::MPI_Send_init(bufferSend.commView(), bufferSend.Size(),
+                realMPI, procSend, 650, data->mygrid->AxisComm, &sendRequest));
 
-  MPI_SAFE_CALL(MPI_Recv_init(bufferRecv.data(), bufferRecv.Size(), realMPI, procRecv,
-                650, data->mygrid->AxisComm, &recvRequest));
+  MPI_SAFE_CALL(idefix::MPI_Recv_init(bufferRecv.commView(), bufferRecv.Size(),
+                realMPI, procRecv, 650, data->mygrid->AxisComm, &recvRequest));
 
   #endif
   idfx::popRegion();
diff --git a/src/fluid/boundary/axis.hpp b/src/fluid/boundary/axis.hpp
index 390a628ce..62a3c1ed3 100644
--- a/src/fluid/boundary/axis.hpp
+++ b/src/fluid/boundary/axis.hpp
@@ -57,8 +57,8 @@ class Axis {
 
   enum {faceTop, faceBot};
 #ifdef WITH_MPI
-  MPI_Request sendRequest;
-  MPI_Request recvRequest;
+  idefix::MPI_Request_1D<real> sendRequest;
+  idefix::MPI_Request_1D<real> recvRequest;
 
   Buffer bufferSend;
   Buffer bufferRecv;
@@ -71,8 +71,8 @@ class Axis {
 #endif
   void InitMPI();
 
-  IdefixArray1D<real> Ex1Avg;
-  IdefixArray2D<real> BAvg;
+  idefix::IdefixCommArray1D<real> Ex1Avg;
+  idefix::IdefixCommArray2D<real> BAvg;
   bool haveCurrent;
   IdefixArray2D<real> JAvg;
   IdefixArray1D<int> symmetryVc;
@@ -155,8 +155,8 @@ Axis::Axis(Boundary<Phys> *boundary) {
     }
     Kokkos::deep_copy(symmetryVs, symmetryVsHost);
 
-    this->Ex1Avg = IdefixArray1D<real>("Axis:Ex1Avg",data->np_tot[IDIR]);
-    this->BAvg = IdefixArray2D<real>("Axis:BxAvg",data->np_tot[IDIR],2);
+    this->Ex1Avg = idefix::IdefixCommArray1D<real>("Axis:Ex1Avg",data->np_tot[IDIR]);
+    this->BAvg = idefix::IdefixCommArray2D<real>("Axis:BxAvg",data->np_tot[IDIR],2);
     if(haveCurrent) {
       this->JAvg = IdefixArray2D<real>("Axis:JAvg",data->np_tot[IDIR],3);
     }
diff --git a/src/fluid/constrainedTransport/EMFexchange.hpp b/src/fluid/constrainedTransport/EMFexchange.hpp
index aad123346..cea40a075 100644
--- a/src/fluid/constrainedTransport/EMFexchange.hpp
+++ b/src/fluid/constrainedTransport/EMFexchange.hpp
@@ -41,7 +41,7 @@ void ConstrainedTransport<Phys>::ExchangeX1(IdefixArray3D<real> ey, IdefixArray3
   MPI_Status recvStatus[2];
 
   double tStart = MPI_Wtime();
-  MPI_SAFE_CALL(MPI_Startall(2, recvRequestX1));
+  MPI_SAFE_CALL(idefix::MPI_Startall(2, recvRequestX1));
   idfx::mpiCallsTimer += MPI_Wtime() - tStart;
 
   BoundaryType lbound = data->lbound[IDIR];
@@ -82,11 +82,10 @@ void ConstrainedTransport<Phys>::ExchangeX1(IdefixArray3D<real> ey, IdefixArray3
   Kokkos::fence();
 
   tStart = MPI_Wtime();
-  MPI_SAFE_CALL(MPI_Startall(2, sendRequestX1));
+  MPI_SAFE_CALL(idefix::MPI_Startall(2, sendRequestX1));
   // Wait for buffers to be received
-
-  MPI_Waitall(2,recvRequestX1,recvStatus);
-  MPI_Waitall(2, sendRequestX1, sendStatus);
+  idefix::MPI_Waitall(2,recvRequestX1,recvStatus);
+  idefix::MPI_Waitall(2, sendRequestX1, sendStatus);
   idfx::mpiCallsTimer += MPI_Wtime() - tStart;
 
   // Unpack
@@ -130,7 +129,7 @@ void ConstrainedTransport<Phys>::ExchangeX2(IdefixArray3D<real> ex, IdefixArray3
   double tStart = MPI_Wtime();
   MPI_Status sendStatus[2];
   MPI_Status recvStatus[2];
-  MPI_SAFE_CALL(MPI_Startall(2, recvRequestX2));
+  MPI_SAFE_CALL(idefix::MPI_Startall(2, recvRequestX2));
   idfx::mpiCallsTimer += MPI_Wtime() - tStart;
 
   BoundaryType lbound = data->lbound[JDIR];
@@ -169,10 +168,10 @@ void ConstrainedTransport<Phys>::ExchangeX2(IdefixArray3D<real> ex, IdefixArray3
   Kokkos::fence();
 
   tStart = MPI_Wtime();
-  MPI_SAFE_CALL(MPI_Startall(2, sendRequestX2));
+  MPI_SAFE_CALL(idefix::MPI_Startall(2, sendRequestX2));
   // Wait for buffers to be received
-  MPI_Waitall(2, recvRequestX2, recvStatus);
-  MPI_Waitall(2, sendRequestX2, sendStatus);
+  idefix::MPI_Waitall(2, recvRequestX2, recvStatus);
+  idefix::MPI_Waitall(2, sendRequestX2, sendStatus);
   idfx::mpiCallsTimer += MPI_Wtime() - tStart;
 
   // Unpack
@@ -216,7 +215,7 @@ void ConstrainedTransport<Phys>::ExchangeX3(IdefixArray3D<real> ex, IdefixArray3
   double tStart = MPI_Wtime();
   MPI_Status sendStatus[2];
   MPI_Status recvStatus[2];
-  MPI_SAFE_CALL(MPI_Startall(2, recvRequestX3));
+  MPI_SAFE_CALL(idefix::MPI_Startall(2, recvRequestX3));
   idfx::mpiCallsTimer += MPI_Wtime() - tStart;
 
   BoundaryType lbound = data->lbound[KDIR];
@@ -256,10 +255,10 @@ void ConstrainedTransport<Phys>::ExchangeX3(IdefixArray3D<real> ex, IdefixArray3
   Kokkos::fence();
 
   tStart = MPI_Wtime();
-  MPI_SAFE_CALL(MPI_Startall(2, sendRequestX3));
+  MPI_SAFE_CALL(idefix::MPI_Startall(2, sendRequestX3));
   // Wait for buffers to be received
-  MPI_Waitall(2, recvRequestX3, recvStatus);
-  MPI_Waitall(2, sendRequestX3, sendStatus);
+  idefix::MPI_Waitall(2, recvRequestX3, recvStatus);
+  idefix::MPI_Waitall(2, sendRequestX3, sendStatus);
   idfx::mpiCallsTimer += MPI_Wtime() - tStart;
 
   // Unpack
diff --git a/src/fluid/constrainedTransport/constrainedTransport.hpp b/src/fluid/constrainedTransport/constrainedTransport.hpp
index ab83b23bc..7734bf5de 100644
--- a/src/fluid/constrainedTransport/constrainedTransport.hpp
+++ b/src/fluid/constrainedTransport/constrainedTransport.hpp
@@ -68,9 +68,9 @@ class ConstrainedTransport {
   IdefixArray3D<real>     Ex3;
 
   // Helper arrays for shearing box
-  IdefixArray2D<real>     sbEyL;
-  IdefixArray2D<real>     sbEyR;
-  IdefixArray2D<real>     sbEyRL;
+  idefix::IdefixCommArray2D<real> sbEyL;
+  idefix::IdefixCommArray2D<real> sbEyR;
+  idefix::IdefixCommArray2D<real> sbEyRL;
 
   // Range of existence
 
@@ -137,12 +137,12 @@ class ConstrainedTransport {
   int bufferSizeX3;
 
   // Requests for MPI persistent communications
-  MPI_Request sendRequestX1[2];
-  MPI_Request sendRequestX2[2];
-  MPI_Request sendRequestX3[2];
-  MPI_Request recvRequestX1[2];
-  MPI_Request recvRequestX2[2];
-  MPI_Request recvRequestX3[2];
+  idefix::MPI_Request_1D<real> sendRequestX1[2];
+  idefix::MPI_Request_1D<real> sendRequestX2[2];
+  idefix::MPI_Request_1D<real> sendRequestX3[2];
+  idefix::MPI_Request_1D<real> recvRequestX1[2];
+  idefix::MPI_Request_1D<real> recvRequestX2[2];
+  idefix::MPI_Request_1D<real> recvRequestX3[2];
 
   Kokkos::Timer timer;    // Internal MPI timer
 #endif
@@ -188,9 +188,9 @@ ConstrainedTransport<Phys>::ConstrainedTransport(Input &input, Fluid<Phys> *hydr
 
   // Allocate shearing box arrays
   if(hydro->haveShearingBox == true) {
-    sbEyL = IdefixArray2D<real>("EMF_sbEyL", data->np_tot[KDIR], data->np_tot[JDIR]);
-    sbEyR = IdefixArray2D<real>("EMF_sbEyR", data->np_tot[KDIR], data->np_tot[JDIR]);
-    sbEyRL = IdefixArray2D<real>("EMF_sbEyRL", data->np_tot[KDIR], data->np_tot[JDIR]);
+    sbEyL = idefix::IdefixCommArray2D<real>("EMF_sbEyL", data->np_tot[KDIR], data->np_tot[JDIR]);
+    sbEyR = idefix::IdefixCommArray2D<real>("EMF_sbEyR", data->np_tot[KDIR], data->np_tot[JDIR]);
+    sbEyRL = idefix::IdefixCommArray2D<real>("EMF_sbEyRL", data->np_tot[KDIR], data->np_tot[JDIR]);
   }
 
   D_EXPAND( ez = IdefixArray3D<real>("EMF_ez",
@@ -331,11 +331,11 @@ ConstrainedTransport<Phys>::ConstrainedTransport(Input &input, Fluid<Phys> *hydr
   if(data->rbound[IDIR] == shearingbox ) procSend = MPI_PROC_NULL;
   if(data->lbound[IDIR] == shearingbox ) procRecv = MPI_PROC_NULL;
 
-  MPI_SAFE_CALL(MPI_Send_init(BufferSendX1[faceRight].data(), bufferSizeX1, realMPI, procSend, 100,
-                mygrid->CartComm, &sendRequestX1[faceRight]));
+  MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX1[faceRight].commView(), bufferSizeX1,
+                realMPI, procSend, 100, mygrid->CartComm, &sendRequestX1[faceRight]));
 
-  MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX1[faceLeft].data(), bufferSizeX1, realMPI, procRecv, 100,
-                mygrid->CartComm, &recvRequestX1[faceLeft]));
+  MPI_SAFE_CALL(idefix::MPI_Recv_init(BufferRecvX1[faceLeft].commView(), bufferSizeX1,
+                realMPI, procRecv, 100, mygrid->CartComm, &recvRequestX1[faceLeft]));
 
   // Send to the left
   // We receive from procRecv, and we send to procSend
@@ -344,52 +344,52 @@ ConstrainedTransport<Phys>::ConstrainedTransport(Input &input, Fluid<Phys> *hydr
   if(data->lbound[IDIR] == shearingbox ) procSend = MPI_PROC_NULL;
   if(data->rbound[IDIR] == shearingbox ) procRecv = MPI_PROC_NULL;
 
-  MPI_SAFE_CALL(MPI_Send_init(BufferSendX1[faceLeft].data(), bufferSizeX1, realMPI, procSend, 101,
-                mygrid->CartComm, &sendRequestX1[faceLeft]));
+  MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX1[faceLeft].commView(), bufferSizeX1,
+                 realMPI, procSend, 101, mygrid->CartComm, &sendRequestX1[faceLeft]));
 
-  MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX1[faceRight].data(), bufferSizeX1, realMPI, procRecv, 101,
-                mygrid->CartComm, &recvRequestX1[faceRight]));
+  MPI_SAFE_CALL(idefix::MPI_Recv_init(BufferRecvX1[faceRight].commView(), bufferSizeX1,
+                 realMPI, procRecv, 101, mygrid->CartComm, &recvRequestX1[faceRight]));
 
   #if DIMENSIONS >= 2
   // We receive from procRecv, and we send to procSend
   MPI_SAFE_CALL(MPI_Cart_shift(mygrid->CartComm,1,1,&procRecv,&procSend ));
 
-  MPI_SAFE_CALL(MPI_Send_init(BufferSendX2[faceRight].data(), bufferSizeX2, realMPI, procSend, 200,
-                mygrid->CartComm, &sendRequestX2[faceRight]));
+  MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX2[faceRight].commView(), bufferSizeX2,
+                 realMPI, procSend, 200, mygrid->CartComm, &sendRequestX2[faceRight]));
 
-  MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX2[faceLeft].data(), bufferSizeX2, realMPI, procRecv, 200,
-                mygrid->CartComm, &recvRequestX2[faceLeft]));
+  MPI_SAFE_CALL(idefix::MPI_Recv_init(BufferRecvX2[faceLeft].commView(), bufferSizeX2,
+                 realMPI, procRecv, 200, mygrid->CartComm, &recvRequestX2[faceLeft]));
 
   // Send to the left
   // We receive from procRecv, and we send to procSend
   MPI_SAFE_CALL(MPI_Cart_shift(mygrid->CartComm,1,-1,&procRecv,&procSend ));
 
-  MPI_SAFE_CALL(MPI_Send_init(BufferSendX2[faceLeft].data(), bufferSizeX2, realMPI, procSend, 201,
-                mygrid->CartComm, &sendRequestX2[faceLeft]));
+  MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX2[faceLeft].commView(), bufferSizeX2,
+                 realMPI, procSend, 201, mygrid->CartComm, &sendRequestX2[faceLeft]));
 
-  MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX2[faceRight].data(), bufferSizeX2, realMPI, procRecv, 201,
-                mygrid->CartComm, &recvRequestX2[faceRight]));
+  MPI_SAFE_CALL(idefix::MPI_Recv_init(BufferRecvX2[faceRight].commView(), bufferSizeX2,
+                 realMPI, procRecv, 201, mygrid->CartComm, &recvRequestX2[faceRight]));
   #endif
 
   #if DIMENSIONS == 3
   // We receive from procRecv, and we send to procSend
   MPI_SAFE_CALL(MPI_Cart_shift(mygrid->CartComm,2,1,&procRecv,&procSend ));
 
-  MPI_SAFE_CALL(MPI_Send_init(BufferSendX3[faceRight].data(), bufferSizeX3, realMPI, procSend, 300,
-                mygrid->CartComm, &sendRequestX3[faceRight]));
+  MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX3[faceRight].commView(), bufferSizeX3,
+                 realMPI, procSend, 300, mygrid->CartComm, &sendRequestX3[faceRight]));
 
-  MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX3[faceLeft].data(), bufferSizeX3, realMPI, procRecv, 300,
-                mygrid->CartComm, &recvRequestX3[faceLeft]));
+  MPI_SAFE_CALL(idefix::MPI_Recv_init(BufferRecvX3[faceLeft].commView(), bufferSizeX3,
+                 realMPI, procRecv, 300, mygrid->CartComm, &recvRequestX3[faceLeft]));
 
   // Send to the left
   // We receive from procRecv, and we send to procSend
   MPI_SAFE_CALL(MPI_Cart_shift(mygrid->CartComm,2,-1,&procRecv,&procSend ));
 
-  MPI_SAFE_CALL(MPI_Send_init(BufferSendX3[faceLeft].data(), bufferSizeX3, realMPI, procSend, 301,
-                mygrid->CartComm, &sendRequestX3[faceLeft]));
+  MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX3[faceLeft].commView(), bufferSizeX3,
+                 realMPI, procSend, 301, mygrid->CartComm, &sendRequestX3[faceLeft]));
 
-  MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX3[faceRight].data(), bufferSizeX3, realMPI, procRecv, 301,
-                mygrid->CartComm, &recvRequestX3[faceRight]));
+  MPI_SAFE_CALL(idefix::MPI_Recv_init(BufferRecvX3[faceRight].commView(), bufferSizeX3,
+                 realMPI, procRecv, 301, mygrid->CartComm, &recvRequestX3[faceRight]));
   #endif
 
 #endif  // WITH_MPI
@@ -409,17 +409,17 @@ ConstrainedTransport<Phys>::~ConstrainedTransport() {
     // Properly clean up the mess
     idfx::cout << "Emf: Cleaning up MPI persistent communication channels" << std::endl;
     for(int i=0 ; i< 2; i++) {
-      MPI_Request_free( &sendRequestX1[i]);
-      MPI_Request_free( &recvRequestX1[i]);
+      idefix::MPI_Request_free( &sendRequestX1[i]);
+      idefix::MPI_Request_free( &recvRequestX1[i]);
 
     #if DIMENSIONS >= 2
-      MPI_Request_free( &sendRequestX2[i]);
-      MPI_Request_free( &recvRequestX2[i]);
+      idefix::MPI_Request_free( &sendRequestX2[i]);
+      idefix::MPI_Request_free( &recvRequestX2[i]);
     #endif
 
     #if DIMENSIONS == 3
-      MPI_Request_free( &sendRequestX3[i]);
-      MPI_Request_free( &recvRequestX3[i]);
+      idefix::MPI_Request_free( &sendRequestX3[i]);
+      idefix::MPI_Request_free( &recvRequestX3[i]);
     #endif
     }
     #endif
diff --git a/src/fluid/constrainedTransport/enforceEMFBoundary.hpp b/src/fluid/constrainedTransport/enforceEMFBoundary.hpp
index 9d3d6913f..f7287a8aa 100644
--- a/src/fluid/constrainedTransport/enforceEMFBoundary.hpp
+++ b/src/fluid/constrainedTransport/enforceEMFBoundary.hpp
@@ -108,9 +108,9 @@ void ConstrainedTransport<Phys>::SymmetrizeEMFShearingBox() {
   idfx::pushRegion("Emf::EnforceEMFBoundary");
   #if MHD == YES
 
-    IdefixArray2D<real> sbEyL = this->sbEyL;
-    IdefixArray2D<real> sbEyR = this->sbEyR;
-    IdefixArray2D<real> sbEyRL = this->sbEyRL;
+    auto sbEyL = this->sbEyL.deviceView();
+    auto sbEyR = this->sbEyR.deviceView();
+    auto sbEyRL = this->sbEyRL.deviceView();
 
     IdefixArray3D<real> ey = this->ey;
 
@@ -142,16 +142,16 @@ void ConstrainedTransport<Phys>::SymmetrizeEMFShearingBox() {
           // We send to our left (which, by periodicity, is the right end of the domain)
           // our value of sbEyL and get
           Kokkos::fence();
-          MPI_Sendrecv(sbEyL.data(), size, realMPI, procLeft, 2001,
-                       sbEyR.data(), size, realMPI, procLeft, 2002,
+          idefix::MPI_Sendrecv(this->sbEyL, size, realMPI, procLeft, 2001,
+                       this->sbEyR, size, realMPI, procLeft, 2002,
                        data->mygrid->CartComm, &status );
         }
         if(data->rbound[IDIR]==shearingbox) {
           Kokkos::fence();
           // We send to our right (which, by periodicity, is the left end (=beginning)
           // of the domain) our value of sbEyR and get sbEyL
-          MPI_Sendrecv(sbEyR.data(), size, realMPI, procRight, 2002,
-                       sbEyL.data(), size, realMPI, procRight, 2001,
+          idefix::MPI_Sendrecv(this->sbEyR, size, realMPI, procRight, 2002,
+                       this->sbEyL, size, realMPI, procRight, 2001,
                        data->mygrid->CartComm, &status );
         }
       }
diff --git a/src/mpi/buffer.hpp b/src/mpi/buffer.hpp
index 35164355b..9fbba748b 100644
--- a/src/mpi/buffer.hpp
+++ b/src/mpi/buffer.hpp
@@ -10,6 +10,7 @@
 
 #include "idefix.hpp"
 #include "arrays.hpp"
+#include "mpiView.hpp"
 
 using BoundingBox = std::array<std::array<int,2>,3>;
 
@@ -17,7 +18,9 @@ using BoundingBox = std::array<std::array<int,2>,3>;
 class Buffer {
  public:
   Buffer() = default;
-  explicit Buffer(size_t size): pointer{0}, array{IdefixArray1D<real>("BufferArray",size)} {};
+  explicit Buffer(size_t size)
+    : pointer{0}
+    , array{IdefixArray1D<real>("BufferArray",size)} {}
 
   // Compute the size of a bounding box
   static size_t ComputeBoxSize(BoundingBox box) {
@@ -27,11 +30,15 @@ class Buffer {
     return(ninjnk);
   }
 
-  void* data() {
+  void* deviceData() {
     return(array.data());
   }
 
-  const IdefixArray1D<real> & getArray(void) const {
+  IdefixArray1D<real> & deviceView(void) {
+    return this->array;
+  }
+
+  idefix::IdefixCommArray1D<real> & commView(void) {
     return this->array;
   }
 
@@ -79,7 +86,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("LoadBuffer4D_var",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(var, k,j,i);
@@ -102,7 +109,7 @@ class Buffer {
     const int jend = box[JDIR][1];
     const int kend = box[KDIR][1];
     const int offset = this->pointer;
-    auto arr = this->array;
+    auto arr = this->array.deviceView();
 
     idefix_for("LoadBuffer4D_map",0,map.size(),
                              kbeg,kend,
@@ -128,7 +135,7 @@ class Buffer {
     const int jend = box[JDIR][1];
     const int kend = box[KDIR][1];
     const int offset = this->pointer;
-    auto arr = this->array;
+    auto arr = this->array.deviceView();
 
     idefix_for("UnLoadBuffer3D",kbeg,kend,jbeg,jend,ibeg,iend,
       KOKKOS_LAMBDA (int k, int j, int i) {
@@ -153,7 +160,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("UnLoadBuffer4D_var",kbeg,kend,jbeg,jend,ibeg,iend,
       KOKKOS_LAMBDA (int k, int j, int i) {
         out(var,k,j,i) = arr(i-ibeg + (j-jbeg)*ni + (k-kbeg)*ninj + offset );
@@ -178,7 +185,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("UnLoadBuffer4D_var_sym",kbeg,kend,jbeg,jend,ibeg,iend,
       KOKKOS_LAMBDA (int k, int j, int i) {
         const int jinverted = jend-(j-jbeg)-1;
@@ -204,7 +211,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("UnLoadBuffer4D_map",0,map.size(),
                               kbeg,kend,
                               jbeg,jend,
@@ -232,7 +239,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("UnLoadBuffer4D_map_sym",0,map.size(),
                               kbeg,kend,
                               jbeg,jend,
@@ -247,10 +254,9 @@ class Buffer {
     this->pointer += ninjnk*map.size();
   }
 
-
  private:
-  size_t pointer;
-  IdefixArray1D<real> array;
+  size_t pointer{0};
+  idefix::IdefixCommArray1D<real> array{};
 };
 
 #endif // MPI_BUFFER_HPP_
diff --git a/src/mpi/exchanger.cpp b/src/mpi/exchanger.cpp
index 033ce4041..5823b7a64 100644
--- a/src/mpi/exchanger.cpp
+++ b/src/mpi/exchanger.cpp
@@ -129,22 +129,22 @@ void Exchanger::Init(
   // X1-dir exchanges
   // We receive from procRecv, and we send to procSend
 
-  MPI_Send_init(BufferSend[faceRight].data(), bufferSizeSend[faceRight], realMPI,
+  idefix::MPI_Send_init(BufferSend[faceRight].commView(), bufferSizeSend[faceRight], realMPI,
             procSend[faceRight], thisInstance*2,
             grid->CartComm, &sendRequest[faceRight]);
 
-  MPI_Recv_init(BufferRecv[faceLeft].data(), bufferSizeRecv[faceLeft], realMPI,
+  idefix::MPI_Recv_init(BufferRecv[faceLeft].commView(), bufferSizeRecv[faceLeft], realMPI,
             procRecv[faceLeft],thisInstance*2,
             grid->CartComm, &recvRequest[faceLeft]);
 
   // Send to the left
   // We receive from procRecv, and we send to procSend
 
-  MPI_Send_init(BufferSend[faceLeft].data(), bufferSizeSend[faceLeft], realMPI,
+  idefix::MPI_Send_init(BufferSend[faceLeft].commView(), bufferSizeSend[faceLeft], realMPI,
             procSend[faceLeft],thisInstance*2+1,
             grid->CartComm, &sendRequest[faceLeft]);
 
-  MPI_Recv_init(BufferRecv[faceRight].data(), bufferSizeRecv[faceRight], realMPI,
+  idefix::MPI_Recv_init(BufferRecv[faceRight].commView(), bufferSizeRecv[faceRight], realMPI,
             procRecv[faceRight], thisInstance*2+1,
             grid->CartComm, &recvRequest[faceRight]);
 
@@ -163,8 +163,8 @@ Exchanger::~Exchanger() {
     // Properly clean up the mess
     #ifdef MPI_PERSISTENT
       for(int i=0 ; i< 2; i++) {
-        MPI_Request_free( &sendRequest[i]);
-        MPI_Request_free( &recvRequest[i]);
+        idefix::MPI_Request_free( &sendRequest[i]);
+        idefix::MPI_Request_free( &recvRequest[i]);
       }
     #endif
     isInitialized = false;
@@ -189,7 +189,7 @@ void Exchanger::Exchange(IdefixArray4D<real> Vc, IdefixArray4D<real> Vs) {
   MPI_Status sendStatus[2];
   MPI_Status recvStatus[2];
 
-  MPI_Startall(2, recvRequest);
+  idefix::MPI_Startall(2, recvRequest);
   idfx::mpiCallsTimer += MPI_Wtime() - tStart;
 #endif
   myTimer += MPI_Wtime();
@@ -211,65 +211,68 @@ void Exchanger::Exchange(IdefixArray4D<real> Vc, IdefixArray4D<real> Vs) {
   Kokkos::fence();
   myTimer -= MPI_Wtime();
   tStart = MPI_Wtime();
+
 #ifdef MPI_PERSISTENT
-  MPI_Startall(2, sendRequest);
+  idefix::MPI_Startall(2, sendRequest);
   // Wait for buffers to be received
-  MPI_Waitall(2,recvRequest,recvStatus);
+  idefix::MPI_Waitall(2, recvRequest, recvStatus);
 
 #else
 
   #ifdef MPI_NON_BLOCKING
   MPI_Status sendStatus[2];
   MPI_Status recvStatus[2];
-  MPI_Request sendRequest[2];
-  MPI_Request recvRequest[2];
+  idefix::MPI_Request sendRequest[2];
+  idefix::MPI_Request recvRequest[2];
 
   // We receive from procRecv, and we send to procSend
 
-  MPI_Isend(BufferSend[faceRight].data(), bufferSizeSend[faceRight], realMPI,
+  idefix::MPI_Isend(BufferSend[faceRight].commView(), bufferSizeSend[faceRight], realMPI,
             procSend[faceRight], 100, mygrid->CartComm, &sendRequest[0]);
 
-  MPI_Irecv(BufferRecv[faceLeft].data(), bufferSizeRecv[faceLeft], realMPI,
+  idefix::MPI_Irecv(BufferRecv[faceLeft].commView(), bufferSizeRecv[faceLeft], realMPI,
             procRecv[faceLeft], 100, mygrid->CartComm, &recvRequest[0]);
   // Send to the left
   // We receive from procRecv, and we send to procSend
 
-  MPI_Isend(BufferSend[faceLeft].data(), bufferSizeSend[faceLeft], realMPI,
+  idefix::MPI_Isend(BufferSend[faceLeft].commView(), bufferSizeSend[faceLeft], realMPI,
             procSend[faceLeft], 101, mygrid->CartComm, &sendRequest[1]);
 
-  MPI_Irecv(BufferRecv[faceRight].data(), bufferSizeRecv[faceRight], realMPI,
+  idefix::MPI_Irecv(BufferRecv[faceRight].commView(), bufferSizeRecv[faceRight], realMPI,
             procRecv[faceRight], 101, mygrid->CartComm, &recvRequest[1]);
 
   // Wait for recv to complete (we don't care about the sends)
-  MPI_Waitall(2, recvRequest, recvStatus);
+  idefix::MPI_Waitall(2, recvRequest, recvStatus);
 
   #else
   MPI_Status status;
   // Send to the right
   // We receive from procRecv, and we send to procSend
 
-  MPI_Sendrecv(BufferSend[faceRight].data(), bufferSizeSend[faceRight], realMPI,
+  idefix::MPI_Sendrecv(BufferSend[faceRight].commView(), bufferSizeSend[faceRight], realMPI,
                 procSend[faceRight], 100,
-                BufferRecv[faceLeft].data(), bufferSizeRecv[faceLeft], realMPI,
+                BufferRecv[faceLeft].commView(), bufferSizeRecv[faceLeft], realMPI,
                 procRecv[faceLeft], 100,
                 grid->CartComm, &status);
 
   // Send to the left
   // We receive from procRecv, and we send to procSend
 
-  MPI_Sendrecv(BufferSend[faceLeft].data(), bufferSizeSend[faceLeft], realMPI,
+  idefix::MPI_Sendrecv(BufferSend[faceLeft].commView(), bufferSizeSend[faceLeft], realMPI,
                 procSend[faceLeft], 101,
-                BufferRecv[faceRight].data(), bufferSizeRecv[faceRight], realMPI,
+                BufferRecv[faceRight].commView(), bufferSizeRecv[faceRight], realMPI,
                 procRecv[faceRight], 101,
                 grid->CartComm, &status);
   #endif
 #endif
 myTimer += MPI_Wtime();
-idfx::mpiCallsTimer += MPI_Wtime() - tStart;
+
 // Unpack
 BufferLeft=BufferRecv[faceLeft];
 BufferRight=BufferRecv[faceRight];
 
+idfx::mpiCallsTimer += MPI_Wtime() - tStart;
+
 BufferLeft.ResetPointer();
 BufferRight.ResetPointer();
 
@@ -292,11 +295,11 @@ if(recvRight) {
 myTimer -= MPI_Wtime();
 #ifdef MPI_NON_BLOCKING
   // Wait for the sends if they have not yet completed
-  MPI_Waitall(2, sendRequest, sendStatus);
+  idefix::MPI_Waitall(2, sendRequest, sendStatus);
 #endif
 
 #ifdef MPI_PERSISTENT
-  MPI_Waitall(2, sendRequest, sendStatus);
+  idefix::MPI_Waitall(2, sendRequest, sendStatus);
 #endif
   myTimer += MPI_Wtime();
   bytesSentOrReceived += (bufferSizeRecv[faceLeft]
diff --git a/src/mpi/exchanger.hpp b/src/mpi/exchanger.hpp
index 2e1affcb0..81d6d1640 100644
--- a/src/mpi/exchanger.hpp
+++ b/src/mpi/exchanger.hpp
@@ -72,8 +72,8 @@ class Exchanger {
   bool haveVs{false};
 
   // Requests for MPI persistent communications
-  MPI_Request sendRequest[2];
-  MPI_Request recvRequest[2];
+  idefix::MPI_Request_1D<real> sendRequest[2];
+  idefix::MPI_Request_1D<real> recvRequest[2];
 
   Grid *grid;
 };
diff --git a/src/mpi/mpi.cpp b/src/mpi/mpi.cpp
index d0759199a..a90e22ef5 100644
--- a/src/mpi/mpi.cpp
+++ b/src/mpi/mpi.cpp
@@ -123,14 +123,14 @@ void Mpi::ExchangeX3(IdefixArray4D<real> Vc, IdefixArray4D<real> Vs) {
 void Mpi::CheckConfig() {
   idfx::pushRegion("Mpi::CheckConfig");
   // compile time check
-  #ifdef KOKKOS_ENABLE_CUDA
+  #if defined(KOKKOS_ENABLE_CUDA) && defined(WITH_MPI_GPU_DIRECT)
     #if defined(MPIX_CUDA_AWARE_SUPPORT) && !MPIX_CUDA_AWARE_SUPPORT
       #error Your MPI library is not CUDA Aware (check Idefix requirements).
     #endif
   #endif /* MPIX_CUDA_AWARE_SUPPORT */
 
   // Run-time check that we can do a reduce on device arrays
-  IdefixArray1D<int64_t> src("MPIChecksrc",1);
+  idefix::IdefixCommArray1D<int64_t> src("MPIChecksrc",1);
   IdefixArray1D<int64_t>::host_mirror_type srcHost = Kokkos::create_mirror_view(src);
 
   if(idfx::prank == 0) {
@@ -155,15 +155,15 @@ void Mpi::CheckConfig() {
       MPI_Status status;
       int ierrSend, ierrRecv;
       if(idfx::prank == 0) {
-        ierrSend = MPI_Send(src.data(), 1, MPI_INT64_T, idfx::prank+1, 1, MPI_COMM_WORLD);
-        ierrRecv = MPI_Recv(src.data(), 1, MPI_INT64_T, idfx::psize-1, 1, MPI_COMM_WORLD, &status);
+        ierrSend = idefix::MPI_Send(src, 1, MPI_INT64_T, idfx::prank+1, 1, MPI_COMM_WORLD);
+        ierrRecv = idefix::MPI_Recv(src, 1, MPI_INT64_T, idfx::psize-1, 1, MPI_COMM_WORLD, &status);
       } else {
-        ierrRecv = MPI_Recv(src.data(), 1, MPI_INT64_T, idfx::prank-1, 1, MPI_COMM_WORLD, &status);
+        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);
         srcHost(0) += idfx::prank;
         Kokkos::deep_copy(src, srcHost);
-        ierrSend = MPI_Send(src.data(), 1, MPI_INT64_T, (idfx::prank+1)%idfx::psize, 1,
+        ierrSend = idefix::MPI_Send(src, 1, MPI_INT64_T, (idfx::prank+1)%idfx::psize, 1,
                                                         MPI_COMM_WORLD);
       }
 
diff --git a/src/mpi/mpi.hpp b/src/mpi/mpi.hpp
index b8f929326..8c31e83f8 100644
--- a/src/mpi/mpi.hpp
+++ b/src/mpi/mpi.hpp
@@ -17,6 +17,7 @@
 #include "grid.hpp"
 #include "buffer.hpp"
 #include "exchanger.hpp"
+#include "mpiView.hpp"
 
 
 class DataBlock;
diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp
new file mode 100644
index 000000000..3ee17b051
--- /dev/null
+++ b/src/mpi/mpiView.hpp
@@ -0,0 +1,621 @@
+// ***********************************************************************************
+// Idefix MHD astrophysical code
+// Copyright(C) Geoffroy R. J. Lesur <geoffroy.lesur@univ-grenoble-alpes.fr>
+// and other code contributors
+// Licensed under CeCILL 2.1 License, see COPYING for more information
+// ***********************************************************************************
+
+#ifndef MPI_MPIVIEW_HPP_
+#define MPI_MPIVIEW_HPP_
+
+//To enable to force transfers and validate that it is working in any conditions.
+//#define FORCE_TRANSFERS_ALL_TIME
+
+#include <mpi.h>
+
+#include <cassert>
+#include <vector>
+#include <array>
+
+#include "idefix.hpp"
+#include "buffer.hpp"
+
+namespace idefix {
+
+/**
+ * Provide a specific array used for MPI communications. It embeds the standard
+ * array as usual, and also optionally a communication copy on the host, depending
+ * if we have access to WITH_MPI_GPU_DIRECT or not.
+ */
+template <class T>
+class IdefixCommArrayGpuDirect : public T {
+ public:
+  //inherit constructors
+  using T::T;
+
+  /**
+   * 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;
+  }
+
+  /**
+   * If needed, transfers the data from the device to the host to be ready to make
+   * a communication.
+   */
+  void syncCommData(void) {}
+
+  /**
+   * If needed, transfers the data from the device to the host to be ready to make
+   * a communication.
+   */
+  void syncCommDataAsync(void) {}
+
+  /**
+   * If needed, transerts the data to the device after a communication to be ready
+   * to use it.
+   */
+  void syncDeviceData(void) {}
+
+  /**
+   * If needed, transerts the data to the device after a communication to be ready
+   * to use it.
+   */
+  void syncDeviceDataAsync(void) {}
+
+  /**
+   * Get access to the communication array. It will be either on host or device
+   * depending on the WITH_MPI_GPU_DIRECT configuration.
+   */
+  void * commData(void) {
+      return this->data();
+  }
+};
+
+/**
+ * Provide a specific array used for MPI communications. It embeds the standard
+ * array as usual, and also optionally a communication copy on the host, depending
+ * if we have access to WITH_MPI_GPU_DIRECT or not.
+ */
+template <class T>
+class IdefixCommArrayNoGpuDirect : public T {
+ public:
+  //inherit constructors
+  using T::T;
+
+  /**
+   * 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;
+  }
+
+  /**
+   * 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);
+  }
+
+  /**
+   * 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);
+  }
+
+  /**
+   * 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);
+  }
+
+  /**
+   * 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);
+  }
+
+  /**
+   * Get access to the communication array. It will be either on host or device
+   * depending on the WITH_MPI_GPU_DIRECT configuration.
+   */
+  void * commData(void) {
+    return this->commArray.data();
+  }
+
+ private:
+  static typename T::host_mirror_type initCommArray(T & deviceArray) {
+    #ifdef FORCE_TRANSFERS_ALL_TIME
+      //for validation purpose, force the transfers by using a copy anytime.
+      return Kokkos::create_mirror(deviceArray);
+    #else
+      //if no GPU DIRECT, it creates a copy only if deviceArray is on GPU,
+      //if on host, it just reference it
+      return Kokkos::create_mirror_view(deviceArray);
+    #endif
+  }
+
+ private:
+  /**
+   * 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)};
+};
+
+#ifdef WITH_MPI_GPU_DIRECT
+  template <class T> using IdefixCommArray = IdefixCommArrayGpuDirect<T>;
+#else
+  template <class T> using IdefixCommArray = IdefixCommArrayNoGpuDirect<T>;
+#endif
+
+/** Wrap the idefix 1D array to use it for communication purpose. */
+template <class T> using IdefixCommArray1D = IdefixCommArray< IdefixArray1D<T> >;
+/** Wrap the idefix 2D array to use it for communication purpose. */
+template <class T> using IdefixCommArray2D = IdefixCommArray< IdefixArray2D<T> >;
+/** Wrap the idefix 3D array to use it for communication purpose. */
+template <class T> using IdefixCommArray3D = IdefixCommArray< IdefixArray3D<T> >;
+/** Wrap the idefix 4D array to use it for communication purpose. */
+template <class T> using IdefixCommArray4D = IdefixCommArray< IdefixArray4D<T> >;
+
+template <class T>
+struct Idefix_MPI_Request {
+  IdefixCommArray<T> commArray;
+  MPI_Request request;
+};
+
+/** Wrap the idefix 1D array to use it for communication purpose. */
+template <class T> using MPI_Request_1D = Idefix_MPI_Request< IdefixArray1D<T> >;
+/** Wrap the idefix 2D array to use it for communication purpose. */
+template <class T> using MPI_Request_2D = Idefix_MPI_Request< IdefixArray2D<T> >;
+/** Wrap the idefix 3D array to use it for communication purpose. */
+template <class T> using MPI_Request_3D = Idefix_MPI_Request< IdefixArray3D<T> >;
+/** Wrap the idefix 4D array to use it for communication purpose. */
+template <class T> using MPI_Request_4D = Idefix_MPI_Request< IdefixArray4D<T> >;
+
+/**
+ * Provide a wrapper of the standard MPI_Sendrecv by handling an idefix
+ * communication array. It will automate the process of transfering the data
+ * to the host if required (when WITH_MPI_GPU_DIRECT is disabled).
+ */
+template <class T>
+int MPI_Sendrecv(IdefixCommArray<T> & sendbuf, int sendcount, MPI_Datatype sendtype,
+  int dest, int sendtag, IdefixCommArray<T> &recvbuf, int recvcount,
+  MPI_Datatype recvtype, int source, int recvtag, MPI_Comm comm, MPI_Status *status) {
+  //check size
+  assert(sendbuf.span() == sendcount);
+  assert(sendbuf.span_is_contiguous());
+  assert(recvbuf.span() == sendcount);
+  assert(recvbuf.span_is_contiguous());
+
+  //make a local allocation if needed
+  sendbuf.syncCommData();
+
+  //make the communication
+  const int result = ::MPI_Sendrecv(sendbuf.commData(), sendcount, sendtype, dest, sendtag,
+               recvbuf.commData(), recvcount, recvtype, source, recvtag,
+               comm, status);
+
+  //copy back if needed
+  recvbuf.syncDeviceData();
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Send_init() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Send_init(IdefixCommArray<T> & buf, int count, MPI_Datatype datatype, int dest,
+                  int tag, MPI_Comm comm, Idefix_MPI_Request<T> *request) {
+    //check
+    assert(request != nullptr);
+    assert(buf.span() == count);
+    assert(buf.span_is_contiguous());
+
+    //store the values
+    request->commArray = buf;
+
+    //call MPI
+    const int result = MPI_Send_init(buf.commData(), count, datatype,
+      dest, tag, comm, &request->request);
+
+    //ok
+    return result;
+}
+
+/**
+ * Overlad the MPI_Recv_init() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Recv_init(IdefixCommArray<T> & buf, int count, MPI_Datatype datatype, int
+    source, int tag, MPI_Comm comm, Idefix_MPI_Request<T>* request) {
+  //check
+  assert(request != nullptr);
+  assert(buf.span() == count);
+  assert(buf.span_is_contiguous());
+
+  //store the values
+  request->commArray = buf;
+
+  //call MPI
+  const int result = MPI_Recv_init(buf.commData(), count, datatype,
+    source, tag, comm, &request->request);
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Request_free() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Request_free(Idefix_MPI_Request<T>* request) {
+  //check
+  assert(request != nullptr);
+
+  //call MPI
+  return MPI_Request_free(&request->request);
+}
+
+/**
+ * Overlad the MPI_Startall() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Startall(int count, Idefix_MPI_Request<T>* request) {
+  //check
+  assert(request != nullptr);
+  assert(count >= 0);
+
+  //sync the copies in async mode
+  for (size_t i = 0 ; i < count ; i++) {
+    request[i].commArray.syncCommDataAsync();
+  }
+
+  //wait transfers to be done
+  Kokkos::fence();
+
+  //launch the comms
+  int final_result = 0;
+  for (size_t i = 0 ; i < count ; i++) {
+    const int result = ::MPI_Start(&request[i].request);
+    if (result < 0)
+      final_result = result;
+  }
+
+  //ok
+  return final_result;
+}
+
+/**
+ * Overlad the MPI_Waitall() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Waitall(int count, Idefix_MPI_Request<T>* request,
+  MPI_Status *array_of_statuses) {
+  //check
+  assert(request != nullptr);
+  assert(count >= 0);
+
+  //wait the comms
+  int final_result = 0;
+  for (size_t i = 0 ; i < count ; i++) {
+    const int result = ::MPI_Wait(&request[i].request, &array_of_statuses[i]);
+    if (result < 0)
+      final_result = result;
+  }
+
+  //sync the copies in async mode
+  for (size_t i = 0 ; i < count ; i++) {
+    request[i].commArray.syncDeviceDataAsync();
+  }
+
+  //wait transfers to be done
+  Kokkos::fence();
+
+  //ok
+  return final_result;
+}
+
+/**
+ * Overlad the MPI_Start() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Start(Idefix_MPI_Request<T>* request) {
+  //check
+  assert(request != nullptr);
+
+  //sync the copies in async mode
+  request->commArray.syncCommData();
+
+  //call MPI
+  return ::MPI_Start(&request->request);
+}
+
+/**
+ * Overlad the MPI_Wait() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Wait(Idefix_MPI_Request<T>* request, MPI_Status *array_of_statuses) {
+  //check
+  assert(request != nullptr);
+
+  //wait the comms
+  const int result = ::MPI_Wait(&request->request, array_of_statuses);
+
+  //sync the copies in async mode
+  request->commArray.syncDeviceData();
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Send() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Send(IdefixCommArray<T> & buf, int count, MPI_Datatype datatype, int
+    source, int tag, MPI_Comm comm) {
+  //check
+  assert(buf.span() == count);
+  assert(buf.span_is_contiguous());
+
+  //sync
+  buf.syncCommData();
+
+  //call MPI
+  const int result = ::MPI_Send(buf.commData(), count, datatype,
+    source, tag, comm);
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Recv() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Recv(IdefixCommArray<T> & buf, int count, MPI_Datatype datatype, int
+    source, int tag, MPI_Comm comm, MPI_Status * status) {
+  //check
+  assert(buf.span() == count);
+  assert(buf.span_is_contiguous());
+
+  //call MPI
+  const int result = ::MPI_Recv(buf.commData(), count, datatype,
+    source, tag, comm, status);
+
+  //sync
+  buf.syncDeviceData();
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Recv() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T, int U>
+int MPI_Recv(std::array<T, U> & buf, int count, MPI_Datatype datatype, int
+    source, int tag, MPI_Comm comm, MPI_Status * status) {
+  //call MPI
+  return ::MPI_Recv(buf.data(), count, datatype, source, tag, comm, status);
+}
+
+/**
+ * Overlad the MPI_Isend() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Isend(IdefixCommArray<T> & buf, int count, MPI_Datatype datatype, int
+    source, int tag, MPI_Comm comm, Idefix_MPI_Request<T>* request) {
+  //check
+  assert(request != nullptr);
+  assert(buf.span() == count);
+  assert(buf.span_is_contiguous());
+
+  //store the values
+  request->commArray = buf;
+
+  //sync
+  buf.syncCommData();
+
+  //call MPI
+  const int result = ::MPI_Isend(buf.commData(), count, datatype,
+    source, tag, comm, &request->request);
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Irecv() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Irecv(IdefixCommArray<T> & buf, int count, MPI_Datatype datatype, int
+    source, int tag, MPI_Comm comm, Idefix_MPI_Request<T>* request) {
+  //check
+  assert(request != nullptr);
+  assert(buf.span() == count);
+  assert(buf.span_is_contiguous());
+
+  //store the values
+  request->commArray = buf;
+
+  //call MPI
+  const int result = ::MPI_Irecv(buf.commData(), count, datatype,
+    source, tag, comm, &request->request);
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Bcast() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Bcast(IdefixCommArray<T> & buffer, int count, MPI_Datatype datatype, int root,
+  MPI_Comm comm) {
+  //check
+  assert(buffer.span() == count);
+  assert(buffer.span_is_contiguous());
+
+  //get current rank
+  int rank = -1;
+  const int status = ::MPI_Comm_rank(comm, &rank);
+  assert(status == MPI_SUCCESS);
+
+  //sync
+  if (rank == root)
+    buffer.syncCommData();
+
+  //call MPI
+  const int result = ::MPI_Bcast(buffer.commData(), count, datatype, root, comm);
+
+  //sync back
+  if (rank != root)
+    buffer.syncDeviceData();
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Bcast() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Bcast(IdefixHostArray1D<T> & buffer, int count, MPI_Datatype datatype, int root,
+  MPI_Comm comm) {
+  //check
+  assert(buffer.span() == count);
+  assert(buffer.span_is_contiguous());
+
+  //call MPI
+  const int result = ::MPI_Bcast(buffer.commData(), count, datatype, root, comm);
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Allreduce() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Allreduce(IdefixCommArray<T> & sendbuf, IdefixCommArray<T> & recvbuf, int count,
+    MPI_Datatype datatype, MPI_Op op, MPI_Comm comm) {
+  //check
+  assert(sendbuf.span() == count);
+  assert(sendbuf.span_is_contiguous());
+  assert(recvbuf.span() == count);
+  assert(recvbuf.span_is_contiguous());
+
+  sendbuf.syncCommData();
+
+  //call MPI
+  const int result = ::MPI_Allreduce(sendbuf.commData(), recvbuf.commData(), count, datatype,
+    op, comm);
+
+  //sync back
+  recvbuf.syncDeviceData();
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Allreduce() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Allreduce(void * sendbuf, IdefixHostArray3D<T> & recvbuf, int count,
+    MPI_Datatype datatype, MPI_Op op, MPI_Comm comm) {
+  //check
+  assert(recvbuf.span() == count);
+  assert(recvbuf.span_is_contiguous());
+
+  //call MPI
+  const int result = ::MPI_Allreduce(sendbuf, recvbuf.data(), count, datatype, op, comm);
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Allreduce() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Allreduce(void * sendbuf, IdefixCommArray<T> & recvbuf, int count,
+    MPI_Datatype datatype, MPI_Op op, MPI_Comm comm) {
+  //check
+  assert(recvbuf.span() == count);
+  assert(recvbuf.span_is_contiguous());
+
+  //call MPI
+  const int result = ::MPI_Allreduce(sendbuf, recvbuf.commData(), count, datatype, op, comm);
+
+  //sync back
+  recvbuf.syncDeviceData();
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_File_write_all() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_File_write_all(MPI_File fh, const IdefixHostArray4D<T> & buf, int count,
+    MPI_Datatype datatype, MPI_Status* status) {
+  //check
+  assert(buf.span() == count);
+  assert(buf.span_is_contiguous());
+
+  //call MPI
+  const int result = ::MPI_File_write_all(fh, buf.data(), count, datatype, status);
+
+  //ok
+  return result;
+}
+
+/**
+ * Overlad the MPI_Gather() function to handle an idefix communication array. It will automate the
+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T>
+int MPI_Gather(const void* sendbuf, int sendcount, MPI_Datatype
+    sendtype, std::vector<T> & recvbuf, int recvcount, MPI_Datatype recvtype,
+    int root, MPI_Comm comm) {
+  assert(recvcount <= recvbuf.size());
+  return MPI_Gather(sendbuf, sendcount, sendtype, recvbuf.data(), recvcount, recvtype, root, comm);
+}
+
+} // namespace idefix
+
+#endif // MPI_MPIVIEW_HPP_
diff --git a/src/output/slice.cpp b/src/output/slice.cpp
index 98caa5c21..cb9d8843b 100644
--- a/src/output/slice.cpp
+++ b/src/output/slice.cpp
@@ -162,7 +162,7 @@ void Slice::CheckForWrite(DataBlock &data, bool force) {
         }}}
         #ifdef WITH_MPI
           Kokkos::fence();
-          MPI_Allreduce(MPI_IN_PLACE, arrOut.data(),
+          idefix::MPI_Allreduce(MPI_IN_PLACE, arrOut,
                         arrOut.extent(0)*arrOut.extent(1)*arrOut.extent(2)*arrOut.extent(3),
                         realMPI, MPI_SUM, avgComm);
         #endif
diff --git a/src/output/vtk.cpp b/src/output/vtk.cpp
index 98e4c2577..d1d52f780 100644
--- a/src/output/vtk.cpp
+++ b/src/output/vtk.cpp
@@ -50,7 +50,7 @@ void Vtk::WriteHeaderNodes(IdfxFileHandler fvtk) {
   int size_int = static_cast<int>(size);
   MPI_SAFE_CALL(MPI_File_set_view(fvtk, this->offset, MPI_FLOAT, this->nodeView,
                                   "native", MPI_INFO_NULL));
-  MPI_SAFE_CALL(MPI_File_write_all(fvtk, node_coord.data(), size_int,
+  MPI_SAFE_CALL(idefix::MPI_File_write_all(fvtk, node_coord, size_int,
                                    MPI_FLOAT, MPI_STATUS_IGNORE));
   this->offset += sizeof(float)*(nx1+ioffset)*(nx2+joffset)*(nx3+koffset)*3;
 #else
diff --git a/src/pydefix.cpp b/src/pydefix.cpp
index f7044af71..262ada0e9 100644
--- a/src/pydefix.cpp
+++ b/src/pydefix.cpp
@@ -96,15 +96,15 @@ py::array_t<real, py::array::c_style> GatherIdefixArray(IdefixHostArray3D<real>
         #ifdef WITH_MPI
           MPI_Status status;
           // Fetch the local array size
-          MPI_Recv(np_int.data(), 3, MPI_INT, rank, 010, MPI_COMM_WORLD, &status);
-          MPI_Recv(np_tot.data(), 3, MPI_INT, rank, 011, MPI_COMM_WORLD, &status);
-          MPI_Recv(beg.data(), 3, MPI_INT, rank, 012, MPI_COMM_WORLD, &status);
-          MPI_Recv(gbeg.data(), 3, MPI_INT, rank, 013, MPI_COMM_WORLD, &status);
+          idefix::MPI_Recv(np_int, 3, MPI_INT, rank, 010, MPI_COMM_WORLD, &status);
+          idefix::MPI_Recv(np_tot, 3, MPI_INT, rank, 011, MPI_COMM_WORLD, &status);
+          idefix::MPI_Recv(beg, 3, MPI_INT, rank, 012, MPI_COMM_WORLD, &status);
+          idefix::MPI_Recv(gbeg, 3, MPI_INT, rank, 013, MPI_COMM_WORLD, &status);
 
           buf = IdefixHostArray3D<real>("pydefix::tempArray",
                                          np_tot[KDIR],np_tot[JDIR],np_tot[IDIR]);
           // Fetch data
-          MPI_Recv(buf.data(), np_tot[IDIR]*np_tot[JDIR]*np_tot[KDIR],
+          idefix::MPI_Recv(buf, np_tot[IDIR]*np_tot[JDIR]*np_tot[KDIR],
                    realMPI, rank, 014, MPI_COMM_WORLD,&status);
         #else
           IDEFIX_ERROR("Can't deal with psize>1 without MPI.");
@@ -147,11 +147,11 @@ py::array_t<real, py::array::c_style> GatherIdefixArray(IdefixHostArray3D<real>
     }
     #ifdef WITH_MPI
       // send the local array size
-      MPI_Send(np_int.data(), 3, MPI_INT, 0, 010, MPI_COMM_WORLD);
-      MPI_Send(np_tot.data(), 3, MPI_INT, 0, 011, MPI_COMM_WORLD);
-      MPI_Send(beg.data(), 3, MPI_INT, 0, 012, MPI_COMM_WORLD);
-      MPI_Send(gbeg.data(), 3, MPI_INT, 0, 013, MPI_COMM_WORLD);
-      MPI_Send(in.data(), np_tot[IDIR]*np_tot[JDIR]*np_tot[KDIR], realMPI, 0, 014, MPI_COMM_WORLD);
+      idefix::MPI_Send(np_int, 3, MPI_INT, 0, 010, MPI_COMM_WORLD);
+      idefix::MPI_Send(np_tot, 3, MPI_INT, 0, 011, MPI_COMM_WORLD);
+      idefix::MPI_Send(beg, 3, MPI_INT, 0, 012, MPI_COMM_WORLD);
+      idefix::MPI_Send(gbeg, 3, MPI_INT, 0, 013, MPI_COMM_WORLD);
+      idefix::MPI_Send(in, np_tot[IDIR]*np_tot[JDIR]*np_tot[KDIR], realMPI, 0, 014, MPI_COMM_WORLD);
     #else
       IDEFIX_ERROR("Can't deal with psize>1 without MPI.");
     #endif
@@ -159,7 +159,7 @@ py::array_t<real, py::array::c_style> GatherIdefixArray(IdefixHostArray3D<real>
   // All is transfered
   #ifdef WITH_MPI
     if(broadcast) {
-      MPI_Bcast(out.data(), out.extent(0)*out.extent(1)*out.extent(2), realMPI, 0, MPI_COMM_WORLD);
+      idefix::MPI_Bcast(out, out.extent(0)*out.extent(1)*out.extent(2), realMPI, 0, MPI_COMM_WORLD);
     }
   #endif
 
diff --git a/src/timeIntegrator.cpp b/src/timeIntegrator.cpp
index 6ffb6acfb..35f4887d3 100644
--- a/src/timeIntegrator.cpp
+++ b/src/timeIntegrator.cpp
@@ -199,7 +199,7 @@ double TimeIntegrator::ComputeBalance() {
     #ifdef WITH_MPI
       const double allowedImbalance = 20.0;
       std::vector<double> computeLogPerCore(idfx::psize);
-      MPI_Gather(&computeLastLog, 1, MPI_DOUBLE, computeLogPerCore.data(), 1, MPI_DOUBLE, 0,
+      idefix::MPI_Gather(&computeLastLog, 1, MPI_DOUBLE, computeLogPerCore, 1, MPI_DOUBLE, 0,
                   MPI_COMM_WORLD);
       computeLastLog = 0; // reset timer for all cores
       if(idfx::prank==0) {
diff --git a/src/utils/column.cpp b/src/utils/column.cpp
index 4657abc55..06d083892 100644
--- a/src/utils/column.cpp
+++ b/src/utils/column.cpp
@@ -16,6 +16,7 @@
 #include "grid.hpp"
 #include "dataBlock.hpp"
 #include "dataBlockHost.hpp"
+#include "mpiView.hpp"
 
 Column::Column(int dir, int sign, DataBlock *data)
                 : direction(dir), sign(sign) {
@@ -35,13 +36,13 @@ Column::Column(int dir, int sign, DataBlock *data)
 
   // allocate helper  array
   if(dir == IDIR) {
-    localSum = IdefixArray2D<real>("localSum",np_tot[KDIR], np_tot[JDIR]);
+    localSum = idefix::IdefixCommArray2D<real>("localSum",np_tot[KDIR], np_tot[JDIR]);
   }
   if(dir == JDIR) {
-    localSum = IdefixArray2D<real>("localSum",np_tot[KDIR], np_tot[IDIR]);
+    localSum = idefix::IdefixCommArray2D<real>("localSum",np_tot[KDIR], np_tot[IDIR]);
   }
   if(dir == KDIR) {
-    localSum = IdefixArray2D<real>("localSum",np_tot[JDIR], np_tot[IDIR]);
+    localSum = idefix::IdefixCommArray2D<real>("localSum",np_tot[JDIR], np_tot[IDIR]);
   }
   #ifdef WITH_MPI
   // Create sub-MPI communicator dedicated to scan
@@ -78,7 +79,7 @@ void Column::ComputeColumn(IdefixArray4D<real> in, const int var) {
   auto column = this->ColumnArray;
   auto dV = this->Volume;
   auto A = this->Area;
-  auto localSum = this->localSum;
+  auto localSum = this->localSum.deviceView();
 
   if(direction==IDIR) {
     // Inspired from loop.hpp
@@ -133,7 +134,7 @@ void Column::ComputeColumn(IdefixArray4D<real> in, const int var) {
         MPI_Status status;
         // Get the cumulative sum from previous processes
         Kokkos::fence();
-        MPI_Recv(localSum.data(), size, realMPI, src, 20, ColumnComm, &status);
+        idefix::MPI_Recv(this->localSum, size, realMPI, src, 20, ColumnComm, &status);
         // Add this to our cumulative sum
         idefix_for("Addsum",kb,ke,jb,je,ib,ie,
           KOKKOS_LAMBDA(int k, int j, int i) {
@@ -164,7 +165,7 @@ void Column::ComputeColumn(IdefixArray4D<real> in, const int var) {
         }
         // And send it
         Kokkos::fence();
-        MPI_Send(localSum.data(),size, realMPI, dst, 20, ColumnComm);
+        idefix::MPI_Send(this->localSum,size, realMPI, dst, 20, ColumnComm);
       } // MPIrank small enough
     #endif
     // If we need it backwards
@@ -194,7 +195,7 @@ void Column::ComputeColumn(IdefixArray4D<real> in, const int var) {
 
       #ifdef WITH_MPI
       Kokkos::fence();
-      MPI_Bcast(localSum.data(),size, realMPI, MPIsize-1, ColumnComm);
+      idefix::MPI_Bcast(this->localSum,size, realMPI, MPIsize-1, ColumnComm);
       #endif
       // All substract the local column from the full column
 
diff --git a/src/utils/column.hpp b/src/utils/column.hpp
index 43383cbe6..4d43ef0c3 100644
--- a/src/utils/column.hpp
+++ b/src/utils/column.hpp
@@ -59,7 +59,7 @@ class Column {
   IdefixArray3D<real> Area;
   IdefixArray3D<real> Volume;
 
-    IdefixArray2D<real> localSum;
+    idefix::IdefixCommArray2D<real> localSum;
   #ifdef WITH_MPI
   Mpi mpi;  // Mpi object when WITH_MPI is set
   MPI_Comm ColumnComm;
diff --git a/src/utils/lookupTable.hpp b/src/utils/lookupTable.hpp
index 72d13a838..396ff79e7 100644
--- a/src/utils/lookupTable.hpp
+++ b/src/utils/lookupTable.hpp
@@ -392,10 +392,10 @@ LookupTable<kDim>::LookupTable(std::string filename, char delimiter, bool errOOB
 
   #ifdef WITH_MPI
     // Share with the others
-    MPI_Bcast(xinHost.data(), xinHost.extent(0), realMPI, 0, MPI_COMM_WORLD);
-    MPI_Bcast(dimensionsHost.data(), dimensionsHost.extent(0), MPI_INT, 0, MPI_COMM_WORLD);
-    MPI_Bcast(offsetHost.data(), offsetHost.extent(0), MPI_INT, 0, MPI_COMM_WORLD);
-    MPI_Bcast(dataHost.data(),dataHost.extent(0), realMPI, 0, MPI_COMM_WORLD);
+    idefix::MPI_Bcast(xinHost, xinHost.extent(0), realMPI, 0, MPI_COMM_WORLD);
+    idefix::MPI_Bcast(dimensionsHost, dimensionsHost.extent(0), MPI_INT, 0, MPI_COMM_WORLD);
+    idefix::MPI_Bcast(offsetHost, offsetHost.extent(0), MPI_INT, 0, MPI_COMM_WORLD);
+    idefix::MPI_Bcast(dataHost,dataHost.extent(0), realMPI, 0, MPI_COMM_WORLD);
   #endif
 
   // Copy to target