AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

[TransferBench] Support multiple of 4 byte sizes, changing default GPU timing mechanism (#307)

Przeglądaj źródła 
* Changing default timing mechanism, adjusting CPU bandwidth calc, adding flag to use combined timing
* Adding support for smaller transfers (byte size must be multiple of 4 instead of 128)
This commit is contained in:
gilbertlee-amd
2020-12-04 14:57:13 -07:00
zatwierdzone przez GitHub
rodzic 882d52ad7e
commit b80ae551b1
1 zmienionych plików z 27 dodań i 22 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
tools/TransferBench/TransferBench.cpp
+27 -22
Wyświetl plik
@@ -54,6 +54,7 @@ int main(int argc, char **argv)
bool useSingleSync = getenv("USE_SINGLE_SYNC"); // Perform synchronization only once after all iterations instead of per iteration
bool useInteractive = getenv("USE_INTERACTIVE"); // Pause for user-input before starting transfer loop
bool useSleep = getenv("USE_SLEEP"); // Adds a 100ms sleep after each synchronization
bool combineTiming = getenv("COMBINE_TIMING"); // Combines the timing with kernel launch
bool reuseStreams = getenv("REUSE_STREAMS"); // Re-use streams instead of creating / destroying per test
bool showAddr = getenv("SHOW_ADDR"); // Print out memory addresses for each Link
bool outputToCsv = getenv("OUTPUT_TO_CSV"); // Output in CSV format
@@ -67,9 +68,9 @@ int main(int argc, char **argv)
std::vector<size_t> valuesOfN;
size_t const numBytesPerLink = argc > 2 ? atoll(argv[2]) : DEFAULT_BYTES_PER_LINK;
if (numBytesPerLink % 128)
if (numBytesPerLink % 4)
{
printf("[ERROR] numBytesPerLink (%lu) must be a multiple of 128\n", numBytesPerLink);
printf("[ERROR] numBytesPerLink (%lu) must be a multiple of 4\n", numBytesPerLink);
exit(1);
}
@@ -130,6 +131,9 @@ int main(int argc, char **argv)
printf("%-20s %8s: %s\n",
"USE_SLEEP", useSleep ? "(set)" : "(unset)",
useSleep ? "Add sleep after each sync" : "No sleep per sync");
printf("%-20s %8s: %s\n",
"COMBINE_TIMING", combineTiming ? "(set)" : "(unset)",
combineTiming ? "Using combined timing+launch" : "Using separate timing / launch");
printf("%-20s %8s: %s\n",
"REUSE_STREAMS", reuseStreams ? "(set)" : "(unset)",
reuseStreams ? "Re-using streams per topology" : "Creating/destroying streams per topology");
@@ -200,7 +204,6 @@ int main(int argc, char **argv)
hipStream_t streams[numLinks]; // hipStream to use per Link
hipEvent_t startEvents[numLinks]; // Start event per Link
hipEvent_t stopEvents[numLinks]; // Stop event per Link
hipEvent_t dummyEvents[numLinks]; // Dummy event per Link
std::vector<BlockParam> cpuBlockParams[numLinks]; // CPU copy of block parameters
BlockParam* gpuBlockParams[numLinks]; // GPU copy of block parameters
@@ -251,7 +254,6 @@ int main(int argc, char **argv)
HIP_CALL(hipSetDevice(exeIndex));
HIP_CALL(hipEventCreate(&startEvents[i]));
HIP_CALL(hipEventCreate(&stopEvents[i]));
HIP_CALL(hipEventCreate(&dummyEvents[i]));
HIP_CALL(hipMalloc((void**)&gpuBlockParams[i], sizeof(BlockParam) * numLinks));
if (reuseStreams)
{
@@ -278,21 +280,23 @@ int main(int argc, char **argv)
// Count # of links / total blocks each GPU will be working on
linkCount[exeIndex]++;
// Each block needs to know src/dst pointers and how many elements to transfer
// Figure out the sub-array each block does for this link
// NOTE: Have each sub-array to work on multiple of 32-floats (128-bytes),
// but divide as evenly as possible
// NOTE: N is always a multiple of 32
int blocksWithExtra = (N / 32) % links[i].numBlocksToUse;
int perBlockBaseN = (N / 32) / links[i].numBlocksToUse * 32;
// Figure out the sub-array each block does for this Link
// - Partition N as evenly as posible, but try to keep blocks as multiples of 32,
// except the very last one, for alignment reasons
size_t assigned = 0;
int maxNumBlocksToUse = std::min((N + 31) / 32, (size_t)links[i].numBlocksToUse);
for (int j = 0; j < links[i].numBlocksToUse; j++)
{
BlockParam param;
param.N = perBlockBaseN + ((j < blocksWithExtra) ? 32 : 0);
param.src = linkSrcMem[i] + ((j * perBlockBaseN) + ((j < blocksWithExtra) ?
j : blocksWithExtra) * 32) + initOffset;
param.dst = linkDstMem[i] + ((j * perBlockBaseN) + ((j < blocksWithExtra) ?
j : blocksWithExtra) * 32) + initOffset;
int blocksLeft = std::max(0, maxNumBlocksToUse - j);
size_t leftover = N - assigned;
size_t roundedN = (leftover + 31) / 32;
param.N = blocksLeft ? std::min(leftover, ((roundedN / blocksLeft) * 32)) : 0;
param.src = linkSrcMem[i] + assigned + initOffset;
param.dst = linkDstMem[i] + assigned + initOffset;
assigned += param.N;
cpuBlockParams[i].push_back(param);
}
@@ -345,13 +349,15 @@ int main(int argc, char **argv)
}
else
{
if (!combineTiming && recordStart) HIP_CALL(hipEventRecord(startEvents[i], streams[i]));
hipExtLaunchKernelGGL(useMemset ? MemsetKernel : CopyKernel,
dim3(links[i].numBlocksToUse, 1, 1),
dim3(BLOCKSIZE, 1, 1),
0, streams[i],
recordStart ? startEvents[i] : NULL,
recordStop ? stopEvents[i] : NULL,
(combineTiming && recordStart) ? startEvents[i] : NULL,
(combineTiming && recordStop) ? stopEvents[i] : NULL,
0, gpuBlockParams[i]);
if (!combineTiming & recordStop) HIP_CALL(hipEventRecord(stopEvents[i], streams[i]));
}
}
@@ -381,7 +387,6 @@ int main(int argc, char **argv)
for (int i = 0; i < numLinks; i++)
{
HIP_CALL(hipSetDevice(links[i].exeIndex));
HIP_CALL(hipEventSynchronize(startEvents[i]));
HIP_CALL(hipEventSynchronize(stopEvents[i]));
float gpuDeltaMsec;
HIP_CALL(hipEventElapsedTime(&gpuDeltaMsec, startEvents[i], stopEvents[i]));
@@ -404,12 +409,11 @@ int main(int argc, char **argv)
// Report timings
totalCpuTime = totalCpuTime / (1.0 * numIterations) * 1000;
double totalBandwidthGbs = 0.0;
double totalBandwidthGbs = (numLinks * N * sizeof(float) / 1.0E6) / totalCpuTime;
for (int i = 0; i < numLinks; i++)
{
double linkDurationMsec = totalGpuTime[i] / (1.0 * numIterations);
double linkBandwidthGbs = (N * sizeof(float) / 1.0E9) / linkDurationMsec * 1000.0f;
totalBandwidthGbs += linkBandwidthGbs;
if (!outputToCsv)
{
printf(" Link %02d: %c%02d -> [GPU %02d:%02d] -> %c%02d | %9.3f GB/s | %8.3f ms | %9s |",
@@ -443,7 +447,7 @@ int main(int argc, char **argv)
// Display aggregate statistics
if (!outputToCsv)
{
printf(" Aggregate Bandwidth | %9.3f GB/s | %8.3f ms |\n", totalBandwidthGbs, totalCpuTime);
printf(" Aggregate Bandwidth (CPU timed) | %9.3f GB/s | %8.3f ms |\n", totalBandwidthGbs, totalCpuTime);
}
else
{
@@ -488,7 +492,7 @@ void DisplayUsage(char const* cmdName)
printf(" configFile: File containing Links to execute (see below for format)\n");
printf(" N : (Optional) Number of bytes to transfer per link.\n");
printf(" If not specified, defaults to %lu bytes. Must be a multiple of 128 bytes\n", DEFAULT_BYTES_PER_LINK);
printf(" If not specified, defaults to %lu bytes. Must be a multiple of 4 bytes\n", DEFAULT_BYTES_PER_LINK);
printf(" If 0 is specified, a range of Ns will be benchmarked\n");
printf(" If a negative number is specified, a configFile gets generated with this number as default number of CUs per link\n");
printf("\n");
@@ -538,6 +542,7 @@ void DisplayUsage(char const* cmdName)
printf(" USE_SINGLE_SYNC - Perform synchronization only once after all iterations instead of per iteration\n");
printf(" USE_INTERACTIVE - Pause for user-input before starting transfer loop\n");
printf(" USE_SLEEP - Adds a 100ms sleep after each synchronization\n");
printf(" COMBINE_TIMING - Combines timing with launch (potentially lower timing overhead)\n");
printf(" REUSE_STREAMS - Re-use streams instead of creating / destroying per test\n");
printf(" SHOW_ADDR - Print out memory addresses for each Link\n");
printf(" OUTPUT_TO_CSV - Outputs to CSV format if set\n");