From fec3e60ec7cba280d0ec860851c49009dae30a6b Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Wed, 8 Jul 2026 15:11:12 +0200 Subject: [PATCH 01/17] gpu-direct: handle communications when GPU DIRECT is not there. --- CMakeLists.txt | 6 ++ src/fluid/boundary/axis.cpp | 7 +- .../constrainedTransport/EMFexchange.hpp | 19 +++++- .../constrainedTransport.hpp | 48 ++++++------- src/mpi/buffer.hpp | 67 ++++++++++++++----- src/mpi/exchanger.cpp | 21 ++++-- 6 files changed, 118 insertions(+), 50 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index fad0058ba..6d7c19a96 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_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_GPU_DIRECT) + add_compile_definitions("WITH_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_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..51a487dcb 100644 --- a/src/fluid/boundary/axis.cpp +++ b/src/fluid/boundary/axis.cpp @@ -470,8 +470,10 @@ void Axis::ExchangeMPI(int side) { Kokkos::fence(); tStart = MPI_Wtime(); + bufferSend.syncCommData(); MPI_SAFE_CALL(MPI_Start(&sendRequest)); MPI_Wait(&recvRequest,&recvStatus); + bufferRecv.syncDeviceData(); idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Unpack @@ -526,7 +528,6 @@ void Axis::ExchangeMPI(int side) { } MPI_Wait(&sendRequest, &sendStatus); - idfx::mpiCallsTimer += MPI_Wtime() - tStart; @@ -593,10 +594,10 @@ 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, + MPI_SAFE_CALL(MPI_Send_init(bufferSend.commData(), bufferSend.Size(), realMPI, procSend, 650, data->mygrid->AxisComm, &sendRequest)); - MPI_SAFE_CALL(MPI_Recv_init(bufferRecv.data(), bufferRecv.Size(), realMPI, procRecv, + MPI_SAFE_CALL(MPI_Recv_init(bufferRecv.commData(), bufferRecv.Size(), realMPI, procRecv, 650, data->mygrid->AxisComm, &recvRequest)); #endif diff --git a/src/fluid/constrainedTransport/EMFexchange.hpp b/src/fluid/constrainedTransport/EMFexchange.hpp index aad123346..558ad4a1b 100644 --- a/src/fluid/constrainedTransport/EMFexchange.hpp +++ b/src/fluid/constrainedTransport/EMFexchange.hpp @@ -82,17 +82,20 @@ void ConstrainedTransport::ExchangeX1(IdefixArray3D ey, IdefixArray3 Kokkos::fence(); tStart = MPI_Wtime(); + BufferRight.syncCommData(); MPI_SAFE_CALL(MPI_Startall(2, sendRequestX1)); // Wait for buffers to be received MPI_Waitall(2,recvRequestX1,recvStatus); MPI_Waitall(2, sendRequestX1, sendStatus); - idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Unpack BufferLeft=BufferRecvX1[faceLeft]; BufferRight=BufferRecvX1[faceRight]; + BufferLeft.syncDeviceData(); + idfx::mpiCallsTimer += MPI_Wtime() - tStart; + // Erase the emf with the one coming from the left process //extend by one the end on jdir && take the ghost revc zone on i BoundingBox recvBoxEz = baseBox; @@ -169,16 +172,22 @@ void ConstrainedTransport::ExchangeX2(IdefixArray3D ex, IdefixArray3 Kokkos::fence(); tStart = MPI_Wtime(); + + BufferRight.syncCommData(); + MPI_SAFE_CALL(MPI_Startall(2, sendRequestX2)); // Wait for buffers to be received MPI_Waitall(2, recvRequestX2, recvStatus); MPI_Waitall(2, sendRequestX2, sendStatus); - idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Unpack BufferLeft=BufferRecvX2[faceLeft]; BufferRight=BufferRecvX2[faceRight]; + BufferLeft.syncDeviceData(); + + idfx::mpiCallsTimer += MPI_Wtime() - tStart; + // We average the edge emfs zones //extend by one the end on idir && take the ghost revc zone on j BoundingBox recvBoxEz = baseBox; @@ -256,16 +265,20 @@ void ConstrainedTransport::ExchangeX3(IdefixArray3D ex, IdefixArray3 Kokkos::fence(); tStart = MPI_Wtime(); + BufferRight.syncCommData(); MPI_SAFE_CALL(MPI_Startall(2, sendRequestX3)); // Wait for buffers to be received MPI_Waitall(2, recvRequestX3, recvStatus); MPI_Waitall(2, sendRequestX3, sendStatus); - idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Unpack BufferLeft=BufferRecvX3[faceLeft]; BufferRight=BufferRecvX3[faceRight]; + BufferLeft.syncDeviceData(); + + idfx::mpiCallsTimer += MPI_Wtime() - tStart; + // We average the edge emfs zones //extend by one the end on jdir && take the ghost revc zone on k BoundingBox recvBoxEx = baseBox; diff --git a/src/fluid/constrainedTransport/constrainedTransport.hpp b/src/fluid/constrainedTransport/constrainedTransport.hpp index ab83b23bc..3bec27bb4 100644 --- a/src/fluid/constrainedTransport/constrainedTransport.hpp +++ b/src/fluid/constrainedTransport/constrainedTransport.hpp @@ -331,11 +331,11 @@ ConstrainedTransport::ConstrainedTransport(Input &input, Fluid *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(MPI_Send_init(BufferSendX1[faceRight].commData(), 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(MPI_Recv_init(BufferRecvX1[faceLeft].commData(), 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::ConstrainedTransport(Input &input, Fluid *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(MPI_Send_init(BufferSendX1[faceLeft].commData(), 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(MPI_Recv_init(BufferRecvX1[faceRight].commData(), 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(MPI_Send_init(BufferSendX2[faceRight].commData(), 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(MPI_Recv_init(BufferRecvX2[faceLeft].commData(), 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(MPI_Send_init(BufferSendX2[faceLeft].commData(), 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(MPI_Recv_init(BufferRecvX2[faceRight].commData(), 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(MPI_Send_init(BufferSendX3[faceRight].commData(), 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(MPI_Recv_init(BufferRecvX3[faceLeft].commData(), 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(MPI_Send_init(BufferSendX3[faceLeft].commData(), 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(MPI_Recv_init(BufferRecvX3[faceRight].commData(), bufferSizeX3, realMPI, procRecv, + 301, mygrid->CartComm, &recvRequestX3[faceRight])); #endif #endif // WITH_MPI diff --git a/src/mpi/buffer.hpp b/src/mpi/buffer.hpp index 35164355b..e16988a46 100644 --- a/src/mpi/buffer.hpp +++ b/src/mpi/buffer.hpp @@ -17,7 +17,10 @@ using BoundingBox = std::array,3>; class Buffer { public: Buffer() = default; - explicit Buffer(size_t size): pointer{0}, array{IdefixArray1D("BufferArray",size)} {}; + explicit Buffer(size_t size) + : pointer{0} + , deviceArray{IdefixArray1D("BufferArray",size)} + , commArray{initCommArray(deviceArray)} {} // Compute the size of a bounding box static size_t ComputeBoxSize(BoundingBox box) { @@ -27,16 +30,34 @@ class Buffer { return(ninjnk); } - void* data() { - return(array.data()); + void* deviceData() { + return(deviceArray.data()); } - const IdefixArray1D & getArray(void) const { - return this->array; + void* commData() { + return(commArray.data()); + } + + void syncCommData() { + #ifndef WITH_GPU_DIRECT + //will make a memory transfer only if deviceArray is on device + Kokkos::deep_copy(this->commArray, this->deviceArray); + #endif + } + + void syncDeviceData() { + #ifndef WITH_GPU_DIRECT + //will make a memory transfer only if deviceArray is on device + Kokkos:deep_copy(this->deviceArray, this->commArray); + #endif + } + + const IdefixArray1D & getDeviceArray(void) const { + return this->deviceArray; } int Size() { - return(array.size()); + return(deviceArray.size()); } void ResetPointer() { @@ -55,7 +76,7 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->array; + auto arr = this->deviceArray; 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); @@ -79,7 +100,7 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->array; + auto arr = this->deviceArray; 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 +123,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->deviceArray; idefix_for("LoadBuffer4D_map",0,map.size(), kbeg,kend, @@ -128,7 +149,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->deviceArray; idefix_for("UnLoadBuffer3D",kbeg,kend,jbeg,jend,ibeg,iend, KOKKOS_LAMBDA (int k, int j, int i) { @@ -153,7 +174,7 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->array; + auto arr = this->deviceArray; 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 +199,7 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->array; + auto arr = this->deviceArray; 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 +225,7 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->array; + auto arr = this->deviceArray; idefix_for("UnLoadBuffer4D_map",0,map.size(), kbeg,kend, jbeg,jend, @@ -232,7 +253,7 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->array; + auto arr = this->deviceArray; idefix_for("UnLoadBuffer4D_map_sym",0,map.size(), kbeg,kend, jbeg,jend, @@ -247,10 +268,26 @@ class Buffer { this->pointer += ninjnk*map.size(); } + private: + static IdefixArray1D initCommArray(IdefixArray1D & deviceArray) { + #ifdef WITH_GPU_DIRECT + //in gpu direct, we use directly the device buffer + //if running on host, it will also just use directly the host storage. + return 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: size_t pointer; - IdefixArray1D array; + IdefixArray1D deviceArray; + //array use for MPI communications, it is deviceArray when running on Host or + //using GPU DIRECT. Otherwise it represent a copy which needs to be synced + //before communications. + IdefixArray1D commArray; }; #endif // MPI_BUFFER_HPP_ diff --git a/src/mpi/exchanger.cpp b/src/mpi/exchanger.cpp index 033ce4041..caa5c6cb8 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, + MPI_Send_init(BufferSend[faceRight].commData(), bufferSizeSend[faceRight], realMPI, procSend[faceRight], thisInstance*2, grid->CartComm, &sendRequest[faceRight]); - MPI_Recv_init(BufferRecv[faceLeft].data(), bufferSizeRecv[faceLeft], realMPI, + MPI_Recv_init(BufferRecv[faceLeft].commData(), 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, + MPI_Send_init(BufferSend[faceLeft].commData(), bufferSizeSend[faceLeft], realMPI, procSend[faceLeft],thisInstance*2+1, grid->CartComm, &sendRequest[faceLeft]); - MPI_Recv_init(BufferRecv[faceRight].data(), bufferSizeRecv[faceRight], realMPI, + MPI_Recv_init(BufferRecv[faceRight].commData(), bufferSizeRecv[faceRight], realMPI, procRecv[faceRight], thisInstance*2+1, grid->CartComm, &recvRequest[faceRight]); @@ -211,6 +211,11 @@ void Exchanger::Exchange(IdefixArray4D Vc, IdefixArray4D Vs) { Kokkos::fence(); myTimer -= MPI_Wtime(); tStart = MPI_Wtime(); + + //sync comm arrays + BufferLeft.syncCommData(); + BufferRight.syncCommData(); + #ifdef MPI_PERSISTENT MPI_Startall(2, sendRequest); // Wait for buffers to be received @@ -265,11 +270,17 @@ void Exchanger::Exchange(IdefixArray4D Vc, IdefixArray4D Vs) { #endif #endif myTimer += MPI_Wtime(); -idfx::mpiCallsTimer += MPI_Wtime() - tStart; + // Unpack BufferLeft=BufferRecv[faceLeft]; BufferRight=BufferRecv[faceRight]; +// sync +BufferLeft.syncDeviceData(); +BufferRight.syncDeviceData(); + +idfx::mpiCallsTimer += MPI_Wtime() - tStart; + BufferLeft.ResetPointer(); BufferRight.ResetPointer(); From 7967a1b21e048d0a7e05e9967d8d939443d47647 Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Thu, 9 Jul 2026 10:30:02 +0200 Subject: [PATCH 02/17] gpu-direct: Implement extensions of MPI API using views and provde the IdefixCommArrayXD<> to handle the no gpu-direct transfers. --- src/mpi/mpiView.hpp | 163 ++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 163 insertions(+) create mode 100644 src/mpi/mpiView.hpp diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp new file mode 100644 index 000000000..425046c63 --- /dev/null +++ b/src/mpi/mpiView.hpp @@ -0,0 +1,163 @@ +// *********************************************************************************** +// Idefix MHD astrophysical code +// Copyright(C) Geoffroy R. J. Lesur +// and other code contributors +// Licensed under CeCILL 2.1 License, see COPYING for more information +// *********************************************************************************** + +#ifndef MPI_MPIVIEW_HPP_ +#define MPI_MPIVIEW_HPP_ + +#include + +#include + +#include "idefix.hpp" +#include "buffer.hpp" + +/** + * Provide a specific array used for MPI communications. It embeds the standard + * array as usual, and also carry optionally a communication copy on host depending + * if we have access to WITH_MPI_GPU_DIRECT or not. + */ +template +class IdefixCommArray : public T { + public: + //inherit constructors + using T::T; + + /** + * 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, transerts the data to the device after a communication to be ready + * to use it. + */ + void syncDeviceData(void) { + Kokkos::deep_copy(*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 WITH_MPI_GPU_DIRECT + //in gpu direct, we use directly the device buffer + //if running on host, it will also just use directly the host storage. + return 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)}; +}; + +/** Wrap the idefix 1D array to use it for communication purpose. */ +template using IdefixCommArray1D = IdefixCommArray< IdefixArray1D >; +/** Wrap the idefix 2D array to use it for communication purpose. */ +template using IdefixCommArray2D = IdefixCommArray< IdefixArray2D >; +/** Wrap the idefix 3D array to use it for communication purpose. */ +template using IdefixCommArray3D = IdefixCommArray< IdefixArray3D >; +/** Wrap the idefix 4D array to use it for communication purpose. */ +template using IdefixCommArray4D = IdefixCommArray< IdefixArray4D >; + +/*template +struct IdefixMpiViewRequest +{ + Mpi + MPI_Request request; +};*/ + +/** + * Provide a wrapper of the standard MPI_Sendrecv by handling a kokkkos + * view. Use it with caution as it will allocate/deallocate a temporary + * host buffer if required (when WITH_MPI_GPU_DIRECT is disabled.). + */ +template +int idefix_MPI_View_Sendrecv(Kokkos::View sendbuf, int sendcount, MPI_Datatype sendtype, + int dest, int sendtag, Kokkos::View 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 + #ifdef WITH_MPI_GPU_DIRECT + auto sendHostArray = sendbuf; + auto recvHostArray = recvbuf; + #else + auto sendHostArray = Kokkos::create_mirror_view(sendbuf); + auto recvHostArray = Kokkos::create_mirror_view(recvbuf); + Kokkos::deep_copy(sendHostArray, sendbuf); + #endif + + //make the communication + const int result = MPI_Sendrecv(sendHostArray.data(), sendcount, sendtype, dest, sendtag, + recvHostArray.data(), recvcount, recvtype, source, recvtag, + comm, status); + + //copy back if needed + #ifndef WITH_MPI_GPU_DIRECT + Kokkos::deep_copy(recvbuf, recvHostArray); + #endif + + //ok + return result; +} + +/** + * 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 +int idefix_MPI_View_Sendrecv(IdefixCommArray sendbuf, int sendcount, MPI_Datatype sendtype, + int dest, int sendtag, IdefixCommArray 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 + #ifndef WITH_MPI_GPU_DIRECT + sendbuf.syncCommData(); + #endif + + //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 + #ifndef WITH_MPI_GPU_DIRECT + recvbuf.syncDeviceData(); + #endif + + //ok + return result; +} + +#endif // MPI_MPIVIEW_HPP_ From a15c690ea98f32e638defaa79f1a89085bc8fc3f Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Thu, 9 Jul 2026 13:52:04 +0200 Subject: [PATCH 03/17] gpu-direct: use the MPI view semantic in EMF, axis and exganger --- CMakeLists.txt | 6 +- src/fluid/boundary/axis.cpp | 18 +- src/fluid/boundary/axis.hpp | 4 +- .../constrainedTransport/EMFexchange.hpp | 44 ++--- .../constrainedTransport.hpp | 84 ++++----- .../enforceEMFBoundary.hpp | 14 +- src/mpi/buffer.hpp | 89 +++++----- src/mpi/exchanger.cpp | 28 ++- src/mpi/exchanger.hpp | 4 +- src/mpi/mpi.cpp | 2 +- src/mpi/mpi.hpp | 1 + src/mpi/mpiView.hpp | 161 +++++++++++++++++- 12 files changed, 288 insertions(+), 167 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 6d7c19a96..3870ad485 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -105,8 +105,8 @@ if(Idefix_MPI) add_subdirectory(src/mpi) endif() -if (Idefix_GPU_DIRECT) - add_compile_definitions("WITH_GPU_DIRECT") +if (Idefix_MPI_GPU_DIRECT) + add_compile_definitions("WITH_MPI_GPU_DIRECT") endif() if(Idefix_HDF5) @@ -269,7 +269,7 @@ endif() message(STATUS " MPI: ${Idefix_MPI}") message(STATUS " HDF5: ${Idefix_HDF5}") message(STATUS " Python: ${Idefix_PYTHON}") -message(STATUS " GPU direct: ${Idefix_GPU_DIRECT}") +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 51a487dcb..4aae0dfd8 100644 --- a/src/fluid/boundary/axis.cpp +++ b/src/fluid/boundary/axis.cpp @@ -408,7 +408,7 @@ void Axis::ExchangeMPI(int side) { MPI_Status recvStatus; double tStart = MPI_Wtime(); - MPI_SAFE_CALL(MPI_Start(&recvRequest)); + MPI_SAFE_CALL(idefix_MPI_View_Start(&recvRequest)); idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Coordinates of the ghost region which needs to be transfered @@ -470,10 +470,8 @@ void Axis::ExchangeMPI(int side) { Kokkos::fence(); tStart = MPI_Wtime(); - bufferSend.syncCommData(); - MPI_SAFE_CALL(MPI_Start(&sendRequest)); - MPI_Wait(&recvRequest,&recvStatus); - bufferRecv.syncDeviceData(); + MPI_SAFE_CALL(idefix_MPI_View_Start(&sendRequest)); + idefix_MPI_View_Wait(&recvRequest,&recvStatus); idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Unpack @@ -527,7 +525,7 @@ void Axis::ExchangeMPI(int side) { } // MHD } - MPI_Wait(&sendRequest, &sendStatus); + idefix_MPI_View_Wait(&sendRequest, &sendStatus); idfx::mpiCallsTimer += MPI_Wtime() - tStart; @@ -594,11 +592,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.commData(), bufferSend.Size(), realMPI, procSend, - 650, data->mygrid->AxisComm, &sendRequest)); + MPI_SAFE_CALL(idefix_MPI_View_Send_init(bufferSend.commView(), bufferSend.Size(), + realMPI, procSend, 650, data->mygrid->AxisComm, &sendRequest)); - MPI_SAFE_CALL(MPI_Recv_init(bufferRecv.commData(), bufferRecv.Size(), realMPI, procRecv, - 650, data->mygrid->AxisComm, &recvRequest)); + MPI_SAFE_CALL(idefix_MPI_View_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..e94b10366 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 sendRequest; + Idefix_MPI_Request_1D recvRequest; Buffer bufferSend; Buffer bufferRecv; diff --git a/src/fluid/constrainedTransport/EMFexchange.hpp b/src/fluid/constrainedTransport/EMFexchange.hpp index 558ad4a1b..a0ae39931 100644 --- a/src/fluid/constrainedTransport/EMFexchange.hpp +++ b/src/fluid/constrainedTransport/EMFexchange.hpp @@ -41,7 +41,7 @@ void ConstrainedTransport::ExchangeX1(IdefixArray3D ey, IdefixArray3 MPI_Status recvStatus[2]; double tStart = MPI_Wtime(); - MPI_SAFE_CALL(MPI_Startall(2, recvRequestX1)); + MPI_SAFE_CALL(idefix_MPI_View_Startall(2, recvRequestX1)); idfx::mpiCallsTimer += MPI_Wtime() - tStart; BoundaryType lbound = data->lbound[IDIR]; @@ -82,20 +82,16 @@ void ConstrainedTransport::ExchangeX1(IdefixArray3D ey, IdefixArray3 Kokkos::fence(); tStart = MPI_Wtime(); - BufferRight.syncCommData(); - MPI_SAFE_CALL(MPI_Startall(2, sendRequestX1)); + MPI_SAFE_CALL(idefix_MPI_View_Startall(2, sendRequestX1)); // Wait for buffers to be received - - MPI_Waitall(2,recvRequestX1,recvStatus); - MPI_Waitall(2, sendRequestX1, sendStatus); + idefix_MPI_View_Waitall(2,recvRequestX1,recvStatus); + idefix_MPI_View_Waitall(2, sendRequestX1, sendStatus); + idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Unpack BufferLeft=BufferRecvX1[faceLeft]; BufferRight=BufferRecvX1[faceRight]; - BufferLeft.syncDeviceData(); - idfx::mpiCallsTimer += MPI_Wtime() - tStart; - // Erase the emf with the one coming from the left process //extend by one the end on jdir && take the ghost revc zone on i BoundingBox recvBoxEz = baseBox; @@ -133,7 +129,7 @@ void ConstrainedTransport::ExchangeX2(IdefixArray3D 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_View_Startall(2, recvRequestX2)); idfx::mpiCallsTimer += MPI_Wtime() - tStart; BoundaryType lbound = data->lbound[JDIR]; @@ -172,22 +168,16 @@ void ConstrainedTransport::ExchangeX2(IdefixArray3D ex, IdefixArray3 Kokkos::fence(); tStart = MPI_Wtime(); - - BufferRight.syncCommData(); - - MPI_SAFE_CALL(MPI_Startall(2, sendRequestX2)); + MPI_SAFE_CALL(idefix_MPI_View_Startall(2, sendRequestX2)); // Wait for buffers to be received - MPI_Waitall(2, recvRequestX2, recvStatus); - MPI_Waitall(2, sendRequestX2, sendStatus); + idefix_MPI_View_Waitall(2, recvRequestX2, recvStatus); + idefix_MPI_View_Waitall(2, sendRequestX2, sendStatus); + idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Unpack BufferLeft=BufferRecvX2[faceLeft]; BufferRight=BufferRecvX2[faceRight]; - BufferLeft.syncDeviceData(); - - idfx::mpiCallsTimer += MPI_Wtime() - tStart; - // We average the edge emfs zones //extend by one the end on idir && take the ghost revc zone on j BoundingBox recvBoxEz = baseBox; @@ -225,7 +215,7 @@ void ConstrainedTransport::ExchangeX3(IdefixArray3D 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_View_Startall(2, recvRequestX3)); idfx::mpiCallsTimer += MPI_Wtime() - tStart; BoundaryType lbound = data->lbound[KDIR]; @@ -265,20 +255,16 @@ void ConstrainedTransport::ExchangeX3(IdefixArray3D ex, IdefixArray3 Kokkos::fence(); tStart = MPI_Wtime(); - BufferRight.syncCommData(); - MPI_SAFE_CALL(MPI_Startall(2, sendRequestX3)); + MPI_SAFE_CALL(idefix_MPI_View_Startall(2, sendRequestX3)); // Wait for buffers to be received - MPI_Waitall(2, recvRequestX3, recvStatus); - MPI_Waitall(2, sendRequestX3, sendStatus); + idefix_MPI_View_Waitall(2, recvRequestX3, recvStatus); + idefix_MPI_View_Waitall(2, sendRequestX3, sendStatus); + idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Unpack BufferLeft=BufferRecvX3[faceLeft]; BufferRight=BufferRecvX3[faceRight]; - BufferLeft.syncDeviceData(); - - idfx::mpiCallsTimer += MPI_Wtime() - tStart; - // We average the edge emfs zones //extend by one the end on jdir && take the ghost revc zone on k BoundingBox recvBoxEx = baseBox; diff --git a/src/fluid/constrainedTransport/constrainedTransport.hpp b/src/fluid/constrainedTransport/constrainedTransport.hpp index 3bec27bb4..2fbfa103f 100644 --- a/src/fluid/constrainedTransport/constrainedTransport.hpp +++ b/src/fluid/constrainedTransport/constrainedTransport.hpp @@ -68,9 +68,9 @@ class ConstrainedTransport { IdefixArray3D Ex3; // Helper arrays for shearing box - IdefixArray2D sbEyL; - IdefixArray2D sbEyR; - IdefixArray2D sbEyRL; + IdefixCommArray2D sbEyL; + IdefixCommArray2D sbEyR; + IdefixCommArray2D 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 sendRequestX1[2]; + Idefix_MPI_Request_1D sendRequestX2[2]; + Idefix_MPI_Request_1D sendRequestX3[2]; + Idefix_MPI_Request_1D recvRequestX1[2]; + Idefix_MPI_Request_1D recvRequestX2[2]; + Idefix_MPI_Request_1D recvRequestX3[2]; Kokkos::Timer timer; // Internal MPI timer #endif @@ -188,9 +188,9 @@ ConstrainedTransport::ConstrainedTransport(Input &input, Fluid *hydr // Allocate shearing box arrays if(hydro->haveShearingBox == true) { - sbEyL = IdefixArray2D("EMF_sbEyL", data->np_tot[KDIR], data->np_tot[JDIR]); - sbEyR = IdefixArray2D("EMF_sbEyR", data->np_tot[KDIR], data->np_tot[JDIR]); - sbEyRL = IdefixArray2D("EMF_sbEyRL", data->np_tot[KDIR], data->np_tot[JDIR]); + sbEyL = IdefixCommArray2D("EMF_sbEyL", data->np_tot[KDIR], data->np_tot[JDIR]); + sbEyR = IdefixCommArray2D("EMF_sbEyR", data->np_tot[KDIR], data->np_tot[JDIR]); + sbEyRL = IdefixCommArray2D("EMF_sbEyRL", data->np_tot[KDIR], data->np_tot[JDIR]); } D_EXPAND( ez = IdefixArray3D("EMF_ez", @@ -331,11 +331,11 @@ ConstrainedTransport::ConstrainedTransport(Input &input, Fluid *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].commData(), bufferSizeX1, realMPI, procSend, - 100, mygrid->CartComm, &sendRequestX1[faceRight])); + MPI_SAFE_CALL(idefix_MPI_View_Send_init(BufferSendX1[faceRight].commView(), bufferSizeX1, + realMPI, procSend, 100, mygrid->CartComm, &sendRequestX1[faceRight])); - MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX1[faceLeft].commData(), bufferSizeX1, realMPI, procRecv, - 100, mygrid->CartComm, &recvRequestX1[faceLeft])); + MPI_SAFE_CALL(idefix_MPI_View_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::ConstrainedTransport(Input &input, Fluid *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].commData(), bufferSizeX1, realMPI, procSend, - 101, mygrid->CartComm, &sendRequestX1[faceLeft])); + MPI_SAFE_CALL(idefix_MPI_View_Send_init(BufferSendX1[faceLeft].commView(), bufferSizeX1, + realMPI, procSend, 101, mygrid->CartComm, &sendRequestX1[faceLeft])); - MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX1[faceRight].commData(), bufferSizeX1, realMPI, procRecv, - 101, mygrid->CartComm, &recvRequestX1[faceRight])); + MPI_SAFE_CALL(idefix_MPI_View_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].commData(), bufferSizeX2, realMPI, procSend, - 200, mygrid->CartComm, &sendRequestX2[faceRight])); + MPI_SAFE_CALL(idefix_MPI_View_Send_init(BufferSendX2[faceRight].commView(), bufferSizeX2, + realMPI, procSend, 200, mygrid->CartComm, &sendRequestX2[faceRight])); - MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX2[faceLeft].commData(), bufferSizeX2, realMPI, procRecv, - 200, mygrid->CartComm, &recvRequestX2[faceLeft])); + MPI_SAFE_CALL(idefix_MPI_View_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].commData(), bufferSizeX2, realMPI, procSend, - 201, mygrid->CartComm, &sendRequestX2[faceLeft])); + MPI_SAFE_CALL(idefix_MPI_View_Send_init(BufferSendX2[faceLeft].commView(), bufferSizeX2, + realMPI, procSend, 201, mygrid->CartComm, &sendRequestX2[faceLeft])); - MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX2[faceRight].commData(), bufferSizeX2, realMPI, procRecv, - 201, mygrid->CartComm, &recvRequestX2[faceRight])); + MPI_SAFE_CALL(idefix_MPI_View_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].commData(), bufferSizeX3, realMPI, procSend, - 300, mygrid->CartComm, &sendRequestX3[faceRight])); + MPI_SAFE_CALL(idefix_MPI_View_Send_init(BufferSendX3[faceRight].commView(), bufferSizeX3, + realMPI, procSend, 300, mygrid->CartComm, &sendRequestX3[faceRight])); - MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX3[faceLeft].commData(), bufferSizeX3, realMPI, procRecv, - 300, mygrid->CartComm, &recvRequestX3[faceLeft])); + MPI_SAFE_CALL(idefix_MPI_View_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].commData(), bufferSizeX3, realMPI, procSend, - 301, mygrid->CartComm, &sendRequestX3[faceLeft])); + MPI_SAFE_CALL(idefix_MPI_View_Send_init(BufferSendX3[faceLeft].commView(), bufferSizeX3, + realMPI, procSend, 301, mygrid->CartComm, &sendRequestX3[faceLeft])); - MPI_SAFE_CALL(MPI_Recv_init(BufferRecvX3[faceRight].commData(), bufferSizeX3, realMPI, procRecv, - 301, mygrid->CartComm, &recvRequestX3[faceRight])); + MPI_SAFE_CALL(idefix_MPI_View_Recv_init(BufferRecvX3[faceRight].commView(), bufferSizeX3, + realMPI, procRecv, 301, mygrid->CartComm, &recvRequestX3[faceRight])); #endif #endif // WITH_MPI @@ -409,17 +409,17 @@ ConstrainedTransport::~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_View_Request_free( &sendRequestX1[i]); + idefix_MPI_View_Request_free( &recvRequestX1[i]); #if DIMENSIONS >= 2 - MPI_Request_free( &sendRequestX2[i]); - MPI_Request_free( &recvRequestX2[i]); + idefix_MPI_View_Request_free( &sendRequestX2[i]); + idefix_MPI_View_Request_free( &recvRequestX2[i]); #endif #if DIMENSIONS == 3 - MPI_Request_free( &sendRequestX3[i]); - MPI_Request_free( &recvRequestX3[i]); + idefix_MPI_View_Request_free( &sendRequestX3[i]); + idefix_MPI_View_Request_free( &recvRequestX3[i]); #endif } #endif diff --git a/src/fluid/constrainedTransport/enforceEMFBoundary.hpp b/src/fluid/constrainedTransport/enforceEMFBoundary.hpp index 9d3d6913f..0b19e6cd0 100644 --- a/src/fluid/constrainedTransport/enforceEMFBoundary.hpp +++ b/src/fluid/constrainedTransport/enforceEMFBoundary.hpp @@ -108,9 +108,9 @@ void ConstrainedTransport::SymmetrizeEMFShearingBox() { idfx::pushRegion("Emf::EnforceEMFBoundary"); #if MHD == YES - IdefixArray2D sbEyL = this->sbEyL; - IdefixArray2D sbEyR = this->sbEyR; - IdefixArray2D sbEyRL = this->sbEyRL; + auto sbEyL = this->sbEyL; + auto sbEyR = this->sbEyR; + auto sbEyRL = this->sbEyRL; IdefixArray3D ey = this->ey; @@ -142,16 +142,16 @@ void ConstrainedTransport::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_View_Sendrecv(sbEyL, size, realMPI, procLeft, 2001, + 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_View_Sendrecv(sbEyR, size, realMPI, procRight, 2002, + sbEyL, size, realMPI, procRight, 2001, data->mygrid->CartComm, &status ); } } diff --git a/src/mpi/buffer.hpp b/src/mpi/buffer.hpp index e16988a46..fd980f61c 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,3>; @@ -19,8 +20,7 @@ class Buffer { Buffer() = default; explicit Buffer(size_t size) : pointer{0} - , deviceArray{IdefixArray1D("BufferArray",size)} - , commArray{initCommArray(deviceArray)} {} + , array{IdefixArray1D("BufferArray",size)} {} // Compute the size of a bounding box static size_t ComputeBoxSize(BoundingBox box) { @@ -31,33 +31,19 @@ class Buffer { } void* deviceData() { - return(deviceArray.data()); + return(array.data()); } - void* commData() { - return(commArray.data()); + IdefixArray1D & deviceView(void) { + return this->array; } - void syncCommData() { - #ifndef WITH_GPU_DIRECT - //will make a memory transfer only if deviceArray is on device - Kokkos::deep_copy(this->commArray, this->deviceArray); - #endif - } - - void syncDeviceData() { - #ifndef WITH_GPU_DIRECT - //will make a memory transfer only if deviceArray is on device - Kokkos:deep_copy(this->deviceArray, this->commArray); - #endif - } - - const IdefixArray1D & getDeviceArray(void) const { - return this->deviceArray; + IdefixCommArray1D & commView(void) { + return this->array; } int Size() { - return(deviceArray.size()); + return(array.size()); } void ResetPointer() { @@ -76,7 +62,7 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->deviceArray; + auto arr = this->array; 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); @@ -100,7 +86,7 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->deviceArray; + auto arr = this->array; 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); @@ -123,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->deviceArray; + auto arr = this->array; idefix_for("LoadBuffer4D_map",0,map.size(), kbeg,kend, @@ -149,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->deviceArray; + auto arr = this->array; idefix_for("UnLoadBuffer3D",kbeg,kend,jbeg,jend,ibeg,iend, KOKKOS_LAMBDA (int k, int j, int i) { @@ -174,7 +160,31 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->deviceArray; + auto arr = this->array; + 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 ); + }); + + // Update pointer + this->pointer += ninjnk; + } + + void UnpackAveraged(IdefixArray4D& out, + const int var, + BoundingBox box) { + const int ni = box[IDIR][1]-box[IDIR][0]; + const int ninj = (box[JDIR][1]-box[JDIR][0])*ni; + const int ninjnk = (box[KDIR][1]-box[KDIR][0])*ninj; + const int ibeg = box[IDIR][0]; + const int jbeg = box[JDIR][0]; + const int kbeg = box[KDIR][0]; + const int iend = box[IDIR][1]; + const int jend = box[JDIR][1]; + const int kend = box[KDIR][1]; + const int offset = this->pointer; + + auto arr = this->array; 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 ); @@ -199,7 +209,7 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->deviceArray; + auto arr = this->array; 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; @@ -225,7 +235,7 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->deviceArray; + auto arr = this->array; idefix_for("UnLoadBuffer4D_map",0,map.size(), kbeg,kend, jbeg,jend, @@ -253,7 +263,7 @@ class Buffer { const int kend = box[KDIR][1]; const int offset = this->pointer; - auto arr = this->deviceArray; + auto arr = this->array; idefix_for("UnLoadBuffer4D_map_sym",0,map.size(), kbeg,kend, jbeg,jend, @@ -268,26 +278,9 @@ class Buffer { this->pointer += ninjnk*map.size(); } - private: - static IdefixArray1D initCommArray(IdefixArray1D & deviceArray) { - #ifdef WITH_GPU_DIRECT - //in gpu direct, we use directly the device buffer - //if running on host, it will also just use directly the host storage. - return 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: size_t pointer; - IdefixArray1D deviceArray; - //array use for MPI communications, it is deviceArray when running on Host or - //using GPU DIRECT. Otherwise it represent a copy which needs to be synced - //before communications. - IdefixArray1D commArray; + IdefixCommArray1D array; }; #endif // MPI_BUFFER_HPP_ diff --git a/src/mpi/exchanger.cpp b/src/mpi/exchanger.cpp index caa5c6cb8..cbd5044fc 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].commData(), bufferSizeSend[faceRight], realMPI, + idefix_MPI_View_Send_init(BufferSend[faceRight].commView(), bufferSizeSend[faceRight], realMPI, procSend[faceRight], thisInstance*2, grid->CartComm, &sendRequest[faceRight]); - MPI_Recv_init(BufferRecv[faceLeft].commData(), bufferSizeRecv[faceLeft], realMPI, + idefix_MPI_View_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].commData(), bufferSizeSend[faceLeft], realMPI, + idefix_MPI_View_Send_init(BufferSend[faceLeft].commView(), bufferSizeSend[faceLeft], realMPI, procSend[faceLeft],thisInstance*2+1, grid->CartComm, &sendRequest[faceLeft]); - MPI_Recv_init(BufferRecv[faceRight].commData(), bufferSizeRecv[faceRight], realMPI, + idefix_MPI_View_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_View_Request_free( &sendRequest[i]); + idefix_MPI_View_Request_free( &recvRequest[i]); } #endif isInitialized = false; @@ -189,7 +189,7 @@ void Exchanger::Exchange(IdefixArray4D Vc, IdefixArray4D Vs) { MPI_Status sendStatus[2]; MPI_Status recvStatus[2]; - MPI_Startall(2, recvRequest); + idefix_MPI_View_Startall(2, recvRequest); idfx::mpiCallsTimer += MPI_Wtime() - tStart; #endif myTimer += MPI_Wtime(); @@ -212,14 +212,10 @@ void Exchanger::Exchange(IdefixArray4D Vc, IdefixArray4D Vs) { myTimer -= MPI_Wtime(); tStart = MPI_Wtime(); - //sync comm arrays - BufferLeft.syncCommData(); - BufferRight.syncCommData(); - #ifdef MPI_PERSISTENT - MPI_Startall(2, sendRequest); + idefix_MPI_View_Startall(2, sendRequest); // Wait for buffers to be received - MPI_Waitall(2,recvRequest,recvStatus); + idefix_MPI_View_Waitall(2, recvRequest, recvStatus); #else @@ -275,10 +271,6 @@ myTimer += MPI_Wtime(); BufferLeft=BufferRecv[faceLeft]; BufferRight=BufferRecv[faceRight]; -// sync -BufferLeft.syncDeviceData(); -BufferRight.syncDeviceData(); - idfx::mpiCallsTimer += MPI_Wtime() - tStart; BufferLeft.ResetPointer(); @@ -307,7 +299,7 @@ myTimer -= MPI_Wtime(); #endif #ifdef MPI_PERSISTENT - MPI_Waitall(2, sendRequest, sendStatus); + idefix_MPI_View_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..122796f82 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 sendRequest[2]; + Idefix_MPI_Request_1D recvRequest[2]; Grid *grid; }; diff --git a/src/mpi/mpi.cpp b/src/mpi/mpi.cpp index d0759199a..cb436ec92 100644 --- a/src/mpi/mpi.cpp +++ b/src/mpi/mpi.cpp @@ -123,7 +123,7 @@ void Mpi::ExchangeX3(IdefixArray4D Vc, IdefixArray4D 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 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 index 425046c63..d4ba44f12 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -34,6 +34,14 @@ class IdefixCommArray : public T { 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. @@ -42,6 +50,14 @@ class IdefixCommArray : public T { 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. @@ -80,12 +96,20 @@ template using IdefixCommArray3D = IdefixCommArray< IdefixArray3D > /** Wrap the idefix 4D array to use it for communication purpose. */ template using IdefixCommArray4D = IdefixCommArray< IdefixArray4D >; -/*template -struct IdefixMpiViewRequest -{ - Mpi +template +struct Idefix_MPI_Request { + IdefixCommArray commArray; MPI_Request request; -};*/ +}; + +/** Wrap the idefix 1D array to use it for communication purpose. */ +template using Idefix_MPI_Request_1D = Idefix_MPI_Request< IdefixArray1D >; +/** Wrap the idefix 2D array to use it for communication purpose. */ +template using Idefix_MPI_Request_2D = Idefix_MPI_Request< IdefixArray2D >; +/** Wrap the idefix 3D array to use it for communication purpose. */ +template using Idefix_MPI_Request_3D = Idefix_MPI_Request< IdefixArray3D >; +/** Wrap the idefix 4D array to use it for communication purpose. */ +template using Idefix_MPI_Request_4D = Idefix_MPI_Request< IdefixArray4D >; /** * Provide a wrapper of the standard MPI_Sendrecv by handling a kokkkos @@ -160,4 +184,131 @@ int idefix_MPI_View_Sendrecv(IdefixCommArray sendbuf, int sendcount, MPI_Data return result; } +template +int idefix_MPI_View_Send_init(IdefixCommArray & buf, int count, MPI_Datatype datatype, int dest, + int tag, MPI_Comm comm, Idefix_MPI_Request *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; +} + +template +int idefix_MPI_View_Recv_init(IdefixCommArray & buf, int count, MPI_Datatype datatype, int + source, int tag, MPI_Comm comm, Idefix_MPI_Request* 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; +} + +template +int idefix_MPI_View_Request_free(Idefix_MPI_Request* request) { + //check + assert(request != nullptr); + + //call MPI + return MPI_Request_free(&request->request); +} + +template +int idefix_MPI_View_Startall(int count, Idefix_MPI_Request* 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; +} + +template +int idefix_MPI_View_Waitall(int count, Idefix_MPI_Request* 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; +} + +template +int idefix_MPI_View_Start(Idefix_MPI_Request* request) { + //check + assert(request != nullptr); + + //sync the copies in async mode + request->commArray.syncCommData(); + + //call MPI + return MPI_Start(&request->request); +} + +template +int idefix_MPI_View_Wait(Idefix_MPI_Request* 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; +} + #endif // MPI_MPIVIEW_HPP_ From 4742f742dd85bc9b30fc3618c1f4955450abe15e Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Thu, 9 Jul 2026 16:51:54 +0200 Subject: [PATCH 04/17] gpu-direct: rename the API to use a namespace and fallback on the original function names --- src/fluid/boundary/axis.cpp | 12 ++-- src/fluid/boundary/axis.hpp | 4 +- .../constrainedTransport/EMFexchange.hpp | 24 ++++---- .../constrainedTransport.hpp | 60 +++++++++---------- .../enforceEMFBoundary.hpp | 4 +- src/mpi/buffer.hpp | 4 +- src/mpi/exchanger.cpp | 44 +++++++------- src/mpi/exchanger.hpp | 4 +- src/mpi/mpiView.hpp | 42 +++++++------ 9 files changed, 101 insertions(+), 97 deletions(-) diff --git a/src/fluid/boundary/axis.cpp b/src/fluid/boundary/axis.cpp index 4aae0dfd8..9303747a4 100644 --- a/src/fluid/boundary/axis.cpp +++ b/src/fluid/boundary/axis.cpp @@ -408,7 +408,7 @@ void Axis::ExchangeMPI(int side) { MPI_Status recvStatus; double tStart = MPI_Wtime(); - MPI_SAFE_CALL(idefix_MPI_View_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 +470,8 @@ void Axis::ExchangeMPI(int side) { Kokkos::fence(); tStart = MPI_Wtime(); - MPI_SAFE_CALL(idefix_MPI_View_Start(&sendRequest)); - idefix_MPI_View_Wait(&recvRequest,&recvStatus); + MPI_SAFE_CALL(idefix::MPI_Start(&sendRequest)); + idefix::MPI_Wait(&recvRequest,&recvStatus); idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Unpack @@ -525,7 +525,7 @@ void Axis::ExchangeMPI(int side) { } // MHD } - idefix_MPI_View_Wait(&sendRequest, &sendStatus); + idefix::MPI_Wait(&sendRequest, &sendStatus); idfx::mpiCallsTimer += MPI_Wtime() - tStart; @@ -592,10 +592,10 @@ void Axis::InitMPI() { MPI_SAFE_CALL(MPI_Cart_shift(data->mygrid->AxisComm,0,data->mygrid->nproc[KDIR]/2, &procRecv,&procSend )); - MPI_SAFE_CALL(idefix_MPI_View_Send_init(bufferSend.commView(), bufferSend.Size(), + MPI_SAFE_CALL(idefix::MPI_Send_init(bufferSend.commView(), bufferSend.Size(), realMPI, procSend, 650, data->mygrid->AxisComm, &sendRequest)); - MPI_SAFE_CALL(idefix_MPI_View_Recv_init(bufferRecv.commView(), bufferRecv.Size(), + MPI_SAFE_CALL(idefix::MPI_Recv_init(bufferRecv.commView(), bufferRecv.Size(), realMPI, procRecv, 650, data->mygrid->AxisComm, &recvRequest)); #endif diff --git a/src/fluid/boundary/axis.hpp b/src/fluid/boundary/axis.hpp index e94b10366..039519228 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 - Idefix_MPI_Request_1D sendRequest; - Idefix_MPI_Request_1D recvRequest; + idefix::MPI_Request_1D sendRequest; + idefix::MPI_Request_1D recvRequest; Buffer bufferSend; Buffer bufferRecv; diff --git a/src/fluid/constrainedTransport/EMFexchange.hpp b/src/fluid/constrainedTransport/EMFexchange.hpp index a0ae39931..cea40a075 100644 --- a/src/fluid/constrainedTransport/EMFexchange.hpp +++ b/src/fluid/constrainedTransport/EMFexchange.hpp @@ -41,7 +41,7 @@ void ConstrainedTransport::ExchangeX1(IdefixArray3D ey, IdefixArray3 MPI_Status recvStatus[2]; double tStart = MPI_Wtime(); - MPI_SAFE_CALL(idefix_MPI_View_Startall(2, recvRequestX1)); + MPI_SAFE_CALL(idefix::MPI_Startall(2, recvRequestX1)); idfx::mpiCallsTimer += MPI_Wtime() - tStart; BoundaryType lbound = data->lbound[IDIR]; @@ -82,10 +82,10 @@ void ConstrainedTransport::ExchangeX1(IdefixArray3D ey, IdefixArray3 Kokkos::fence(); tStart = MPI_Wtime(); - MPI_SAFE_CALL(idefix_MPI_View_Startall(2, sendRequestX1)); + MPI_SAFE_CALL(idefix::MPI_Startall(2, sendRequestX1)); // Wait for buffers to be received - idefix_MPI_View_Waitall(2,recvRequestX1,recvStatus); - idefix_MPI_View_Waitall(2, sendRequestX1, sendStatus); + idefix::MPI_Waitall(2,recvRequestX1,recvStatus); + idefix::MPI_Waitall(2, sendRequestX1, sendStatus); idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Unpack @@ -129,7 +129,7 @@ void ConstrainedTransport::ExchangeX2(IdefixArray3D ex, IdefixArray3 double tStart = MPI_Wtime(); MPI_Status sendStatus[2]; MPI_Status recvStatus[2]; - MPI_SAFE_CALL(idefix_MPI_View_Startall(2, recvRequestX2)); + MPI_SAFE_CALL(idefix::MPI_Startall(2, recvRequestX2)); idfx::mpiCallsTimer += MPI_Wtime() - tStart; BoundaryType lbound = data->lbound[JDIR]; @@ -168,10 +168,10 @@ void ConstrainedTransport::ExchangeX2(IdefixArray3D ex, IdefixArray3 Kokkos::fence(); tStart = MPI_Wtime(); - MPI_SAFE_CALL(idefix_MPI_View_Startall(2, sendRequestX2)); + MPI_SAFE_CALL(idefix::MPI_Startall(2, sendRequestX2)); // Wait for buffers to be received - idefix_MPI_View_Waitall(2, recvRequestX2, recvStatus); - idefix_MPI_View_Waitall(2, sendRequestX2, sendStatus); + idefix::MPI_Waitall(2, recvRequestX2, recvStatus); + idefix::MPI_Waitall(2, sendRequestX2, sendStatus); idfx::mpiCallsTimer += MPI_Wtime() - tStart; // Unpack @@ -215,7 +215,7 @@ void ConstrainedTransport::ExchangeX3(IdefixArray3D ex, IdefixArray3 double tStart = MPI_Wtime(); MPI_Status sendStatus[2]; MPI_Status recvStatus[2]; - MPI_SAFE_CALL(idefix_MPI_View_Startall(2, recvRequestX3)); + MPI_SAFE_CALL(idefix::MPI_Startall(2, recvRequestX3)); idfx::mpiCallsTimer += MPI_Wtime() - tStart; BoundaryType lbound = data->lbound[KDIR]; @@ -255,10 +255,10 @@ void ConstrainedTransport::ExchangeX3(IdefixArray3D ex, IdefixArray3 Kokkos::fence(); tStart = MPI_Wtime(); - MPI_SAFE_CALL(idefix_MPI_View_Startall(2, sendRequestX3)); + MPI_SAFE_CALL(idefix::MPI_Startall(2, sendRequestX3)); // Wait for buffers to be received - idefix_MPI_View_Waitall(2, recvRequestX3, recvStatus); - idefix_MPI_View_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 2fbfa103f..7734bf5de 100644 --- a/src/fluid/constrainedTransport/constrainedTransport.hpp +++ b/src/fluid/constrainedTransport/constrainedTransport.hpp @@ -68,9 +68,9 @@ class ConstrainedTransport { IdefixArray3D Ex3; // Helper arrays for shearing box - IdefixCommArray2D sbEyL; - IdefixCommArray2D sbEyR; - IdefixCommArray2D sbEyRL; + idefix::IdefixCommArray2D sbEyL; + idefix::IdefixCommArray2D sbEyR; + idefix::IdefixCommArray2D sbEyRL; // Range of existence @@ -137,12 +137,12 @@ class ConstrainedTransport { int bufferSizeX3; // Requests for MPI persistent communications - Idefix_MPI_Request_1D sendRequestX1[2]; - Idefix_MPI_Request_1D sendRequestX2[2]; - Idefix_MPI_Request_1D sendRequestX3[2]; - Idefix_MPI_Request_1D recvRequestX1[2]; - Idefix_MPI_Request_1D recvRequestX2[2]; - Idefix_MPI_Request_1D recvRequestX3[2]; + idefix::MPI_Request_1D sendRequestX1[2]; + idefix::MPI_Request_1D sendRequestX2[2]; + idefix::MPI_Request_1D sendRequestX3[2]; + idefix::MPI_Request_1D recvRequestX1[2]; + idefix::MPI_Request_1D recvRequestX2[2]; + idefix::MPI_Request_1D recvRequestX3[2]; Kokkos::Timer timer; // Internal MPI timer #endif @@ -188,9 +188,9 @@ ConstrainedTransport::ConstrainedTransport(Input &input, Fluid *hydr // Allocate shearing box arrays if(hydro->haveShearingBox == true) { - sbEyL = IdefixCommArray2D("EMF_sbEyL", data->np_tot[KDIR], data->np_tot[JDIR]); - sbEyR = IdefixCommArray2D("EMF_sbEyR", data->np_tot[KDIR], data->np_tot[JDIR]); - sbEyRL = IdefixCommArray2D("EMF_sbEyRL", data->np_tot[KDIR], data->np_tot[JDIR]); + sbEyL = idefix::IdefixCommArray2D("EMF_sbEyL", data->np_tot[KDIR], data->np_tot[JDIR]); + sbEyR = idefix::IdefixCommArray2D("EMF_sbEyR", data->np_tot[KDIR], data->np_tot[JDIR]); + sbEyRL = idefix::IdefixCommArray2D("EMF_sbEyRL", data->np_tot[KDIR], data->np_tot[JDIR]); } D_EXPAND( ez = IdefixArray3D("EMF_ez", @@ -331,10 +331,10 @@ ConstrainedTransport::ConstrainedTransport(Input &input, Fluid *hydr if(data->rbound[IDIR] == shearingbox ) procSend = MPI_PROC_NULL; if(data->lbound[IDIR] == shearingbox ) procRecv = MPI_PROC_NULL; - MPI_SAFE_CALL(idefix_MPI_View_Send_init(BufferSendX1[faceRight].commView(), bufferSizeX1, + MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX1[faceRight].commView(), bufferSizeX1, realMPI, procSend, 100, mygrid->CartComm, &sendRequestX1[faceRight])); - MPI_SAFE_CALL(idefix_MPI_View_Recv_init(BufferRecvX1[faceLeft].commView(), bufferSizeX1, + MPI_SAFE_CALL(idefix::MPI_Recv_init(BufferRecvX1[faceLeft].commView(), bufferSizeX1, realMPI, procRecv, 100, mygrid->CartComm, &recvRequestX1[faceLeft])); // Send to the left @@ -344,30 +344,30 @@ ConstrainedTransport::ConstrainedTransport(Input &input, Fluid *hydr if(data->lbound[IDIR] == shearingbox ) procSend = MPI_PROC_NULL; if(data->rbound[IDIR] == shearingbox ) procRecv = MPI_PROC_NULL; - MPI_SAFE_CALL(idefix_MPI_View_Send_init(BufferSendX1[faceLeft].commView(), bufferSizeX1, + MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX1[faceLeft].commView(), bufferSizeX1, realMPI, procSend, 101, mygrid->CartComm, &sendRequestX1[faceLeft])); - MPI_SAFE_CALL(idefix_MPI_View_Recv_init(BufferRecvX1[faceRight].commView(), bufferSizeX1, + 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(idefix_MPI_View_Send_init(BufferSendX2[faceRight].commView(), bufferSizeX2, + MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX2[faceRight].commView(), bufferSizeX2, realMPI, procSend, 200, mygrid->CartComm, &sendRequestX2[faceRight])); - MPI_SAFE_CALL(idefix_MPI_View_Recv_init(BufferRecvX2[faceLeft].commView(), bufferSizeX2, + 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(idefix_MPI_View_Send_init(BufferSendX2[faceLeft].commView(), bufferSizeX2, + MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX2[faceLeft].commView(), bufferSizeX2, realMPI, procSend, 201, mygrid->CartComm, &sendRequestX2[faceLeft])); - MPI_SAFE_CALL(idefix_MPI_View_Recv_init(BufferRecvX2[faceRight].commView(), bufferSizeX2, + MPI_SAFE_CALL(idefix::MPI_Recv_init(BufferRecvX2[faceRight].commView(), bufferSizeX2, realMPI, procRecv, 201, mygrid->CartComm, &recvRequestX2[faceRight])); #endif @@ -375,20 +375,20 @@ ConstrainedTransport::ConstrainedTransport(Input &input, Fluid *hydr // We receive from procRecv, and we send to procSend MPI_SAFE_CALL(MPI_Cart_shift(mygrid->CartComm,2,1,&procRecv,&procSend )); - MPI_SAFE_CALL(idefix_MPI_View_Send_init(BufferSendX3[faceRight].commView(), bufferSizeX3, + MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX3[faceRight].commView(), bufferSizeX3, realMPI, procSend, 300, mygrid->CartComm, &sendRequestX3[faceRight])); - MPI_SAFE_CALL(idefix_MPI_View_Recv_init(BufferRecvX3[faceLeft].commView(), bufferSizeX3, + 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(idefix_MPI_View_Send_init(BufferSendX3[faceLeft].commView(), bufferSizeX3, + MPI_SAFE_CALL(idefix::MPI_Send_init(BufferSendX3[faceLeft].commView(), bufferSizeX3, realMPI, procSend, 301, mygrid->CartComm, &sendRequestX3[faceLeft])); - MPI_SAFE_CALL(idefix_MPI_View_Recv_init(BufferRecvX3[faceRight].commView(), bufferSizeX3, + MPI_SAFE_CALL(idefix::MPI_Recv_init(BufferRecvX3[faceRight].commView(), bufferSizeX3, realMPI, procRecv, 301, mygrid->CartComm, &recvRequestX3[faceRight])); #endif @@ -409,17 +409,17 @@ ConstrainedTransport::~ConstrainedTransport() { // Properly clean up the mess idfx::cout << "Emf: Cleaning up MPI persistent communication channels" << std::endl; for(int i=0 ; i< 2; i++) { - idefix_MPI_View_Request_free( &sendRequestX1[i]); - idefix_MPI_View_Request_free( &recvRequestX1[i]); + idefix::MPI_Request_free( &sendRequestX1[i]); + idefix::MPI_Request_free( &recvRequestX1[i]); #if DIMENSIONS >= 2 - idefix_MPI_View_Request_free( &sendRequestX2[i]); - idefix_MPI_View_Request_free( &recvRequestX2[i]); + idefix::MPI_Request_free( &sendRequestX2[i]); + idefix::MPI_Request_free( &recvRequestX2[i]); #endif #if DIMENSIONS == 3 - idefix_MPI_View_Request_free( &sendRequestX3[i]); - idefix_MPI_View_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 0b19e6cd0..6f5e008d6 100644 --- a/src/fluid/constrainedTransport/enforceEMFBoundary.hpp +++ b/src/fluid/constrainedTransport/enforceEMFBoundary.hpp @@ -142,7 +142,7 @@ void ConstrainedTransport::SymmetrizeEMFShearingBox() { // We send to our left (which, by periodicity, is the right end of the domain) // our value of sbEyL and get Kokkos::fence(); - idefix_MPI_View_Sendrecv(sbEyL, size, realMPI, procLeft, 2001, + idefix::MPI_Sendrecv(sbEyL, size, realMPI, procLeft, 2001, sbEyR, size, realMPI, procLeft, 2002, data->mygrid->CartComm, &status ); } @@ -150,7 +150,7 @@ void ConstrainedTransport::SymmetrizeEMFShearingBox() { 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 - idefix_MPI_View_Sendrecv(sbEyR, size, realMPI, procRight, 2002, + idefix::MPI_Sendrecv(sbEyR, size, realMPI, procRight, 2002, sbEyL, size, realMPI, procRight, 2001, data->mygrid->CartComm, &status ); } diff --git a/src/mpi/buffer.hpp b/src/mpi/buffer.hpp index fd980f61c..869daee77 100644 --- a/src/mpi/buffer.hpp +++ b/src/mpi/buffer.hpp @@ -38,7 +38,7 @@ class Buffer { return this->array; } - IdefixCommArray1D & commView(void) { + idefix::IdefixCommArray1D & commView(void) { return this->array; } @@ -280,7 +280,7 @@ class Buffer { private: size_t pointer; - IdefixCommArray1D array; + idefix::IdefixCommArray1D array; }; #endif // MPI_BUFFER_HPP_ diff --git a/src/mpi/exchanger.cpp b/src/mpi/exchanger.cpp index cbd5044fc..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 - idefix_MPI_View_Send_init(BufferSend[faceRight].commView(), bufferSizeSend[faceRight], realMPI, + idefix::MPI_Send_init(BufferSend[faceRight].commView(), bufferSizeSend[faceRight], realMPI, procSend[faceRight], thisInstance*2, grid->CartComm, &sendRequest[faceRight]); - idefix_MPI_View_Recv_init(BufferRecv[faceLeft].commView(), 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 - idefix_MPI_View_Send_init(BufferSend[faceLeft].commView(), bufferSizeSend[faceLeft], realMPI, + idefix::MPI_Send_init(BufferSend[faceLeft].commView(), bufferSizeSend[faceLeft], realMPI, procSend[faceLeft],thisInstance*2+1, grid->CartComm, &sendRequest[faceLeft]); - idefix_MPI_View_Recv_init(BufferRecv[faceRight].commView(), 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++) { - idefix_MPI_View_Request_free( &sendRequest[i]); - idefix_MPI_View_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 Vc, IdefixArray4D Vs) { MPI_Status sendStatus[2]; MPI_Status recvStatus[2]; - idefix_MPI_View_Startall(2, recvRequest); + idefix::MPI_Startall(2, recvRequest); idfx::mpiCallsTimer += MPI_Wtime() - tStart; #endif myTimer += MPI_Wtime(); @@ -213,54 +213,54 @@ void Exchanger::Exchange(IdefixArray4D Vc, IdefixArray4D Vs) { tStart = MPI_Wtime(); #ifdef MPI_PERSISTENT - idefix_MPI_View_Startall(2, sendRequest); + idefix::MPI_Startall(2, sendRequest); // Wait for buffers to be received - idefix_MPI_View_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 @@ -295,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 - idefix_MPI_View_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 122796f82..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 - Idefix_MPI_Request_1D sendRequest[2]; - Idefix_MPI_Request_1D recvRequest[2]; + idefix::MPI_Request_1D sendRequest[2]; + idefix::MPI_Request_1D recvRequest[2]; Grid *grid; }; diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp index d4ba44f12..be2b4d4ae 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -15,6 +15,8 @@ #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 carry optionally a communication copy on host depending @@ -103,13 +105,13 @@ struct Idefix_MPI_Request { }; /** Wrap the idefix 1D array to use it for communication purpose. */ -template using Idefix_MPI_Request_1D = Idefix_MPI_Request< IdefixArray1D >; +template using MPI_Request_1D = Idefix_MPI_Request< IdefixArray1D >; /** Wrap the idefix 2D array to use it for communication purpose. */ -template using Idefix_MPI_Request_2D = Idefix_MPI_Request< IdefixArray2D >; +template using MPI_Request_2D = Idefix_MPI_Request< IdefixArray2D >; /** Wrap the idefix 3D array to use it for communication purpose. */ -template using Idefix_MPI_Request_3D = Idefix_MPI_Request< IdefixArray3D >; +template using MPI_Request_3D = Idefix_MPI_Request< IdefixArray3D >; /** Wrap the idefix 4D array to use it for communication purpose. */ -template using Idefix_MPI_Request_4D = Idefix_MPI_Request< IdefixArray4D >; +template using MPI_Request_4D = Idefix_MPI_Request< IdefixArray4D >; /** * Provide a wrapper of the standard MPI_Sendrecv by handling a kokkkos @@ -117,7 +119,7 @@ template using Idefix_MPI_Request_4D = Idefix_MPI_Request< IdefixArray * host buffer if required (when WITH_MPI_GPU_DIRECT is disabled.). */ template -int idefix_MPI_View_Sendrecv(Kokkos::View sendbuf, int sendcount, MPI_Datatype sendtype, +int MPI_Sendrecv(Kokkos::View sendbuf, int sendcount, MPI_Datatype sendtype, int dest, int sendtag, Kokkos::View recvbuf, int recvcount, MPI_Datatype recvtype, int source, int recvtag, MPI_Comm comm, MPI_Status *status) { //check size @@ -137,7 +139,7 @@ int idefix_MPI_View_Sendrecv(Kokkos::View sendbuf, int sendcount, MPI_D #endif //make the communication - const int result = MPI_Sendrecv(sendHostArray.data(), sendcount, sendtype, dest, sendtag, + const int result = ::MPI_Sendrecv(sendHostArray.data(), sendcount, sendtype, dest, sendtag, recvHostArray.data(), recvcount, recvtype, source, recvtag, comm, status); @@ -156,7 +158,7 @@ int idefix_MPI_View_Sendrecv(Kokkos::View sendbuf, int sendcount, MPI_D * to the host if required (when WITH_MPI_GPU_DIRECT is disabled). */ template -int idefix_MPI_View_Sendrecv(IdefixCommArray sendbuf, int sendcount, MPI_Datatype sendtype, +int MPI_Sendrecv(IdefixCommArray sendbuf, int sendcount, MPI_Datatype sendtype, int dest, int sendtag, IdefixCommArray recvbuf, int recvcount, MPI_Datatype recvtype, int source, int recvtag, MPI_Comm comm, MPI_Status *status) { //check size @@ -171,7 +173,7 @@ int idefix_MPI_View_Sendrecv(IdefixCommArray sendbuf, int sendcount, MPI_Data #endif //make the communication - const int result = MPI_Sendrecv(sendbuf.commData(), sendcount, sendtype, dest, sendtag, + const int result = ::MPI_Sendrecv(sendbuf.commData(), sendcount, sendtype, dest, sendtag, recvbuf.commData(), recvcount, recvtype, source, recvtag, comm, status); @@ -185,7 +187,7 @@ int idefix_MPI_View_Sendrecv(IdefixCommArray sendbuf, int sendcount, MPI_Data } template -int idefix_MPI_View_Send_init(IdefixCommArray & buf, int count, MPI_Datatype datatype, int dest, +int MPI_Send_init(IdefixCommArray & buf, int count, MPI_Datatype datatype, int dest, int tag, MPI_Comm comm, Idefix_MPI_Request *request) { //check assert(request != nullptr); @@ -204,7 +206,7 @@ int idefix_MPI_View_Send_init(IdefixCommArray & buf, int count, MPI_Datatype } template -int idefix_MPI_View_Recv_init(IdefixCommArray & buf, int count, MPI_Datatype datatype, int +int MPI_Recv_init(IdefixCommArray & buf, int count, MPI_Datatype datatype, int source, int tag, MPI_Comm comm, Idefix_MPI_Request* request) { //check assert(request != nullptr); @@ -223,7 +225,7 @@ int idefix_MPI_View_Recv_init(IdefixCommArray & buf, int count, MPI_Datatype } template -int idefix_MPI_View_Request_free(Idefix_MPI_Request* request) { +int MPI_Request_free(Idefix_MPI_Request* request) { //check assert(request != nullptr); @@ -232,7 +234,7 @@ int idefix_MPI_View_Request_free(Idefix_MPI_Request* request) { } template -int idefix_MPI_View_Startall(int count, Idefix_MPI_Request* request) { +int MPI_Startall(int count, Idefix_MPI_Request* request) { //check assert(request != nullptr); assert(count >= 0); @@ -248,7 +250,7 @@ int idefix_MPI_View_Startall(int count, Idefix_MPI_Request* request) { //launch the comms int final_result = 0; for (size_t i = 0 ; i < count ; i++) { - const int result = MPI_Start(&request[i].request); + const int result = ::MPI_Start(&request[i].request); if (result < 0) final_result = result; } @@ -258,7 +260,7 @@ int idefix_MPI_View_Startall(int count, Idefix_MPI_Request* request) { } template -int idefix_MPI_View_Waitall(int count, Idefix_MPI_Request* request, +int MPI_Waitall(int count, Idefix_MPI_Request* request, MPI_Status *array_of_statuses) { //check assert(request != nullptr); @@ -267,7 +269,7 @@ int idefix_MPI_View_Waitall(int count, Idefix_MPI_Request* request, //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]); + const int result = ::MPI_Wait(&request[i].request, &array_of_statuses[i]); if (result < 0) final_result = result; } @@ -285,7 +287,7 @@ int idefix_MPI_View_Waitall(int count, Idefix_MPI_Request* request, } template -int idefix_MPI_View_Start(Idefix_MPI_Request* request) { +int MPI_Start(Idefix_MPI_Request* request) { //check assert(request != nullptr); @@ -293,16 +295,16 @@ int idefix_MPI_View_Start(Idefix_MPI_Request* request) { request->commArray.syncCommData(); //call MPI - return MPI_Start(&request->request); + return ::MPI_Start(&request->request); } template -int idefix_MPI_View_Wait(Idefix_MPI_Request* request, MPI_Status *array_of_statuses) { +int MPI_Wait(Idefix_MPI_Request* request, MPI_Status *array_of_statuses) { //check assert(request != nullptr); //wait the comms - const int result = MPI_Wait(&request->request, array_of_statuses); + const int result = ::MPI_Wait(&request->request, array_of_statuses); //sync the copies in async mode request->commArray.syncDeviceData(); @@ -311,4 +313,6 @@ int idefix_MPI_View_Wait(Idefix_MPI_Request* request, MPI_Status *array_of_st return result; } +} // namespace idefix + #endif // MPI_MPIVIEW_HPP_ From 2af6af638ebeb88ef2ffd6c4747c5351b037956a Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Thu, 9 Jul 2026 17:48:37 +0200 Subject: [PATCH 05/17] gpu-direct: fix some extra parts of idefix using MPI --- src/mpi/mpiView.hpp | 154 ++++++++++++++++++++++++++++++++++++++++++- src/utils/column.cpp | 13 ++-- src/utils/column.hpp | 2 +- 3 files changed, 161 insertions(+), 8 deletions(-) diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp index be2b4d4ae..15ed3416d 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -19,7 +19,7 @@ namespace idefix { /** * Provide a specific array used for MPI communications. It embeds the standard - * array as usual, and also carry optionally a communication copy on host depending + * 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 @@ -186,6 +186,10 @@ int MPI_Sendrecv(IdefixCommArray sendbuf, int sendcount, MPI_Datatype sendtyp 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 int MPI_Send_init(IdefixCommArray & buf, int count, MPI_Datatype datatype, int dest, int tag, MPI_Comm comm, Idefix_MPI_Request *request) { @@ -205,6 +209,10 @@ int MPI_Send_init(IdefixCommArray & buf, int count, MPI_Datatype datatype, in 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 int MPI_Recv_init(IdefixCommArray & buf, int count, MPI_Datatype datatype, int source, int tag, MPI_Comm comm, Idefix_MPI_Request* request) { @@ -224,6 +232,10 @@ int MPI_Recv_init(IdefixCommArray & buf, int count, MPI_Datatype datatype, in 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 int MPI_Request_free(Idefix_MPI_Request* request) { //check @@ -233,6 +245,10 @@ int MPI_Request_free(Idefix_MPI_Request* request) { 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 int MPI_Startall(int count, Idefix_MPI_Request* request) { //check @@ -259,6 +275,10 @@ int MPI_Startall(int count, Idefix_MPI_Request* request) { 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 int MPI_Waitall(int count, Idefix_MPI_Request* request, MPI_Status *array_of_statuses) { @@ -286,6 +306,10 @@ int MPI_Waitall(int count, Idefix_MPI_Request* request, 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 int MPI_Start(Idefix_MPI_Request* request) { //check @@ -298,6 +322,10 @@ int MPI_Start(Idefix_MPI_Request* request) { 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 int MPI_Wait(Idefix_MPI_Request* request, MPI_Status *array_of_statuses) { //check @@ -313,6 +341,130 @@ int MPI_Wait(Idefix_MPI_Request* request, MPI_Status *array_of_statuses) { 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 +int MPI_Send(IdefixCommArray & 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 +int MPI_Recv(IdefixCommArray & 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_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 +int MPI_Isend(IdefixCommArray & buf, int count, MPI_Datatype datatype, int + source, int tag, MPI_Comm comm, Idefix_MPI_Request* 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 +int MPI_Irecv(IdefixCommArray & buf, int count, MPI_Datatype datatype, int + source, int tag, MPI_Comm comm, Idefix_MPI_Request* 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 +int MPI_Bcast(IdefixCommArray & 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; +} + } // namespace idefix #endif // MPI_MPIVIEW_HPP_ diff --git a/src/utils/column.cpp b/src/utils/column.cpp index 4657abc55..fb26a4596 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("localSum",np_tot[KDIR], np_tot[JDIR]); + localSum = idefix::IdefixCommArray2D("localSum",np_tot[KDIR], np_tot[JDIR]); } if(dir == JDIR) { - localSum = IdefixArray2D("localSum",np_tot[KDIR], np_tot[IDIR]); + localSum = idefix::IdefixCommArray2D("localSum",np_tot[KDIR], np_tot[IDIR]); } if(dir == KDIR) { - localSum = IdefixArray2D("localSum",np_tot[JDIR], np_tot[IDIR]); + localSum = idefix::IdefixCommArray2D("localSum",np_tot[JDIR], np_tot[IDIR]); } #ifdef WITH_MPI // Create sub-MPI communicator dedicated to scan @@ -133,7 +134,7 @@ void Column::ComputeColumn(IdefixArray4D 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(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 in, const int var) { } // And send it Kokkos::fence(); - MPI_Send(localSum.data(),size, realMPI, dst, 20, ColumnComm); + idefix::MPI_Send(localSum,size, realMPI, dst, 20, ColumnComm); } // MPIrank small enough #endif // If we need it backwards @@ -194,7 +195,7 @@ void Column::ComputeColumn(IdefixArray4D in, const int var) { #ifdef WITH_MPI Kokkos::fence(); - MPI_Bcast(localSum.data(),size, realMPI, MPIsize-1, ColumnComm); + idefix::MPI_Bcast(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 Area; IdefixArray3D Volume; - IdefixArray2D localSum; + idefix::IdefixCommArray2D localSum; #ifdef WITH_MPI Mpi mpi; // Mpi object when WITH_MPI is set MPI_Comm ColumnComm; From f38f1069474a9dad9fbc7da369ab9da3df6f1411 Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 10:29:45 +0200 Subject: [PATCH 06/17] gpu-direct: wrap all the internal MPI_.+ using .data() internally with idefix::MPI_ API. --- src/fluid/boundary/axis.cpp | 12 ++-- src/fluid/boundary/axis.hpp | 8 +-- src/mpi/mpi.cpp | 10 +-- src/mpi/mpiView.hpp | 125 ++++++++++++++++++++++++++++++++++++ src/output/slice.cpp | 2 +- src/output/vtk.cpp | 2 +- src/pydefix.cpp | 22 +++---- src/timeIntegrator.cpp | 2 +- src/utils/lookupTable.hpp | 8 +-- 9 files changed, 158 insertions(+), 33 deletions(-) diff --git a/src/fluid/boundary/axis.cpp b/src/fluid/boundary/axis.cpp index 9303747a4..206951516 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 Ex1) { #if DIMENSIONS == 3 - IdefixArray1D Ex1Avg = this->Ex1Avg; + idefix::IdefixCommArray1D Ex1Avg = this->Ex1Avg; idefix_for("Ex1_ini",0,data->np_tot[IDIR], KOKKOS_LAMBDA(int i) { @@ -40,7 +40,7 @@ void Axis::SymmetrizeEx1Side(int jref, IdefixArray3D 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, Ex1Avg, data->np_tot[IDIR], realMPI, MPI_SUM, data->mygrid->AxisComm); #endif } @@ -88,7 +88,7 @@ void Axis::RegularizeCurrentSide(int side) { jc = data->end[JDIR]-1; sign = -1; } - IdefixArray1D BAvg = this->Ex1Avg; + idefix::IdefixCommArray1D BAvg = this->Ex1Avg; IdefixArray1D x1 = data->x[IDIR]; IdefixArray1D dx3 = data->dx[KDIR]; IdefixArray1D dx2 = data->dx[JDIR]; @@ -107,7 +107,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, BAvg, data->np_tot[IDIR], realMPI, MPI_SUM, data->mygrid->AxisComm); #endif } @@ -169,7 +169,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 Vs = this->Vs; - IdefixArray2D BAvg = this->BAvg; + idefix::IdefixCommArray2D BAvg = this->BAvg; IdefixArray1D phi = data->x[KDIR]; int jin = 0; @@ -208,7 +208,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, BAvg, 2*data->np_tot[IDIR], realMPI, MPI_SUM, data->mygrid->AxisComm); #endif } diff --git a/src/fluid/boundary/axis.hpp b/src/fluid/boundary/axis.hpp index 039519228..62a3c1ed3 100644 --- a/src/fluid/boundary/axis.hpp +++ b/src/fluid/boundary/axis.hpp @@ -71,8 +71,8 @@ class Axis { #endif void InitMPI(); - IdefixArray1D Ex1Avg; - IdefixArray2D BAvg; + idefix::IdefixCommArray1D Ex1Avg; + idefix::IdefixCommArray2D BAvg; bool haveCurrent; IdefixArray2D JAvg; IdefixArray1D symmetryVc; @@ -155,8 +155,8 @@ Axis::Axis(Boundary *boundary) { } Kokkos::deep_copy(symmetryVs, symmetryVsHost); - this->Ex1Avg = IdefixArray1D("Axis:Ex1Avg",data->np_tot[IDIR]); - this->BAvg = IdefixArray2D("Axis:BxAvg",data->np_tot[IDIR],2); + this->Ex1Avg = idefix::IdefixCommArray1D("Axis:Ex1Avg",data->np_tot[IDIR]); + this->BAvg = idefix::IdefixCommArray2D("Axis:BxAvg",data->np_tot[IDIR],2); if(haveCurrent) { this->JAvg = IdefixArray2D("Axis:JAvg",data->np_tot[IDIR],3); } diff --git a/src/mpi/mpi.cpp b/src/mpi/mpi.cpp index cb436ec92..a90e22ef5 100644 --- a/src/mpi/mpi.cpp +++ b/src/mpi/mpi.cpp @@ -130,7 +130,7 @@ void Mpi::CheckConfig() { #endif /* MPIX_CUDA_AWARE_SUPPORT */ // Run-time check that we can do a reduce on device arrays - IdefixArray1D src("MPIChecksrc",1); + idefix::IdefixCommArray1D src("MPIChecksrc",1); IdefixArray1D::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/mpiView.hpp b/src/mpi/mpiView.hpp index 15ed3416d..5a647525a 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -11,6 +11,8 @@ #include #include +#include +#include #include "idefix.hpp" #include "buffer.hpp" @@ -385,6 +387,17 @@ int MPI_Recv(IdefixCommArray & buf, int count, MPI_Datatype datatype, int 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 +int MPI_Recv(std::array & 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. @@ -465,6 +478,118 @@ int MPI_Bcast(IdefixCommArray & buffer, int count, MPI_Datatype datatype, int 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 +int MPI_Bcast(IdefixHostArray1D & 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 +int MPI_Allreduce(IdefixCommArray & sendbuf, IdefixCommArray & 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 +int MPI_Allreduce(void * sendbuf, IdefixHostArray3D & 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 +int MPI_Allreduce(void * sendbuf, IdefixCommArray & 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 +int MPI_File_write_all(MPI_File fh, const IdefixHostArray4D & 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 +int MPI_Gather(const void* sendbuf, int sendcount, MPI_Datatype + sendtype, std::vector & recvbuf, int recvcount, MPI_Datatype recvtype, + int root, MPI_Comm comm) { + return MPI_Gather(sendbuf, sendcount, sendtype, recvbuf, 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(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 GatherIdefixArray(IdefixHostArray3D #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("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 GatherIdefixArray(IdefixHostArray3D } #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 GatherIdefixArray(IdefixHostArray3D // 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 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/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::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 From 7b89da6ddb7457c764df58b8e20f33ef83652d31 Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 10:33:59 +0200 Subject: [PATCH 07/17] gpu-direct: Add a macro to activate to validate by forcing the transfers all time for validation purpose --- src/mpi/mpiView.hpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp index 5a647525a..b8a6cf9aa 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -8,6 +8,8 @@ #ifndef MPI_MPIVIEW_HPP_ #define MPI_MPIVIEW_HPP_ +//#define FORCE_TRANSFERS_ALL_TIME + #include #include @@ -76,6 +78,9 @@ class IdefixCommArray : public T { //in gpu direct, we use directly the device buffer //if running on host, it will also just use directly the host storage. return deviceArray; + #elif defined(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 From 9b58b1c83b46bef3946a8bc7f0306995dba9188d Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 10:49:40 +0200 Subject: [PATCH 08/17] gpu-direct: fix an issue in MPI_Gather --- src/mpi/mpiView.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp index b8a6cf9aa..0c38cbd35 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -8,6 +8,7 @@ #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 @@ -592,7 +593,7 @@ template int MPI_Gather(const void* sendbuf, int sendcount, MPI_Datatype sendtype, std::vector & recvbuf, int recvcount, MPI_Datatype recvtype, int root, MPI_Comm comm) { - return MPI_Gather(sendbuf, sendcount, sendtype, recvbuf, recvcount, recvtype, root, comm); + return MPI_Gather(sendbuf, sendcount, sendtype, recvbuf.data(), recvcount, recvtype, root, comm); } } // namespace idefix From 75a4c3331335edae64dd4ce14d13750eaddc2c43 Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 11:44:33 +0200 Subject: [PATCH 09/17] gpu-direct: Fix some compilation warnings with CUDA --- src/fluid/boundary/axis.cpp | 2 +- src/mpi/mpiView.hpp | 9 ++++++++- 2 files changed, 9 insertions(+), 2 deletions(-) diff --git a/src/fluid/boundary/axis.cpp b/src/fluid/boundary/axis.cpp index 206951516..c0568aace 100644 --- a/src/fluid/boundary/axis.cpp +++ b/src/fluid/boundary/axis.cpp @@ -392,7 +392,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 map = this->mapVars; diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp index 0c38cbd35..aa373ce5e 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -38,6 +38,11 @@ class IdefixCommArray : public T { * a communication. */ void syncCommData(void) { + //do lazy init (cannot do in constructor without having a warning under CUDA build) + if (this->commArray.is_allocated() == false) + this->commArray = initCommArray(*this); + + //transfer the data from device to host. Kokkos::deep_copy(this->commArray, *this); } @@ -70,6 +75,7 @@ class IdefixCommArray : public T { * depending on the WITH_MPI_GPU_DIRECT configuration. */ void * commData(void) { + assert(this->commArray.is_allocated()); return this->commArray.data(); } @@ -94,7 +100,7 @@ class IdefixCommArray : public T { * 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; }; /** Wrap the idefix 1D array to use it for communication purpose. */ @@ -593,6 +599,7 @@ template int MPI_Gather(const void* sendbuf, int sendcount, MPI_Datatype sendtype, std::vector & 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); } From 433c9c13cc1bdd806d9a3bab1b5a82f722891bdc Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 12:07:09 +0200 Subject: [PATCH 10/17] gpu-direct: fix mistake in lazy-init, should be done on all entry points --- src/mpi/mpiView.hpp | 18 ++++++++++++------ 1 file changed, 12 insertions(+), 6 deletions(-) diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp index aa373ce5e..5b45750f8 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -38,11 +38,7 @@ class IdefixCommArray : public T { * a communication. */ void syncCommData(void) { - //do lazy init (cannot do in constructor without having a warning under CUDA build) - if (this->commArray.is_allocated() == false) - this->commArray = initCommArray(*this); - - //transfer the data from device to host. + this->lazyInit(); Kokkos::deep_copy(this->commArray, *this); } @@ -51,6 +47,7 @@ class IdefixCommArray : public T { * a communication. */ void syncCommDataAsync(void) { + this->lazyInit(); Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), this->commArray, *this); } @@ -59,6 +56,7 @@ class IdefixCommArray : public T { * to use it. */ void syncDeviceData(void) { + this->lazyInit(); Kokkos::deep_copy(*this, this->commArray); } @@ -67,6 +65,7 @@ class IdefixCommArray : public T { * to use it. */ void syncDeviceDataAsync(void) { + this->lazyInit(); Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), *this, this->commArray); } @@ -75,10 +74,17 @@ class IdefixCommArray : public T { * depending on the WITH_MPI_GPU_DIRECT configuration. */ void * commData(void) { - assert(this->commArray.is_allocated()); + this->lazyInit(); return this->commArray.data(); } + private: + void lazyInit(void) { + //do lazy init (cannot do in constructor without having a warning under CUDA build) + if (this->commArray.is_allocated() == false) + this->commArray = initCommArray(*this); + } + private: static typename T::host_mirror_type initCommArray(T & deviceArray) { #ifdef WITH_MPI_GPU_DIRECT From 17e0c762c85231693d62f1f148da5de3ed6f0c6b Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 13:37:10 +0200 Subject: [PATCH 11/17] gpu-direct: Ooops, no cannot lazy init due to the copy semantic --- src/mpi/mpiView.hpp | 21 ++++++++------------- 1 file changed, 8 insertions(+), 13 deletions(-) diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp index 5b45750f8..5f7565f54 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -33,12 +33,18 @@ class IdefixCommArray : public T { //inherit constructors using T::T; + /** + * Overload the copy constructor + */ + IdefixCommArray(const IdefixCommArray & orig) + :T(orig) + ,commArray(orig.commArray) {} + /** * If needed, transfers the data from the device to the host to be ready to make * a communication. */ void syncCommData(void) { - this->lazyInit(); Kokkos::deep_copy(this->commArray, *this); } @@ -47,7 +53,6 @@ class IdefixCommArray : public T { * a communication. */ void syncCommDataAsync(void) { - this->lazyInit(); Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), this->commArray, *this); } @@ -56,7 +61,6 @@ class IdefixCommArray : public T { * to use it. */ void syncDeviceData(void) { - this->lazyInit(); Kokkos::deep_copy(*this, this->commArray); } @@ -65,7 +69,6 @@ class IdefixCommArray : public T { * to use it. */ void syncDeviceDataAsync(void) { - this->lazyInit(); Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), *this, this->commArray); } @@ -74,17 +77,9 @@ class IdefixCommArray : public T { * depending on the WITH_MPI_GPU_DIRECT configuration. */ void * commData(void) { - this->lazyInit(); return this->commArray.data(); } - private: - void lazyInit(void) { - //do lazy init (cannot do in constructor without having a warning under CUDA build) - if (this->commArray.is_allocated() == false) - this->commArray = initCommArray(*this); - } - private: static typename T::host_mirror_type initCommArray(T & deviceArray) { #ifdef WITH_MPI_GPU_DIRECT @@ -106,7 +101,7 @@ class IdefixCommArray : public T { * 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; + typename T::host_mirror_type commArray{initCommArray(*this)}; }; /** Wrap the idefix 1D array to use it for communication purpose. */ From 411866f7fd265fdaf63fad865e570fd8c14fdb05 Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 13:58:43 +0200 Subject: [PATCH 12/17] gpu-direct: try to avoid the warning in normal way to run, need to find a solution for the NO GPU DIRECT case --- src/mpi/mpiView.hpp | 67 ++++++++++++++++++++++++--------------------- 1 file changed, 36 insertions(+), 31 deletions(-) diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp index 5f7565f54..ed4a4485e 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -33,19 +33,14 @@ class IdefixCommArray : public T { //inherit constructors using T::T; - /** - * Overload the copy constructor - */ - IdefixCommArray(const IdefixCommArray & orig) - :T(orig) - ,commArray(orig.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); + #ifndef WITH_MPI_GPU_DIRECT + Kokkos::deep_copy(this->commArray, *this); + #endif //WITH_MPI_GPU_DIRECT } /** @@ -53,7 +48,9 @@ class IdefixCommArray : public T { * a communication. */ void syncCommDataAsync(void) { - Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), this->commArray, *this); + #ifndef WITH_MPI_GPU_DIRECT + Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), this->commArray, *this); + #endif //WITH_MPI_GPU_DIRECT } /** @@ -61,7 +58,9 @@ class IdefixCommArray : public T { * to use it. */ void syncDeviceData(void) { - Kokkos::deep_copy(*this, this->commArray); + #ifndef WITH_MPI_GPU_DIRECT + Kokkos::deep_copy(*this, this->commArray); + #endif //WITH_MPI_GPU_DIRECT } /** @@ -69,7 +68,9 @@ class IdefixCommArray : public T { * to use it. */ void syncDeviceDataAsync(void) { - Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), *this, this->commArray); + #ifndef WITH_MPI_GPU_DIRECT + Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), *this, this->commArray); + #endif //WITH_MPI_GPU_DIRECT } /** @@ -77,31 +78,35 @@ class IdefixCommArray : public T { * 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 WITH_MPI_GPU_DIRECT - //in gpu direct, we use directly the device buffer - //if running on host, it will also just use directly the host storage. - return deviceArray; - #elif defined(FORCE_TRANSFERS_ALL_TIME) - //for validation purpose, force the transfers by using a copy anytime. - return Kokkos::create_mirror(deviceArray); + return this->data(); #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 + return this->commArray.data(); + #endif //WITH_MPI_GPU_DIRECT } 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)}; + #ifndef WITH_MPI_GPU_DIRECT + 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 + } + #endif //WITH_MPI_GPU_DIRECT + + private: + #ifndef WITH_MPI_GPU_DIRECT + /** + * 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)}; + #endif //WITH_MPI_GPU_DIRECT }; /** Wrap the idefix 1D array to use it for communication purpose. */ From 4a5b620f006f5eff7136c2d230749334e6a540ec Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 14:14:42 +0200 Subject: [PATCH 13/17] gpu-direct: Cleanup the way to switch the class ON/OFF by using a single #if/else --- src/mpi/mpiView.hpp | 108 +++++++++++++++++++++++++++++--------------- 1 file changed, 71 insertions(+), 37 deletions(-) diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp index ed4a4485e..0c902058e 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -28,7 +28,51 @@ namespace idefix { * if we have access to WITH_MPI_GPU_DIRECT or not. */ template -class IdefixCommArray : public T { +class IdefixCommArrayGpuDirect : public T { + public: + //inherit constructors + using T::T; + + /** + * 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 IdefixCommArrayNoGpuDirect : public T { public: //inherit constructors using T::T; @@ -38,9 +82,7 @@ class IdefixCommArray : public T { * a communication. */ void syncCommData(void) { - #ifndef WITH_MPI_GPU_DIRECT - Kokkos::deep_copy(this->commArray, *this); - #endif //WITH_MPI_GPU_DIRECT + Kokkos::deep_copy(this->commArray, *this); } /** @@ -48,9 +90,7 @@ class IdefixCommArray : public T { * a communication. */ void syncCommDataAsync(void) { - #ifndef WITH_MPI_GPU_DIRECT - Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), this->commArray, *this); - #endif //WITH_MPI_GPU_DIRECT + Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), this->commArray, *this); } /** @@ -58,9 +98,7 @@ class IdefixCommArray : public T { * to use it. */ void syncDeviceData(void) { - #ifndef WITH_MPI_GPU_DIRECT - Kokkos::deep_copy(*this, this->commArray); - #endif //WITH_MPI_GPU_DIRECT + Kokkos::deep_copy(*this, this->commArray); } /** @@ -68,9 +106,7 @@ class IdefixCommArray : public T { * to use it. */ void syncDeviceDataAsync(void) { - #ifndef WITH_MPI_GPU_DIRECT - Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), *this, this->commArray); - #endif //WITH_MPI_GPU_DIRECT + Kokkos::deep_copy(Kokkos::DefaultExecutionSpace(), *this, this->commArray); } /** @@ -78,37 +114,35 @@ class IdefixCommArray : public T { * depending on the WITH_MPI_GPU_DIRECT configuration. */ void * commData(void) { - #ifdef WITH_MPI_GPU_DIRECT - return this->data(); - #else - return this->commArray.data(); - #endif //WITH_MPI_GPU_DIRECT + return this->commArray.data(); } private: - #ifndef WITH_MPI_GPU_DIRECT - 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 - } - #endif //WITH_MPI_GPU_DIRECT + 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: - #ifndef WITH_MPI_GPU_DIRECT - /** - * 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)}; - #endif //WITH_MPI_GPU_DIRECT + /** + * 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 using IdefixCommArray = IdefixCommArrayGpuDirect; +#else + template using IdefixCommArray = IdefixCommArrayNoGpuDirect; +#endif + /** Wrap the idefix 1D array to use it for communication purpose. */ template using IdefixCommArray1D = IdefixCommArray< IdefixArray1D >; /** Wrap the idefix 2D array to use it for communication purpose. */ From afa7a4e974f3d36213f995b23113a84681309e00 Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 14:24:09 +0200 Subject: [PATCH 14/17] gpu-direct: Opps wrong name in option declaration --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 3870ad485..84974d3b4 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -16,7 +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_GPU_DIRECT "Enable usage of GPU direct to avoid CPU-GPU copies for MPI messages." ON) +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") From 03e0f40163051d3d06fa347666c89bdf450b1338 Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 15:26:37 +0200 Subject: [PATCH 15/17] gpu-direct: remove a useless scafolding --- src/mpi/buffer.hpp | 24 ------------------------ 1 file changed, 24 deletions(-) diff --git a/src/mpi/buffer.hpp b/src/mpi/buffer.hpp index 869daee77..4d4bc3904 100644 --- a/src/mpi/buffer.hpp +++ b/src/mpi/buffer.hpp @@ -170,30 +170,6 @@ class Buffer { this->pointer += ninjnk; } - void UnpackAveraged(IdefixArray4D& out, - const int var, - BoundingBox box) { - const int ni = box[IDIR][1]-box[IDIR][0]; - const int ninj = (box[JDIR][1]-box[JDIR][0])*ni; - const int ninjnk = (box[KDIR][1]-box[KDIR][0])*ninj; - const int ibeg = box[IDIR][0]; - const int jbeg = box[JDIR][0]; - const int kbeg = box[KDIR][0]; - const int iend = box[IDIR][1]; - const int jend = box[JDIR][1]; - const int kend = box[KDIR][1]; - const int offset = this->pointer; - - auto arr = this->array; - 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 ); - }); - - // Update pointer - this->pointer += ninjnk; - } - void UnpackJDirSymetric(IdefixArray4D& out, const int var, const int symMultiplier, From 76d46ffb6b4b51cb538bdfa97e07f850dde9400b Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 16:24:21 +0200 Subject: [PATCH 16/17] gpu-direct: try to implement the copy/move constructor & operator= to solve the Axis issue --- src/mpi/buffer.hpp | 4 +- src/mpi/mpiView.hpp | 89 +++++++++++++++++++++++++-------------------- 2 files changed, 52 insertions(+), 41 deletions(-) diff --git a/src/mpi/buffer.hpp b/src/mpi/buffer.hpp index 4d4bc3904..5a3bc876d 100644 --- a/src/mpi/buffer.hpp +++ b/src/mpi/buffer.hpp @@ -255,8 +255,8 @@ class Buffer { } private: - size_t pointer; - idefix::IdefixCommArray1D array; + size_t pointer{0}; + idefix::IdefixCommArray1D array{}; }; #endif // MPI_BUFFER_HPP_ diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp index 0c902058e..dfa2d9424 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -16,6 +16,7 @@ #include #include #include +#include #include "idefix.hpp" #include "buffer.hpp" @@ -33,6 +34,28 @@ class IdefixCommArrayGpuDirect : public T { //inherit constructors using T::T; + IdefixCommArrayGpuDirect(const IdefixCommArrayGpuDirect & orig) { + T & tmp = *this; + tmp = (T&)orig; + } + + IdefixCommArrayGpuDirect(IdefixCommArrayGpuDirect && orig) { + T & tmp = *this; + tmp = std::move((T&)orig); + } + + IdefixCommArrayGpuDirect & operator=(const IdefixCommArrayGpuDirect & orig) { + T & tmp = *this; + tmp = (T&)orig; + return *this; + } + + IdefixCommArrayGpuDirect & operator=(IdefixCommArrayGpuDirect && orig) { + T & tmp = *this; + tmp = std::move((T&)orig); + return *this; + } + /** * If needed, transfers the data from the device to the host to be ready to make * a communication. @@ -77,6 +100,33 @@ class IdefixCommArrayNoGpuDirect : public T { //inherit constructors using T::T; + IdefixCommArrayNoGpuDirect(const IdefixCommArrayNoGpuDirect & orig) { + T & tmp = *this; + tmp = (T&)orig; + this->commArray = orig.commArray; + } + + IdefixCommArrayNoGpuDirect(IdefixCommArrayNoGpuDirect && orig) { + T & tmp = *this; + tmp = std::move((T&)orig); + this->commArray = std::move(orig.commArray); + } + + IdefixCommArrayNoGpuDirect & operator=(IdefixCommArrayNoGpuDirect & orig) { + T & tmp = *this; + tmp = (T&)orig; + this->commArray = orig.commArray; + return *this; + } + + IdefixCommArrayNoGpuDirect & operator=(IdefixCommArrayNoGpuDirect && orig) { + T & tmp = *this; + tmp = std::move((T&)orig); + this->commArray = std::move(orig.commArray); + return *this; + } + + /** * If needed, transfers the data from the device to the host to be ready to make * a communication. @@ -167,45 +217,6 @@ template using MPI_Request_3D = Idefix_MPI_Request< IdefixArray3D > /** Wrap the idefix 4D array to use it for communication purpose. */ template using MPI_Request_4D = Idefix_MPI_Request< IdefixArray4D >; -/** - * Provide a wrapper of the standard MPI_Sendrecv by handling a kokkkos - * view. Use it with caution as it will allocate/deallocate a temporary - * host buffer if required (when WITH_MPI_GPU_DIRECT is disabled.). - */ -template -int MPI_Sendrecv(Kokkos::View sendbuf, int sendcount, MPI_Datatype sendtype, - int dest, int sendtag, Kokkos::View 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 - #ifdef WITH_MPI_GPU_DIRECT - auto sendHostArray = sendbuf; - auto recvHostArray = recvbuf; - #else - auto sendHostArray = Kokkos::create_mirror_view(sendbuf); - auto recvHostArray = Kokkos::create_mirror_view(recvbuf); - Kokkos::deep_copy(sendHostArray, sendbuf); - #endif - - //make the communication - const int result = ::MPI_Sendrecv(sendHostArray.data(), sendcount, sendtype, dest, sendtag, - recvHostArray.data(), recvcount, recvtype, source, recvtag, - comm, status); - - //copy back if needed - #ifndef WITH_MPI_GPU_DIRECT - Kokkos::deep_copy(recvbuf, recvHostArray); - #endif - - //ok - return result; -} - /** * Provide a wrapper of the standard MPI_Sendrecv by handling an idefix * communication array. It will automate the process of transfering the data From aab2e8a1635b3701a263bd3c5f3b5bd2b4164375 Mon Sep 17 00:00:00 2001 From: Sebastien Valat Date: Fri, 10 Jul 2026 16:48:25 +0200 Subject: [PATCH 17/17] gpu-direct: solve most of the compiler warnings when compiling with CUDA --- src/fluid/boundary/axis.cpp | 13 ++-- .../enforceEMFBoundary.hpp | 14 ++-- src/mpi/buffer.hpp | 14 ++-- src/mpi/mpiView.hpp | 66 ++++--------------- src/utils/column.cpp | 8 +-- 5 files changed, 39 insertions(+), 76 deletions(-) diff --git a/src/fluid/boundary/axis.cpp b/src/fluid/boundary/axis.cpp index c0568aace..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 Ex1) { #if DIMENSIONS == 3 - idefix::IdefixCommArray1D 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 Ex1) { #ifdef WITH_MPI Kokkos::fence(); // sum along all of the processes on the same r - idefix::MPI_Allreduce(MPI_IN_PLACE, Ex1Avg, 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; } - idefix::IdefixCommArray1D BAvg = this->Ex1Avg; + auto BAvg = this->Ex1Avg.deviceView(); + auto BAvgComm = this->Ex1Avg; IdefixArray1D x1 = data->x[IDIR]; IdefixArray1D dx3 = data->dx[KDIR]; IdefixArray1D 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->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 Vs = this->Vs; - idefix::IdefixCommArray2D BAvg = this->BAvg; + auto BAvg = this->BAvg.deviceView(); IdefixArray1D 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 - idefix::MPI_Allreduce(MPI_IN_PLACE, BAvg, 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 } diff --git a/src/fluid/constrainedTransport/enforceEMFBoundary.hpp b/src/fluid/constrainedTransport/enforceEMFBoundary.hpp index 6f5e008d6..f7287a8aa 100644 --- a/src/fluid/constrainedTransport/enforceEMFBoundary.hpp +++ b/src/fluid/constrainedTransport/enforceEMFBoundary.hpp @@ -108,9 +108,9 @@ void ConstrainedTransport::SymmetrizeEMFShearingBox() { idfx::pushRegion("Emf::EnforceEMFBoundary"); #if MHD == YES - auto sbEyL = this->sbEyL; - auto sbEyR = this->sbEyR; - auto sbEyRL = this->sbEyRL; + auto sbEyL = this->sbEyL.deviceView(); + auto sbEyR = this->sbEyR.deviceView(); + auto sbEyRL = this->sbEyRL.deviceView(); IdefixArray3D ey = this->ey; @@ -142,16 +142,16 @@ void ConstrainedTransport::SymmetrizeEMFShearingBox() { // We send to our left (which, by periodicity, is the right end of the domain) // our value of sbEyL and get Kokkos::fence(); - idefix::MPI_Sendrecv(sbEyL, size, realMPI, procLeft, 2001, - sbEyR, 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 - idefix::MPI_Sendrecv(sbEyR, size, realMPI, procRight, 2002, - sbEyL, 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 5a3bc876d..9fbba748b 100644 --- a/src/mpi/buffer.hpp +++ b/src/mpi/buffer.hpp @@ -86,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); @@ -109,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, @@ -135,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) { @@ -160,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 ); @@ -185,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; @@ -211,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, @@ -239,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, diff --git a/src/mpi/mpiView.hpp b/src/mpi/mpiView.hpp index dfa2d9424..3ee17b051 100644 --- a/src/mpi/mpiView.hpp +++ b/src/mpi/mpiView.hpp @@ -16,7 +16,6 @@ #include #include #include -#include #include "idefix.hpp" #include "buffer.hpp" @@ -34,25 +33,11 @@ class IdefixCommArrayGpuDirect : public T { //inherit constructors using T::T; - IdefixCommArrayGpuDirect(const IdefixCommArrayGpuDirect & orig) { - T & tmp = *this; - tmp = (T&)orig; - } - - IdefixCommArrayGpuDirect(IdefixCommArrayGpuDirect && orig) { - T & tmp = *this; - tmp = std::move((T&)orig); - } - - IdefixCommArrayGpuDirect & operator=(const IdefixCommArrayGpuDirect & orig) { - T & tmp = *this; - tmp = (T&)orig; - return *this; - } - - IdefixCommArrayGpuDirect & operator=(IdefixCommArrayGpuDirect && orig) { - T & tmp = *this; - tmp = std::move((T&)orig); + /** + * 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; } @@ -100,33 +85,14 @@ class IdefixCommArrayNoGpuDirect : public T { //inherit constructors using T::T; - IdefixCommArrayNoGpuDirect(const IdefixCommArrayNoGpuDirect & orig) { - T & tmp = *this; - tmp = (T&)orig; - this->commArray = orig.commArray; - } - - IdefixCommArrayNoGpuDirect(IdefixCommArrayNoGpuDirect && orig) { - T & tmp = *this; - tmp = std::move((T&)orig); - this->commArray = std::move(orig.commArray); - } - - IdefixCommArrayNoGpuDirect & operator=(IdefixCommArrayNoGpuDirect & orig) { - T & tmp = *this; - tmp = (T&)orig; - this->commArray = orig.commArray; - return *this; - } - - IdefixCommArrayNoGpuDirect & operator=(IdefixCommArrayNoGpuDirect && orig) { - T & tmp = *this; - tmp = std::move((T&)orig); - this->commArray = std::move(orig.commArray); + /** + * 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. @@ -223,8 +189,8 @@ template using MPI_Request_4D = Idefix_MPI_Request< IdefixArray4D > * to the host if required (when WITH_MPI_GPU_DIRECT is disabled). */ template -int MPI_Sendrecv(IdefixCommArray sendbuf, int sendcount, MPI_Datatype sendtype, - int dest, int sendtag, IdefixCommArray recvbuf, int recvcount, +int MPI_Sendrecv(IdefixCommArray & sendbuf, int sendcount, MPI_Datatype sendtype, + int dest, int sendtag, IdefixCommArray &recvbuf, int recvcount, MPI_Datatype recvtype, int source, int recvtag, MPI_Comm comm, MPI_Status *status) { //check size assert(sendbuf.span() == sendcount); @@ -233,9 +199,7 @@ int MPI_Sendrecv(IdefixCommArray sendbuf, int sendcount, MPI_Datatype sendtyp assert(recvbuf.span_is_contiguous()); //make a local allocation if needed - #ifndef WITH_MPI_GPU_DIRECT - sendbuf.syncCommData(); - #endif + sendbuf.syncCommData(); //make the communication const int result = ::MPI_Sendrecv(sendbuf.commData(), sendcount, sendtype, dest, sendtag, @@ -243,9 +207,7 @@ int MPI_Sendrecv(IdefixCommArray sendbuf, int sendcount, MPI_Datatype sendtyp comm, status); //copy back if needed - #ifndef WITH_MPI_GPU_DIRECT - recvbuf.syncDeviceData(); - #endif + recvbuf.syncDeviceData(); //ok return result; diff --git a/src/utils/column.cpp b/src/utils/column.cpp index fb26a4596..06d083892 100644 --- a/src/utils/column.cpp +++ b/src/utils/column.cpp @@ -79,7 +79,7 @@ void Column::ComputeColumn(IdefixArray4D 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 @@ -134,7 +134,7 @@ void Column::ComputeColumn(IdefixArray4D in, const int var) { MPI_Status status; // Get the cumulative sum from previous processes Kokkos::fence(); - idefix::MPI_Recv(localSum, 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) { @@ -165,7 +165,7 @@ void Column::ComputeColumn(IdefixArray4D in, const int var) { } // And send it Kokkos::fence(); - idefix::MPI_Send(localSum,size, realMPI, dst, 20, ColumnComm); + idefix::MPI_Send(this->localSum,size, realMPI, dst, 20, ColumnComm); } // MPIrank small enough #endif // If we need it backwards @@ -195,7 +195,7 @@ void Column::ComputeColumn(IdefixArray4D in, const int var) { #ifdef WITH_MPI Kokkos::fence(); - idefix::MPI_Bcast(localSum,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