AI/rocm-systems
2
0 Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

gtest: add scatter to combined calls and use loops (#303)

Browse Source 
* gtest: add scatter to combined calls and use loops

* gtest: run validation inside loop

* gtest: revert small element count to 2520

* gtest: fix memory leak in validation

(cherry picked from commit 36935cfbee)

* Fix combined call UT

* Fix memory leak

* Fix alltoallv test


[ROCm/rccl commit: b33a2cac8b]
This commit is contained in:
Wenkai Du authored and Wenkai Du committed 2021-01-14 19:28:01 -05:00
1 parent 2c49121171
commit 560224fe9f
8 files changed
+61 -37
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

projects/rccl/test/test_AllToAllv.cpp
+5 -5
View File
@@ -22,7 +22,7 @@ namespace CorrectnessTests
size_t chunksize = numElements*2/numDevices;
#define MAX_ALLTOALLV_RANKS 16
static size_t sendcounts[MAX_ALLTOALLV_RANKS], recvcounts[MAX_ALLTOALLV_RANKS], sdispls[MAX_ALLTOALLV_RANKS], rdispls[MAX_ALLTOALLV_RANKS];
static size_t sendcounts[MAX_ALLTOALLV_RANKS*MAX_ALLTOALLV_RANKS], recvcounts[MAX_ALLTOALLV_RANKS*MAX_ALLTOALLV_RANKS], sdispls[MAX_ALLTOALLV_RANKS*MAX_ALLTOALLV_RANKS], rdispls[MAX_ALLTOALLV_RANKS*MAX_ALLTOALLV_RANKS];
// Launch the reduction (1 thread per GPU)
ncclGroupStart();
for (int r = 0; r < numDevices; r++) {
@@ -31,12 +31,12 @@ namespace CorrectnessTests
size_t scount = ((i+r)%numDevices)*chunksize;
if (i+r == numDevices-1)
scount += (numElements*numDevices-chunksize*(numDevices-1)*numDevices/2);
sendcounts[i] = recvcounts[i] = scount;
sdispls[i] = rdispls[i] = disp;
sendcounts[i+r*MAX_ALLTOALLV_RANKS] = recvcounts[i+r*MAX_ALLTOALLV_RANKS] = scount;
sdispls[i+r*MAX_ALLTOALLV_RANKS] = rdispls[i+r*MAX_ALLTOALLV_RANKS] = disp;
disp += scount;
}
ncclAllToAllv((char*)dataset.inputs[r], sendcounts, sdispls,
(char*)dataset.outputs[r], recvcounts, rdispls, dataType, comms[r], streams[r]);
ncclAllToAllv((char*)dataset.inputs[r], sendcounts+r*MAX_ALLTOALLV_RANKS, sdispls+r*MAX_ALLTOALLV_RANKS,
(char*)dataset.outputs[r], recvcounts+r*MAX_ALLTOALLV_RANKS, rdispls+r*MAX_ALLTOALLV_RANKS, dataType, comms[r], streams[r]);
}
ncclGroupEnd();
// Wait for reduction to complete