From cb84597d5e90fc5930aeea1f1e5a3400fc2e2c15 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 12 Feb 2016 21:30:43 -0600 Subject: [PATCH 1/5] Add Bus Bandwidth test, leveraged from SHOC. [ROCm/hip-tests commit: 8cb885f03ae902b5caefeaaa44c7e7896ada49cc] --- .../1_Utils/hipBusBandwidth/LICENSE.txt | 27 + .../samples/1_Utils/hipBusBandwidth/Makefile | 16 + .../hipBusBandwidth/ResultDatabase.cpp | 520 ++++++++++++++++++ .../1_Utils/hipBusBandwidth/ResultDatabase.h | 100 ++++ .../hipBusBandwidth/hipBusBandwidth.cpp | 170 ++++++ 5 files changed, 833 insertions(+) create mode 100644 projects/hip-tests/samples/1_Utils/hipBusBandwidth/LICENSE.txt create mode 100644 projects/hip-tests/samples/1_Utils/hipBusBandwidth/Makefile create mode 100644 projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp create mode 100644 projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.h create mode 100644 projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/LICENSE.txt b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/LICENSE.txt new file mode 100644 index 0000000000..5d0d603232 --- /dev/null +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/LICENSE.txt @@ -0,0 +1,27 @@ + +Copyright (c) 2011, UT-Battelle, LLC +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are met: + +* Redistributions of source code must retain the above copyright + notice, this list of conditions and the following disclaimer. +* Redistributions in binary form must reproduce the above copyright + notice, this list of conditions and the following disclaimer in the + documentation and/or other materials provided with the distribution. +* Neither the name of Oak Ridge National Laboratory, nor UT-Battelle, LLC, nor + the names of its contributors may be used to endorse or promote products + derived from this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND +ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED +WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/Makefile b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/Makefile new file mode 100644 index 0000000000..d233216313 --- /dev/null +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/Makefile @@ -0,0 +1,16 @@ +HIP_PATH?=$(shell hipconfig -p) +HIPCC=$(HIP_PATH)/bin/hipcc + +EXE=hipBusBandwidth + +all: install + +$(EXE): hipBusBandwidth.cpp ResultDatabase.cpp + $(HIPCC) $^ -o $@ + +install: $(EXE) + cp $(EXE) $(HIP_PATH)/bin + + +clean: + rm -f *.o $(EXE) diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp new file mode 100644 index 0000000000..f57aed11be --- /dev/null +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp @@ -0,0 +1,520 @@ +#include "ResultDatabase.h" + +#include +#include +#include +#include + +using namespace std; + +bool ResultDatabase::Result::operator<(const Result &rhs) const +{ + if (test < rhs.test) + return true; + if (test > rhs.test) + return false; + if (atts < rhs.atts) + return true; + if (atts > rhs.atts) + return false; + return false; // less-operator returns false on equal +} + +double ResultDatabase::Result::GetMin() const +{ + double r = FLT_MAX; + for (int i=0; i= 100) + return value[n-1]; + + double index = ((n + 1.) * q / 100.) - 1; + + vector sorted = value; + sort(sorted.begin(), sorted.end()); + + if (n == 2) + return (sorted[0] * (1 - q/100.) + sorted[1] * (q/100.)); + + int index_lo = int(index); + double frac = index - index_lo; + if (frac == 0) + return sorted[index_lo]; + + double lo = sorted[index_lo]; + double hi = sorted[index_lo + 1]; + return lo + (hi-lo)*frac; +} + +double ResultDatabase::Result::GetMean() const +{ + double r = 0; + for (int i=0; i &values) +{ + for (int i=0; i= results.size()) + { + Result r; + r.test = test; + r.atts = atts; + r.unit = unit; + results.push_back(r); + } + + results[index].value.push_back(value); +} + +// **************************************************************************** +// Method: ResultDatabase::DumpDetailed +// +// Purpose: +// Writes the full results, including all trials. +// +// Arguments: +// out where to print +// +// Programmer: Jeremy Meredith +// Creation: August 14, 2009 +// +// Modifications: +// Jeremy Meredith, Wed Nov 10 14:25:17 EST 2010 +// Renamed to DumpDetailed to make room for a DumpSummary. +// +// Jeremy Meredith, Thu Nov 11 11:39:57 EST 2010 +// Added note about (*) missing value tag. +// +// Jeremy Meredith, Tue Nov 23 13:57:02 EST 2010 +// Changed note about missing values to be worded a little better. +// +// **************************************************************************** +void ResultDatabase::DumpDetailed(ostream &out) +{ + vector sorted(results); + + sort(sorted.begin(), sorted.end()); + + int maxtrials = 1; + for (int i=0; i maxtrials) + maxtrials = sorted[i].value.size(); + } + + // TODO: in big parallel runs, the "trials" are the procs + // and we really don't want to print them all out.... + out << "test\t" + << "atts\t" + << "units\t" + << "median\t" + << "mean\t" + << "stddev\t" + << "min\t" + << "max\t"; + for (int i=0; i sorted(results); + + sort(sorted.begin(), sorted.end()); + + out << std::fixed << right << std::setw(9) << std::setprecision(4); + + // TODO: in big parallel runs, the "trials" are the procs + // and we really don't want to print them all out.... + out << "test\t" + << "atts\t" + << "units\t" + << "median\t" + << "mean\t" + << "stddev\t" + << "min\t" + << "max\t"; + out << endl; + + for (int i=0; i sorted(results); + + sort(sorted.begin(), sorted.end()); + + //Check to see if the file is empty - if so, add the headers + emptyFile = this->IsFileEmpty(fileName); + + //Open file and append by default + ofstream out; + out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app); + + //Add headers only for empty files + if(emptyFile) + { + // TODO: in big parallel runs, the "trials" are the procs + // and we really don't want to print them all out.... + out << "test, " + << "atts, " + << "units, " + << "median, " + << "mean, " + << "stddev, " + << "min, " + << "max, "; + out << endl; + } + + for (int i=0; i +ResultDatabase::GetResultsForTest(const string &test) +{ + // get only the given test results + vector retval; + for (int i=0; i & +ResultDatabase::GetResults() const +{ + return results; +} diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.h b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.h new file mode 100644 index 0000000000..4b63a02a1f --- /dev/null +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.h @@ -0,0 +1,100 @@ +#ifndef RESULT_DATABASE_H +#define RESULT_DATABASE_H + +#include +#include +#include +#include +#include +using std::string; +using std::vector; +using std::ostream; +using std::ofstream; +using std::ifstream; + + +// **************************************************************************** +// Class: ResultDatabase +// +// Purpose: +// Track numerical results as they are generated. +// Print statistics of raw results. +// +// Programmer: Jeremy Meredith +// Creation: June 12, 2009 +// +// Modifications: +// Jeremy Meredith, Wed Nov 10 14:20:47 EST 2010 +// Split timing reports into detailed and summary. E.g. for serial code, +// we might report all trial values, but skip them in parallel. +// +// Jeremy Meredith, Thu Nov 11 11:40:18 EST 2010 +// Added check for missing value tag. +// +// Jeremy Meredith, Mon Nov 22 13:37:10 EST 2010 +// Added percentile statistic. +// +// Jeremy Meredith, Fri Dec 3 16:30:31 EST 2010 +// Added a method to extract a subset of results based on test name. Also, +// the Result class is now public, so that clients can use them directly. +// Added a GetResults method as well, and made several functions const. +// +// **************************************************************************** +class ResultDatabase +{ + public: + // + // A performance result for a single SHOC benchmark run. + // + struct Result + { + string test; // e.g. "readback" + string atts; // e.g. "pagelocked 4k^2" + string unit; // e.g. "MB/sec" + vector value; // e.g. "837.14" + double GetMin() const; + double GetMax() const; + double GetMedian() const; + double GetPercentile(double q) const; + double GetMean() const; + double GetStdDev() const; + + bool operator<(const Result &rhs) const; + + bool HadAnyFLTMAXValues() const + { + for (int i=0; i= FLT_MAX) + return true; + } + return false; + } + }; + + protected: + vector results; + + public: + void AddResult(const string &test, + const string &atts, + const string &unit, + double value); + void AddResults(const string &test, + const string &atts, + const string &unit, + const vector &values); + vector GetResultsForTest(const string &test); + const vector &GetResults() const; + void ClearAllResults(); + void DumpDetailed(ostream&); + void DumpSummary(ostream&); + void DumpCsv(string fileName); + + private: + bool IsFileEmpty(string fileName); + +}; + + +#endif diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp new file mode 100644 index 0000000000..8481476fc8 --- /dev/null +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -0,0 +1,170 @@ +#include +#include +#include + +#include "ResultDatabase.h" + +// Cmdline parms: +const bool p_verbose = false; +const bool p_pinned = true; +const unsigned int p_iters = 10; + +#define CHECK_HIP_ERROR() \ +{ \ + hipError_t err = hipGetLastError(); \ + if (err != hipSuccess) \ + { \ + printf("error=%d name=%s at " \ + "ln: %d\n ",err,hipGetErrorString(err),__LINE__); \ + exit(EXIT_FAILURE); \ + } \ +} + + +// **************************************************************************** +// Function: runBenchmark +// +// Purpose: +// Measures the bandwidth of the bus connecting the host processor to the +// OpenCL device. This benchmark repeatedly transfers data chunks of various +// sizes across the bus to the OpenCL device, and calculates the bandwidth. +// +// +// Arguments: +// +// Returns: nothing +// +// Programmer: Jeremy Meredith +// Creation: September 08, 2009 +// +// Modifications: +// Jeremy Meredith, Wed Dec 1 17:05:27 EST 2010 +// Added calculation of latency estimate. +// Ben Sander - moved to standalone test +// +// **************************************************************************** +void RunBenchmark(ResultDatabase &resultDB) +{ + // Sizes are in kb + int nSizes = 20; + int sizes[20] = {1,2,4,8,16,32,64,128,256,512,1024,2048,4096,8192,16384, + 32768,65536,131072,262144,524288}; + long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + + // Create some host memory pattern + float *hostMem = NULL; + if (p_pinned) + { + hipMallocHost((void**)&hostMem, sizeof(float) * numMaxFloats); + while (hipGetLastError() != hipSuccess) + { + // 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; + hipMallocHost((void**)&hostMem, sizeof(float) * numMaxFloats); + } + } + else + { + hostMem = new float[numMaxFloats]; + } + + for (int i = 0; i < numMaxFloats; i++) + { + hostMem[i] = i % 77; + } + + float *device; + hipMalloc((void**)&device, sizeof(float) * numMaxFloats); + while (hipGetLastError() != hipSuccess) + { + // 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; + hipMalloc((void**)&device, sizeof(float) * numMaxFloats); + } + + + hipEvent_t start, stop; + hipEventCreate(&start); + hipEventCreate(&stop); + CHECK_HIP_ERROR(); + + // Three passes, forward and backward both + for (int pass = 0; pass < p_iters; 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 = sizes[sizeIndex] * 1024; + + hipEventRecord(start, 0); + hipMemcpy(device, hostMem, nbytes, hipMemcpyHostToDevice); + 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 " << sizes[sizeIndex] << "k took " << t << + " ms\n"; + } + + double speed = (double(sizes[sizeIndex]) * 1024. / (1000*1000)) / t; + char sizeStr[256]; + sprintf(sizeStr, "% 7dkB", sizes[sizeIndex]); + resultDB.AddResult("DownloadSpeed", sizeStr, "GB/sec", speed); + resultDB.AddResult("DownloadTime", sizeStr, "ms", t); + } + } + + // Cleanup + hipFree((void*)device); + CHECK_HIP_ERROR(); + if (p_pinned) + { + hipFreeHost((void*)hostMem); + CHECK_HIP_ERROR(); + } + else + { + delete[] hostMem; + } + hipEventDestroy(start); + hipEventDestroy(stop); +} + + + +int main(int argc, char *argv[]) +{ + ResultDatabase resultDB; + RunBenchmark(resultDB); + + resultDB.DumpSummary(std::cout); + + resultDB.DumpDetailed(std::cout); +} From 4b0bf5c1a123da34ab41d7af1e29803531f03d75 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 12 Feb 2016 22:46:34 -0600 Subject: [PATCH 2/5] Add D2H test [ROCm/hip-tests commit: 42257c5d71e4cb95332979e62c71dbaafa78d191] --- .../hipBusBandwidth/hipBusBandwidth.cpp | 235 +++++++++++++++++- 1 file changed, 226 insertions(+), 9 deletions(-) diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 8481476fc8..c908fa655e 100644 --- a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -5,9 +5,15 @@ #include "ResultDatabase.h" // Cmdline parms: -const bool p_verbose = false; -const bool p_pinned = true; -const unsigned int p_iters = 10; +bool p_verbose = false; +bool p_pinned = true; +int p_iterations = 10; +int p_device = 0; +int p_detailed = 0; + +bool p_h2d = true; +bool p_d2h = true; + #define CHECK_HIP_ERROR() \ { \ @@ -43,7 +49,7 @@ const unsigned int p_iters = 10; // Ben Sander - moved to standalone test // // **************************************************************************** -void RunBenchmark(ResultDatabase &resultDB) +void RunBenchmark_H2D(ResultDatabase &resultDB) { // Sizes are in kb int nSizes = 20; @@ -51,6 +57,8 @@ void RunBenchmark(ResultDatabase &resultDB) 32768,65536,131072,262144,524288}; long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + hipSetDevice(p_device); + // Create some host memory pattern float *hostMem = NULL; if (p_pinned) @@ -103,7 +111,7 @@ void RunBenchmark(ResultDatabase &resultDB) CHECK_HIP_ERROR(); // Three passes, forward and backward both - for (int pass = 0; pass < p_iters; pass++) + for (int pass = 0; pass < p_iterations; pass++) { // store the times temporarily to estimate latency //float times[nSizes]; @@ -158,13 +166,222 @@ void RunBenchmark(ResultDatabase &resultDB) } +void RunBenchmark_D2H(ResultDatabase &resultDB) +{ + + // Sizes are in kb + int nSizes = 20; + int sizes[20] = {1,2,4,8,16,32,64,128,256,512,1024,2048,4096,8192,16384, + 32768,65536,131072,262144,524288}; + long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + + // Create some host memory pattern + float *hostMem1; + float *hostMem2; + if (p_pinned) + { + hipMallocHost((void**)&hostMem1, sizeof(float)*numMaxFloats); + hipError_t err1 = hipGetLastError(); + hipMallocHost((void**)&hostMem2, sizeof(float)*numMaxFloats); + hipError_t err2 = hipGetLastError(); + while (err1 != hipSuccess || err2 != hipSuccess) + { + // free the first buffer if only the second failed + if (err1 == hipSuccess) + hipFreeHost((void*)hostMem1); + + // 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; + hipMallocHost((void**)&hostMem1, sizeof(float)*numMaxFloats); + err1 = hipGetLastError(); + hipMallocHost((void**)&hostMem2, sizeof(float)*numMaxFloats); + err2 = hipGetLastError(); + } + } + else + { + hostMem1 = new float[numMaxFloats]; + hostMem2 = new float[numMaxFloats]; + } + for (int i=0; i= argc || !parseInt(argv[i], &p_iterations)) { + failed("Bad iterations argument"); + } + } else if (!strcmp(arg, "--device") || (!strcmp(arg, "-d"))) { + if (++i >= argc || !parseInt(argv[i], &p_device)) { + failed("Bad device argument"); + } + } else if (!strcmp(arg, "--unpinned")) { + p_pinned = 0; + } else if (!strcmp(arg, "--h2d")) { + p_h2d = true; + p_d2h = false; + + } else if (!strcmp(arg, "--d2h")) { + p_h2d = false; + p_d2h = true; + + } else if (!strcmp(arg, "--help") || (!strcmp(arg, "-h"))) { + help(); + + } else if (!strcmp(arg, "--verbose")) { + p_verbose = 1; + } else if (!strcmp(arg, "--detailed")) { + p_detailed = 1; + } else { + failed("Bad argument '%s'", arg); + } + } + + return 0; +}; + + int main(int argc, char *argv[]) { - ResultDatabase resultDB; - RunBenchmark(resultDB); + parseStandardArguments(argc, argv); - resultDB.DumpSummary(std::cout); + if (p_h2d) { + ResultDatabase resultDB; + RunBenchmark_H2D(resultDB); - resultDB.DumpDetailed(std::cout); + resultDB.DumpSummary(std::cout); + + if (p_detailed) { + resultDB.DumpDetailed(std::cout); + } + } + + if (p_d2h) { + ResultDatabase resultDB; + RunBenchmark_D2H(resultDB); + + resultDB.DumpSummary(std::cout); + + if (p_detailed) { + resultDB.DumpDetailed(std::cout); + } + } } From 19c6c95b8e727a9abef4b91901afa79ca70037c8 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 12 Feb 2016 22:47:26 -0600 Subject: [PATCH 3/5] Add D2H test [ROCm/hip-tests commit: 63d79554ab5010dd2b138f7e465348fb75bdbe23] --- .../samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index c908fa655e..b847f8db40 100644 --- a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -144,8 +144,8 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) double speed = (double(sizes[sizeIndex]) * 1024. / (1000*1000)) / t; char sizeStr[256]; sprintf(sizeStr, "% 7dkB", sizes[sizeIndex]); - resultDB.AddResult("DownloadSpeed", sizeStr, "GB/sec", speed); - resultDB.AddResult("DownloadTime", sizeStr, "ms", t); + resultDB.AddResult("H2D_Bandwidth", sizeStr, "GB/sec", speed); + resultDB.AddResult("H2D_Time", sizeStr, "ms", t); } } @@ -273,8 +273,8 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) double speed = (double(sizes[sizeIndex]) * 1024. / (1000*1000)) / t; char sizeStr[256]; sprintf(sizeStr, "% 7dkB", sizes[sizeIndex]); - resultDB.AddResult("ReadbackSpeed", sizeStr, "GB/sec", speed); - resultDB.AddResult("ReadbackTime", sizeStr, "ms", t); + resultDB.AddResult("D2H_Bandwidth", sizeStr, "GB/sec", speed); + resultDB.AddResult("D2H_Time", sizeStr, "ms", t); } //resultDB.AddResult("ReadbackLatencyEstimate", "1-2kb", "ms", times[0]-(times[1]-times[0])/1.); //resultDB.AddResult("ReadbackLatencyEstimate", "1-4kb", "ms", times[0]-(times[2]-times[0])/3.); From 931a24bad1481a3c2bd826638a204125c0cb5d1d Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 13 Feb 2016 01:14:01 -0600 Subject: [PATCH 4/5] Result formatting [ROCm/hip-tests commit: b5743abae250e668d808a6062a8b52756605bb3d] --- .../samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp index f57aed11be..7d2f3aef84 100644 --- a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp @@ -278,13 +278,16 @@ void ResultDatabase::DumpSummary(ostream &out) { vector sorted(results); + int testW = 15 ; + const int fieldW = 9; + sort(sorted.begin(), sorted.end()); - out << std::fixed << right << std::setw(9) << std::setprecision(4); + out << std::fixed << right << std::setprecision(4); // TODO: in big parallel runs, the "trials" are the procs // and we really don't want to print them all out.... - out << "test\t" + out << setw(testW) << "test\t" << setw(fieldW) << "atts\t" << "units\t" << "median\t" @@ -297,7 +300,7 @@ void ResultDatabase::DumpSummary(ostream &out) for (int i=0; i Date: Sat, 13 Feb 2016 03:17:42 -0600 Subject: [PATCH 5/5] Enable -O3, style points on array size [ROCm/hip-tests commit: f12a1e95816adbfaffb86a2cd809150e3bf49b8d] --- projects/hip-tests/samples/1_Utils/hipBusBandwidth/Makefile | 3 ++- .../samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp | 6 +++--- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/Makefile b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/Makefile index d233216313..77a92fb1a6 100644 --- a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/Makefile +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/Makefile @@ -2,11 +2,12 @@ HIP_PATH?=$(shell hipconfig -p) HIPCC=$(HIP_PATH)/bin/hipcc EXE=hipBusBandwidth +CXXFLAGS = -O3 -g all: install $(EXE): hipBusBandwidth.cpp ResultDatabase.cpp - $(HIPCC) $^ -o $@ + $(HIPCC) $(CXXFLAGS) $^ -o $@ install: $(EXE) cp $(EXE) $(HIP_PATH)/bin diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index b847f8db40..d276725921 100644 --- a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -52,9 +52,9 @@ bool p_d2h = true; void RunBenchmark_H2D(ResultDatabase &resultDB) { // Sizes are in kb - int nSizes = 20; - int sizes[20] = {1,2,4,8,16,32,64,128,256,512,1024,2048,4096,8192,16384, - 32768,65536,131072,262144,524288}; + int sizes[] = {1,2,4,8,16,32,64,128,256,512,1024,2048,4096,8192,16384, 32768,65536,131072,262144,524288}; + int nSizes = sizeof(sizes) / sizeof(int); + long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; hipSetDevice(p_device);