diff --git a/test/CorrectnessTest.hpp b/test/CorrectnessTest.hpp index 95010ee701..3c7d6722b9 100644 --- a/test/CorrectnessTest.hpp +++ b/test/CorrectnessTest.hpp @@ -121,7 +121,7 @@ namespace CorrectnessTests // Explicit memory release to avoid double-free from subDatasets void Release() { - for (int i = 0; i < outputs.size(); i++) + for (int i = 0; i < numDevices; i++) { if (!inPlace) hipFree(outputs[i]); hipFree(inputs[i]); @@ -443,6 +443,7 @@ namespace CorrectnessTests } ASSERT_EQ(isMatch, true); } + free(outputI1); } // Passed in parameters from TestTuple diff --git a/test/test_AllGather.hpp b/test/test_AllGather.hpp index 615616f36d..06a0297c3d 100644 --- a/test/test_AllGather.hpp +++ b/test/test_AllGather.hpp @@ -25,6 +25,8 @@ namespace CorrectnessTests for (int i = 0; i < dataset.numDevices; i++) memcpy(dataset.expected[i], result, dataset.NumBytes()); + + free(result); } }; } diff --git a/test/test_AllReduceAbort.cpp b/test/test_AllReduceAbort.cpp index ba63ea9004..23d7129814 100644 --- a/test/test_AllReduceAbort.cpp +++ b/test/test_AllReduceAbort.cpp @@ -114,6 +114,7 @@ namespace CorrectnessTests if (idle) pthread_yield(); } + free(done); HIPCHECK(hipHostFree(fake_head)); HIPCHECK(hipStreamDestroy(stream)); dataset.Release(); diff --git a/test/test_AllToAllv.cpp b/test/test_AllToAllv.cpp index 6f2e317f44..fe78b0f352 100644 --- a/test/test_AllToAllv.cpp +++ b/test/test_AllToAllv.cpp @@ -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 diff --git a/test/test_BroadcastAbort.cpp b/test/test_BroadcastAbort.cpp index 96f1de6ccf..ddab18a8e9 100644 --- a/test/test_BroadcastAbort.cpp +++ b/test/test_BroadcastAbort.cpp @@ -116,6 +116,7 @@ namespace CorrectnessTests if (idle) pthread_yield(); } + free(done); HIPCHECK(hipHostFree(fake_head)); HIPCHECK(hipStreamDestroy(stream)); dataset.Release(); diff --git a/test/test_CombinedCalls.cpp b/test/test_CombinedCalls.cpp index db4f938fc0..e7be7385a1 100644 --- a/test/test_CombinedCalls.cpp +++ b/test/test_CombinedCalls.cpp @@ -10,6 +10,7 @@ #include "test_Broadcast.hpp" #include "test_Reduce.hpp" #include "test_ReduceScatter.hpp" +#include "test_Scatter.hpp" namespace CorrectnessTests { @@ -25,6 +26,10 @@ namespace CorrectnessTests FillDatasetWithPattern(datasets[i]); } + Dataset scatter_dataset; + scatter_dataset.Initialize(numDevices, numElements, dataType, inPlace, ncclCollScatter); + FillDatasetWithPattern(scatter_dataset); + // Compute expected results for each dataset in combined int const root = 0; AllGatherCorrectnessTest::ComputeExpectedResults(datasets[0]); @@ -32,46 +37,60 @@ namespace CorrectnessTests BroadcastCorrectnessTest::ComputeExpectedResults(datasets[2], root); ReduceCorrectnessTest::ComputeExpectedResults(datasets[3], op, root); ReduceScatterCorrectnessTest::ComputeExpectedResults(datasets[4], op); + ScatterCorrectnessTest::ComputeExpectedResults(scatter_dataset, root); size_t const byteCount = datasets[0].NumBytes() / numDevices; size_t const elemCount = numElements / numDevices; - ncclGroupStart(); - for (int i = 0; i < numDevices; i++) + for (int j = 0; j < 10; j++) { - ncclAllGather((int8_t *)datasets[0].inputs[i] + (i * byteCount), - datasets[0].outputs[i], elemCount, - dataType, comms[i], streams[i]); + ncclGroupStart(); + for (int i = 0; i < numDevices; i++) + { + ncclScatter(scatter_dataset.inputs[i], + scatter_dataset.outputs[i], + numElements, dataType, + root, comms[i], streams[i]); + } + ncclGroupEnd(); + ncclGroupStart(); + for (int i = 0; i < numDevices; i++) + { + ncclAllGather((int8_t *)datasets[0].inputs[i] + (i * byteCount), + datasets[0].outputs[i], elemCount, + dataType, comms[i], streams[i]); - ncclAllReduce(datasets[1].inputs[i], datasets[1].outputs[i], - numElements, dataType, op, comms[i], streams[i]); + ncclAllReduce(datasets[1].inputs[i], datasets[1].outputs[i], + numElements, dataType, op, comms[i], streams[i]); - ncclBroadcast(datasets[2].inputs[i], - datasets[2].outputs[i], - numElements, dataType, - root, comms[i], streams[i]); + ncclBroadcast(datasets[2].inputs[i], + datasets[2].outputs[i], + numElements, dataType, + root, comms[i], streams[i]); - ncclReduce(datasets[3].inputs[i], - datasets[3].outputs[i], - numElements, dataType, op, - root, comms[i], streams[i]); + ncclReduce(datasets[3].inputs[i], + datasets[3].outputs[i], + numElements, dataType, op, + root, comms[i], streams[i]); - ncclReduceScatter(datasets[4].inputs[i], - (int8_t *)datasets[4].outputs[i] + (i * byteCount), - elemCount, dataType, op, - comms[i], streams[i]); + ncclReduceScatter(datasets[4].inputs[i], + (int8_t *)datasets[4].outputs[i] + (i * byteCount), + elemCount, dataType, op, + comms[i], streams[i]); + } + ncclGroupEnd(); + // Wait for reduction to complete + Synchronize(); + // Check results for each collective in the combined + for (int i = 0; i < 5; i++) + ValidateResults(datasets[i]); + + ValidateResults(scatter_dataset); } - ncclGroupEnd(); - // Wait for reduction to complete - Synchronize(); - - // Check results for each collective in the combined for (int i = 0; i < 5; i++) - { - ValidateResults(datasets[i]); datasets[i].Release(); - } + scatter_dataset.Release(); } INSTANTIATE_TEST_SUITE_P(CombinedCallsCorrectnessSweep, @@ -95,7 +114,7 @@ namespace CorrectnessTests // Number of devices testing::Values(2,3,4,5,6,7,8), // In-place or not - testing::Values(false, true), - testing::Values("")), + testing::Values(false), + testing::Values("RCCL_ALLTOALL_KERNEL_DISABLE=1")), CorrectnessTest::PrintToStringParamName()); } // namespace diff --git a/test/test_Gather.cpp b/test/test_Gather.cpp index 8169127966..bf917c739d 100644 --- a/test/test_Gather.cpp +++ b/test/test_Gather.cpp @@ -66,6 +66,6 @@ namespace CorrectnessTests testing::Values(2,3,4,5,6,7,8), // In-place or not testing::Values(false), - testing::Values("RCCL_ALLTOALL_KERNEL_DISABLE=0", "RCCL_ALLTOALL_KERNEL_DISABLE=1")), + testing::Values("RCCL_ALLTOALL_KERNEL_DISABLE=1")), CorrectnessTest::PrintToStringParamName()); } // namespace diff --git a/test/test_Scatter.cpp b/test/test_Scatter.cpp index 78cfa01ec3..aaca748217 100644 --- a/test/test_Scatter.cpp +++ b/test/test_Scatter.cpp @@ -66,6 +66,6 @@ namespace CorrectnessTests testing::Values(2,3,4,5,6,7,8), // In-place or not testing::Values(false), - testing::Values("RCCL_ALLTOALL_KERNEL_DISABLE=0", "RCCL_ALLTOALL_KERNEL_DISABLE=1")), + testing::Values("RCCL_ALLTOALL_KERNEL_DISABLE=1")), CorrectnessTest::PrintToStringParamName()); } // namespace