AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Add Bidir copy test and help.

Browse Source
This commit is contained in:
Ben Sander
2016-03-14 14:39:23 -05:00
parent ac6ed35ba0
commit 5606bee076
1 changed files with 198 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files
samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp
+198 -11
View File
@@ -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);
}
}
}