Skip to content

gpu-direct: make it optional and have fallback solution with device-host explicit transfers#389

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
svalat wants to merge 28 commits into
base:develop
Choose a base branch
from
Draft

gpu-direct: make it optional and have fallback solution with device-host explicit transfers #389

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
28 commits
Select commit Hold shift + click to select a range
fec3e60
gpu-direct: handle communications when GPU DIRECT is not there.
svalat Jul 8, 2026
7967a1b
gpu-direct: Implement extensions of MPI API using views and provde th…
svalat Jul 9, 2026
a15c690
gpu-direct: use the MPI view semantic in EMF, axis and exganger
svalat Jul 9, 2026
4742f74
gpu-direct: rename the API to use a namespace and fallback on the ori…
svalat Jul 9, 2026
2af6af6
gpu-direct: fix some extra parts of idefix using MPI
svalat Jul 9, 2026
f38f106
gpu-direct: wrap all the internal MPI_.+ using .data() internally wit…
svalat Jul 10, 2026
7b89da6
gpu-direct: Add a macro to activate to validate by forcing the transf…
svalat Jul 10, 2026
9b58b1c
gpu-direct: fix an issue in MPI_Gather
svalat Jul 10, 2026
75a4c33
gpu-direct: Fix some compilation warnings with CUDA
svalat Jul 10, 2026
433c9c1
gpu-direct: fix mistake in lazy-init, should be done on all entry points
svalat Jul 10, 2026
17e0c76
gpu-direct: Ooops, no cannot lazy init due to the copy semantic
svalat Jul 10, 2026
411866f
gpu-direct: try to avoid the warning in normal way to run, need to fi…
svalat Jul 10, 2026
4a5b620
gpu-direct: Cleanup the way to switch the class ON/OFF by using a sin…
svalat Jul 10, 2026
afa7a4e
gpu-direct: Opps wrong name in option declaration
svalat Jul 10, 2026
03e0f40
gpu-direct: remove a useless scafolding
svalat Jul 10, 2026
76d46ff
gpu-direct: try to implement the copy/move constructor & operator= to…
svalat Jul 10, 2026
aab2e8a
gpu-direct: solve most of the compiler warnings when compiling with CUDA
svalat Jul 10, 2026
61b78f8
gpu-direct: Fix build of test/HD/KHI due to cyclic includes + fix mis…
svalat Aug 14, 2026
2e23ff7
gpu-direct: Fix pyidefix issues by adding the missing MPI wrappers
svalat Aug 14, 2026
7382ba1
gpu-direct: Fix a build issue due to assert variable names in debug mode
svalat Aug 14, 2026
bf3c87f
gpu-direct: fix an issue coming from the previous refactoring about a…
svalat Aug 14, 2026
9ab9f27
gpu-direct: Fix issue with test utils/dumpImage due to missing include
svalat Aug 14, 2026
3d9b17d
test: Add the possibility to test the MPI modes via test options (and…
svalat Aug 17, 2026
15c59bc
gpu-direct: fix none building case in exchanger about MPI_NON_BLOCKIN…
svalat Aug 17, 2026
5834964
gpu-direct: add -mpiGpuForceCopy on test script to run by disabling G…
svalat Aug 17, 2026
643e0c4
gpu-direct: Add testing of the NO_GPU_DIRECT feature by forceing the …
svalat Aug 17, 2026
5431e11
gpu-direct: remove the display of the MPI MODE status in the output o…
svalat Aug 18, 2026
93e5794
Merge branch 'develop' into feature/no-gpu-direct-support
svalat Aug 18, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
26 changes: 26 additions & 0 deletions CMakeLists.txt
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -16,6 +16,9 @@ option(Idefix_DEBUG "Enable Idefix debug features (makes the code very slow)" OF
option(Idefix_RUNTIME_CHECKS "Enable runtime sanity checks" OFF)
option(Idefix_WERROR "Treat compiler warnings as errors" OFF)
option(Idefix_PYTHON "Enable python bindings (requires pybind11)" OFF)
option(Idefix_MPI_GPU_DIRECT "Enable usage of GPU direct to avoid CPU-GPU copies for MPI messages." ON)
option(Idefix_MPI_GPU_FORCE_COPY "When GPU_DIRECT is disabled, force using transfers for validation purpose." OFF)
mark_as_advanced(Idefix_MPI_GPU_FORCE_TRANSFERS)
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")
Expand All@@ -35,6 +38,10 @@ set_property(CACHE Idefix_PRECISION PROPERTY STRINGS Double Single)
set(Idefix_LOOP_PATTERN "Default" CACHE STRING "Loop pattern for idefix_for")
set_property(CACHE Idefix_LOOP_PATTERN PROPERTY STRINGS Default SIMD Range MDRange TeamPolicy TeamPolicyInnerVector)

set(Idefix_MPI_MODE "Persistent" CACHE STRING "MPI communication mode to use.")
set_property(CACHE Idefix_MPI_MODE PROPERTY STRINGS Persistent Blocking NonBlocking)
mark_as_advanced(Idefix_MPI_MODE)

# load git revision tools
list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/")
include(GetGitRevisionDescription)
Expand DownExpand Up@@ -104,6 +111,10 @@ if(Idefix_MPI)
add_subdirectory(src/mpi)
endif()

if (Idefix_MPI_GPU_DIRECT)
add_compile_definitions("WITH_MPI_GPU_DIRECT")
endif()

if(Idefix_HDF5)
add_compile_definitions("WITH_HDF5")
if(Idefix_MPI)
Expand DownExpand Up@@ -227,6 +238,20 @@ if(${Idefix_PRECISION} STREQUAL "Single")
add_compile_definitions("SINGLE_PRECISION")
endif()

# MPI mode
if(${Idefix_MPI_MODE} STREQUAL "Persistent")
add_compile_definitions("COMMUNICATION_MODE_PERSISTENT")
elseif(${Idefix_MPI_MODE} STREQUAL "NonBlocking")
add_compile_definitions("COMMUNICATION_MODE_NON_BLOCKING")
else(NOT ${Idefix_MPI_MODE} STREQUAL "Blocking")
message(ERROR "Unknown MPI communication mode : Idefix_MPI_MODE=${Idefix_MPI_MODE}")
endif()

# GPU direct disabled and force transfers
if (${Idefix_MPI_GPU_FORCE_COPY})
add_compile_definitions("WITH_GPU_FORCE_COPY")
endif()

target_include_directories(idefix PUBLIC
"${Idefix_PROBLEM_DIR_ABS}"
)
Expand DownExpand Up@@ -264,6 +289,7 @@ endif()
message(STATUS " MPI: ${Idefix_MPI}")
message(STATUS " HDF5: ${Idefix_HDF5}")
message(STATUS " Python: ${Idefix_PYTHON}")
message(STATUS " GPU direct: ${Idefix_MPI_GPU_DIRECT}")
message(STATUS " Reconstruction: ${Idefix_RECONSTRUCTION}")
message(STATUS " Precision: ${Idefix_PRECISION}")
message(STATUS " Version: ${Idefix_VERSION}")
Expand Down
21 changes: 21 additions & 0 deletions pytools/idfx_test.py
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -92,6 +92,20 @@ def __init__(self, current_test_file, name=""):

parser.add_argument("-mpi", help="Enable MPI", action="store_true")

parser.add_argument(
"-mpiMode",
default="Persistent",
choices=["Default", "Persistent", "NonBlocking", "Blocking"],
help="MPI communication mode to use (Default)",
type=str,
)

parser.add_argument(
"-mpiGpuForceCopy",
action="store_true",
help="Disable MPI GPU direct and force memory transfers for validation purpose.",
)

parser.add_argument(
"-all",
help="Do all test suite (otherwise, just do the test with the current configuration)",
Expand DownExpand Up@@ -268,6 +282,13 @@ def _genCmakeCommand(self, definitionFile=""):
else:
comm.append("-DIdefix_MPI=OFF")

if self.mpiMode != "Default":
comm.append(f"-DIdefix_MPI_MODE={self.mpiMode}")

if self.mpiGpuForceCopy:
comm.append("-DIdefix_MPI_GPU_DIRECT=OFF")
comm.append("-DIdefix_MPI_GPU_FORCE_COPY=ON")

if self.reconstruction == 2:
comm.append("-DIdefix_RECONSTRUCTION=Linear")
elif self.reconstruction == 3:
Expand Down
31 changes: 15 additions & 16 deletions src/fluid/boundary/axis.cpp
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -25,7 +25,7 @@ void Axis::ShowConfig() {
void Axis::SymmetrizeEx1Side(int jref, IdefixArray3D<real> Ex1) {
#if DIMENSIONS == 3

IdefixArray1D<real> Ex1Avg = this->Ex1Avg;
auto Ex1Avg = this->Ex1Avg.deviceView();

idefix_for("Ex1_ini",0,data->np_tot[IDIR],
KOKKOS_LAMBDA(int i) {
Expand All@@ -40,7 +40,7 @@ void Axis::SymmetrizeEx1Side(int jref, IdefixArray3D<real> Ex1) {
#ifdef WITH_MPI
Kokkos::fence();
// sum along all of the processes on the same r
MPI_Allreduce(MPI_IN_PLACE, Ex1Avg.data(), data->np_tot[IDIR], realMPI,
idefix::MPI_Allreduce(MPI_IN_PLACE, this->Ex1Avg, data->np_tot[IDIR], realMPI,
MPI_SUM, data->mygrid->AxisComm);
#endif
}
Expand DownExpand Up@@ -88,7 +88,8 @@ void Axis::RegularizeCurrentSide(int side) {
jc = data->end[JDIR]-1;
sign = -1;
}
IdefixArray1D<real> BAvg = this->Ex1Avg;
auto BAvg = this->Ex1Avg.deviceView();
auto BAvgComm = this->Ex1Avg;
IdefixArray1D<real> x1 = data->x[IDIR];
IdefixArray1D<real> dx3 = data->dx[KDIR];
IdefixArray1D<real> dx2 = data->dx[JDIR];
Expand All@@ -107,7 +108,7 @@ void Axis::RegularizeCurrentSide(int side) {
#ifdef WITH_MPI
Kokkos::fence();
// sum along all of the processes on the same r
MPI_Allreduce(MPI_IN_PLACE, BAvg.data(), data->np_tot[IDIR], realMPI,
MPI_Allreduce(MPI_IN_PLACE, BAvgComm, data->np_tot[IDIR], realMPI,
MPI_SUM, data->mygrid->AxisComm);
#endif
}
Expand DownExpand Up@@ -169,7 +170,7 @@ void Axis::FixBx2sAxis(int side) {
// Compute the values of Bx and By that are consistent with BX2 along the axis
#if DIMENSIONS == 3
IdefixArray4D<real> Vs = this->Vs;
IdefixArray2D<real> BAvg = this->BAvg;
auto BAvg = this->BAvg.deviceView();
IdefixArray1D<real> phi = data->x[KDIR];

int jin = 0;
Expand DownExpand Up@@ -208,7 +209,7 @@ void Axis::FixBx2sAxis(int side) {
Kokkos::fence();
#ifdef WITH_MPI
// sum along all of the processes on the same r
MPI_Allreduce(MPI_IN_PLACE, BAvg.data(), 2*data->np_tot[IDIR], realMPI,
idefix::MPI_Allreduce(MPI_IN_PLACE, this->BAvg, 2*data->np_tot[IDIR], realMPI,
MPI_SUM, data->mygrid->AxisComm);
#endif
}
Expand DownExpand Up@@ -392,7 +393,6 @@ void Axis::ExchangeMPI(int side) {
idfx::pushRegion("Axis::ExchangeMPI");
#ifdef WITH_MPI
// Load the buffers with data
[[maybe_unused]] int ibeg,iend,jbeg,jend,kbeg,kend;
int offset;
int ny;
Buffer bufferSend = this->bufferSend;
Expand All@@ -409,7 +409,7 @@ void Axis::ExchangeMPI(int side) {
MPI_Status recvStatus;

double tStart = MPI_Wtime();
MPI_SAFE_CALL(MPI_Start(&recvRequest));
MPI_SAFE_CALL(idefix::MPI_Start(&recvRequest));
idfx::mpiCallsTimer += MPI_Wtime() - tStart;

// Coordinates of the ghost region which needs to be transfered
Expand DownExpand Up@@ -471,8 +471,8 @@ void Axis::ExchangeMPI(int side) {
Kokkos::fence();

tStart = MPI_Wtime();
MPI_SAFE_CALL(MPI_Start(&sendRequest));
MPI_Wait(&recvRequest,&recvStatus);
MPI_SAFE_CALL(idefix::MPI_Start(&sendRequest));
idefix::MPI_Wait(&recvRequest,&recvStatus);
idfx::mpiCallsTimer += MPI_Wtime() - tStart;

// Unpack
Expand DownExpand Up@@ -526,8 +526,7 @@ void Axis::ExchangeMPI(int side) {
} // MHD
}

MPI_Wait(&sendRequest, &sendStatus);

idefix::MPI_Wait(&sendRequest, &sendStatus);
idfx::mpiCallsTimer += MPI_Wtime() - tStart;


Expand DownExpand Up@@ -594,11 +593,11 @@ void Axis::InitMPI() {
MPI_SAFE_CALL(MPI_Cart_shift(data->mygrid->AxisComm,0,data->mygrid->nproc[KDIR]/2,
&procRecv,&procSend ));

MPI_SAFE_CALL(MPI_Send_init(bufferSend.data(), bufferSend.Size(), realMPI, procSend,
650, data->mygrid->AxisComm, &sendRequest));
MPI_SAFE_CALL(idefix::MPI_Send_init(bufferSend.commView(), bufferSend.Size(),
realMPI, procSend, 650, data->mygrid->AxisComm, &sendRequest));

MPI_SAFE_CALL(MPI_Recv_init(bufferRecv.data(), bufferRecv.Size(), realMPI, procRecv,
650, data->mygrid->AxisComm, &recvRequest));
MPI_SAFE_CALL(idefix::MPI_Recv_init(bufferRecv.commView(), bufferRecv.Size(),
realMPI, procRecv, 650, data->mygrid->AxisComm, &recvRequest));

#endif
idfx::popRegion();
Expand Down
12 changes: 6 additions & 6 deletions src/fluid/boundary/axis.hpp
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -57,8 +57,8 @@ class Axis {

enum {faceTop, faceBot};
#ifdef WITH_MPI
MPI_Request sendRequest;
MPI_Request recvRequest;
idefix::MPI_Request_1D<real> sendRequest;
idefix::MPI_Request_1D<real> recvRequest;

Buffer bufferSend;
Buffer bufferRecv;
Expand All@@ -71,8 +71,8 @@ class Axis {
#endif
void InitMPI();

IdefixArray1D<real> Ex1Avg;
IdefixArray2D<real> BAvg;
idefix::IdefixCommArray1D<real> Ex1Avg;
idefix::IdefixCommArray2D<real> BAvg;
bool haveCurrent;
IdefixArray2D<real> JAvg;
IdefixArray1D<int> symmetryVc;
Expand DownExpand Up@@ -155,8 +155,8 @@ Axis::Axis(Boundary<Phys> *boundary) {
}
Kokkos::deep_copy(symmetryVs, symmetryVsHost);

this->Ex1Avg = IdefixArray1D<real>("Axis:Ex1Avg",data->np_tot[IDIR]);
this->BAvg = IdefixArray2D<real>("Axis:BxAvg",data->np_tot[IDIR],2);
this->Ex1Avg = idefix::IdefixCommArray1D<real>("Axis:Ex1Avg",data->np_tot[IDIR]);
this->BAvg = idefix::IdefixCommArray2D<real>("Axis:BxAvg",data->np_tot[IDIR],2);
if(haveCurrent) {
this->JAvg = IdefixArray2D<real>("Axis:JAvg",data->np_tot[IDIR],3);
}
Expand Down
25 changes: 12 additions & 13 deletions src/fluid/constrainedTransport/EMFexchange.hpp
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -41,7 +41,7 @@ void ConstrainedTransport<Phys>::ExchangeX1(IdefixArray3D<real> ey, IdefixArray3
MPI_Status recvStatus[2];

double tStart = MPI_Wtime();
MPI_SAFE_CALL(MPI_Startall(2, recvRequestX1));
MPI_SAFE_CALL(idefix::MPI_Startall(2, recvRequestX1));
idfx::mpiCallsTimer += MPI_Wtime() - tStart;

BoundaryType lbound = data->lbound[IDIR];
Expand DownExpand Up@@ -82,11 +82,10 @@ void ConstrainedTransport<Phys>::ExchangeX1(IdefixArray3D<real> ey, IdefixArray3
Kokkos::fence();

tStart = MPI_Wtime();
MPI_SAFE_CALL(MPI_Startall(2, sendRequestX1));
MPI_SAFE_CALL(idefix::MPI_Startall(2, sendRequestX1));
// Wait for buffers to be received

MPI_Waitall(2,recvRequestX1,recvStatus);
MPI_Waitall(2, sendRequestX1, sendStatus);
idefix::MPI_Waitall(2,recvRequestX1,recvStatus);
idefix::MPI_Waitall(2, sendRequestX1, sendStatus);
idfx::mpiCallsTimer += MPI_Wtime() - tStart;

// Unpack
Expand DownExpand Up@@ -130,7 +129,7 @@ void ConstrainedTransport<Phys>::ExchangeX2(IdefixArray3D<real> ex, IdefixArray3
double tStart = MPI_Wtime();
MPI_Status sendStatus[2];
MPI_Status recvStatus[2];
MPI_SAFE_CALL(MPI_Startall(2, recvRequestX2));
MPI_SAFE_CALL(idefix::MPI_Startall(2, recvRequestX2));
idfx::mpiCallsTimer += MPI_Wtime() - tStart;

BoundaryType lbound = data->lbound[JDIR];
Expand DownExpand Up@@ -169,10 +168,10 @@ void ConstrainedTransport<Phys>::ExchangeX2(IdefixArray3D<real> ex, IdefixArray3
Kokkos::fence();

tStart = MPI_Wtime();
MPI_SAFE_CALL(MPI_Startall(2, sendRequestX2));
MPI_SAFE_CALL(idefix::MPI_Startall(2, sendRequestX2));
// Wait for buffers to be received
MPI_Waitall(2, recvRequestX2, recvStatus);
MPI_Waitall(2, sendRequestX2, sendStatus);
idefix::MPI_Waitall(2, recvRequestX2, recvStatus);
idefix::MPI_Waitall(2, sendRequestX2, sendStatus);
idfx::mpiCallsTimer += MPI_Wtime() - tStart;

// Unpack
Expand DownExpand Up@@ -216,7 +215,7 @@ void ConstrainedTransport<Phys>::ExchangeX3(IdefixArray3D<real> ex, IdefixArray3
double tStart = MPI_Wtime();
MPI_Status sendStatus[2];
MPI_Status recvStatus[2];
MPI_SAFE_CALL(MPI_Startall(2, recvRequestX3));
MPI_SAFE_CALL(idefix::MPI_Startall(2, recvRequestX3));
idfx::mpiCallsTimer += MPI_Wtime() - tStart;

BoundaryType lbound = data->lbound[KDIR];
Expand DownExpand Up@@ -256,10 +255,10 @@ void ConstrainedTransport<Phys>::ExchangeX3(IdefixArray3D<real> ex, IdefixArray3
Kokkos::fence();

tStart = MPI_Wtime();
MPI_SAFE_CALL(MPI_Startall(2, sendRequestX3));
MPI_SAFE_CALL(idefix::MPI_Startall(2, sendRequestX3));
// Wait for buffers to be received
MPI_Waitall(2, recvRequestX3, recvStatus);
MPI_Waitall(2, sendRequestX3, sendStatus);
idefix::MPI_Waitall(2, recvRequestX3, recvStatus);
idefix::MPI_Waitall(2, sendRequestX3, sendStatus);
idfx::mpiCallsTimer += MPI_Wtime() - tStart;

// Unpack
Expand Down
Loading
Loading