From 5606bee07658e1ce4e34818a299091c66012b160 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 14 Mar 2016 14:39:23 -0500 Subject: [PATCH] Add Bidir copy test and help. --- .../hipBusBandwidth/hipBusBandwidth.cpp | 209 +++++++++++++++++- 1 file changed, 198 insertions(+), 11 deletions(-) diff --git a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 369fa9e37a..88d547e2c5 100644 --- a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -12,9 +12,11 @@ bool p_pinned = true; int p_iterations = 10; int p_device = 0; int p_detailed = 0; +bool p_async = 0; -bool p_h2d = true; -bool p_d2h = true; +bool p_h2d = true; +bool p_d2h = true; +bool p_bidir = true; #define CHECK_HIP_ERROR() \ @@ -42,6 +44,7 @@ std::string sizeToString(int size) using namespace std; stringstream ss; if (size < 0) { + // char (01) sorts before " " so will cause Byte values to be displayed before kB. ss << char(0x1) << setfill('0') << setw(3) << -size << "B"; } else { ss << size << "kB"; @@ -50,6 +53,18 @@ std::string sizeToString(int size) } +// **************************************************************************** +hipError_t memcopy(void * dst, const void *src, size_t sizeBytes, enum hipMemcpyKind kind) +{ + if (p_async) { + return hipMemcpyAsync(dst, src, sizeBytes, kind, NULL); + } else { + return hipMemcpy(dst, src, sizeBytes, kind); + } +} + + + // **************************************************************************** // -sizes are in bytes, +sizes are in kb, last size must be largest int sizes[] = {-64, -256, -512, 1,2,4,8,16,32,64,128,256,512,1024,2048,4096,8192,16384, 32768,65536,131072,262144,524288}; @@ -152,7 +167,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) int nbytes = sizeToBytes(sizes[sizeIndex]); hipEventRecord(start, 0); - hipMemcpy(device, hostMem, nbytes, hipMemcpyHostToDevice); + memcopy(device, hostMem, nbytes, hipMemcpyHostToDevice); hipEventRecord(stop, 0); hipEventSynchronize(stop); float t = 0; @@ -251,8 +266,7 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) hipMalloc((void**)&device, sizeof(float) * numMaxFloats); } - hipMemcpy(device, hostMem1, - numMaxFloats*sizeof(float), hipMemcpyHostToDevice); + hipMemcpy(device, hostMem1, numMaxFloats*sizeof(float), hipMemcpyHostToDevice); hipDeviceSynchronize(); hipEvent_t start, stop; @@ -277,8 +291,7 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) int nbytes = sizeToBytes(sizes[sizeIndex]); hipEventRecord(start, 0); - hipMemcpy(hostMem2, device, - nbytes, hipMemcpyDeviceToHost); + memcopy(hostMem2, device, nbytes, hipMemcpyDeviceToHost); hipEventRecord(stop, 0); hipEventSynchronize(stop); float t = 0; @@ -323,6 +336,147 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) } +void RunBenchmark_Bidir(ResultDatabase &resultDB) +{ + long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + + hipSetDevice(p_device); + + hipStream_t stream[2]; + + + // Create some host memory pattern + float *hostMem[2] = {NULL, NULL}; + if (p_pinned) + { + while (1) + { + hipError_t e1 = hipMallocHost((void**)&hostMem[0], sizeof(float) * numMaxFloats); + hipError_t e2 = hipMallocHost((void**)&hostMem[1], sizeof(float) * numMaxFloats); + + if ((e1 == hipSuccess) && (e2 == hipSuccess)) { + break; + } else { + // drop the size and try again + if (p_verbose) std::cout << " - dropping size allocating pinned mem\n"; + --nSizes; + if (nSizes < 1) + { + std::cerr << "Error: Couldn't allocated any pinned buffer\n"; + return; + } + numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + } + } + } + else + { + hostMem[0] = new float[numMaxFloats]; + hostMem[1] = new float[numMaxFloats]; + } + + for (int i = 0; i < numMaxFloats; i++) + { + hostMem[0][i] = i % 77; + } + + float *deviceMem[2]; + while (1) { + hipError_t e1 = hipMalloc((void**)&deviceMem[0], sizeof(float) * numMaxFloats); + hipError_t e2 = hipMalloc((void**)&deviceMem[1], sizeof(float) * numMaxFloats); + + if ((e1 == hipSuccess) && (e2 == hipSuccess)) { + break; + } else { + if (e1) { + // First alloc succeeded, so free it before trying again + hipFree(&deviceMem[0]); + } + // drop the size and try again + if (p_verbose) std::cout << " - dropping size allocating device mem\n"; + --nSizes; + if (nSizes < 1) + { + std::cerr << "Error: Couldn't allocated any device buffer\n"; + return; + } + numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + } + }; + + + hipMemset(deviceMem[1], 0xFA, numMaxFloats); + + + hipEvent_t start, stop; + hipEventCreate(&start); + hipEventCreate(&stop); + CHECK_HIP_ERROR(); + hipStreamCreate(&stream[0]); + hipStreamCreate(&stream[1]); + + // Three passes, forward and backward both + for (int pass = 0; pass < p_iterations; pass++) + { + // store the times temporarily to estimate latency + //float times[nSizes]; + // Step through sizes forward on even passes and backward on odd + for (int i = 0; i < nSizes; i++) + { + int sizeIndex; + if ((pass % 2) == 0) + sizeIndex = i; + else + sizeIndex = (nSizes - 1) - i; + + int nbytes = sizeToBytes(sizes[sizeIndex]); + + hipEventRecord(start, 0); + hipMemcpyAsync(deviceMem[0], hostMem[0], nbytes, hipMemcpyHostToDevice, stream[0]); + hipMemcpyAsync(hostMem[1], deviceMem[1], nbytes, hipMemcpyDeviceToHost, stream[1]); + hipEventRecord(stop, 0); + hipEventSynchronize(stop); + float t = 0; + hipEventElapsedTime(&t, start, stop); + //times[sizeIndex] = t; + + // Convert to GB/sec + if (p_verbose) + { + std::cerr << "size " << sizeToString(sizes[sizeIndex]) << " took " << t << + " ms\n"; + } + + double speed = (double(sizeToBytes(sizes[sizeIndex])) / (1000*1000)) / t; + char sizeStr[256]; + sprintf(sizeStr, "%9s", sizeToString(sizes[sizeIndex]).c_str()); + resultDB.AddResult("Bidir_Bandwidth", sizeStr, "GB/sec", speed); + resultDB.AddResult("Bidir_Time", sizeStr, "ms", t); + } + } + + // Cleanup + hipFree((void*)deviceMem[0]); + hipFree((void*)deviceMem[1]); + CHECK_HIP_ERROR(); + if (p_pinned) + { + hipFreeHost((void*)hostMem[0]); + hipFreeHost((void*)hostMem[1]); + CHECK_HIP_ERROR(); + } + else + { + delete[] hostMem[0]; + delete[] hostMem[1]; + } + hipEventDestroy(start); + hipEventDestroy(stop); + hipStreamDestroy(stream[0]); + hipStreamDestroy(stream[1]); +} + + #define failed(...) \ printf ("error: ");\ printf (__VA_ARGS__);\ @@ -337,6 +491,17 @@ int parseInt(const char *str, int *output) } void help() { + printf ("Usage: hipBusBandwidth [OPTIONS]\n"); + printf (" --iterations, -i : Number of copy iterations to run.\n"); + printf (" --device, -d : Device ID to use (0..numDevices).\n"); + printf (" --unpinned : Use unpinned host memory.\n"); + printf (" --d2h : Run only device-to-host test.\n"); + printf (" --h2d : Run only host-to-device test.\n"); + printf (" --bidir : Run only bidir copy test.\n"); + printf (" --verbose : Print verbose status messages as test is run.\n"); + printf (" --detailed : Print detailed report (including all trials).\n"); + printf (" --async : Use hipMemcpyAsync(with NULL stream) for H2D/D2H. Default uses hipMemcpy.\n"); + }; int parseStandardArguments(int argc, char *argv[]) @@ -357,18 +522,28 @@ int parseStandardArguments(int argc, char *argv[]) } else if (!strcmp(arg, "--unpinned")) { p_pinned = 0; } else if (!strcmp(arg, "--h2d")) { - p_h2d = true; - p_d2h = false; + p_h2d = true; + p_d2h = false; + p_bidir = false; } else if (!strcmp(arg, "--d2h")) { - p_h2d = false; - p_d2h = true; + p_h2d = false; + p_d2h = true; + p_bidir = false; + + } else if (!strcmp(arg, "--bidir")) { + p_h2d = false; + p_d2h = false; + p_bidir = true; } else if (!strcmp(arg, "--help") || (!strcmp(arg, "-h"))) { help(); + exit(EXIT_SUCCESS); } else if (!strcmp(arg, "--verbose")) { p_verbose = 1; + } else if (!strcmp(arg, "--async")) { + p_async = 1; } else if (!strcmp(arg, "--detailed")) { p_detailed = 1; } else { @@ -406,4 +581,16 @@ int main(int argc, char *argv[]) resultDB.DumpDetailed(std::cout); } } + + + if (p_bidir) { + ResultDatabase resultDB; + RunBenchmark_Bidir(resultDB); + + resultDB.DumpSummary(std::cout); + + if (p_detailed) { + resultDB.DumpDetailed(std::cout); + } + } }