From b80ae551b1491a60e3833c870740ad3e6c237763 Mon Sep 17 00:00:00 2001 From: gilbertlee-amd <44450918+gilbertlee-amd@users.noreply.github.com> Date: Fri, 4 Dec 2020 14:57:13 -0700 Subject: [PATCH] [TransferBench] Support multiple of 4 byte sizes, changing default GPU timing mechanism (#307) * 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) --- tools/TransferBench/TransferBench.cpp | 49 +++++++++++++++------------ 1 file changed, 27 insertions(+), 22 deletions(-) diff --git a/tools/TransferBench/TransferBench.cpp b/tools/TransferBench/TransferBench.cpp index 732a405ee7..531589695d 100644 --- a/tools/TransferBench/TransferBench.cpp +++ b/tools/TransferBench/TransferBench.cpp @@ -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 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 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");