Allreduce algorithm, performance and codepath issue on ZE gpus

[kaushikvelusamy]

1. We observe a sudden abnormal increase ( 2-3x) in the collective communications with all reduce from 1MB and beyond with GPUs. You can reproduce this issue by measuring the time taken to complete 100 iterations of MPI_Allreduce() excluding the first and the second iterations in order to eliminate the initialization cost and include from 3-100 to get the measurement for only pure communications.

This issue is persistent irrespective of scale, different all reduce algorithms, CPU / NIC bindings , with and w/o MPI_Barriers, variability, C Vs SYCL and mpich versions.

2. I do see setting the --env MPIR_CVAR_ALLREDUCE_INTRA_ALGORITHM=recursive_doubling change the MPIR_CVAR_ALLREDUCE_INTRA_ALGORITHM=0 to MPIR_CVAR_ALLREDUCE_INTRA_ALGORITHM=3 but do not see any difference in the performance. Is there a possibility if the CVARs are not getting passed down or if any CVAR is modified during runtime?

3. Could you confirm if this mpich upstream is following the codepath in which it is using the Xelinks?

[raffenet]

#7070 was to address this issue, but currently the CVAR default value does not change the existing behavior. The question remains whether any MPICH module on Aurora will set the reduction threshold to an appropriate value as to not cause the poor performance observed in your tests.

[cbelusar]

Seems like we still need to close on if this should be default on Aurora for new builds of MPICH. If there are no perceived downsides, we can make this recommendation for future builds, but I'd like to confirm that there are no concerns before recommending that.

[raffenet]

Default collective selection has been improved since this issue. We should check the large message allreduce case again to see if still bad.

Overhaul of MPICH selection will address this, too.

[hzhou]

Finally revisited this issue. It turns out the culprit is the release_gather path selected by ch4-layer composition algorithms. The sudden drop of performance at 1MB is because ch4 composition performs a buffer swap for < 1MB messages, bypassing the issue. The fix in #7070, currently enabled at aurora_test, effectively does buffer swap for all message path, thus bypass the release_gather issue.

Actually, the correct algorithm for large GPU buffer should be the reduce_scatter_allgather (Rabenseifner's Algorithm) without buffer swapping.

Anyway, the algorithm currently developed in #7748 will out perform the reduce_scatter_allgather algorithm by 2x.

[colleeneb]

This issue seems similar to #7752 -- any issues with us closing this one to consolidate the "slow allreduce" ones?

[colleeneb]

I'm closing this to group allreduce issues into #7752, but we can reopen if needed or if anyone disagrees with the grouping.