From 3277453ea570045c9f46f8a9fafd35a7fc4ab29e Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 12 Mar 2018 11:29:03 +0530 Subject: [PATCH] Apply .clangformat to all repo source files Change-Id: I7e79c6058f0303f9a98911e3b7dd2e8596079344 [ROCm/hip-tests commit: 6b09bde6752112b1865317e07410faf6a4f662f5] --- .../0_Intro/bit_extract/bit_extract.cpp | 67 +- .../hcc_dialects/vadd_amp_arrayview.cpp | 37 +- .../0_Intro/hcc_dialects/vadd_hc_am.cpp | 52 +- .../0_Intro/hcc_dialects/vadd_hc_array.cpp | 35 +- .../hcc_dialects/vadd_hc_arrayview.cpp | 30 +- .../samples/0_Intro/hcc_dialects/vadd_hip.cpp | 28 +- .../0_Intro/module_api/defaultDriver.cpp | 20 +- .../0_Intro/module_api/launchKernelHcc.cpp | 39 +- .../samples/0_Intro/module_api/runKernel.cpp | 38 +- .../0_Intro/module_api/vcpy_kernel.cpp | 3 +- .../0_Intro/module_api_global/runKernel.cpp | 59 +- .../0_Intro/module_api_global/vcpy_kernel.cpp | 12 +- .../samples/0_Intro/square/square.hipref.cpp | 67 +- .../hipBusBandwidth/ResultDatabase.cpp | 269 +++---- .../1_Utils/hipBusBandwidth/ResultDatabase.h | 55 +- .../hipBusBandwidth/hipBusBandwidth.cpp | 705 ++++++++--------- .../1_Utils/hipCommander/ResultDatabase.cpp | 267 +++---- .../1_Utils/hipCommander/ResultDatabase.h | 55 +- .../1_Utils/hipCommander/hipCommander.cpp | 744 ++++++++---------- .../1_Utils/hipCommander/nullkernel.hip.cpp | 2 +- .../samples/1_Utils/hipCommander/testcase.cpp | 19 +- .../hipDispatchLatency/ResultDatabase.cpp | 269 +++---- .../hipDispatchLatency/ResultDatabase.h | 55 +- .../hipDispatchLatency/hipDispatchLatency.cpp | 87 +- .../samples/1_Utils/hipInfo/hipInfo.cpp | 138 ++-- .../0_MatrixTranspose/MatrixTranspose.cpp | 147 ++-- .../2_Cookbook/10_inline_asm/inline_asm.cpp | 226 +++--- .../11_texture_driver/tex2dKernel.cpp | 11 +- .../11_texture_driver/texture2dDrv.cpp | 122 +-- .../MatrixTranspose.cpp | 147 ++-- .../2_Cookbook/1_hipEvent/hipEvent.cpp | 221 +++--- .../2_Cookbook/2_Profiler/MatrixTranspose.cpp | 331 ++++---- .../3_shared_memory/sharedMemory.cpp | 151 ++-- .../samples/2_Cookbook/4_shfl/shfl.cpp | 154 ++-- .../samples/2_Cookbook/5_2dshfl/2dshfl.cpp | 152 ++-- .../6_dynamic_shared/dynamic_shared.cpp | 149 ++-- .../samples/2_Cookbook/7_streams/stream.cpp | 84 +- .../2_Cookbook/8_peer2peer/peer2peer.cpp | 152 ++-- .../samples/2_Cookbook/9_unroll/unroll.cpp | 154 ++-- 39 files changed, 2386 insertions(+), 2967 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp b/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp index a30f2d052d..be7c5b020f 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp +++ b/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp @@ -28,79 +28,76 @@ THE SOFTWARE. #endif -#define CHECK(cmd) \ -{\ - hipError_t error = cmd;\ - if (error != hipSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \ - exit(EXIT_FAILURE);\ - }\ -} +#define CHECK(cmd) \ + { \ + hipError_t error = cmd; \ + if (error != hipSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, \ + __FILE__, __LINE__); \ + exit(EXIT_FAILURE); \ + } \ + } -__global__ void -bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t N) -{ +__global__ void bit_extract_kernel(hipLaunchParm lp, uint32_t* C_d, const uint32_t* A_d, size_t N) { size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); - size_t stride = hipBlockDim_x * hipGridDim_x ; + size_t stride = hipBlockDim_x * hipGridDim_x; - for (size_t i=offset; i> 8); + C_d[i] = ((A_d[i] & 0xf00) >> 8); #endif } } -int main(int argc, char *argv[]) -{ +int main(int argc, char* argv[]) { uint32_t *A_d, *C_d; uint32_t *A_h, *C_h; size_t N = 1000000; size_t Nbytes = N * sizeof(uint32_t); int deviceId; - CHECK (hipGetDevice(&deviceId)); + CHECK(hipGetDevice(&deviceId)); hipDeviceProp_t props; CHECK(hipGetDeviceProperties(&props, deviceId)); - printf ("info: running on device #%d %s\n", deviceId, props.name); + printf("info: running on device #%d %s\n", deviceId, props.name); - printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); + printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); A_h = (uint32_t*)malloc(Nbytes); - CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); + CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess); C_h = (uint32_t*)malloc(Nbytes); - CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); + CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess); - for (size_t i=0; i> 8); if (C_h[i] != Agold) { - fprintf (stderr, "mismatch detected.\n"); - printf ("%zu: %08x =? %08x (Ain=%08x)\n", i, C_h[i], Agold, A_h[i]); + fprintf(stderr, "mismatch detected.\n"); + printf("%zu: %08x =? %08x (Ain=%08x)\n", i, C_h[i], Agold, A_h[i]); CHECK(hipErrorUnknown); } } - printf ("PASSED!\n"); + printf("PASSED!\n"); } diff --git a/projects/hip-tests/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp b/projects/hip-tests/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp index a3162bccb9..bb0a4157f8 100644 --- a/projects/hip-tests/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp +++ b/projects/hip-tests/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp @@ -26,8 +26,8 @@ THE SOFTWARE. // will automatically copy data to and from the host, without the user needing // to manually perform such copies. This is an excellent mode for developers // new to GPU programming and matches the memory models provided by recent systems where -// CPU and GPU share the same memory pool. Advanced programmers may prefer -// more explicit control over the data movement - shown in the other vadd_hc_array and +// CPU and GPU share the same memory pool. Advanced programmers may prefer +// more explicit control over the data movement - shown in the other vadd_hc_array and // vadd_hc_am examples. // This example shows the similarity between C++AMP and and HC for simple cases where // implicit data transfer is used - really the only difference is the namespace. @@ -35,8 +35,7 @@ THE SOFTWARE. #include -int main(int argc, char *argv[]) -{ +int main(int argc, char* argv[]) { int sizeElements = 1000000; bool pass = true; @@ -46,28 +45,30 @@ int main(int argc, char *argv[]) concurrency::array_view C(sizeElements); // Initialize host data - for (int i=0; i (sizeElements), - [=] (concurrency::index<1> idx) restrict(amp) { - int i = idx[0]; - C[i] = A[i] + B[i]; - }); + // The HCC runtime will ensure that A and B are available on the accelerator before launching + // the kernel. + concurrency::parallel_for_each(concurrency::extent<1>(sizeElements), + [=](concurrency::index<1> idx) restrict(amp) { + int i = idx[0]; + C[i] = A[i] + B[i]; + }); - for (int i=0; i #include -int main(int argc, char *argv[]) -{ +int main(int argc, char* argv[]) { int sizeElements = 1000000; size_t sizeBytes = sizeElements * sizeof(float); bool pass = true; // Allocate host memory - float *A_h = (float*)malloc(sizeBytes); - float *B_h = (float*)malloc(sizeBytes); - float *C_h = (float*)malloc(sizeBytes); + float* A_h = (float*)malloc(sizeBytes); + float* B_h = (float*)malloc(sizeBytes); + float* C_h = (float*)malloc(sizeBytes); // Allocate device pointers: // Unlike array_view, these must be explicitly managed by user: @@ -51,36 +50,37 @@ int main(int argc, char *argv[]) C_d = hc::am_alloc(sizeBytes, acc, 0); // Initialize host data - for (int i=0; i (sizeElements), - [=] (hc::index<1> idx) [[hc]] { - int i = idx[0]; - C_d[i] = A_d[i] + B_d[i]; - }); - - - // This copy is in same AV as the kernel and thus will wait for the kernel to finish before executing. - av.copy(C_d, C_h, sizeBytes); // C++ copy D2H + // and we don't need additional synchronization to ensure the copies complete before the PFE + // begins. + hc::completion_future cf = + hc::parallel_for_each(av, hc::extent<1>(sizeElements), [=](hc::index<1> idx)[[hc]] { + int i = idx[0]; + C_d[i] = A_d[i] + B_d[i]; + }); - for (int i=0; i -int main(int argc, char *argv[]) -{ +int main(int argc, char* argv[]) { int sizeElements = 1000000; size_t sizeBytes = sizeElements * sizeof(float); bool pass = true; // Allocate host memory - float *A_h = (float*)malloc(sizeBytes); - float *B_h = (float*)malloc(sizeBytes); - float *C_h = (float*)malloc(sizeBytes); + float* A_h = (float*)malloc(sizeBytes); + float* B_h = (float*)malloc(sizeBytes); + float* C_h = (float*)malloc(sizeBytes); // Allocate device arrays<> // Unlike array_view, these must be explicitly managed by user: @@ -47,32 +46,32 @@ int main(int argc, char *argv[]) hc::array C_d(sizeElements); // Initialize host data - for (int i=0; i types are not implicitly copied, so we performed copies above. - hc::parallel_for_each(hc::extent<1> (sizeElements), - [&] (hc::index<1> idx) [[hc]] { + hc::parallel_for_each(hc::extent<1>(sizeElements), [&](hc::index<1> idx)[[hc]] { int i = idx[0]; C_d[i] = A_d[i] + B_d[i]; }); - // HCC runtime knows that C_d depends on previous PFE and will force the copy to wait for the PFE to complte. - hc::copy(C_d, C_h); // C++ copy D2H + // HCC runtime knows that C_d depends on previous PFE and will force the copy to wait for the + // PFE to complte. + hc::copy(C_d, C_h); // C++ copy D2H - for (int i=0; i -int main(int argc, char *argv[]) -{ +int main(int argc, char* argv[]) { int sizeElements = 1000000; bool pass = true; @@ -46,28 +45,29 @@ int main(int argc, char *argv[]) hc::array_view C(sizeElements); // Initialize host data - for (int i=0; i (sizeElements), - [=] (hc::index<1> idx) [[hc]] { + // The HCC runtime will ensure that A and B are available on the accelerator before launching + // the kernel. + hc::parallel_for_each(hc::extent<1>(sizeElements), [=](hc::index<1> idx)[[hc]] { int i = idx[0]; C[i] = A[i] + B[i]; }); - for (int i=0; i -#include -#include +#include +#include +#include #define LEN 64 -#define SIZE LEN<<2 +#define SIZE LEN << 2 #define fileName "test.co" #define kernel_name "vadd" -int main(){ +int main() { float *A, *B, *C; hipDeviceptr_t Ad, Bd, Cd; A = new float[LEN]; B = new float[LEN]; C = new float[LEN]; - for(uint32_t i=0;i @@ -33,22 +32,25 @@ THE SOFTWARE. #endif #define LEN 64 -#define SIZE LEN<<2 +#define SIZE LEN << 2 #define fileName "vcpy_kernel.code.adipose" #define kernel_name "hello_world" -#define HIP_CHECK(status) \ -if(status != hipSuccess) {std::cout<<"Got Status: "< #define LEN 64 -#define SIZE LEN*sizeof(float) +#define SIZE LEN * sizeof(float) #define fileName "vcpy_kernel.code.adipose" float myDeviceGlobal; float myDeviceGlobalArray[16]; -#define HIP_CHECK(cmd) \ -{\ - hipError_t status = cmd;\ - if(status != hipSuccess) {std::cout<<"error: #"< #include "hip/hip_runtime.h" -#define CHECK(cmd) \ -{\ - hipError_t error = cmd;\ - if (error != hipSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \ - exit(EXIT_FAILURE);\ - }\ -} +#define CHECK(cmd) \ + { \ + hipError_t error = cmd; \ + if (error != hipSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, \ + __FILE__, __LINE__); \ + exit(EXIT_FAILURE); \ + } \ + } /* * Square each element in the array A and write to array C. */ template -__global__ void -vector_square(hipLaunchParm lp, T *C_d, const T *A_d, size_t N) -{ +__global__ void vector_square(hipLaunchParm lp, T* C_d, const T* A_d, size_t N) { size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); - size_t stride = hipBlockDim_x * hipGridDim_x ; + size_t stride = hipBlockDim_x * hipGridDim_x; - for (size_t i=offset; i rhs.test) - return false; -#if (SORT_RETAIN_ATTS_ORDER == 0) +bool ResultDatabase::Result::operator<(const Result& rhs) const { + if (test < rhs.test) return true; + if (test > rhs.test) return false; +#if (SORT_RETAIN_ATTS_ORDER == 0) // For ties, sort by the value of the attribute: - if (atts < rhs.atts) - return true; - if (atts > rhs.atts) - return false; + if (atts < rhs.atts) return true; + if (atts > rhs.atts) return false; #endif - return false; // less-operator returns false on equal + return false; // less-operator returns false on equal } -double ResultDatabase::Result::GetMin() const -{ +double ResultDatabase::Result::GetMin() const { double r = FLT_MAX; - for (int i=0; i= 100) - return value[n-1]; + if (q <= 0) return value[0]; + if (q >= 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.)); + 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]; + if (frac == 0) return sorted[index_lo]; double lo = sorted[index_lo]; double hi = sorted[index_lo + 1]; - return lo + (hi-lo)*frac; + return lo + (hi - lo) * frac; } -double ResultDatabase::Result::GetMean() const -{ +double ResultDatabase::Result::GetMean() const { double r = 0; - for (int i=0; i &values) -{ - for (int i=0; i& values) { + for (int i = 0; i < values.size(); i++) { AddResult(test, atts, unit, values[i]); } } -static string RemoveAllButLeadingSpaces(const string &a) -{ +static string RemoveAllButLeadingSpaces(const string& a) { string b; int n = a.length(); int i = 0; - while (i= results.size()) - { + if (index >= results.size()) { Result r; r.test = test; r.atts = atts; @@ -192,41 +152,33 @@ void ResultDatabase::AddResult(const string &test_orig, // Changed note about missing values to be worded a little better. // // **************************************************************************** -void ResultDatabase::DumpDetailed(ostream &out) -{ +void ResultDatabase::DumpDetailed(ostream& out) { vector sorted(results); stable_sort(sorted.begin(), sorted.end()); - const int testNameW = 24 ; + const int testNameW = 24; const int attW = 12; const int fieldW = 11; out << std::fixed << right << std::setprecision(4); int maxtrials = 1; - for (int i=0; i maxtrials) - maxtrials = sorted[i].value.size(); + for (int i = 0; i < sorted.size(); i++) { + if (sorted[i].value.size() > 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 << setw(testNameW) << "test\t" - << setw(attW) << "atts\t" - << setw(fieldW) - << "median\t" + out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "median\t" << "mean\t" << "stddev\t" << "min\t" << "max\t"; - for (int i=0; i sorted(results); stable_sort(sorted.begin(), sorted.end()); - const int testNameW = 24 ; + const int testNameW = 24; const int attW = 12; const int fieldW = 9; 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 << setw(testNameW) << "test\t" - << setw(attW) << "atts\t" - << setw(fieldW) - << "units\t" + out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "units\t" << "median\t" << "mean\t" << "stddev\t" @@ -309,9 +256,8 @@ void ResultDatabase::DumpSummary(ostream &out) << "max\t"; out << endl; - for (int i=0; i sorted(results); stable_sort(sorted.begin(), sorted.end()); - //Check to see if the file is empty - if so, add the headers + // Check to see if the file is empty - if so, add the headers emptyFile = this->IsFileEmpty(fileName); - //Open file and append by default + // Open file and append by default ofstream out; - out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app); + 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; + // 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) -{ +vector ResultDatabase::GetResultsForTest(const string& test) { // get only the given test results vector retval; - for (int i=0; i & -ResultDatabase::GetResults() const -{ - return results; -} +const vector& 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 index 4b63a02a1f..ca6a00fc91 100644 --- a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.h +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.h @@ -6,11 +6,11 @@ #include #include #include +using std::ifstream; +using std::ofstream; +using std::ostream; using std::string; using std::vector; -using std::ostream; -using std::ofstream; -using std::ifstream; // **************************************************************************** @@ -40,18 +40,16 @@ using std::ifstream; // Added a GetResults method as well, and made several functions const. // // **************************************************************************** -class ResultDatabase -{ - public: +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" + 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; @@ -59,41 +57,32 @@ class ResultDatabase double GetMean() const; double GetStdDev() const; - bool operator<(const Result &rhs) const; + bool operator<(const Result& rhs) const; - bool HadAnyFLTMAXValues() const - { - for (int i=0; i= FLT_MAX) - return true; + bool HadAnyFLTMAXValues() const { + for (int i = 0; i < value.size(); ++i) { + if (value[i] >= FLT_MAX) return true; } return false; } }; - protected: + 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; + 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: + private: bool IsFileEmpty(string fileName); - }; diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 4bd1bc0163..bbeddf2f48 100644 --- a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -6,64 +6,67 @@ #include "ResultDatabase.h" -enum MallocMode {MallocPinned, MallocUnpinned, MallocRegistered}; +enum MallocMode { MallocPinned, MallocUnpinned, MallocRegistered }; // Cmdline parms: -bool p_verbose = false; -MallocMode p_malloc_mode = MallocPinned; -int p_numa_ctl = -1; -int p_iterations = 10; -int p_beatsperiteration=1; -int p_device = 0; -int p_detailed = 0; -bool p_async = 0; -int p_alignedhost = 0; // align host allocs to this granularity, in bytes. 64 or 4096 are good values to try. -int p_onesize = 0; +bool p_verbose = false; +MallocMode p_malloc_mode = MallocPinned; +int p_numa_ctl = -1; +int p_iterations = 10; +int p_beatsperiteration = 1; +int p_device = 0; +int p_detailed = 0; +bool p_async = 0; +int p_alignedhost = + 0; // align host allocs to this granularity, in bytes. 64 or 4096 are good values to try. +int p_onesize = 0; -bool p_h2d = true; -bool p_d2h = true; -bool p_bidir = true; -bool p_p2p = false; +bool p_h2d = true; +bool p_d2h = true; +bool p_bidir = true; +bool p_p2p = false; //#define NO_CHECK -#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); \ - } \ -} +#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); \ + } \ + } std::string mallocModeString(int mallocMode) { switch (mallocMode) { - case MallocPinned : return "pinned"; - case MallocUnpinned: return "unpinned"; - case MallocRegistered: return "registered"; - default: return "mallocmode-UNKNOWN"; + case MallocPinned: + return "pinned"; + case MallocUnpinned: + return "unpinned"; + case MallocRegistered: + return "registered"; + default: + return "mallocmode-UNKNOWN"; }; }; // **************************************************************************** -int sizeToBytes(int size) { - return (size < 0) ? -size : size * 1024; -} +int sizeToBytes(int size) { return (size < 0) ? -size : size * 1024; } // **************************************************************************** -std::string sizeToString(int size) -{ +std::string sizeToString(int size) { using namespace std; stringstream ss; if (size < 0) { // char (-) lexically sorts before " " so will cause Byte values to be displayed before kB. - ss << "+" << setfill('0') << setw(3) << -size << "By"; + ss << "+" << setfill('0') << setw(3) << -size << "By"; } else { ss << size << "kB"; } @@ -72,8 +75,7 @@ std::string sizeToString(int size) // **************************************************************************** -hipError_t memcopy(void * dst, const void *src, size_t sizeBytes, enum hipMemcpyKind kind ) -{ +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 { @@ -82,11 +84,11 @@ hipError_t memcopy(void * dst, const void *src, size_t sizeBytes, enum hipMemcpy } - // **************************************************************************** // -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}; -int nSizes = sizeof(sizes) / sizeof(int); +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}; +int nSizes = sizeof(sizes) / sizeof(int); // **************************************************************************** @@ -108,74 +110,61 @@ int nSizes = sizeof(sizes) / sizeof(int); // Modifications: // Jeremy Meredith, Wed Dec 1 17:05:27 EST 2010 // Added calculation of latency estimate. -// Ben Sander - moved to standalone test +// Ben Sander - moved to standalone test // // **************************************************************************** -void RunBenchmark_H2D(ResultDatabase &resultDB) -{ - long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; +void RunBenchmark_H2D(ResultDatabase& resultDB) { + long long numMaxFloats = 1024 * (sizes[nSizes - 1]) / 4; hipSetDevice(p_device); // Create some host memory pattern - float *hostMem = NULL; - if (p_malloc_mode == MallocPinned) - { + float* hostMem = NULL; + if (p_malloc_mode == MallocPinned) { hipHostMalloc((void**)&hostMem, sizeof(float) * numMaxFloats); - while (hipGetLastError() != hipSuccess) - { + while (hipGetLastError() != hipSuccess) { // drop the size and try again if (p_verbose) std::cout << " - dropping size allocating pinned mem\n"; --nSizes; - if (nSizes < 1) - { + if (nSizes < 1) { std::cerr << "Error: Couldn't allocate any pinned buffer\n"; - return; + return; } - numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + numMaxFloats = 1024 * (sizes[nSizes - 1]) / 4; hipHostMalloc((void**)&hostMem, sizeof(float) * numMaxFloats); } - } - else if (p_malloc_mode == MallocUnpinned) - { + } else if (p_malloc_mode == MallocUnpinned) { if (p_alignedhost) { - hostMem = (float*)aligned_alloc(p_alignedhost, numMaxFloats*sizeof(float)); + hostMem = (float*)aligned_alloc(p_alignedhost, numMaxFloats * sizeof(float)); } else { hostMem = new float[numMaxFloats]; } - } - else if (p_malloc_mode == MallocRegistered) - { + } else if (p_malloc_mode == MallocRegistered) { if (p_numa_ctl == -1) { - hostMem = (float*)malloc(numMaxFloats*sizeof(float)); + hostMem = (float*)malloc(numMaxFloats * sizeof(float)); } hipHostRegister(hostMem, numMaxFloats * sizeof(float), 0); CHECK_HIP_ERROR(); - } - else - { + } else { assert(0); } - for (int i = 0; i < numMaxFloats; i++) - { + for (int i = 0; i < numMaxFloats; i++) { hostMem[i] = i % 77; } - float *device; + float* device; hipMalloc((void**)&device, sizeof(float) * numMaxFloats); - while (hipGetLastError() != hipSuccess) - { + while (hipGetLastError() != hipSuccess) { // drop the size and try again if (p_verbose) std::cout << " - dropping size allocating device mem\n"; --nSizes; - if (nSizes < 1) - { + if (nSizes < 1) { std::cerr << "Error: Couldn't allocate any device buffer\n"; return; } - numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + numMaxFloats = 1024 * (sizes[nSizes - 1]) / 4; hipMalloc((void**)&device, sizeof(float) * numMaxFloats); } @@ -186,13 +175,11 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) CHECK_HIP_ERROR(); // Three passes, forward and backward both - for (int pass = 0; pass < p_iterations; pass++) - { + for (int pass = 0; pass < p_iterations; pass++) { // store the times temporarily to estimate latency - //float times[nSizes]; + // float times[nSizes]; // Step through sizes forward on even passes and backward on odd - for (int i = 0; i < nSizes; i++) - { + for (int i = 0; i < nSizes; i++) { int sizeIndex; if ((pass % 2) == 0) sizeIndex = i; @@ -203,30 +190,32 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) const int nbytes = sizeToBytes(thisSize); hipEventRecord(start, 0); - for (int j=0;j1) { + if (p_beatsperiteration > 1) { sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), p_beatsperiteration); } else { sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str()); } - resultDB.AddResult(std::string("H2D_Bandwidth") + "_" + mallocModeString(p_malloc_mode), sizeStr, "GB/sec", speed); - resultDB.AddResult(std::string("H2D_Time") + mallocModeString(p_malloc_mode), sizeStr, "ms", t); + resultDB.AddResult(std::string("H2D_Bandwidth") + "_" + mallocModeString(p_malloc_mode), + sizeStr, "GB/sec", speed); + resultDB.AddResult(std::string("H2D_Time") + mallocModeString(p_malloc_mode), sizeStr, + "ms", t); if (p_onesize) { break; @@ -240,17 +229,16 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) #ifndef NO_CHECK - // Check. First reset the host memory, then copy-back result. Then compare against original ref value. - for (int i = 0; i < numMaxFloats; i++) - { + // Check. First reset the host memory, then copy-back result. Then compare against original + // ref value. + for (int i = 0; i < numMaxFloats; i++) { hostMem[i] = 0; } - hipMemcpy(hostMem, device, numMaxFloats*sizeof(float), hipMemcpyDeviceToHost); - for (int i = 0; i < numMaxFloats; i++) - { + hipMemcpy(hostMem, device, numMaxFloats * sizeof(float), hipMemcpyDeviceToHost); + for (int i = 0; i < numMaxFloats; i++) { float ref = i % 77; if (ref != hostMem[i]) { - printf ("error: H2D. i=%d reference:%6.f != copyback:%6.2f\n", i, ref, hostMem[i]); + printf("error: H2D. i=%d reference:%6.f != copyback:%6.2f\n", i, ref, hostMem[i]); } } #endif @@ -260,26 +248,26 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) hipFree((void*)device); CHECK_HIP_ERROR(); switch (p_malloc_mode) { - case MallocPinned: - hipHostFree((void*)hostMem); - CHECK_HIP_ERROR(); - break; + case MallocPinned: + hipHostFree((void*)hostMem); + CHECK_HIP_ERROR(); + break; - case MallocUnpinned: - if (p_alignedhost) { - delete[] hostMem; - } else { + case MallocUnpinned: + if (p_alignedhost) { + delete[] hostMem; + } else { + free(hostMem); + } + break; + + case MallocRegistered: + hipHostUnregister(hostMem); + CHECK_HIP_ERROR(); free(hostMem); - } - break; - - case MallocRegistered: - hipHostUnregister(hostMem); - CHECK_HIP_ERROR(); - free(hostMem); - break; - default: - assert(0); + break; + default: + assert(0); } @@ -288,85 +276,70 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) } - // **************************************************************************** -void RunBenchmark_D2H(ResultDatabase &resultDB) -{ - long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; +void RunBenchmark_D2H(ResultDatabase& resultDB) { + long long numMaxFloats = 1024 * (sizes[nSizes - 1]) / 4; // Create some host memory pattern - float *hostMem1; - float *hostMem2; - if (p_malloc_mode == MallocPinned) - { - hipHostMalloc((void**)&hostMem1, sizeof(float)*numMaxFloats); + float* hostMem1; + float* hostMem2; + if (p_malloc_mode == MallocPinned) { + hipHostMalloc((void**)&hostMem1, sizeof(float) * numMaxFloats); hipError_t err1 = hipGetLastError(); - hipHostMalloc((void**)&hostMem2, sizeof(float)*numMaxFloats); + hipHostMalloc((void**)&hostMem2, sizeof(float) * numMaxFloats); hipError_t err2 = hipGetLastError(); - while (err1 != hipSuccess || err2 != hipSuccess) - { + while (err1 != hipSuccess || err2 != hipSuccess) { // free the first buffer if only the second failed - if (err1 == hipSuccess) - hipHostFree((void*)hostMem1); + if (err1 == hipSuccess) hipHostFree((void*)hostMem1); // drop the size and try again if (p_verbose) std::cout << " - dropping size allocating pinned mem\n"; --nSizes; - if (nSizes < 1) - { + if (nSizes < 1) { std::cerr << "Error: Couldn't allocate any pinned buffer\n"; - return; + return; } - numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; - hipHostMalloc((void**)&hostMem1, sizeof(float)*numMaxFloats); - err1 = hipGetLastError(); - hipHostMalloc((void**)&hostMem2, sizeof(float)*numMaxFloats); - err2 = hipGetLastError(); + numMaxFloats = 1024 * (sizes[nSizes - 1]) / 4; + hipHostMalloc((void**)&hostMem1, sizeof(float) * numMaxFloats); + err1 = hipGetLastError(); + hipHostMalloc((void**)&hostMem2, sizeof(float) * numMaxFloats); + err2 = hipGetLastError(); } - } - else if (p_malloc_mode == MallocUnpinned) - { + } else if (p_malloc_mode == MallocUnpinned) { hostMem1 = new float[numMaxFloats]; hostMem2 = new float[numMaxFloats]; - } - else if (p_malloc_mode == MallocRegistered) - { + } else if (p_malloc_mode == MallocRegistered) { if (p_numa_ctl == -1) { - hostMem1 = (float*)malloc(numMaxFloats*sizeof(float)); - hostMem2 = (float*)malloc(numMaxFloats*sizeof(float)); + hostMem1 = (float*)malloc(numMaxFloats * sizeof(float)); + hostMem2 = (float*)malloc(numMaxFloats * sizeof(float)); } hipHostRegister(hostMem1, numMaxFloats * sizeof(float), 0); CHECK_HIP_ERROR(); hipHostRegister(hostMem2, numMaxFloats * sizeof(float), 0); CHECK_HIP_ERROR(); - } - else - { + } else { assert(0); } - for (int i=0; i1) { + if (p_beatsperiteration > 1) { sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), p_beatsperiteration); } else { sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str()); } - resultDB.AddResult(std::string("D2H_Bandwidth") +"_" + mallocModeString(p_malloc_mode) , sizeStr, "GB/sec", speed); - resultDB.AddResult(std::string("D2H_Time") +"_" + mallocModeString(p_malloc_mode) , sizeStr, "ms", t); + resultDB.AddResult(std::string("D2H_Bandwidth") + "_" + mallocModeString(p_malloc_mode), + sizeStr, "GB/sec", speed); + resultDB.AddResult(std::string("D2H_Time") + "_" + mallocModeString(p_malloc_mode), + sizeStr, "ms", t); if (p_onesize) { break; } @@ -427,12 +398,12 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) if (p_onesize) { numMaxFloats = sizeToBytes(p_onesize) / sizeof(float); } - // Check. First reset the host memory, then copy-back result. Then compare against original ref value. - for (int i = 0; i < numMaxFloats; i++) - { + // Check. First reset the host memory, then copy-back result. Then compare against original + // ref value. + for (int i = 0; i < numMaxFloats; i++) { float ref = i % 77; if (ref != hostMem2[i]) { - printf ("error: D2H. i=%d reference:%6.f != copyback:%6.2f\n", i, ref, hostMem2[i]); + printf("error: D2H. i=%d reference:%6.f != copyback:%6.2f\n", i, ref, hostMem2[i]); } } @@ -441,25 +412,25 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) CHECK_HIP_ERROR(); switch (p_malloc_mode) { - case MallocPinned: - hipHostFree((void*)hostMem1); - CHECK_HIP_ERROR(); - hipHostFree((void*)hostMem2); - CHECK_HIP_ERROR(); - break; - case MallocUnpinned: - delete[] hostMem1; - delete[] hostMem2; - break; - case MallocRegistered: - hipHostUnregister(hostMem1); - CHECK_HIP_ERROR(); - free(hostMem1); - hipHostUnregister(hostMem2); - free(hostMem2); - break; - default: - assert(0); + case MallocPinned: + hipHostFree((void*)hostMem1); + CHECK_HIP_ERROR(); + hipHostFree((void*)hostMem2); + CHECK_HIP_ERROR(); + break; + case MallocUnpinned: + delete[] hostMem1; + delete[] hostMem2; + break; + case MallocRegistered: + hipHostUnregister(hostMem1); + CHECK_HIP_ERROR(); + free(hostMem1); + hipHostUnregister(hostMem2); + free(hostMem2); + break; + default: + assert(0); } hipEventDestroy(start); @@ -467,9 +438,8 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) } -void RunBenchmark_Bidir(ResultDatabase &resultDB) -{ - long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; +void RunBenchmark_Bidir(ResultDatabase& resultDB) { + long long numMaxFloats = 1024 * (sizes[nSizes - 1]) / 4; hipSetDevice(p_device); @@ -477,11 +447,9 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB) // Create some host memory pattern - float *hostMem[2] = {NULL, NULL}; - if (p_malloc_mode == MallocPinned) - { - while (1) - { + float* hostMem[2] = {NULL, NULL}; + if (p_malloc_mode == MallocPinned) { + while (1) { hipError_t e1 = hipHostMalloc((void**)&hostMem[0], sizeof(float) * numMaxFloats); hipError_t e2 = hipHostMalloc((void**)&hostMem[1], sizeof(float) * numMaxFloats); @@ -491,43 +459,35 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB) // drop the size and try again if (p_verbose) std::cout << " - dropping size allocating pinned mem\n"; --nSizes; - if (nSizes < 1) - { + if (nSizes < 1) { std::cerr << "Error: Couldn't allocate any pinned buffer\n"; - return; + return; } - numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + numMaxFloats = 1024 * (sizes[nSizes - 1]) / 4; } } - } - else if (p_malloc_mode == MallocUnpinned) - { + } else if (p_malloc_mode == MallocUnpinned) { hostMem[0] = new float[numMaxFloats]; hostMem[1] = new float[numMaxFloats]; - } - else if (p_malloc_mode == MallocRegistered) - { + } else if (p_malloc_mode == MallocRegistered) { if (p_numa_ctl == -1) { - hostMem[0] = (float*)malloc(numMaxFloats*sizeof(float)); - hostMem[1] = (float*)malloc(numMaxFloats*sizeof(float)); + hostMem[0] = (float*)malloc(numMaxFloats * sizeof(float)); + hostMem[1] = (float*)malloc(numMaxFloats * sizeof(float)); } hipHostRegister(hostMem[0], numMaxFloats * sizeof(float), 0); CHECK_HIP_ERROR(); hipHostRegister(hostMem[1], numMaxFloats * sizeof(float), 0); CHECK_HIP_ERROR(); - } - else - { + } else { assert(0); } - for (int i = 0; i < numMaxFloats; i++) - { + for (int i = 0; i < numMaxFloats; i++) { hostMem[0][i] = i % 77; } - float *deviceMem[2]; + 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); @@ -542,12 +502,11 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB) // drop the size and try again if (p_verbose) std::cout << " - dropping size allocating device mem\n"; --nSizes; - if (nSizes < 1) - { + if (nSizes < 1) { std::cerr << "Error: Couldn't allocate any device buffer\n"; return; } - numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + numMaxFloats = 1024 * (sizes[nSizes - 1]) / 4; } }; @@ -563,13 +522,11 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB) hipStreamCreate(&stream[1]); // Three passes, forward and backward both - for (int pass = 0; pass < p_iterations; pass++) - { + for (int pass = 0; pass < p_iterations; pass++) { // store the times temporarily to estimate latency - //float times[nSizes]; + // float times[nSizes]; // Step through sizes forward on even passes and backward on odd - for (int i = 0; i < nSizes; i++) - { + for (int i = 0; i < nSizes; i++) { int sizeIndex; if ((pass % 2) == 0) sizeIndex = i; @@ -580,25 +537,26 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB) const int nbytes = sizeToBytes(thisSize); hipEventRecord(start, 0); - hipMemcpyAsync(deviceMem[0], hostMem[0], nbytes, hipMemcpyHostToDevice, stream[0]); - hipMemcpyAsync(hostMem[1], deviceMem[1], nbytes, hipMemcpyDeviceToHost, stream[1]); + 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); // Convert to GB/sec - if (p_verbose) - { - std::cerr << "size " << sizeToString(thisSize) << " took " << t << - " ms\n"; + if (p_verbose) { + std::cerr << "size " << sizeToString(thisSize) << " took " << t << " ms\n"; } - double speed = (double(sizeToBytes(2*thisSize)) / (1000*1000)) / t; + double speed = (double(sizeToBytes(2 * thisSize)) / (1000 * 1000)) / t; char sizeStr[256]; sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str()); - resultDB.AddResult(std::string("Bidir_Bandwidth") + "_" + mallocModeString(p_malloc_mode), sizeStr, "GB/sec", speed); - resultDB.AddResult(std::string("Bidir_Time") + "_" + mallocModeString(p_malloc_mode), sizeStr, "ms", t); + resultDB.AddResult( + std::string("Bidir_Bandwidth") + "_" + mallocModeString(p_malloc_mode), sizeStr, + "GB/sec", speed); + resultDB.AddResult(std::string("Bidir_Time") + "_" + mallocModeString(p_malloc_mode), + sizeStr, "ms", t); } } @@ -607,24 +565,24 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB) hipFree((void*)deviceMem[1]); CHECK_HIP_ERROR(); switch (p_malloc_mode) { - case MallocPinned: - hipHostFree((void*)hostMem[0]); - hipHostFree((void*)hostMem[1]); - CHECK_HIP_ERROR(); - break; - case MallocUnpinned: - delete[] hostMem[0]; - delete[] hostMem[1]; - break; - case MallocRegistered: - for (int i=0; i<2; i++) { - hipHostUnregister(hostMem[i]); + case MallocPinned: + hipHostFree((void*)hostMem[0]); + hipHostFree((void*)hostMem[1]); CHECK_HIP_ERROR(); - free(hostMem[i]); - } - break; - default: - assert(0); + break; + case MallocUnpinned: + delete[] hostMem[0]; + delete[] hostMem[1]; + break; + case MallocRegistered: + for (int i = 0; i < 2; i++) { + hipHostUnregister(hostMem[i]); + CHECK_HIP_ERROR(); + free(hostMem[i]); + } + break; + default: + assert(0); }; hipEventDestroy(start); @@ -634,37 +592,31 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB) } - - - -#define failed(...) \ - printf ("error: ");\ - printf (__VA_ARGS__);\ - printf ("\n");\ +#define failed(...) \ + printf("error: "); \ + printf(__VA_ARGS__); \ + printf("\n"); \ exit(EXIT_FAILURE); -int parseInt(const char *str, int *output) -{ - char *next; +int parseInt(const char* str, int* output) { + char* next; *output = strtol(str, &next, 0); return !strlen(next); } -void checkPeer2PeerSupport() -{ +void checkPeer2PeerSupport() { int deviceCnt; hipGetDeviceCount(&deviceCnt); std::cout << "Total no. of available gpu #" << deviceCnt << "\n" << std::endl; - for(int deviceId=0; deviceIdhost then host-->GPU2)\n\n" << std::endl; + std::cout << "\nNote: For non-supported peer2peer devices, memcopy will use/follow the normal " + "behaviour (GPU1-->host then host-->GPU2)\n\n" + << std::endl; } -void enablePeer2Peer(int currentGpu, int peerGpu) -{ +void enablePeer2Peer(int currentGpu, int peerGpu) { int canAccessPeer; hipSetDevice(currentGpu); hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu); - if(canAccessPeer==1){ + if (canAccessPeer == 1) { hipDeviceEnablePeerAccess(peerGpu, 0); } } -void disablePeer2Peer(int currentGpu, int peerGpu) -{ +void disablePeer2Peer(int currentGpu, int peerGpu) { int canAccessPeer; hipSetDevice(currentGpu); hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu); - if(canAccessPeer==1){ + if (canAccessPeer == 1) { hipDeviceDisablePeerAccess(peerGpu); } } -std::string gpuIDToString(int gpuID) -{ +std::string gpuIDToString(int gpuID) { using namespace std; stringstream ss; - ss << gpuID; + ss << gpuID; return ss.str(); } -void RunBenchmark_P2P_Unidir(ResultDatabase &resultDB) -{ +void RunBenchmark_P2P_Unidir(ResultDatabase& resultDB) { int gpuCount; hipGetDeviceCount(&gpuCount); int currentGpu, peerGpu; - long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + long long numMaxFloats = 1024 * (sizes[nSizes - 1]) / 4; - for (currentGpu=0; currentGpu1) { - sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), p_beatsperiteration); + if (p_beatsperiteration > 1) { + sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), + p_beatsperiteration); } else { sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str()); } @@ -809,14 +757,17 @@ void RunBenchmark_P2P_Unidir(ResultDatabase &resultDB) cGpu = gpuIDToString(currentGpu); pGpu = gpuIDToString(peerGpu); - resultDB.AddResult(std::string("p2p_uni") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "GB/sec", speed); - resultDB.AddResult(std::string("P2P_uni") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "ms", t); + resultDB.AddResult(std::string("p2p_uni") + "_gpu" + std::string(cGpu) + + "_gpu" + std::string(pGpu), + sizeStr, "GB/sec", speed); + resultDB.AddResult(std::string("P2P_uni") + "_gpu" + std::string(cGpu) + + "_gpu" + std::string(pGpu), + sizeStr, "ms", t); if (p_onesize) { break; } } - } if (p_onesize) { @@ -839,13 +790,10 @@ void RunBenchmark_P2P_Unidir(ResultDatabase &resultDB) hipSetDevice(currentGpu); hipDeviceReset(); } - } - } -void RunBenchmark_P2P_Bidir(ResultDatabase &resultDB) { - +void RunBenchmark_P2P_Bidir(ResultDatabase& resultDB) { int gpuCount; hipGetDeviceCount(&gpuCount); @@ -853,14 +801,11 @@ void RunBenchmark_P2P_Bidir(ResultDatabase &resultDB) { int currentGpu, peerGpu; - long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + long long numMaxFloats = 1024 * (sizes[nSizes - 1]) / 4; - for (currentGpu=0; currentGpu1) { - sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), p_beatsperiteration); + if (p_beatsperiteration > 1) { + sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), + p_beatsperiteration); } else { sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str()); } @@ -934,14 +881,17 @@ void RunBenchmark_P2P_Bidir(ResultDatabase &resultDB) { cGpu = gpuIDToString(currentGpu); pGpu = gpuIDToString(peerGpu); - resultDB.AddResult(std::string("p2p_bi") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "GB/sec", speed); - resultDB.AddResult(std::string("P2P_bi") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "ms", t); + resultDB.AddResult(std::string("p2p_bi") + "_gpu" + std::string(cGpu) + "_gpu" + + std::string(pGpu), + sizeStr, "GB/sec", speed); + resultDB.AddResult(std::string("P2P_bi") + "_gpu" + std::string(cGpu) + "_gpu" + + std::string(pGpu), + sizeStr, "ms", t); if (p_onesize) { break; } } - } if (p_onesize) { @@ -953,7 +903,7 @@ void RunBenchmark_P2P_Bidir(ResultDatabase &resultDB) { hipEventDestroy(start); hipEventDestroy(stop); - for (int i=0; i<2; i++) { + for (int i = 0; i < 2; i++) { hipStreamDestroy(stream[i]); hipFree((void*)currentGpuMem[i]); @@ -975,75 +925,81 @@ void printConfig() { hipDeviceProp_t props; hipGetDeviceProperties(&props, p_device); - printf ("Device:%s Mem=%.1fGB #CUs=%d Freq=%.0fMhz MallocMode=%s\n", props.name, props.totalGlobalMem/1024.0/1024.0/1024.0, props.multiProcessorCount, props.clockRate/1000.0, mallocModeString(p_malloc_mode).c_str()); + printf("Device:%s Mem=%.1fGB #CUs=%d Freq=%.0fMhz MallocMode=%s\n", props.name, + props.totalGlobalMem / 1024.0 / 1024.0 / 1024.0, props.multiProcessorCount, + props.clockRate / 1000.0, mallocModeString(p_malloc_mode).c_str()); } void help() { - printf ("Usage: hipBusBandwidth [OPTIONS]\n"); - printf (" --iterations, -i : Number of copy iterations to run.\n"); - printf (" --beatsperiterations, -b : Number of beats (back-to-back copies of same size) per iteration 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 (" --p2p : Run only peer2peer unidir and bidir copy tests.\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"); - printf (" --onesize, -o : Only run one measurement, at specified size (in KB, or if negative in bytes)\n"); - + printf("Usage: hipBusBandwidth [OPTIONS]\n"); + printf(" --iterations, -i : Number of copy iterations to run.\n"); + printf( + " --beatsperiterations, -b : Number of beats (back-to-back copies of same size) per " + "iteration 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(" --p2p : Run only peer2peer unidir and bidir copy tests.\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"); + printf( + " --onesize, -o : Only run one measurement, at specified size (in KB, or if " + "negative in bytes)\n"); }; -int parseStandardArguments(int argc, char *argv[]) -{ +int parseStandardArguments(int argc, char* argv[]) { for (int i = 1; i < argc; i++) { - const char *arg = argv[i]; + const char* arg = argv[i]; if (!strcmp(arg, " ")) { // skip NULL args. } else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) { if (++i >= argc || !parseInt(argv[i], &p_iterations)) { - failed("Bad iterations argument"); + failed("Bad iterations argument"); } } else if (!strcmp(arg, "--beatsperiteration") || (!strcmp(arg, "-b"))) { if (++i >= argc || !parseInt(argv[i], &p_beatsperiteration)) { - failed("Bad beatsperiteration argument"); + failed("Bad beatsperiteration argument"); } } else if (!strcmp(arg, "--device") || (!strcmp(arg, "-d"))) { if (++i >= argc || !parseInt(argv[i], &p_device)) { - failed("Bad device argument"); + failed("Bad device argument"); } } else if (!strcmp(arg, "--onesize") || (!strcmp(arg, "-o"))) { if (++i >= argc || !parseInt(argv[i], &p_onesize)) { - failed("Bad onesize argument"); + failed("Bad onesize argument"); } } else if (!strcmp(arg, "--unpinned")) { p_malloc_mode = MallocUnpinned; } else if (!strcmp(arg, "--registered")) { p_malloc_mode = MallocRegistered; } 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_h2d = false; + p_d2h = false; p_bidir = true; } else if (!strcmp(arg, "--p2p")) { - p_h2d = false; - p_d2h = false; + p_h2d = false; + p_d2h = false; p_bidir = false; - p_p2p = true; + p_p2p = true; - } else if (!strcmp(arg, "--help") || (!strcmp(arg, "-h"))) { + } else if (!strcmp(arg, "--help") || (!strcmp(arg, "-h"))) { help(); exit(EXIT_SUCCESS); @@ -1057,15 +1013,13 @@ int parseStandardArguments(int argc, char *argv[]) } else { failed("Bad argument '%s'", arg); } - } + } return 0; }; - -int main(int argc, char *argv[]) -{ +int main(int argc, char* argv[]) { parseStandardArguments(argc, argv); if (p_p2p) { @@ -1083,8 +1037,7 @@ int main(int argc, char *argv[]) resultDB_Unidir.DumpDetailed(std::cout); resultDB_Bidir.DumpDetailed(std::cout); } - } - else { + } else { printConfig(); if (p_h2d) { diff --git a/projects/hip-tests/samples/1_Utils/hipCommander/ResultDatabase.cpp b/projects/hip-tests/samples/1_Utils/hipCommander/ResultDatabase.cpp index 2ec686f260..1b1ee3a70d 100644 --- a/projects/hip-tests/samples/1_Utils/hipCommander/ResultDatabase.cpp +++ b/projects/hip-tests/samples/1_Utils/hipCommander/ResultDatabase.cpp @@ -7,93 +7,69 @@ 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 +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 ResultDatabase::Result::GetMin() const { double r = FLT_MAX; - for (int i=0; i= 100) - return value[n-1]; + if (q <= 0) return value[0]; + if (q >= 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.)); + 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]; + if (frac == 0) return sorted[index_lo]; double lo = sorted[index_lo]; double hi = sorted[index_lo + 1]; - return lo + (hi-lo)*frac; + return lo + (hi - lo) * frac; } -double ResultDatabase::Result::GetMean() const -{ +double ResultDatabase::Result::GetMean() const { double r = 0; - for (int i=0; i &values) -{ - for (int i=0; i& values) { + for (int i = 0; i < values.size(); i++) { AddResult(test, atts, unit, values[i]); } } -static string RemoveAllButLeadingSpaces(const string &a) -{ +static string RemoveAllButLeadingSpaces(const string& a) { string b; int n = a.length(); int i = 0; - while (i= results.size()) - { + if (index >= results.size()) { Result r; r.test = test; r.atts = atts; @@ -186,40 +146,32 @@ void ResultDatabase::AddResult(const string &test_orig, // Changed note about missing values to be worded a little better. // // **************************************************************************** -void ResultDatabase::DumpDetailed(ostream &out) -{ +void ResultDatabase::DumpDetailed(ostream& out) { vector sorted(results); sort(sorted.begin(), sorted.end()); - const int testNameW = 24 ; + const int testNameW = 24; const int attW = 12; const int fieldW = 11; out << std::fixed << right << std::setprecision(4); int maxtrials = 1; - for (int i=0; i maxtrials) - maxtrials = sorted[i].value.size(); + for (int i = 0; i < sorted.size(); i++) { + if (sorted[i].value.size() > 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 << setw(testNameW) << "test\t" - << setw(attW) << "atts\t" - << setw(fieldW) - << "median\t" + out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "median\t" << "mean\t" << "stddev\t" << "min\t" << "max\t"; - for (int i=0; i sorted(results); sort(sorted.begin(), sorted.end()); - const int testNameW = 24 ; + const int testNameW = 24; const int attW = 12; const int fieldW = 9; 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 << setw(testNameW) << "test\t" - << setw(attW) << "atts\t" - << setw(fieldW) - << "units\t" + out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "units\t" << "median\t" << "mean\t" << "stddev\t" @@ -301,9 +248,8 @@ void ResultDatabase::DumpSummary(ostream &out) << "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 + // Check to see if the file is empty - if so, add the headers emptyFile = this->IsFileEmpty(fileName); - //Open file and append by default + // Open file and append by default ofstream out; - out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app); + 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; + // 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) -{ +vector ResultDatabase::GetResultsForTest(const string& test) { // get only the given test results vector retval; - for (int i=0; i & -ResultDatabase::GetResults() const -{ - return results; -} +const vector& ResultDatabase::GetResults() const { return results; } diff --git a/projects/hip-tests/samples/1_Utils/hipCommander/ResultDatabase.h b/projects/hip-tests/samples/1_Utils/hipCommander/ResultDatabase.h index 4b63a02a1f..ca6a00fc91 100644 --- a/projects/hip-tests/samples/1_Utils/hipCommander/ResultDatabase.h +++ b/projects/hip-tests/samples/1_Utils/hipCommander/ResultDatabase.h @@ -6,11 +6,11 @@ #include #include #include +using std::ifstream; +using std::ofstream; +using std::ostream; using std::string; using std::vector; -using std::ostream; -using std::ofstream; -using std::ifstream; // **************************************************************************** @@ -40,18 +40,16 @@ using std::ifstream; // Added a GetResults method as well, and made several functions const. // // **************************************************************************** -class ResultDatabase -{ - public: +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" + 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; @@ -59,41 +57,32 @@ class ResultDatabase double GetMean() const; double GetStdDev() const; - bool operator<(const Result &rhs) const; + bool operator<(const Result& rhs) const; - bool HadAnyFLTMAXValues() const - { - for (int i=0; i= FLT_MAX) - return true; + bool HadAnyFLTMAXValues() const { + for (int i = 0; i < value.size(); ++i) { + if (value[i] >= FLT_MAX) return true; } return false; } }; - protected: + 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; + 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: + private: bool IsFileEmpty(string fileName); - }; diff --git a/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp b/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp index 4b93180b18..345c2d5006 100644 --- a/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp @@ -21,55 +21,46 @@ bool g_printedTiming = false; // Cmdline parms: -int p_device = 0; -const char* p_command = "setstream(1); H2D; NullKernel; D2H;"; -const char* p_file = nullptr; -unsigned p_verbose = 0x0; -unsigned p_db = 0x0; -unsigned p_blockingSync = 0x0; +int p_device = 0; +const char* p_command = "setstream(1); H2D; NullKernel; D2H;"; +const char* p_file = nullptr; +unsigned p_verbose = 0x0; +unsigned p_db = 0x0; +unsigned p_blockingSync = 0x0; //--- -int p_iterations = 1; +int p_iterations = 1; -#define KNRM "\x1B[0m" -#define KRED "\x1B[31m" -#define KGRN "\x1B[32m" +#define KNRM "\x1B[0m" +#define KRED "\x1B[31m" +#define KGRN "\x1B[32m" -#define failed(...) \ - printf ("error: ");\ - printf (__VA_ARGS__);\ - printf ("\n");\ +#define failed(...) \ + printf("error: "); \ + printf(__VA_ARGS__); \ + printf("\n"); \ abort(); -#define HIPCHECK(error) \ -{\ - hipError_t localError = error; \ - if (localError != hipSuccess) { \ - printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \ - KRED,hipGetErrorString(localError), localError,\ - #error,\ - __FILE__, __LINE__,KNRM); \ - failed("API returned error code.");\ - }\ -} -#define HIPASSERT(condition, msg) \ - if (! (condition) ) { \ - failed("%sassertion %s at %s:%d: %s%s\n", \ - KRED, #condition,\ - __FILE__, __LINE__,msg, KNRM); \ +#define HIPCHECK(error) \ + { \ + hipError_t localError = error; \ + if (localError != hipSuccess) { \ + printf("%serror: '%s'(%d) from %s at %s:%d%s\n", KRED, hipGetErrorString(localError), \ + localError, #error, __FILE__, __LINE__, KNRM); \ + failed("API returned error code."); \ + } \ + } +#define HIPASSERT(condition, msg) \ + if (!(condition)) { \ + failed("%sassertion %s at %s:%d: %s%s\n", KRED, #condition, __FILE__, __LINE__, msg, \ + KNRM); \ } - - - - - -int parseInt(const char *str, int *output) -{ - char *next; +int parseInt(const char* str, int* output) { + char* next; *output = strtol(str, &next, 0); return !strlen(next); } @@ -79,50 +70,50 @@ void printConfig() { hipDeviceProp_t props; HIPCHECK(hipGetDeviceProperties(&props, p_device)); - printf ("Device:%s Mem=%.1fGB #CUs=%d Freq=%.0fMhz\n", props.name, props.totalGlobalMem/1024.0/1024.0/1024.0, props.multiProcessorCount, props.clockRate/1000.0); + printf("Device:%s Mem=%.1fGB #CUs=%d Freq=%.0fMhz\n", props.name, + props.totalGlobalMem / 1024.0 / 1024.0 / 1024.0, props.multiProcessorCount, + props.clockRate / 1000.0); } - - void help() { - printf ("Usage: hipBusBandwidth [OPTIONS]\n"); - printf (" --file, -f : Read string of commands from file\n"); - printf (" --command, -c : String specifying commands to run.\n"); - printf (" --iterations, -i : Number of copy iterations to run.\n"); - printf (" --device, -d : Device ID to use (0..numDevices).\n"); - printf (" --verbose, -v : Verbose printing of status. Fore more info, combine with HIP_TRACE_API on ROCm\n"); + printf("Usage: hipBusBandwidth [OPTIONS]\n"); + printf(" --file, -f : Read string of commands from file\n"); + printf(" --command, -c : String specifying commands to run.\n"); + printf(" --iterations, -i : Number of copy iterations to run.\n"); + printf(" --device, -d : Device ID to use (0..numDevices).\n"); + printf( + " --verbose, -v : Verbose printing of status. Fore more info, combine with " + "HIP_TRACE_API on ROCm\n"); }; - -int parseStandardArguments(int argc, char *argv[]) -{ +int parseStandardArguments(int argc, char* argv[]) { for (int i = 1; i < argc; i++) { - const char *arg = argv[i]; + const char* arg = argv[i]; if (!strcmp(arg, " ")) { // skip NULL args. } else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) { if (++i >= argc || !parseInt(argv[i], &p_iterations)) { - failed("Bad --iterations argument"); + failed("Bad --iterations argument"); } } else if (!strcmp(arg, "--device") || (!strcmp(arg, "-d"))) { if (++i >= argc || !parseInt(argv[i], &p_device)) { - failed("Bad --device argument"); + failed("Bad --device argument"); } } else if (!strcmp(arg, "--file") || (!strcmp(arg, "-f"))) { if (++i >= argc) { - failed("Bad --file argument"); + failed("Bad --file argument"); } else { p_file = argv[i]; } } else if (!strcmp(arg, "--commands") || (!strcmp(arg, "-c"))) { if (++i >= argc) { - failed("Bad --commands argument"); + failed("Bad --commands argument"); } else { p_command = argv[i]; } @@ -134,21 +125,20 @@ int parseStandardArguments(int argc, char *argv[]) p_blockingSync = 1; - } else if (!strcmp(arg, "--help") || (!strcmp(arg, "-h"))) { + } else if (!strcmp(arg, "--help") || (!strcmp(arg, "-h"))) { help(); exit(EXIT_SUCCESS); } else { failed("Bad argument '%s'", arg); } - } + } return 0; }; // Returns the current system time in microseconds -inline long long get_time() -{ +inline long long get_time() { struct timeval tv; gettimeofday(&tv, 0); return (tv.tv_sec * 1000000) + tv.tv_usec; @@ -161,142 +151,132 @@ class Command; //================================================================================================= // A stream of commands , specified as a string. class CommandStream { -public: + public: // State that is inherited by sub-blocks: struct CommandStreamState { - hipStream_t _currentStream; - std::vector _streams; - vector _subBlocks; + hipStream_t _currentStream; + std::vector _streams; + vector _subBlocks; }; -public: + + public: CommandStream(std::string commandStreamString, int iterations); ~CommandStream(); hipStream_t currentStream() const { return _state._currentStream; }; - void print(const std::string &indent="") const; - void printBrief(std::ostream &s=std::cout) const ; + void print(const std::string& indent = "") const; + void printBrief(std::ostream& s = std::cout) const; void run(); void recordTime(); - void printTiming(int iterations=0); + void printTiming(int iterations = 0); - CommandStream *currentCommandStream() { + CommandStream* currentCommandStream() { return _parseInSubBlock ? _state._subBlocks.back() : this; }; - void enterSubBlock(CommandStream *commandStream) { + void enterSubBlock(CommandStream* commandStream) { _parseInSubBlock = true; _state._subBlocks.push_back(commandStream); }; - void exitSubBlock() { - _parseInSubBlock = false; - }; + void exitSubBlock() { _parseInSubBlock = false; }; - void setParent(CommandStream *parentCmdStream) - { - _parentCommandStream = parentCmdStream; - _state = parentCmdStream->_state; + void setParent(CommandStream* parentCmdStream) { + _parentCommandStream = parentCmdStream; + _state = parentCmdStream->_state; }; - CommandStream * getParent() { return _parentCommandStream; }; + CommandStream* getParent() { return _parentCommandStream; }; void setStream(int streamIndex); - CommandStreamState &getState() { return _state; }; + CommandStreamState& getState() { return _state; }; -private: - static void tokenize(const std::string &s, char delim, std::vector &tokens); + private: + static void tokenize(const std::string& s, char delim, std::vector& tokens); void parse(const std::string fullCmd); -protected: - CommandStreamState _state; -private: - + protected: + CommandStreamState _state; + private: // List of commands to run in this stream: - std::vector _commands; - + std::vector _commands; // Number of iterations to run the command loop - int _iterations; - - + int _iterations; // Us to run the the command-stream. Only valid after run is called. - long long _startTime; - double _elapsedUs; + long long _startTime; + double _elapsedUs; // Track nested loop of command streams: - CommandStream *_parentCommandStream; + CommandStream* _parentCommandStream; // Track if we are parsing commands in the subblock. - bool _parseInSubBlock; - + bool _parseInSubBlock; }; //================================================================================================= class Command { -public: - + public: // @p minArgs : Minimum arguments for command. -1 = don't check. - // @p maxArgs : Minimum arguments for command. 0 means min=max, ie exact #arguments expected. -1 = don't check max. - Command(CommandStream *cmdStream, const std::vector &args, int minArgs=0, int maxArgs=0) - : _commandStream(cmdStream), - _args(args) - { + // @p maxArgs : Minimum arguments for command. 0 means min=max, ie exact #arguments expected. + // -1 = don't check max. + Command(CommandStream* cmdStream, const std::vector& args, int minArgs = 0, + int maxArgs = 0) + : _commandStream(cmdStream), _args(args) { int numArgs = args.size() - 1; - if ((minArgs != -1 ) && (numArgs < minArgs)) { + if ((minArgs != -1) && (numArgs < minArgs)) { // TODO - print full command here. - failed ("Not enough arguments for command %s. (Expected %d, got %d)", args[0].c_str(), minArgs, numArgs); + failed("Not enough arguments for command %s. (Expected %d, got %d)", args[0].c_str(), + minArgs, numArgs); } // Check for an exact number of arguments: if (maxArgs == 0) { maxArgs = minArgs; } - if ((maxArgs != -1 ) && (numArgs > maxArgs)) { - failed ("Too many arguments for command %s. (Expected %d, got %d)", args[0].c_str(), maxArgs, numArgs); + if ((maxArgs != -1) && (numArgs > maxArgs)) { + failed("Too many arguments for command %s. (Expected %d, got %d)", args[0].c_str(), + maxArgs, numArgs); } }; - void printBrief(std::ostream &s=std::cout) const - { - s << _args[0]; - } + void printBrief(std::ostream& s = std::cout) const { s << _args[0]; } - virtual ~Command() {}; + virtual ~Command(){}; - virtual void print(const std::string &indent = "") const { + virtual void print(const std::string& indent = "") const { std::cout << indent << "["; - std::for_each(_args.begin(), _args.end(), [] (const std::string &s) { - std::cout << s; - }); + std::for_each(_args.begin(), _args.end(), [](const std::string& s) { std::cout << s; }); std::cout << "]"; }; virtual void run() = 0; -protected: - int readIntArg(int argIndex, const std::string &argName) - { + protected: + int readIntArg(int argIndex, const std::string& argName) { // TODO - catch references to non-existant arguments here. int argVal; try { argVal = std::stoi(_args[argIndex]); } catch (std::invalid_argument) { - failed ("Command %s has bad %s argument ('%s')", _args[0].c_str(), argName.c_str(), _args[argIndex].c_str()); + failed("Command %s has bad %s argument ('%s')", _args[0].c_str(), argName.c_str(), + _args[argIndex].c_str()); } return argVal; } -protected: - CommandStream *_commandStream; - std::vector _args; + + protected: + CommandStream* _commandStream; + std::vector _args; }; @@ -307,33 +287,30 @@ protected: #ifdef __HIP_PLATFORM_HCC__ //================================================================================================= // Use Aql to launch the NULL kernel. -class AqlKernelCommand : public Command -{ -public: - AqlKernelCommand(CommandStream *cmdStream, const std::vector args) : - Command(cmdStream, args) - { - hc::accelerator_view *av; +class AqlKernelCommand : public Command { + public: + AqlKernelCommand(CommandStream* cmdStream, const std::vector args) + : Command(cmdStream, args) { + hc::accelerator_view* av; HIPCHECK(hipHccGetAcceleratorView(cmdStream->currentStream(), &av)); hc::accelerator acc = av->get_accelerator(); hsa_region_t systemRegion = *(hsa_region_t*)acc.get_hsa_am_system_region(); - _hsaAgent = *(hsa_agent_t*) acc.get_hsa_agent(); + _hsaAgent = *(hsa_agent_t*)acc.get_hsa_agent(); std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); std::streamsize fsize = file.tellg(); file.seekg(0, std::ios::beg); std::vector buffer(fsize); - if (file.read(buffer.data(), fsize)) - { + if (file.read(buffer.data(), fsize)) { uint64_t elfSize = ElfSize(&buffer[0]); assert(fsize == elfSize); - //TODO - replace module load code with explicit module load and unload. + // TODO - replace module load code with explicit module load and unload. hipModule_t module; HIPCHECK(hipModuleLoadData(&module, &buffer[0])); @@ -344,56 +321,53 @@ public: } }; - ~AqlKernelCommand() {}; + ~AqlKernelCommand(){}; void run() override { #define LEN 64 - uint32_t len = LEN; - uint32_t one = 1; + uint32_t len = LEN; + uint32_t one = 1; - float *Ad = NULL; + float* Ad = NULL; size_t argSize = 36; char argBuffer[argSize]; - *(uint32_t*) (&argBuffer[0]) = len; - *(uint32_t*) (&argBuffer[4]) = one; - *(uint32_t*) (&argBuffer[8]) = one; - *(uint32_t*) (&argBuffer[12]) = len; - *(uint32_t*) (&argBuffer[16]) = one; - *(uint32_t*) (&argBuffer[20]) = one; - *(float**) (&argBuffer[24]) = Ad; // Ad pointer argument + *(uint32_t*)(&argBuffer[0]) = len; + *(uint32_t*)(&argBuffer[4]) = one; + *(uint32_t*)(&argBuffer[8]) = one; + *(uint32_t*)(&argBuffer[12]) = len; + *(uint32_t*)(&argBuffer[16]) = one; + *(uint32_t*)(&argBuffer[20]) = one; + *(float**)(&argBuffer[24]) = Ad; // Ad pointer argument - void *config[] = { - HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], - HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, - HIP_LAUNCH_PARAM_END - }; + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], + HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, HIP_LAUNCH_PARAM_END}; - hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); + hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); }; -public: - hsa_queue_t _hsaQueue; - hsa_agent_t _hsaAgent; + public: + hsa_queue_t _hsaQueue; + hsa_agent_t _hsaAgent; hipFunction_t _function; -private: - static uint64_t ElfSize(const void *emi){ - const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi; - const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff); + private: + static uint64_t ElfSize(const void* emi) { + const Elf64_Ehdr* ehdr = (const Elf64_Ehdr*)emi; + const Elf64_Shdr* shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff); uint64_t max_offset = ehdr->e_shoff; uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum; - for(uint16_t i=0;i < ehdr->e_shnum;++i){ + for (uint16_t i = 0; i < ehdr->e_shnum; ++i) { uint64_t cur_offset = static_cast(shdr[i].sh_offset); - if(max_offset < cur_offset){ + if (max_offset < cur_offset) { max_offset = cur_offset; total_size = max_offset; - if(SHT_NOBITS != shdr[i].sh_type){ + if (SHT_NOBITS != shdr[i].sh_type) { total_size += static_cast(shdr[i].sh_size); } } @@ -405,122 +379,110 @@ private: //================================================================================================= // HCC optimizes away fully NULL kernel calls, so run one that is nearly null: -class ModuleKernelCommand : public Command -{ -public: - ModuleKernelCommand(CommandStream *cmdStream, const std::vector args) : - Command(cmdStream, args), - _stream (cmdStream->currentStream()) { - - hipModule_t module; - HIPCHECK(hipModuleLoad(&module, FILENAME)); - HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME)); - }; - ~ModuleKernelCommand() {}; +class ModuleKernelCommand : public Command { + public: + ModuleKernelCommand(CommandStream* cmdStream, const std::vector args) + : Command(cmdStream, args), _stream(cmdStream->currentStream()) { + hipModule_t module; + HIPCHECK(hipModuleLoad(&module, FILENAME)); + HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME)); + }; + ~ModuleKernelCommand(){}; void run() override { #define LEN 64 - uint32_t len = LEN; - uint32_t one = 1; + uint32_t len = LEN; + uint32_t one = 1; - float *Ad = NULL; + float* Ad = NULL; size_t argSize = 36; char argBuffer[argSize]; - *(uint32_t*) (&argBuffer[0]) = len; - *(uint32_t*) (&argBuffer[4]) = one; - *(uint32_t*) (&argBuffer[8]) = one; - *(uint32_t*) (&argBuffer[12]) = len; - *(uint32_t*) (&argBuffer[16]) = one; - *(uint32_t*) (&argBuffer[20]) = one; - *(float**) (&argBuffer[24]) = Ad; // Ad pointer argument + *(uint32_t*)(&argBuffer[0]) = len; + *(uint32_t*)(&argBuffer[4]) = one; + *(uint32_t*)(&argBuffer[8]) = one; + *(uint32_t*)(&argBuffer[12]) = len; + *(uint32_t*)(&argBuffer[16]) = one; + *(uint32_t*)(&argBuffer[20]) = one; + *(float**)(&argBuffer[24]) = Ad; // Ad pointer argument - void *config[] = { - HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], - HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, - HIP_LAUNCH_PARAM_END - }; + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], + HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, HIP_LAUNCH_PARAM_END}; - hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); + hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); }; -public: + public: hipFunction_t _function; hipStream_t _stream; }; -class KernelCommand : public Command -{ -public: - enum Type {Null, VectorAdd}; - KernelCommand(CommandStream *cmdStream, const std::vector args, Type kind) : - Command(cmdStream, args), - _kind(kind), - _stream (cmdStream->currentStream()) { - }; - ~KernelCommand() {}; +class KernelCommand : public Command { + public: + enum Type { Null, VectorAdd }; + KernelCommand(CommandStream* cmdStream, const std::vector args, Type kind) + : Command(cmdStream, args), _kind(kind), _stream(cmdStream->currentStream()){}; + ~KernelCommand(){}; void run() override { - static const int gridX = 64; + static const int gridX = 64; static const int groupX = 64; switch (_kind) { case Null: - hipLaunchKernel(NullKernel, dim3(gridX/groupX), dim3(gridX), 0, _stream, nullptr); + hipLaunchKernel(NullKernel, dim3(gridX / groupX), dim3(gridX), 0, _stream, nullptr); break; case VectorAdd: - assert(0); // TODO + assert(0); // TODO break; }; } -private: - Type _kind; + + private: + Type _kind; hipStream_t _stream; }; #ifdef __HIP_PLATFORM_HCC__ //================================================================================================= -class PfeCommand : public Command -{ -public: - PfeCommand(CommandStream *cmdStream, const std::vector args, hipStream_t stream = 0) : - Command(cmdStream, args) - { +class PfeCommand : public Command { + public: + PfeCommand(CommandStream* cmdStream, const std::vector args, + hipStream_t stream = 0) + : Command(cmdStream, args) { HIPCHECK(hipHccGetAcceleratorView(stream, &_av)); } - ~PfeCommand() { } + ~PfeCommand() {} void run() override { - static const int gridX = 64; + static const int gridX = 64; static const int groupX = 64; auto cf = hc::parallel_for_each(*_av, hc::extent<1>(gridX).tile(groupX), - [=](hc::index<1>& idx) __HC__ { - }); + [=](hc::index<1>& idx) __HC__ {}); } -private: - hc::accelerator_view *_av; + + private: + hc::accelerator_view* _av; }; #endif - //================================================================================================= -class CopyCommand : public Command -{ -enum MemType {PinnedHost, UnpinnedHost, Device} ; +class CopyCommand : public Command { + enum MemType { PinnedHost, UnpinnedHost, Device }; -public: - CopyCommand(CommandStream *cmdStream, const std::vector &args, hipMemcpyKind kind, bool isAsync, bool isPinnedHost) ; + public: + CopyCommand(CommandStream* cmdStream, const std::vector& args, hipMemcpyKind kind, + bool isAsync, bool isPinnedHost); - ~CopyCommand() - { + ~CopyCommand() { if (_dst) { dealloc(_dst, _dstType); _dst = NULL; @@ -541,9 +503,9 @@ public: } }; -private: - void * alloc(size_t size, MemType memType) { - void * p; + private: + void* alloc(size_t size, MemType memType) { + void* p; if (memType == Device) { HIPCHECK(hipMalloc(&p, size)); @@ -562,7 +524,7 @@ private: }; - void dealloc(void *p, MemType memType) { + void dealloc(void* p, MemType memType) { if (memType == Device) { HIPCHECK(hipFree(p)); } else if (memType == PinnedHost) { @@ -575,53 +537,42 @@ private: } -private: - bool _isAsync; - hipStream_t _stream; + private: + bool _isAsync; + hipStream_t _stream; hipMemcpyKind _kind; - size_t _sizeBytes; - void *_dst; + size_t _sizeBytes; + void* _dst; MemType _dstType; - void *_src; + void* _src; MemType _srcType; }; //================================================================================================= -class DeviceSyncCommand : public Command -{ -public: - DeviceSyncCommand(CommandStream *cmdStream, const std::vector &args) : - Command(cmdStream, args) {}; +class DeviceSyncCommand : public Command { + public: + DeviceSyncCommand(CommandStream* cmdStream, const std::vector& args) + : Command(cmdStream, args){}; - void run() override { - HIPCHECK(hipDeviceSynchronize()); - }; + void run() override { HIPCHECK(hipDeviceSynchronize()); }; }; //================================================================================================= -class StreamSyncCommand : public Command -{ -public: - StreamSyncCommand(CommandStream *cmdStream, const std::vector &args) : - Command(cmdStream, args), - _stream(cmdStream->currentStream()) - {}; +class StreamSyncCommand : public Command { + public: + StreamSyncCommand(CommandStream* cmdStream, const std::vector& args) + : Command(cmdStream, args), _stream(cmdStream->currentStream()){}; - const char *help() { - return "synchronizes the current stream"; - }; + const char* help() { return "synchronizes the current stream"; }; + void run() override { HIPCHECK(hipStreamSynchronize(_stream)); }; - void run() override { - HIPCHECK(hipStreamSynchronize(_stream)); - }; - -private: + private: hipStream_t _stream; }; @@ -629,53 +580,45 @@ private: //================================================================================================= //================================================================================================= -class LoopCommand : public Command -{ -public: - LoopCommand(CommandStream *parentCmdStream, const std::vector &args) : - Command(parentCmdStream, args, 1) - { +class LoopCommand : public Command { + public: + LoopCommand(CommandStream* parentCmdStream, const std::vector& args) + : Command(parentCmdStream, args, 1) { int loopCnt; try { loopCnt = std::stoi(args[1]); } catch (std::invalid_argument) { - failed ("bad LOOP_CNT=%s", args[1].c_str()); + failed("bad LOOP_CNT=%s", args[1].c_str()); } _commandStream = new CommandStream("", loopCnt); _commandStream->setParent(parentCmdStream); parentCmdStream->enterSubBlock(_commandStream); - }; - void print(const std::string &indent = "") const override { + void print(const std::string& indent = "") const override { Command::print(); - _commandStream->print (indent + " "); + _commandStream->print(indent + " "); }; - void run() override { - _commandStream->run(); - }; + void run() override { _commandStream->run(); }; }; //================================================================================================= -class EndBlockCommand : public Command -{ -public: - EndBlockCommand(CommandStream *blockCmdStream, CommandStream *parentCmdStream, const std::vector &args) : - Command(parentCmdStream, args, 0, 1), - _blockCmdStream(blockCmdStream), - _printTiming(0) - { - int argCnt = args.size()-1; - if (argCnt >= 1 ) { +class EndBlockCommand : public Command { + public: + EndBlockCommand(CommandStream* blockCmdStream, CommandStream* parentCmdStream, + const std::vector& args) + : Command(parentCmdStream, args, 0, 1), _blockCmdStream(blockCmdStream), _printTiming(0) { + int argCnt = args.size() - 1; + if (argCnt >= 1) { _printTiming = readIntArg(1, "PRINT_TIMING"); } if (parentCmdStream == nullptr) { - failed ("%s without corresponding command to start block", args[0].c_str()); + failed("%s without corresponding command to start block", args[0].c_str()); } parentCmdStream->exitSubBlock(); }; @@ -684,11 +627,10 @@ public: if (_printTiming) { _blockCmdStream->printTiming(); } - }; -private: - CommandStream *_blockCmdStream; + private: + CommandStream* _blockCmdStream; // print the stream when loop exits. int _printTiming; @@ -696,69 +638,60 @@ private: //================================================================================================= -class SetStreamCommand : public Command -{ -public: - SetStreamCommand(CommandStream *cmdStream, const std::vector &args) : - Command(cmdStream, args, 1) - { +class SetStreamCommand : public Command { + public: + SetStreamCommand(CommandStream* cmdStream, const std::vector& args) + : Command(cmdStream, args, 1) { int streamIndex = readIntArg(1, "STREAM_INDEX"); cmdStream->setStream(streamIndex); - }; - void run() override { - }; + void run() override{}; }; //================================================================================================= -class PrintTimingCommand : public Command -{ -public: - PrintTimingCommand(CommandStream *cmdStream, const std::vector &args) - : Command(cmdStream, args, 1) - { +class PrintTimingCommand : public Command { + public: + PrintTimingCommand(CommandStream* cmdStream, const std::vector& args) + : Command(cmdStream, args, 1) { _iterations = readIntArg(1, "ITERATIONS"); }; - void run() override { - _commandStream->printTiming(_iterations); - }; + void run() override { _commandStream->printTiming(_iterations); }; -private: - int _iterations; + private: + int _iterations; }; //================================================================================================= -CopyCommand::CopyCommand(CommandStream *cmdStream, const std::vector &args, - hipMemcpyKind kind, bool isAsync, bool isPinnedHost) : - Command(cmdStream, args) , - _isAsync(isAsync), - _kind(kind), - _stream(cmdStream->currentStream()) - { - switch (kind) { - case hipMemcpyDeviceToHost: - _srcType = Device; - _dstType = isPinnedHost ? PinnedHost : UnpinnedHost; - break; - case hipMemcpyHostToDevice: - _srcType = isPinnedHost ? PinnedHost : UnpinnedHost; - _dstType = Device; - break; - default: - HIPASSERT(0, "Unknown hipMemcpyKind"); - }; - - _sizeBytes = 64; //TODO, support reading from arg. - - _dst = alloc(_sizeBytes, _dstType); - _src = alloc(_sizeBytes, _srcType); +CopyCommand::CopyCommand(CommandStream* cmdStream, const std::vector& args, + hipMemcpyKind kind, bool isAsync, bool isPinnedHost) + : Command(cmdStream, args), + _isAsync(isAsync), + _kind(kind), + _stream(cmdStream->currentStream()) { + switch (kind) { + case hipMemcpyDeviceToHost: + _srcType = Device; + _dstType = isPinnedHost ? PinnedHost : UnpinnedHost; + break; + case hipMemcpyHostToDevice: + _srcType = isPinnedHost ? PinnedHost : UnpinnedHost; + _dstType = Device; + break; + default: + HIPASSERT(0, "Unknown hipMemcpyKind"); }; + _sizeBytes = 64; // TODO, support reading from arg. + + _dst = alloc(_sizeBytes, _dstType); + _src = alloc(_sizeBytes, _srcType); +}; + //================================================================================================= //================================================================================================= @@ -771,41 +704,31 @@ CommandStream::CommandStream(std::string commandStreamString, int iterations) _startTime(0), _elapsedUs(0.0), _parentCommandStream(nullptr), - _parseInSubBlock(false) -{ + _parseInSubBlock(false) { std::vector tokens; tokenize(commandStreamString, ';', tokens); - std::for_each(tokens.begin(), tokens.end(), [&] (const std::string s) { - this->parse(s); - }); + std::for_each(tokens.begin(), tokens.end(), [&](const std::string s) { this->parse(s); }); setStream(0); } -CommandStream::~CommandStream() -{ - std::for_each(_state._streams.begin(), _state._streams.end(), [&] (hipStream_t s) { +CommandStream::~CommandStream() { + std::for_each(_state._streams.begin(), _state._streams.end(), [&](hipStream_t s) { if (s) { HIPCHECK(hipStreamDestroy(s)); } }); - std::for_each(_commands.begin(), _commands.end(), [&] (Command *c) { - delete c; - }); - - + std::for_each(_commands.begin(), _commands.end(), [&](Command* c) { delete c; }); } -void CommandStream::setStream(int streamIndex) -{ - +void CommandStream::setStream(int streamIndex) { if (streamIndex >= _state._streams.size()) { - _state._streams.resize(streamIndex+1); + _state._streams.resize(streamIndex + 1); } if (streamIndex && (_state._streams[streamIndex] == nullptr)) { @@ -816,51 +739,45 @@ void CommandStream::setStream(int streamIndex) _state._currentStream = stream; } else { // Use existing stream: - + _state._currentStream = _state._streams[streamIndex]; } - } -void CommandStream::tokenize(const std::string &s, char delim, std::vector &tokens) -{ +void CommandStream::tokenize(const std::string& s, char delim, std::vector& tokens) { std::stringstream ss; ss.str(s); std::string item; while (getline(ss, item, delim)) { - item.erase (std::remove (item.begin(), item.end(), ' '), item.end()); // remove whitespace. + item.erase(std::remove(item.begin(), item.end(), ' '), item.end()); // remove whitespace. tokens.push_back(item); } } -void trim(std::string *s) -{ +void trim(std::string* s) { // trim whitespace from begin and end: - const char *t = "\t\n\r\f\v"; + const char* t = "\t\n\r\f\v"; s->erase(0, s->find_first_not_of(t)); - s->erase(s->find_last_not_of(t)+1); + s->erase(s->find_last_not_of(t) + 1); } -void ltrim(std::string *s) -{ +void ltrim(std::string* s) { // trim whitespace from begin and end: - const char *t = "\t\n\r\f\v"; + const char* t = "\t\n\r\f\v"; s->erase(0, s->find_first_not_of(t)); } -void CommandStream::parse(std::string fullCmd) -{ - //convert to lower-case: +void CommandStream::parse(std::string fullCmd) { + // convert to lower-case: std::transform(fullCmd.begin(), fullCmd.end(), fullCmd.begin(), ::tolower); trim(&fullCmd); if (p_db) { - printf ("parse: <%s>\n", fullCmd.c_str()); + printf("parse: <%s>\n", fullCmd.c_str()); } - std::string c; std::vector args; size_t leftParenZ = fullCmd.find_first_of('('); @@ -871,40 +788,38 @@ void CommandStream::parse(std::string fullCmd) c = fullCmd.substr(0, leftParenZ); args.push_back(c); size_t rightParenZ = fullCmd.find_first_of(')', leftParenZ); - std::string argStr = fullCmd.substr(leftParenZ+1, rightParenZ-leftParenZ-1); - //printf ("c=%s argstr='%s' leftParenZ=%zu rightParenZ=%zu\n", c.c_str(), argStr.c_str(), leftParenZ, rightParenZ); + std::string argStr = fullCmd.substr(leftParenZ + 1, rightParenZ - leftParenZ - 1); + // printf ("c=%s argstr='%s' leftParenZ=%zu rightParenZ=%zu\n", c.c_str(), argStr.c_str(), + // leftParenZ, rightParenZ); tokenize(argStr, ',', args); - } - - - - if ((args.size()==0) || (fullCmd.c_str()[0] == '#') ) { + if ((args.size() == 0) || (fullCmd.c_str()[0] == '#')) { if (p_db) { - printf (" skip comment\n"); + printf(" skip comment\n"); } return; - } + } + Command* cmd = NULL; + CommandStream* cmdStream = currentCommandStream(); - Command *cmd = NULL; - CommandStream *cmdStream = currentCommandStream(); - if (c == "h2d") { - cmd = new CopyCommand(cmdStream, args, hipMemcpyHostToDevice, true/*isAsync*/, true/*isPinned*/); + cmd = new CopyCommand(cmdStream, args, hipMemcpyHostToDevice, true /*isAsync*/, + true /*isPinned*/); //= h2d //= Performs an async host-to-device copy of array A_h to A_d. //= The size of these arrays may be set with the datasize command. } else if (c == "d2h") { - cmd = new CopyCommand(cmdStream, args, hipMemcpyDeviceToHost, true/*isAsync*/, true/*isPinned*/); + cmd = new CopyCommand(cmdStream, args, hipMemcpyDeviceToHost, true /*isAsync*/, + true /*isPinned*/); //= d2h //= Performs an async device-to-host copy of array A_d to A_h. //= The size of these arrays may be set with the datasize command. - + } else if (c == "modulekernel") { cmd = new ModuleKernelCommand(cmdStream, args); @@ -929,20 +844,21 @@ void CommandStream::parse(std::string fullCmd) } else if (c == "streamsync") { //= streamsync - //= Execute hipStreamSynchronize. + //= Execute hipStreamSynchronize. //= This will cause the host thread to wait until the current stream //= completes all pending operations. cmd = new StreamSyncCommand(cmdStream, args); } else if (c == "setstream") { //= setstream(STREAM_INDEX); - //= Set current stream used by subsequent commands. - //= STREAM_INDEX is index starting from 0...N. - //= This function will create new stream on first call to setstream or re-use previous + //= Set current stream used by subsequent commands. + //= STREAM_INDEX is index starting from 0...N. + //= This function will create new stream on first call to setstream or re-use previous //= stream if setstream has already been called with STREAM_INDEX. - //= STREAM_INDEX=0 will use the default "null" stream associated with the device, and will not create a new stream. - //= The default stream has special, conservative synchronization properties. - + //= STREAM_INDEX=0 will use the default "null" stream associated with the device, and will + //not create a new stream. = The default stream has special, conservative synchronization + //properties. + cmd = new SetStreamCommand(cmdStream, args); } else if (c == "printtiming") { @@ -952,15 +868,15 @@ void CommandStream::parse(std::string fullCmd) //= loop(LOOP_CNT) //= Loop over next set of commands (until 'endloop' command) for LOOP_CNT iterations. //= Loops can be nested. - + cmd = new LoopCommand(cmdStream, args); } else if (c == "endloop") { //= endloop - //= End a looped sequence. Must be paired with a preceding loop command. - //= Command between the `loop` and `endloop` must be executed - - CommandStream * parentCmdStream = cmdStream->getParent() ; + //= End a looped sequence. Must be paired with a preceding loop command. + //= Command between the `loop` and `endloop` must be executed + + CommandStream* parentCmdStream = cmdStream->getParent(); cmd = new EndBlockCommand(cmdStream, parentCmdStream, args); cmdStream = parentCmdStream; @@ -975,28 +891,23 @@ void CommandStream::parse(std::string fullCmd) } - - -void CommandStream::print(const std::string &indent) const -{ +void CommandStream::print(const std::string& indent) const { for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) { (*cmdI)->print(indent); }; } -void CommandStream::printBrief(std::ostream &s) const -{ +void CommandStream::printBrief(std::ostream& s) const { for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) { (*cmdI)->printBrief(s); s << ";"; }; } -void CommandStream::run() -{ +void CommandStream::run() { _startTime = get_time(); - for (int i=0; i<_iterations; i++) { + for (int i = 0; i < _iterations; i++) { for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) { if (p_verbose) { (*cmdI)->print(); @@ -1009,8 +920,7 @@ void CommandStream::run() recordTime(); }; -void CommandStream::recordTime() -{ +void CommandStream::recordTime() { if (_elapsedUs == 0.0) { auto stopTime = get_time(); _elapsedUs = stopTime - _startTime; @@ -1018,11 +928,9 @@ void CommandStream::recordTime() } -void CommandStream::printTiming(int iterations) -{ - - if ((_state._subBlocks.size() == 1) && (_commands.size()==1)) { - //printf ("print just the loop\n"); +void CommandStream::printTiming(int iterations) { + if ((_state._subBlocks.size() == 1) && (_commands.size() == 1)) { + // printf ("print just the loop\n"); _state._subBlocks.front()->printTiming(iterations); } else { g_printedTiming = true; @@ -1031,33 +939,31 @@ void CommandStream::printTiming(int iterations) if (iterations == 0) { iterations = _iterations; } - std::cout << "command<"; printBrief(std::cout); - std::cout << ">," ; - printf (" iterations,%d, total_time,%6.3f, time/iteration,%6.3f\n", iterations, _elapsedUs, _elapsedUs/iterations); + std::cout << "command<"; + printBrief(std::cout); + std::cout << ">,"; + printf(" iterations,%d, total_time,%6.3f, time/iteration,%6.3f\n", iterations, + _elapsedUs, _elapsedUs / iterations); } }; - - - //================================================================================================= -int main(int argc, char *argv[]) -{ +int main(int argc, char* argv[]) { parseStandardArguments(argc, argv); printConfig(); - CommandStream *cs; + CommandStream* cs; if (p_blockingSync) { #ifdef __HIP_PLATFORM_HCC__ - printf ("setting BlockingSync for AMD\n"); + printf("setting BlockingSync for AMD\n"); setenv("HIP_BLOCKING_SYNC", "1", 1); #endif #ifdef __HIP_PLATFORM_NVCC__ - printf ("setting cudaDeviceBlockingSync\n"); + printf("setting cudaDeviceBlockingSync\n"); HIPCHECK(hipSetDeviceFlags(cudaDeviceBlockingSync)); #endif }; @@ -1068,10 +974,9 @@ int main(int argc, char *argv[]) std::ifstream file(p_file); std::string str; std::string file_contents; - while (std::getline(file, str)) - { + while (std::getline(file, str)) { file_contents += str; - } + } cs = new CommandStream(file_contents, p_iterations); @@ -1080,7 +985,7 @@ int main(int argc, char *argv[]) } cs->print(); - printf ("------\n"); + printf("------\n"); cs->run(); if (!g_printedTiming) { @@ -1091,5 +996,4 @@ int main(int argc, char *argv[]) } - // TODO - add error checking for arguments. diff --git a/projects/hip-tests/samples/1_Utils/hipCommander/nullkernel.hip.cpp b/projects/hip-tests/samples/1_Utils/hipCommander/nullkernel.hip.cpp index 890e9bdc1e..410796c7e6 100644 --- a/projects/hip-tests/samples/1_Utils/hipCommander/nullkernel.hip.cpp +++ b/projects/hip-tests/samples/1_Utils/hipCommander/nullkernel.hip.cpp @@ -1,6 +1,6 @@ #include "hip/hip_runtime.h" -extern "C" __global__ void NullKernel(hipLaunchParm lp, float* Ad){ +extern "C" __global__ void NullKernel(hipLaunchParm lp, float* Ad) { if (Ad) { Ad[0] = 42; } diff --git a/projects/hip-tests/samples/1_Utils/hipCommander/testcase.cpp b/projects/hip-tests/samples/1_Utils/hipCommander/testcase.cpp index 0eeae7c50f..93ebcf40c1 100644 --- a/projects/hip-tests/samples/1_Utils/hipCommander/testcase.cpp +++ b/projects/hip-tests/samples/1_Utils/hipCommander/testcase.cpp @@ -1,20 +1,17 @@ #include -static const int BLOCKSIZEX=32; -static const int BLOCKSIZEY=16; +static const int BLOCKSIZEX = 32; +static const int BLOCKSIZEY = 16; -__global__ void fails(hipLaunchParm lp, float* pErrorI) -{ - if(pErrorI!=0) - { - pErrorI[0]=1; +__global__ void fails(hipLaunchParm lp, float* pErrorI) { + if (pErrorI != 0) { + pErrorI[0] = 1; } } -int main() -{ - dim3 blocks(1,1); - dim3 threads(BLOCKSIZEX,BLOCKSIZEY); +int main() { + dim3 blocks(1, 1); + dim3 threads(BLOCKSIZEX, BLOCKSIZEY); float error; hipLaunchKernel(HIP_KERNEL_NAME(fails), blocks, threads, 0, 0, &error); diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.cpp b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.cpp index f6f2fab709..b769ca4b32 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.cpp +++ b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.cpp @@ -11,96 +11,72 @@ using namespace std; #define SORT_RETAIN_ATTS_ORDER 1 -bool ResultDatabase::Result::operator<(const Result &rhs) const -{ - if (test < rhs.test) - return true; - if (test > rhs.test) - return false; -#if (SORT_RETAIN_ATTS_ORDER == 0) +bool ResultDatabase::Result::operator<(const Result& rhs) const { + if (test < rhs.test) return true; + if (test > rhs.test) return false; +#if (SORT_RETAIN_ATTS_ORDER == 0) // For ties, sort by the value of the attribute: - if (atts < rhs.atts) - return true; - if (atts > rhs.atts) - return false; + if (atts < rhs.atts) return true; + if (atts > rhs.atts) return false; #endif - return false; // less-operator returns false on equal + return false; // less-operator returns false on equal } -double ResultDatabase::Result::GetMin() const -{ +double ResultDatabase::Result::GetMin() const { double r = FLT_MAX; - for (int i=0; i= 100) - return value[n-1]; + if (q <= 0) return value[0]; + if (q >= 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.)); + 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]; + if (frac == 0) return sorted[index_lo]; double lo = sorted[index_lo]; double hi = sorted[index_lo + 1]; - return lo + (hi-lo)*frac; + return lo + (hi - lo) * frac; } -double ResultDatabase::Result::GetMean() const -{ +double ResultDatabase::Result::GetMean() const { double r = 0; - for (int i=0; i &values) -{ - for (int i=0; i& values) { + for (int i = 0; i < values.size(); i++) { AddResult(test, atts, unit, values[i]); } } -static string RemoveAllButLeadingSpaces(const string &a) -{ +static string RemoveAllButLeadingSpaces(const string& a) { string b; int n = a.length(); int i = 0; - while (i= results.size()) - { + if (index >= results.size()) { Result r; r.test = test; r.atts = atts; @@ -193,43 +153,35 @@ void ResultDatabase::AddResult(const string &test_orig, // Changed note about missing values to be worded a little better. // // **************************************************************************** -void ResultDatabase::DumpDetailed(ostream &out) -{ +void ResultDatabase::DumpDetailed(ostream& out) { vector sorted(results); #if SORT_BY_NAME stable_sort(sorted.begin(), sorted.end()); #endif - const int testNameW = 24 ; + const int testNameW = 24; const int attW = 12; const int fieldW = 11; out << std::fixed << right << std::setprecision(4); int maxtrials = 1; - for (int i=0; i maxtrials) - maxtrials = sorted[i].value.size(); + for (int i = 0; i < sorted.size(); i++) { + if (sorted[i].value.size() > 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 << setw(testNameW) << "test\t" - << setw(attW) << "atts\t" - << setw(fieldW) - << "median\t" + out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "median\t" << "mean\t" << "stddev\t" << "min\t" << "max\t"; - for (int i=0; i sorted(results); #if SORT_BY_NAME stable_sort(sorted.begin(), sorted.end()); #endif - const int testNameW = 32 ; + const int testNameW = 32; const int attW = 12; const int fieldW = 9; out << std::fixed << right << std::setprecision(2); // TODO: in big parallel runs, the "trials" are the procs // and we really don't want to print them all out.... - out << setw(testNameW) << "test\t" - << setw(attW) << "atts\t" - << setw(fieldW) - << "units\t" + out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "units\t" << "median\t" << "mean\t" << "stddev\t" @@ -316,9 +263,8 @@ void ResultDatabase::DumpSummary(ostream &out) << "max\t"; out << endl; - for (int i=0; i sorted(results); @@ -398,32 +340,30 @@ void ResultDatabase::DumpCsv(string fileName) stable_sort(sorted.begin(), sorted.end()); #endif - //Check to see if the file is empty - if so, add the headers + // Check to see if the file is empty - if so, add the headers emptyFile = this->IsFileEmpty(fileName); - //Open file and append by default + // Open file and append by default ofstream out; - out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app); + 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; + // 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) -{ +vector ResultDatabase::GetResultsForTest(const string& test) { // get only the given test results vector retval; - for (int i=0; i & -ResultDatabase::GetResults() const -{ - return results; -} +const vector& ResultDatabase::GetResults() const { return results; } diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.h b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.h index 4b63a02a1f..ca6a00fc91 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.h +++ b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.h @@ -6,11 +6,11 @@ #include #include #include +using std::ifstream; +using std::ofstream; +using std::ostream; using std::string; using std::vector; -using std::ostream; -using std::ofstream; -using std::ifstream; // **************************************************************************** @@ -40,18 +40,16 @@ using std::ifstream; // Added a GetResults method as well, and made several functions const. // // **************************************************************************** -class ResultDatabase -{ - public: +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" + 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; @@ -59,41 +57,32 @@ class ResultDatabase double GetMean() const; double GetStdDev() const; - bool operator<(const Result &rhs) const; + bool operator<(const Result& rhs) const; - bool HadAnyFLTMAXValues() const - { - for (int i=0; i= FLT_MAX) - return true; + bool HadAnyFLTMAXValues() const { + for (int i = 0; i < value.size(); ++i) { + if (value[i] >= FLT_MAX) return true; } return false; } }; - protected: + 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; + 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: + private: bool IsFileEmpty(string fileName); - }; diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index 2a4f6ff649..7aa8fa4992 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -21,35 +21,34 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#include -#include -#include"ResultDatabase.h" +#include +#include +#include "ResultDatabase.h" #define PRINT_PROGRESS 0 -#define check(cmd) \ -{\ - hipError_t status = cmd;\ - if(status != hipSuccess){ \ - printf("error: '%s'(%d) from %s at %s:%d\n", \ - hipGetErrorString(status), status, #cmd,\ - __FILE__, __LINE__); \ - abort(); \ - }\ -} +#define check(cmd) \ + { \ + hipError_t status = cmd; \ + if (status != hipSuccess) { \ + printf("error: '%s'(%d) from %s at %s:%d\n", hipGetErrorString(status), status, #cmd, \ + __FILE__, __LINE__); \ + abort(); \ + } \ + } -#define LEN 1024*1024 +#define LEN 1024 * 1024 #define NUM_GROUPS 1 #define GROUP_SIZE 64 -#define TEST_ITERS 20 +#define TEST_ITERS 20 #define DISPATCHES_PER_TEST 100 const unsigned p_tests = 0xfffffff; // HCC optimizes away fully NULL kernel calls, so run one that is nearly null: -__global__ void NearlyNull(hipLaunchParm lp, float* Ad){ +__global__ void NearlyNull(hipLaunchParm lp, float* Ad) { if (Ad) { Ad[0] = 42; } @@ -59,38 +58,35 @@ __global__ void NearlyNull(hipLaunchParm lp, float* Ad){ ResultDatabase resultDB; -void stopTest(hipEvent_t start, hipEvent_t stop, const char *msg, int iters) -{ - float mS = 0; +void stopTest(hipEvent_t start, hipEvent_t stop, const char* msg, int iters) { + float mS = 0; check(hipEventRecord(stop)); check(hipDeviceSynchronize()); check(hipEventElapsedTime(&mS, start, stop)); - resultDB.AddResult(std::string(msg), "", "uS", mS*1000/iters); - if (PRINT_PROGRESS & 0x1 ) { - std::cout<< msg <<"\t\t"< #include "hip/hip_runtime.h" -#define KNRM "\x1B[0m" -#define KRED "\x1B[31m" -#define KGRN "\x1B[32m" -#define KYEL "\x1B[33m" -#define KBLU "\x1B[34m" -#define KMAG "\x1B[35m" -#define KCYN "\x1B[36m" -#define KWHT "\x1B[37m" +#define KNRM "\x1B[0m" +#define KRED "\x1B[31m" +#define KGRN "\x1B[32m" +#define KYEL "\x1B[33m" +#define KBLU "\x1B[34m" +#define KMAG "\x1B[35m" +#define KCYN "\x1B[36m" +#define KWHT "\x1B[37m" -#define failed(...) \ - printf ("%serror: ", KRED);\ - printf (__VA_ARGS__);\ - printf ("\n");\ - printf ("error: TEST FAILED\n%s", KNRM );\ +#define failed(...) \ + printf("%serror: ", KRED); \ + printf(__VA_ARGS__); \ + printf("\n"); \ + printf("error: TEST FAILED\n%s", KNRM); \ exit(EXIT_FAILURE); -#define HIPCHECK(error) \ - if (error != hipSuccess) { \ - printf("%serror: '%s'(%d) at %s:%d%s\n", \ - KRED, hipGetErrorString(error), error,\ - __FILE__, __LINE__,KNRM);\ - failed("API returned error code.");\ +#define HIPCHECK(error) \ + if (error != hipSuccess) { \ + printf("%serror: '%s'(%d) at %s:%d%s\n", KRED, hipGetErrorString(error), error, __FILE__, \ + __LINE__, KNRM); \ + failed("API returned error code."); \ } -void printCompilerInfo () -{ +void printCompilerInfo() { #ifdef __HCC__ - printf ("compiler: hcc version=%s, workweek (YYWWD) = %u\n", __hcc_version__, __hcc_workweek__); + printf("compiler: hcc version=%s, workweek (YYWWD) = %u\n", __hcc_version__, __hcc_workweek__); #endif #ifdef __NVCC__ - printf ("compiler: nvcc\n"); + printf("compiler: nvcc\n"); #endif } -double bytesToGB(size_t s) -{ - return (double)s / (1024.0*1024.0*1024.0); -} +double bytesToGB(size_t s) { return (double)s / (1024.0 * 1024.0 * 1024.0); } -#define printLimit(w1, limit, units) \ -{\ - size_t val;\ - cudaDeviceGetLimit(&val, limit);\ - std::cout << setw(w1) << #limit": " << val << " " << units << std::endl;\ -} +#define printLimit(w1, limit, units) \ + { \ + size_t val; \ + cudaDeviceGetLimit(&val, limit); \ + std::cout << setw(w1) << #limit ": " << val << " " << units << std::endl; \ + } -void printDeviceProp (int deviceId) -{ +void printDeviceProp(int deviceId) { using namespace std; const int w1 = 34; cout << left; - cout << setw(w1) << "--------------------------------------------------------------------------------" << endl; + cout << setw(w1) + << "--------------------------------------------------------------------------------" + << endl; cout << setw(w1) << "device#" << deviceId << endl; hipDeviceProp_t props; @@ -88,16 +84,22 @@ void printDeviceProp (int deviceId) cout << setw(w1) << "pciBusID: " << props.pciBusID << endl; cout << setw(w1) << "pciDeviceID: " << props.pciDeviceID << endl; cout << setw(w1) << "multiProcessorCount: " << props.multiProcessorCount << endl; - cout << setw(w1) << "maxThreadsPerMultiProcessor: " << props.maxThreadsPerMultiProcessor << endl; + cout << setw(w1) << "maxThreadsPerMultiProcessor: " << props.maxThreadsPerMultiProcessor + << endl; cout << setw(w1) << "isMultiGpuBoard: " << props.isMultiGpuBoard << endl; cout << setw(w1) << "clockRate: " << (float)props.clockRate / 1000.0 << " Mhz" << endl; - cout << setw(w1) << "memoryClockRate: " << (float)props.memoryClockRate / 1000.0 << " Mhz" << endl; + cout << setw(w1) << "memoryClockRate: " << (float)props.memoryClockRate / 1000.0 << " Mhz" + << endl; cout << setw(w1) << "memoryBusWidth: " << props.memoryBusWidth << endl; - cout << setw(w1) << "clockInstructionRate: " << (float)props.clockInstructionRate / 1000.0 << " Mhz" << endl; - cout << setw(w1) << "totalGlobalMem: " << fixed << setprecision(2) << bytesToGB(props.totalGlobalMem) << " GB" << endl; - cout << setw(w1) << "maxSharedMemoryPerMultiProcessor: " << fixed << setprecision(2) << bytesToGB(props.maxSharedMemoryPerMultiProcessor) << " GB" << endl; + cout << setw(w1) << "clockInstructionRate: " << (float)props.clockInstructionRate / 1000.0 + << " Mhz" << endl; + cout << setw(w1) << "totalGlobalMem: " << fixed << setprecision(2) + << bytesToGB(props.totalGlobalMem) << " GB" << endl; + cout << setw(w1) << "maxSharedMemoryPerMultiProcessor: " << fixed << setprecision(2) + << bytesToGB(props.maxSharedMemoryPerMultiProcessor) << " GB" << endl; cout << setw(w1) << "totalConstMem: " << props.totalConstMem << endl; - cout << setw(w1) << "sharedMemPerBlock: " << (float)props.sharedMemPerBlock / 1024.0 << " KB" << endl; + cout << setw(w1) << "sharedMemPerBlock: " << (float)props.sharedMemPerBlock / 1024.0 << " KB" + << endl; cout << setw(w1) << "regsPerBlock: " << props.regsPerBlock << endl; cout << setw(w1) << "warpSize: " << props.warpSize << endl; cout << setw(w1) << "l2CacheSize: " << props.l2CacheSize << endl; @@ -112,29 +114,31 @@ void printDeviceProp (int deviceId) cout << setw(w1) << "major: " << props.major << endl; cout << setw(w1) << "minor: " << props.minor << endl; cout << setw(w1) << "concurrentKernels: " << props.concurrentKernels << endl; - cout << setw(w1) << "arch.hasGlobalInt32Atomics: " << props.arch.hasGlobalInt32Atomics << endl; - cout << setw(w1) << "arch.hasGlobalFloatAtomicExch: " << props.arch.hasGlobalFloatAtomicExch << endl; - cout << setw(w1) << "arch.hasSharedInt32Atomics: " << props.arch.hasSharedInt32Atomics << endl; - cout << setw(w1) << "arch.hasSharedFloatAtomicExch: " << props.arch.hasSharedFloatAtomicExch << endl; - cout << setw(w1) << "arch.hasFloatAtomicAdd: " << props.arch.hasFloatAtomicAdd << endl; - cout << setw(w1) << "arch.hasGlobalInt64Atomics: " << props.arch.hasGlobalInt64Atomics << endl; - cout << setw(w1) << "arch.hasSharedInt64Atomics: " << props.arch.hasSharedInt64Atomics << endl; - cout << setw(w1) << "arch.hasDoubles: " << props.arch.hasDoubles << endl; - cout << setw(w1) << "arch.hasWarpVote: " << props.arch.hasWarpVote << endl; - cout << setw(w1) << "arch.hasWarpBallot: " << props.arch.hasWarpBallot << endl; - cout << setw(w1) << "arch.hasWarpShuffle: " << props.arch.hasWarpShuffle << endl; - cout << setw(w1) << "arch.hasFunnelShift: " << props.arch.hasFunnelShift << endl; - cout << setw(w1) << "arch.hasThreadFenceSystem: " << props.arch.hasThreadFenceSystem << endl; - cout << setw(w1) << "arch.hasSyncThreadsExt: " << props.arch.hasSyncThreadsExt << endl; - cout << setw(w1) << "arch.hasSurfaceFuncs: " << props.arch.hasSurfaceFuncs << endl; - cout << setw(w1) << "arch.has3dGrid: " << props.arch.has3dGrid << endl; - cout << setw(w1) << "arch.hasDynamicParallelism: " << props.arch.hasDynamicParallelism << endl; - cout << setw(w1) << "gcnArch: " << props.gcnArch << endl; + cout << setw(w1) << "arch.hasGlobalInt32Atomics: " << props.arch.hasGlobalInt32Atomics << endl; + cout << setw(w1) << "arch.hasGlobalFloatAtomicExch: " << props.arch.hasGlobalFloatAtomicExch + << endl; + cout << setw(w1) << "arch.hasSharedInt32Atomics: " << props.arch.hasSharedInt32Atomics << endl; + cout << setw(w1) << "arch.hasSharedFloatAtomicExch: " << props.arch.hasSharedFloatAtomicExch + << endl; + cout << setw(w1) << "arch.hasFloatAtomicAdd: " << props.arch.hasFloatAtomicAdd << endl; + cout << setw(w1) << "arch.hasGlobalInt64Atomics: " << props.arch.hasGlobalInt64Atomics << endl; + cout << setw(w1) << "arch.hasSharedInt64Atomics: " << props.arch.hasSharedInt64Atomics << endl; + cout << setw(w1) << "arch.hasDoubles: " << props.arch.hasDoubles << endl; + cout << setw(w1) << "arch.hasWarpVote: " << props.arch.hasWarpVote << endl; + cout << setw(w1) << "arch.hasWarpBallot: " << props.arch.hasWarpBallot << endl; + cout << setw(w1) << "arch.hasWarpShuffle: " << props.arch.hasWarpShuffle << endl; + cout << setw(w1) << "arch.hasFunnelShift: " << props.arch.hasFunnelShift << endl; + cout << setw(w1) << "arch.hasThreadFenceSystem: " << props.arch.hasThreadFenceSystem << endl; + cout << setw(w1) << "arch.hasSyncThreadsExt: " << props.arch.hasSyncThreadsExt << endl; + cout << setw(w1) << "arch.hasSurfaceFuncs: " << props.arch.hasSurfaceFuncs << endl; + cout << setw(w1) << "arch.has3dGrid: " << props.arch.has3dGrid << endl; + cout << setw(w1) << "arch.hasDynamicParallelism: " << props.arch.hasDynamicParallelism << endl; + cout << setw(w1) << "gcnArch: " << props.gcnArch << endl; int deviceCnt; hipGetDeviceCount(&deviceCnt); cout << setw(w1) << "peers: "; - for (int i=0; i +#include // hip header file #include "hip/hip_runtime.h" -#define WIDTH 1024 +#define WIDTH 1024 -#define NUM (WIDTH*WIDTH) +#define NUM (WIDTH * WIDTH) -#define THREADS_PER_BLOCK_X 4 -#define THREADS_PER_BLOCK_Y 4 -#define THREADS_PER_BLOCK_Z 1 +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 // Device (Kernel) function, it must be void // hipLaunchParm provides the execution configuration -__global__ void matrixTranspose(hipLaunchParm lp, - float *out, - float *in, - const int width) -{ +__global__ void matrixTranspose(hipLaunchParm lp, float* out, float* in, const int width) { int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; @@ -49,88 +45,79 @@ __global__ void matrixTranspose(hipLaunchParm lp, } // CPU implementation of matrix transpose -void matrixTransposeCPUReference( - float * output, - float * input, - const unsigned int width) -{ - for(unsigned int j=0; j < width; j++) - { - for(unsigned int i=0; i < width; i++) - { - output[i*width + j] = input[j*width + i]; +void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) { + for (unsigned int j = 0; j < width; j++) { + for (unsigned int i = 0; i < width; i++) { + output[i * width + j] = input[j * width + i]; } } } int main() { + float* Matrix; + float* TransposeMatrix; + float* cpuTransposeMatrix; - float* Matrix; - float* TransposeMatrix; - float* cpuTransposeMatrix; + float* gpuMatrix; + float* gpuTransposeMatrix; - float* gpuMatrix; - float* gpuTransposeMatrix; + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); - hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + std::cout << "Device name " << devProp.name << std::endl; - std::cout << "Device name " << devProp.name << std::endl; + int i; + int errors; - int i; - int errors; + Matrix = (float*)malloc(NUM * sizeof(float)); + TransposeMatrix = (float*)malloc(NUM * sizeof(float)); + cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float)); - Matrix = (float*)malloc(NUM * sizeof(float)); - TransposeMatrix = (float*)malloc(NUM * sizeof(float)); - cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float)); - - // initialize the input data - for (i = 0; i < NUM; i++) { - Matrix[i] = (float)i*10.0f; - } - - // allocate the memory on the device side - hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); - - // Memory transfer from host to device - hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice); - - // Lauching kernel from host - hipLaunchKernel(matrixTranspose, - dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y), - dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), - 0, 0, - gpuTransposeMatrix , gpuMatrix, WIDTH); - - // Memory transfer from device to host - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost); - - // CPU MatrixTranspose computation - matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); - - // verify the results - errors = 0; - double eps = 1.0E-6; - for (i = 0; i < NUM; i++) { - if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) { - errors++; + // initialize the input data + for (i = 0; i < NUM; i++) { + Matrix[i] = (float)i * 10.0f; } - } - if (errors!=0) { - printf("FAILED: %d errors\n",errors); - } else { - printf ("PASSED!\n"); - } - //free the resources on device side - hipFree(gpuMatrix); - hipFree(gpuTransposeMatrix); + // allocate the memory on the device side + hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); + hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); - //free the resources on host side - free(Matrix); - free(TransposeMatrix); - free(cpuTransposeMatrix); + // Memory transfer from host to device + hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); - return errors; + // Lauching kernel from host + hipLaunchKernel(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, + gpuMatrix, WIDTH); + + // Memory transfer from device to host + hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + + // CPU MatrixTranspose computation + matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); + + // verify the results + errors = 0; + double eps = 1.0E-6; + for (i = 0; i < NUM; i++) { + if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) { + errors++; + } + } + if (errors != 0) { + printf("FAILED: %d errors\n", errors); + } else { + printf("PASSED!\n"); + } + + // free the resources on device side + hipFree(gpuMatrix); + hipFree(gpuTransposeMatrix); + + // free the resources on host side + free(Matrix); + free(TransposeMatrix); + free(cpuTransposeMatrix); + + return errors; } diff --git a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/inline_asm.cpp b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/inline_asm.cpp index 2b4fc3de90..d9aee9c0a8 100644 --- a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/inline_asm.cpp +++ b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/inline_asm.cpp @@ -20,155 +20,141 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include +#include // hip header file #include "hip/hip_runtime.h" -#define WIDTH 1024 +#define WIDTH 1024 -#define NUM (WIDTH*WIDTH) +#define NUM (WIDTH * WIDTH) -#define THREADS_PER_BLOCK_X 4 -#define THREADS_PER_BLOCK_Y 4 -#define THREADS_PER_BLOCK_Z 1 +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 // Device (Kernel) function, it must be void // hipLaunchParm provides the execution configuration -__global__ void matrixTranspose(hipLaunchParm lp, - float *out, - float *in, - const int width) -{ - +__global__ void matrixTranspose(hipLaunchParm lp, float* out, float* in, const int width) { int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; - asm volatile ("v_mov_b32_e32 %0, %1" : "=v" (out[x*width + y]) : "v" (in[y*width + x])); + asm volatile("v_mov_b32_e32 %0, %1" : "=v"(out[x * width + y]) : "v"(in[y * width + x])); } // CPU implementation of matrix transpose -void matrixTransposeCPUReference( - float * output, - float * input, - const unsigned int width) -{ - for(unsigned int j=0; j < width; j++) - { - for(unsigned int i=0; i < width; i++) - { - output[i*width + j] = input[j*width + i]; +void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) { + for (unsigned int j = 0; j < width; j++) { + for (unsigned int i = 0; i < width; i++) { + output[i * width + j] = input[j * width + i]; } } } int main() { + float* Matrix; + float* TransposeMatrix; + float* cpuTransposeMatrix; - float* Matrix; - float* TransposeMatrix; - float* cpuTransposeMatrix; + float* gpuMatrix; + float* gpuTransposeMatrix; - float* gpuMatrix; - float* gpuTransposeMatrix; + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); - hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + std::cout << "Device name " << devProp.name << std::endl; - std::cout << "Device name " << devProp.name << std::endl; + hipEvent_t start, stop; + hipEventCreate(&start); + hipEventCreate(&stop); + float eventMs = 1.0f; - hipEvent_t start, stop; - hipEventCreate(&start); - hipEventCreate(&stop); - float eventMs = 1.0f; + int i; + int errors; - int i; - int errors; + Matrix = (float*)malloc(NUM * sizeof(float)); + TransposeMatrix = (float*)malloc(NUM * sizeof(float)); + cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float)); - Matrix = (float*)malloc(NUM * sizeof(float)); - TransposeMatrix = (float*)malloc(NUM * sizeof(float)); - cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float)); - - // initialize the input data - for (i = 0; i < NUM; i++) { - Matrix[i] = (float)i*10.0f; - } - - // allocate the memory on the device side - hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); - - // Record the start event - hipEventRecord(start, NULL); - - // Memory transfer from host to device - hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice); - - // Record the stop event - hipEventRecord(stop, NULL); - hipEventSynchronize(stop); - - hipEventElapsedTime(&eventMs, start, stop); - - printf ("hipMemcpyHostToDevice time taken = %6.3fms\n", eventMs); - - // Record the start event - hipEventRecord(start, NULL); - - // Lauching kernel from host - hipLaunchKernel(matrixTranspose, - dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y), - dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), - 0, 0, - gpuTransposeMatrix , gpuMatrix, WIDTH); - - // Record the stop event - hipEventRecord(stop, NULL); - hipEventSynchronize(stop); - - hipEventElapsedTime(&eventMs, start, stop); - - printf ("kernel Execution time = %6.3fms\n", eventMs); - - // Record the start event - hipEventRecord(start, NULL); - - // Memory transfer from device to host - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost); - - // Record the stop event - hipEventRecord(stop, NULL); - hipEventSynchronize(stop); - - hipEventElapsedTime(&eventMs, start, stop); - - printf ("hipMemcpyDeviceToHost time taken = %6.3fms\n", eventMs); - - // CPU MatrixTranspose computation - matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); - - // verify the results - errors = 0; - double eps = 1.0E-6; - for (i = 0; i < NUM; i++) { - if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) { - printf("gpu%f cpu %f \n",TransposeMatrix[i],cpuTransposeMatrix[i]); - errors++; + // initialize the input data + for (i = 0; i < NUM; i++) { + Matrix[i] = (float)i * 10.0f; } - } - if (errors!=0) { - printf("FAILED: %d errors\n",errors); - } else { - printf ("PASSED!\n"); - } - //free the resources on device side - hipFree(gpuMatrix); - hipFree(gpuTransposeMatrix); + // allocate the memory on the device side + hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); + hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); - //free the resources on host side - free(Matrix); - free(TransposeMatrix); - free(cpuTransposeMatrix); + // Record the start event + hipEventRecord(start, NULL); - return errors; + // Memory transfer from host to device + hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + + // Record the stop event + hipEventRecord(stop, NULL); + hipEventSynchronize(stop); + + hipEventElapsedTime(&eventMs, start, stop); + + printf("hipMemcpyHostToDevice time taken = %6.3fms\n", eventMs); + + // Record the start event + hipEventRecord(start, NULL); + + // Lauching kernel from host + hipLaunchKernel(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, + gpuMatrix, WIDTH); + + // Record the stop event + hipEventRecord(stop, NULL); + hipEventSynchronize(stop); + + hipEventElapsedTime(&eventMs, start, stop); + + printf("kernel Execution time = %6.3fms\n", eventMs); + + // Record the start event + hipEventRecord(start, NULL); + + // Memory transfer from device to host + hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + + // Record the stop event + hipEventRecord(stop, NULL); + hipEventSynchronize(stop); + + hipEventElapsedTime(&eventMs, start, stop); + + printf("hipMemcpyDeviceToHost time taken = %6.3fms\n", eventMs); + + // CPU MatrixTranspose computation + matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); + + // verify the results + errors = 0; + double eps = 1.0E-6; + for (i = 0; i < NUM; i++) { + if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) { + printf("gpu%f cpu %f \n", TransposeMatrix[i], cpuTransposeMatrix[i]); + errors++; + } + } + if (errors != 0) { + printf("FAILED: %d errors\n", errors); + } else { + printf("PASSED!\n"); + } + + // free the resources on device side + hipFree(gpuMatrix); + hipFree(gpuTransposeMatrix); + + // free the resources on host side + free(Matrix); + free(TransposeMatrix); + free(cpuTransposeMatrix); + + return errors; } diff --git a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp index 17ed911808..c817f4e42c 100644 --- a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp +++ b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp @@ -23,11 +23,8 @@ THE SOFTWARE. #include "hip/hip_runtime.h" extern texture tex; -__global__ void tex2dKernel(hipLaunchParm lp, float* outputData, - int width, - int height) -{ - int x = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x; - int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y; - outputData[y*width + x] = tex2D(tex, x, y); +__global__ void tex2dKernel(hipLaunchParm lp, float* outputData, int width, int height) { + int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + outputData[y * width + x] = tex2D(tex, x, y); } diff --git a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp index a19f4376c3..322bd1370e 100644 --- a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -32,111 +32,113 @@ THE SOFTWARE. texture tex; bool testResult = false; -#define HIP_CHECK(cmd) \ -{\ - hipError_t status = cmd;\ - if(status != hipSuccess) {std::cout<<"error: #"<