AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad

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

Explorar el Código 
* 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 b0853ccd51)

* Fix combined call UT

* Fix memory leak

* Fix alltoallv test
Este commit está contenido en:
Wenkai Du
2020-11-13 17:57:44 -08:00
cometido por Wenkai Du
padre f4d5d3d620
commit b33a2cac8b
Se han modificado 8 ficheros con 61 adiciones y 37 borrados
Descargar archivo de parche Descargar archivo de diferencias Expandir todos los archivos Contraer todos los archivos
test/CorrectnessTest.hpp
+2 -1
Ver fichero
@@ -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
test/test_AllGather.hpp
+2
Ver fichero
@@ -25,6 +25,8 @@ namespace CorrectnessTests
for (int i = 0; i < dataset.numDevices; i++)
memcpy(dataset.expected[i], result, dataset.NumBytes());
free(result);
}
};
}
test/test_AllReduceAbort.cpp
+1
Ver fichero
@@ -114,6 +114,7 @@ namespace CorrectnessTests
if (idle) pthread_yield();
}
free(done);
HIPCHECK(hipHostFree(fake_head));
HIPCHECK(hipStreamDestroy(stream));
dataset.Release();
test/test_AllToAllv.cpp
+5 -5
Ver fichero
@@ -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
test/test_BroadcastAbort.cpp
+1
Ver fichero
@@ -116,6 +116,7 @@ namespace CorrectnessTests
if (idle) pthread_yield();
}
free(done);
HIPCHECK(hipHostFree(fake_head));
HIPCHECK(hipStreamDestroy(stream));
dataset.Release();
test/test_CombinedCalls.cpp
+48 -29
Ver fichero
@@ -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
test/test_Gather.cpp
+1 -1
Ver fichero
@@ -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
test/test_Scatter.cpp
+1 -1
Ver fichero
@@ -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