gpu-direct: make it optional and have fallback solution with device-host explicit transfers by svalat · Pull Request #389 · idefix-code/idefix

svalat · 2026-07-10T15:20:35Z

Handle when there is no-gpu-direct

The patch mostly wrap the MPI interface to make the transfers manually between the host & device when GPU direct is not available.

In order to avoid allocating an host buffer for every transfers, the IdefixArrayXD classes has been extended as IdefixCommArrayXD which contains the standard array on device, and optionally (if GPU direct is disabled and if running on GPU) a host buffer to make the communications.

In order to apply in the code, all the MPI communication calls (MPI_Send, MPI_Recv, MPI_Start....) need to be replaced by the idefix::MPI_* corresponding frunctions which expect an IdefixCommArray as parameter.

It means that we should not continue to use .data() to feed the MPI calls (the compiler will complain for the idefix::MPI_* functions).

When the wrapper is called it ensure to sync the communication buffer in & out during the communication process so it is transparent for the calling code.

One care

You should not use directly the IdefixCommArrayXDin a GPU loop (get a compiler warning in case due to host array in use). This is way the patch uses the .deviceView() to get it casted in a friendly way for the GPU.

Note: in case you still keep it that way, it should not be a problem appart the warning in the compilation output because the host array will not be used in the kernel because not accessible (private).

Note: this semantic is only used if we communicate with the buffer, all the rest of the code using standard IdefixArray is not concerned.

Case of persitent comms

For the persistent comms and async comms, an new idefix::MPI_Request_XD has been defined to carry the information on how to sync the communication buffer with its device counter part when required. Again if you use idefix::MPI_* just follow the compiler which will force you to type correctly using idefix::MPI_Request_XD instead of MPI_Request.

New CMake flags :

Idefix_MPI_GPU_DIRECT : Enable the fallback imlpementation at compile time (will apply only if at runtime we cannot GPU direct and if we run on GPU.

Variable marked as advanced (not seen by default in ccmake) :

Idefix_MPI_GPU_FORCE_COPY : For validation purpose, force the copy in any situation (even on CPU). It should come in conjunition with Idefix_MPI_GPU_DIRECT to have an effect.
Idefix_MPI_MODE: To select a compile time via cmake which communication mode to use. Can be Persistent, Blocking, NonBlocking.

Validation

In order to validate the implementation in any condition, we can :

Disable GPU direct with the CMake variable : Idefix_MPI_GPU_DIRECT.
Force the copy on communications in any case via : Idefix_MPI_GPU_FORCE_COPY.

This is done on the test suite by using the -mpiGpuForceCopy option.

PYTHONPATH=$PWD ./test.py -subdir ./test/MHD/ -k "mpi" -mpiGpuForceCopy -cuda

It can be played in conjuntion with the -mpiMode option which can be either : Persistent, NonBlocking, BLocking to validate the various paths.

Validation in github actions

The test suite has been extended to test automatically those combinations on the test : MHD/AxisFluxTube.

When having GPU direct

in this case, the IdefixCommArrayGpuDirect implementation is used, which is just an empty stub which just does nothing.

Questions

When making a MPI_waitall with multiple comms, I start all the host-device transfers asynchronously and using Kokkos::fence(). I can also avoid it to get simpler code as we have max 2 comms at this location.
I currently named the MPI_* API the same inside the idefix:: namespace, we can also name them all MPI_View_ to clearly separate the concept. To be discussed what is best.
I patched Mpi::CheckConfig() so it uses the copies when not having GPU Direct, is it what we want ?

TODO

Fix the issue on Axis cases
Validate all configuration (mpi sync/async/persistent)
Validate on a non gpu-direct condition.
Some minor cleanup

…e IdefixCommArrayXD<> to handle the no gpu-direct transfers.

…ginal function names

…h idefix::MPI_ API.

…ers all time for validation purpose

…nd a solution for the NO GPU DIRECT case

…gle #if/else

… solve the Axis issue

…sing #if WITH_MPI when MPI is disabled

…ccessing a GPU buffer on CPU

… cmake one) : Persistent, Blocking, NonBlocking

…G comms

…PU_DIRECT and forcing copy in any case for validation

…copy in some tests

…f idefix

…FF in related error messages

glesur

a few relatively minor things.

What is unclear at this stage is whether we want to accept all possible input data for MPI function, and not just the ones involving IdefixCommArray?

glesur · 2026-08-24T08:30:26Z

+
+#include "idefix.hpp"
+
+namespace idefix {


Suggested change

namespace idefix {

namespace idfx {

glesur · 2026-08-24T08:31:01Z

-  IdefixArray2D<real>     sbEyL;
-  IdefixArray2D<real>     sbEyR;
-  IdefixArray2D<real>     sbEyRL;
+  idefix::IdefixCommArray2D<real> sbEyL;


name space should be idfx:: (like cout, and other global functionalities), or idfxMPI:: (to highlight this is our own implementation)

glesur · 2026-08-26T08:00:44Z

      #endif
+      #if defined(KOKKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
+        errmsg << "You can look on Idefix cmake option -DIdefix_MPI_GPU_DIRECT=OFF to not"
+                  "use it (will be slower)." << std::endl;


Suggested change

"use it (will be slower)." << std::endl;

glesur · 2026-08-26T08:03:43Z

  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,


cette réécriture n'est pas utile car node_coord est déjà sur l'hote

glesur · 2026-08-26T08:10:58Z

      const double allowedImbalance = 20.0;
      std::vector<double> computeLogPerCore(idfx::psize);
-      MPI_Gather(&computeLastLog, 1, MPI_DOUBLE, computeLogPerCore.data(), 1, MPI_DOUBLE, 0,
+      idefix::MPI_Gather(&computeLastLog, 1, MPI_DOUBLE, computeLogPerCore, 1, MPI_DOUBLE, 0,


comprends pas: computeLogPerCore est un std::vector, pas un IdefixArray

glesur · 2026-08-26T08:12:57Z

-          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);


tout ces tableaux sont des vecteurs sur l'hote, je ne comprends même pas comment ça peut marcher (mais cette routine n'est probablement pas testée)

glesur · 2026-08-26T08:16:28Z

+ * CPU / GPU transfers if the array in on GPU and WITH_MPI_GPU_DIRECT is disabled.
+ */
+template <class T, size_t U>
+int MPI_Send(const std::array<T, U> & buf, int count, MPI_Datatype datatype,


pas sur de comprendre l'intéret de cette fonction. Tout encapsuler dans ton API, même pour des fonctions qui n'utilisent pas le GPU aware?

Co-authored-by: Geoffroy Lesur <geoffroy.lesur@univ-grenoble-alpes.fr>

svalat added 17 commits July 8, 2026 15:34

gpu-direct: handle communications when GPU DIRECT is not there.

fec3e60

gpu-direct: Implement extensions of MPI API using views and provde th…

7967a1b

…e IdefixCommArrayXD<> to handle the no gpu-direct transfers.

gpu-direct: use the MPI view semantic in EMF, axis and exganger

a15c690

gpu-direct: rename the API to use a namespace and fallback on the ori…

4742f74

…ginal function names

gpu-direct: fix some extra parts of idefix using MPI

2af6af6

gpu-direct: wrap all the internal MPI_.+ using .data() internally wit…

f38f106

…h idefix::MPI_ API.

gpu-direct: Add a macro to activate to validate by forcing the transf…

7b89da6

…ers all time for validation purpose

gpu-direct: fix an issue in MPI_Gather

9b58b1c

gpu-direct: Fix some compilation warnings with CUDA

75a4c33

gpu-direct: fix mistake in lazy-init, should be done on all entry points

433c9c1

gpu-direct: Ooops, no cannot lazy init due to the copy semantic

17e0c76

gpu-direct: try to avoid the warning in normal way to run, need to fi…

411866f

…nd a solution for the NO GPU DIRECT case

gpu-direct: Cleanup the way to switch the class ON/OFF by using a sin…

4a5b620

…gle #if/else

gpu-direct: Opps wrong name in option declaration

afa7a4e

gpu-direct: remove a useless scafolding

03e0f40

gpu-direct: try to implement the copy/move constructor & operator= to…

76d46ff

… solve the Axis issue

gpu-direct: solve most of the compiler warnings when compiling with CUDA

aab2e8a

svalat marked this pull request as draft July 10, 2026 15:21

svalat added 9 commits August 14, 2026 11:21

gpu-direct: Fix build of test/HD/KHI due to cyclic includes + fix mis…

61b78f8

…sing #if WITH_MPI when MPI is disabled

gpu-direct: Fix pyidefix issues by adding the missing MPI wrappers

2e23ff7

gpu-direct: Fix a build issue due to assert variable names in debug mode

7382ba1

gpu-direct: fix an issue coming from the previous refactoring about a…

bf3c87f

…ccessing a GPU buffer on CPU

gpu-direct: Fix issue with test utils/dumpImage due to missing include

9ab9f27

test: Add the possibility to test the MPI modes via test options (and…

3d9b17d

… cmake one) : Persistent, Blocking, NonBlocking

gpu-direct: fix none building case in exchanger about MPI_NON_BLOCKIN…

15c59bc

…G comms

gpu-direct: add -mpiGpuForceCopy on test script to run by disabling G…

5834964

…PU_DIRECT and forcing copy in any case for validation

gpu-direct: Add testing of the NO_GPU_DIRECT feature by forceing the …

643e0c4

…copy in some tests

glesur self-requested a review August 18, 2026 07:28

gpu-direct: remove the display of the MPI MODE status in the output o…

5431e11

…f idefix

svalat force-pushed the feature/no-gpu-direct-support branch from 1e989f8 to 5431e11 Compare August 18, 2026 08:28

Merge branch 'develop' into feature/no-gpu-direct-support

93e5794

svalat force-pushed the feature/no-gpu-direct-support branch from 04913f8 to f24de7d Compare August 19, 2026 09:05

svalat added 2 commits August 19, 2026 11:05

gpu-direct: fix the build issues with Intel compiler.

f24de7d

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

f6bf6e2

svalat marked this pull request as ready for review August 20, 2026 08:12

svalat force-pushed the feature/no-gpu-direct-support branch from 001c28a to e5b4522 Compare August 20, 2026 13:13

gpu-direct: Update the doc and inform about -DIdefix_MPI_GPU_DIRECT=O…

43d4040

…FF in related error messages

svalat force-pushed the feature/no-gpu-direct-support branch from e5b4522 to 43d4040 Compare August 20, 2026 13:13

gpu-direct: fix pre-commit issue

a29cbd1

svalat force-pushed the feature/no-gpu-direct-support branch from c3991ca to a29cbd1 Compare August 20, 2026 13:24

glesur requested changes Aug 26, 2026

View reviewed changes

Apply suggestions from code review

7b21d3d

Co-authored-by: Geoffroy Lesur <geoffroy.lesur@univ-grenoble-alpes.fr>

Navigation Menu

Uh oh!

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

gpu-direct: make it optional and have fallback solution with device-host explicit transfers#389
svalat wants to merge 33 commits into
idefix-code:developfrom
svalat:feature/no-gpu-direct-support

svalat commented Jul 10, 2026 •

edited

Loading

Uh oh!

glesur left a comment

Uh oh!

Uh oh!

Uh oh!

glesur Aug 24, 2026

Uh oh!

glesur Aug 24, 2026

Uh oh!

Uh oh!

glesur Aug 26, 2026

Uh oh!

glesur Aug 26, 2026

Uh oh!

glesur Aug 26, 2026

Uh oh!

glesur Aug 26, 2026

Uh oh!

glesur Aug 26, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

svalat commented Jul 10, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Handle when there is no-gpu-direct

One care

Case of persitent comms

New CMake flags :

Validation

Validation in github actions

When having GPU direct

Questions

TODO

Uh oh!

glesur left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

glesur Aug 24, 2026

Choose a reason for hiding this comment

Uh oh!

glesur Aug 24, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

glesur Aug 26, 2026

Choose a reason for hiding this comment

Uh oh!

glesur Aug 26, 2026

Choose a reason for hiding this comment

Uh oh!

glesur Aug 26, 2026

Choose a reason for hiding this comment

Uh oh!

glesur Aug 26, 2026

Choose a reason for hiding this comment

Uh oh!

glesur Aug 26, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

svalat commented Jul 10, 2026 •

edited

Loading