From f6f37fb73f9ad7e48f256bcaab21d2f01e2e3090 Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Sun, 19 Jul 2020 14:55:10 -0400 Subject: [PATCH 1/5] Support numa if libnume-dev is installed Let rocclr decide numa dependence Change-Id: I65bdfba7ec0d06b550f86632318bcfd1f765cfa9 --- perftests/memory/hipHostNumaAlloc.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/perftests/memory/hipHostNumaAlloc.cpp b/perftests/memory/hipHostNumaAlloc.cpp index 75d6edf0cf..38401c8046 100644 --- a/perftests/memory/hipHostNumaAlloc.cpp +++ b/perftests/memory/hipHostNumaAlloc.cpp @@ -34,12 +34,13 @@ THE SOFTWARE. #include #include "hip/hip_runtime.h" /* HIT_START - * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * BUILD_CMD: hipHostNumaAlloc %hc -I%S/../../src %S/%s %S/../../src/test_common.cpp -lnuma -o %T/%t EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ -// To run it correctly, we must not export HIP_VISIBLE_DEVICES +// To run it correctly, we must not export HIP_VISIBLE_DEVICES. +// And we must explicitly link libnuma because of numa api move_pages(). #define NUM_PAGES 4 char *h = nullptr; char *d_h = nullptr; @@ -127,6 +128,7 @@ bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) { printf("\n"); HIPCHECK(hipHostFree((void* )h)); + hipHostUnregister(m); free(m); if (cpuId >= 0 && (numaMode == MPOL_BIND || numaMode == MPOL_PREFERRED)) { @@ -149,8 +151,7 @@ bool runTest(const int &cpuCount, const int &gpuCount, for (int i = 0; i < cpuCount; i++) { for (int j = 0; j < gpuCount; j++) { - if (!test(i, j, mode[m], - hipHostMallocDefault | hipHostMallocNumaUser)) { + if (!test(i, j, mode[m], hostMallocflags)) { return false; } } From 9e7e04c875504ca54c66b6ad59a215356e5bb052 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Fri, 7 Aug 2020 09:57:11 -0400 Subject: [PATCH 2/5] [HipPerf] add a test for measuring shared memory read speed SWDEV-245290/SWDEV-247330 Change-Id: If78fa0879ff58aab84775e412a86665c7e5959e7 --- .../memory/hipPerfSharedMemReadSpeed.cpp | 250 ++++++++++++++++++ 1 file changed, 250 insertions(+) create mode 100644 perftests/memory/hipPerfSharedMemReadSpeed.cpp diff --git a/perftests/memory/hipPerfSharedMemReadSpeed.cpp b/perftests/memory/hipPerfSharedMemReadSpeed.cpp new file mode 100644 index 0000000000..86adc5f354 --- /dev/null +++ b/perftests/memory/hipPerfSharedMemReadSpeed.cpp @@ -0,0 +1,250 @@ +/* + Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + */ + +/* HIT_START + * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include +#include +#include "test_common.h" + +using namespace std; + +#define sharedMemSize1 2048 +#define sharedMemSize2 256 + +__global__ void sharedMemReadSpeed1(float *outBuf, ulong N) { + + size_t gid = (blockIdx.x * blockDim.x + threadIdx.x); + size_t lid = threadIdx.x; + __shared__ float local[sharedMemSize1]; + + float val1 = 0; + float val2 = 0; + float val3 = 0; + float val4 = 0; + + for (int i = 0; i < (sharedMemSize1 / 64); i++) { + local[lid + i * 64] = lid; + } + + __syncthreads(); + + val1 += local[lid]; + val2 += local[lid + 64]; + val3 += local[lid + 128]; + val4 += local[lid + 192]; + val1 += local[lid + 256]; + val2 += local[lid + 320]; + val3 += local[lid + 384]; + val4 += local[lid + 448]; + val1 += local[lid + 512]; + val2 += local[lid + 576]; + val3 += local[lid + 640]; + val4 += local[lid + 704]; + val1 += local[lid + 768]; + val2 += local[lid + 832]; + val3 += local[lid + 896]; + val4 += local[lid + 960]; + val1 += local[lid + 1024]; + val2 += local[lid + 1088]; + val3 += local[lid + 1152]; + val4 += local[lid + 1216]; + val1 += local[lid + 1280]; + val2 += local[lid + 1344]; + val3 += local[lid + 1408]; + val4 += local[lid + 1472]; + val1 += local[lid + 1536]; + val2 += local[lid + 1600]; + val3 += local[lid + 1664]; + val4 += local[lid + 1728]; + val1 += local[lid + 1792]; + val2 += local[lid + 1856]; + val3 += local[lid + 1920]; + val4 += local[lid + 1984]; + + if (gid < N) { + outBuf[gid] = val1 + val2 + val3 + val4; + } +}; + +__global__ void sharedMemReadSpeed2(float *outBuf, ulong N) { + size_t gid = (blockIdx.x * blockDim.x + threadIdx.x); + size_t lid = threadIdx.x; + __shared__ float local[sharedMemSize2]; + + float val0 = 0.0f; + float val1 = 0.0f; + + for (int i = 0; i < (sharedMemSize2 / 64); i++) { + local[lid + i * 64] = lid; + } + + __syncthreads(); + +#pragma nounroll + for (uint i = 0; i < 32; i++) { + val0 += local[8 * i + 0]; + val1 += local[8 * i + 1]; + val0 += local[8 * i + 2]; + val1 += local[8 * i + 3]; + val0 += local[8 * i + 4]; + val1 += local[8 * i + 5]; + val0 += local[8 * i + 6]; + val1 += local[8 * i + 7]; + } + + if (gid < N) { + outBuf[gid] = val0 + val1; + } +}; + +int main(int argc, char *argv[]) { + float *dDst; + float *hDst; + hipStream_t stream; + constexpr uint numSizes = 4; + constexpr uint Sizes[numSizes] = {262144, 1048576, 4194304, 16777216}; + uint numReads1 = 32; + uint numReads2 = 256; + uint sharedMemSizeBytes1 = sharedMemSize1 * sizeof(float); + uint sharedMemSizeBytes2 = sharedMemSize2 * sizeof(float); + int nIter = 1000; + const unsigned threadsPerBlock = 64; + + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + cout << "info: didn't find any GPU! skipping the test!\n"; + passed(); + return 0; + } + + static int device = 0; + HIPCHECK(hipSetDevice(device)); + hipDeviceProp_t props; + HIPCHECK(hipGetDeviceProperties(&props, device)); + cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name + << " with " << props.multiProcessorCount << " CUs" << endl; + + HIPCHECK(hipStreamCreate(&stream)); + + for (int nTest = 0; nTest < numSizes; nTest++) { + uint nBytes = Sizes[nTest % numSizes]; + ulong N = nBytes / sizeof(float); + const unsigned blocks = N / threadsPerBlock; + + hDst = new float[nBytes]; + HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess); + memset(hDst, 0, nBytes); + + HIPCHECK(hipMalloc(&dDst, nBytes)); + HIPCHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), dim3(threadsPerBlock), + 0, stream, dDst, N); + HIPCHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost)); + hipDeviceSynchronize(); + + int tmp = 0; + for (int i = 0; i < N; i++) { + if (i % threadsPerBlock == 0) { + tmp = 0; + } + if (hDst[i] != tmp) { + cout << "info: Data validation failed for warm up run!" << endl; + cout << "info: expected " << tmp << " got " << hDst[i] << endl; + HIPCHECK (hipErrorUnknown); + } + tmp += threadsPerBlock / 2; + } + + auto all_start = chrono::steady_clock::now(); + for (int i = 0; i < nIter; i++) { + hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), + dim3(threadsPerBlock), 0, stream, dDst, N); + } + hipDeviceSynchronize(); + + auto all_end = chrono::steady_clock::now(); + chrono::duration all_kernel_time = all_end - all_start; + + // read speed in GB/s + double perf = ((double) blocks * threadsPerBlock + * (numReads1 * sizeof(float) + sharedMemSizeBytes1 / 64) * nIter + * (double) (1e-09)) / all_kernel_time.count(); + + cout << "info: read speed = " << setw(8) << perf << " GB/s for " + << sharedMemSizeBytes1 / 1024 << " KB shared memory" + " with " << setw(8) << blocks * threadsPerBlock << " threads, " + << setw(4) << numReads1 << " reads in sharedMemReadSpeed1 kernel" << endl; + + delete[] hDst; + hipFree(dDst); + } + + + for (int nTest = 0; nTest < numSizes; nTest++) { + uint nBytes = Sizes[nTest % numSizes]; + ulong N = nBytes / sizeof(float); + const unsigned blocks = N / threadsPerBlock; + + hDst = new float[nBytes]; + HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess); + memset(hDst, 0, nBytes); + + HIPCHECK(hipMalloc(&dDst, nBytes)); + HIPCHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), dim3(threadsPerBlock), + 0, stream, dDst, N); + HIPCHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost)); + hipDeviceSynchronize(); + + auto all_start = chrono::steady_clock::now(); + for (int i = 0; i < nIter; i++) { + hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), + dim3(threadsPerBlock), 0, stream, dDst, N); + } + hipDeviceSynchronize(); + + auto all_end = chrono::steady_clock::now(); + chrono::duration all_kernel_time = all_end - all_start; + + // read speed in GB/s + double perf = ((double) blocks * threadsPerBlock + * (numReads2 * sizeof(float) + sharedMemSizeBytes2 / 64) * nIter + * (double) (1e-09)) / all_kernel_time.count(); + + cout << "info: read speed = " << setw(8) << perf << " GB/s for " + << sharedMemSizeBytes2 / 1024 << " KB shared memory" + " with " << setw(8) << blocks * threadsPerBlock << " threads, " + << setw(4) << numReads2 << " reads in sharedMemReadSpeed2 kernel" << endl; + + delete[] hDst; + hipFree(dDst); + } + + HIPCHECK(hipStreamDestroy(stream)); + + passed(); +} From 87558f64482bb685437630119146124b5cea1080 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Wed, 29 Jul 2020 21:26:10 -0400 Subject: [PATCH 3/5] [HipPerf] add two subtests for measuring maximum device memory read/write speed SWDEV-245290 / SWDEV-246220 [HIPPerf] Port OCLPerfDevMemWriteSpeed/OCLPerfDevMemReadSpeed into hip performance subtests Change-Id: I5dc323c75cebbc17596dcb4ed9492e18c5246868 --- perftests/memory/hipPerfDevMemReadSpeed.cpp | 165 +++++++++++++++++++ perftests/memory/hipPerfDevMemWriteSpeed.cpp | 155 +++++++++++++++++ 2 files changed, 320 insertions(+) create mode 100644 perftests/memory/hipPerfDevMemReadSpeed.cpp create mode 100644 perftests/memory/hipPerfDevMemWriteSpeed.cpp diff --git a/perftests/memory/hipPerfDevMemReadSpeed.cpp b/perftests/memory/hipPerfDevMemReadSpeed.cpp new file mode 100644 index 0000000000..f740d50ace --- /dev/null +++ b/perftests/memory/hipPerfDevMemReadSpeed.cpp @@ -0,0 +1,165 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include +#include +#include "test_common.h" + +using namespace std; + +#define arraySize 16 + +typedef struct d_uint16 { + uint data[arraySize]; +} d_uint16; + +__global__ void read_kernel(d_uint16 *src, ulong N, uint *dst) { + + size_t idx = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x ; + + uint tmp = 0; + for (size_t i = idx; i < N; i += stride) { + for (size_t j = 0; j < arraySize; j++) { + tmp += src[i].data[j]; + } + } + + atomicAdd(dst, tmp); +} + +int main(int argc, char* argv[]) { + d_uint16 *dSrc; + d_uint16 *hSrc; + uint *dDst; + uint *hDst; + hipStream_t stream; + ulong N = 4 * 1024 * 1024; + uint nBytes = N * sizeof(d_uint16); + + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + cout << "info: didn't find any GPU! skipping the test!\n"; + passed(); + return 0; + } + + static int device = 0; + HIPCHECK(hipSetDevice(device)); + hipDeviceProp_t props; + HIPCHECK(hipGetDeviceProperties(&props, device)); + cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name << + " with " << props.multiProcessorCount << " CUs" << endl; + + const unsigned threadsPerBlock = 64; + const unsigned blocks = props.multiProcessorCount * 4; + + uint inputData = 0x1; + int nIter = 1000; + + hSrc = new d_uint16[nBytes]; + HIPCHECK(hSrc == 0 ? hipErrorOutOfMemory : hipSuccess); + hDst = new uint; + hDst[0] = 0; + HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess); + for (size_t i = 0; i < N; i++) { + for (int j = 0; j < arraySize; j++) { + hSrc[i].data[j] = inputData; + } + } + + HIPCHECK(hipMalloc(&dSrc, nBytes)); + HIPCHECK(hipMalloc(&dDst, sizeof(uint))); + + HIPCHECK(hipStreamCreate(&stream)); + + HIPCHECK(hipMemcpy(dSrc, hSrc, nBytes, hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(dDst, hDst, sizeof(uint), hipMemcpyHostToDevice)); + + cout << "info: warm up launch for 'read_kernel' on the stream " << stream << endl; + + hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst); + HIPCHECK(hipMemcpy(hDst, dDst, sizeof(uint), hipMemcpyDeviceToHost)); + hipDeviceSynchronize(); + + if (hDst[0] != (nBytes / sizeof(uint))) { + cout << "info: Data validation failed for warm up run!" << endl; + cout << "info: expected " << nBytes / sizeof(uint) << " got " << hDst[0] << endl; + HIPCHECK(hipErrorUnknown); + } + + cout << "info: data validated for warm up launch for 'read_kernel'" << endl; + cout << "info: launching 'read_kernel' on the stream " << stream << " for "<< nIter << " iterations"<< endl; + + // measure performance based on host time + auto all_start = chrono::steady_clock::now(); + + for(int i = 0; i < nIter; i++) { + hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst); + } + hipDeviceSynchronize(); + + auto all_end = chrono::steady_clock::now(); + chrono::duration all_kernel_time = all_end - all_start; + + // read speed in GB/s + double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count(); + + cout << "info: average read speed of " << perf << " GB/s " << "achieved for memory size of " << + nBytes / (1024 * 1024) << " MB, calculated based on host time" << endl; + + // measure performance based on events time + hipEvent_t start, stop; + HIPCHECK(hipEventCreate(&start)); + HIPCHECK(hipEventCreate(&stop)); + float allEventMs = 0; + for(int i = 0; i < nIter; i++) { + HIPCHECK(hipEventRecord(start, NULL)); + + hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst); + + HIPCHECK(hipEventRecord(stop, NULL)); + HIPCHECK(hipEventSynchronize(stop)); + + float eventMs = 1.0f; + HIPCHECK(hipEventElapsedTime(&eventMs, start, stop)); + + allEventMs += eventMs; + + } + + double perfe = ((double)nBytes * nIter * (double)(1e-06)) / allEventMs; + cout << "info: average read speed of " << perfe << " GB/s " << "achieved for memory size of " << + nBytes / (1024 * 1024) << " MB, calculated based on events time" << endl; + + delete [] hSrc; + delete hDst; + hipFree(dSrc); + hipFree(dDst); + HIPCHECK(hipStreamDestroy(stream)); + + passed(); +} diff --git a/perftests/memory/hipPerfDevMemWriteSpeed.cpp b/perftests/memory/hipPerfDevMemWriteSpeed.cpp new file mode 100644 index 0000000000..9760a161c7 --- /dev/null +++ b/perftests/memory/hipPerfDevMemWriteSpeed.cpp @@ -0,0 +1,155 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include +#include +#include "test_common.h" + +using namespace std; + +#define arraySize 16 + +typedef struct d_uint16 { + uint data[arraySize]; +} d_uint16; + +__global__ void write_kernel(d_uint16 *dst, ulong N, d_uint16 pval) { + size_t idx = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + for (size_t i = idx; i < N; i += stride) { + dst[i] = pval; + } +}; + +int main(int argc, char* argv[]) { + d_uint16 *dDst; + d_uint16 *hDst; + hipStream_t stream; + ulong N = 4 * 1024 * 1024; + uint nBytes = N * sizeof(d_uint16); + d_uint16 pval; + + for (int i = 0; i < arraySize; i++) { + pval.data[i] = 0xabababab; + } + + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + cout << "info: didn't find any GPU! skipping the test!\n"; + passed(); + return 0; + } + + static int device = 0; + HIPCHECK(hipSetDevice(device)); + hipDeviceProp_t props; + HIPCHECK(hipGetDeviceProperties(&props, device)); + cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name << + " with " << props.multiProcessorCount << " CUs" << endl; + + size_t threadsPerBlock = 64; + size_t blocks = props.multiProcessorCount * 4; + + uint inputData = 0xabababab; + int nIter = 1000; + + hDst = new d_uint16[nBytes]; + HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess); + for (size_t i = 0; i < N; i++) { + for (size_t j = 0; j < arraySize; j++) { + hDst[i].data[j] = 0; + } + } + + HIPCHECK(hipMalloc(&dDst, nBytes)); + + HIPCHECK(hipStreamCreate(&stream)); + + + cout << "info: warm up launch for 'write_kernel' on the stream " << stream << endl; + + hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval); + HIPCHECK(hipMemcpy(hDst, dDst, nBytes , hipMemcpyDeviceToHost)); + hipDeviceSynchronize(); + + for (uint i = 0; i < N; i++) { + for (uint j = 0; j < arraySize; j++) { + if (hDst[i].data[j] != inputData) { + cout << "info: Data validation failed for warm up run! " << endl; + cout << "at index i: " << i << " element j: " << j << endl; + cout << hex << "expected 0x" << inputData << " but got 0x" << hDst[i].data[j] << endl; + HIPCHECK(hipErrorUnknown); + } + } + } + + cout << "info: data validated for warm up launch for 'write_kernel" << endl; + cout << "info: launching 'write_kernel' on the stream " << stream << " for "<< nIter << " iterations"<< endl; + + auto all_start = chrono::steady_clock::now(); + for(int i = 0; i < nIter; i++) { + hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval); + } + hipDeviceSynchronize(); + auto all_end = chrono::steady_clock::now(); + chrono::duration all_kernel_time = all_end - all_start; + + // read speed in GB/s + double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count(); + + cout << "info: average write speed of " << perf << " GB/s " << "achieved for memory size of " << + nBytes / (1024 * 1024) << " MB, calculated based on host time" << endl; + + // measure performance based on events time + hipEvent_t start, stop; + HIPCHECK(hipEventCreate(&start)); + HIPCHECK(hipEventCreate(&stop)); + float allEventMs = 0; + for(int i = 0; i < nIter; i++) { + HIPCHECK(hipEventRecord(start, NULL)); + + hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval); + + HIPCHECK(hipEventRecord(stop, NULL)); + HIPCHECK(hipEventSynchronize(stop)); + + float eventMs = 1.0f; + HIPCHECK(hipEventElapsedTime(&eventMs, start, stop)); + + allEventMs += eventMs; + + } + + double perfe = ((double)nBytes * nIter * (double)(1e-06)) / allEventMs; + cout << "info: average write speed of " << perfe << " GB/s " << "achieved for memory size of " << + nBytes / (1024 * 1024) << " MB, calculated based on events time" << endl; + + delete [] hDst; + hipFree(dDst); + HIPCHECK(hipStreamDestroy(stream)); + + passed(); +} From f7570dce0354526e4ee53e52913db99dcbeb7e2a Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Thu, 13 Aug 2020 14:23:58 -0400 Subject: [PATCH 4/5] [HipPerf] report performance based on wall time only for hipPerfDevMemReadSpeed/hipPerfDevMemWriteSpeed Change-Id: I1fda2ec76da6fad6852d328e0a3fc39e28af57bb --- perftests/memory/hipPerfDevMemReadSpeed.cpp | 31 +------------------- perftests/memory/hipPerfDevMemWriteSpeed.cpp | 31 +------------------- 2 files changed, 2 insertions(+), 60 deletions(-) diff --git a/perftests/memory/hipPerfDevMemReadSpeed.cpp b/perftests/memory/hipPerfDevMemReadSpeed.cpp index f740d50ace..181cd37f24 100644 --- a/perftests/memory/hipPerfDevMemReadSpeed.cpp +++ b/perftests/memory/hipPerfDevMemReadSpeed.cpp @@ -99,8 +99,6 @@ int main(int argc, char* argv[]) { HIPCHECK(hipMemcpy(dSrc, hSrc, nBytes, hipMemcpyHostToDevice)); HIPCHECK(hipMemcpy(dDst, hDst, sizeof(uint), hipMemcpyHostToDevice)); - cout << "info: warm up launch for 'read_kernel' on the stream " << stream << endl; - hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst); HIPCHECK(hipMemcpy(hDst, dDst, sizeof(uint), hipMemcpyDeviceToHost)); hipDeviceSynchronize(); @@ -111,9 +109,6 @@ int main(int argc, char* argv[]) { HIPCHECK(hipErrorUnknown); } - cout << "info: data validated for warm up launch for 'read_kernel'" << endl; - cout << "info: launching 'read_kernel' on the stream " << stream << " for "<< nIter << " iterations"<< endl; - // measure performance based on host time auto all_start = chrono::steady_clock::now(); @@ -129,31 +124,7 @@ int main(int argc, char* argv[]) { double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count(); cout << "info: average read speed of " << perf << " GB/s " << "achieved for memory size of " << - nBytes / (1024 * 1024) << " MB, calculated based on host time" << endl; - - // measure performance based on events time - hipEvent_t start, stop; - HIPCHECK(hipEventCreate(&start)); - HIPCHECK(hipEventCreate(&stop)); - float allEventMs = 0; - for(int i = 0; i < nIter; i++) { - HIPCHECK(hipEventRecord(start, NULL)); - - hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst); - - HIPCHECK(hipEventRecord(stop, NULL)); - HIPCHECK(hipEventSynchronize(stop)); - - float eventMs = 1.0f; - HIPCHECK(hipEventElapsedTime(&eventMs, start, stop)); - - allEventMs += eventMs; - - } - - double perfe = ((double)nBytes * nIter * (double)(1e-06)) / allEventMs; - cout << "info: average read speed of " << perfe << " GB/s " << "achieved for memory size of " << - nBytes / (1024 * 1024) << " MB, calculated based on events time" << endl; + nBytes / (1024 * 1024) << " MB" << endl; delete [] hSrc; delete hDst; diff --git a/perftests/memory/hipPerfDevMemWriteSpeed.cpp b/perftests/memory/hipPerfDevMemWriteSpeed.cpp index 9760a161c7..4d706cdde9 100644 --- a/perftests/memory/hipPerfDevMemWriteSpeed.cpp +++ b/perftests/memory/hipPerfDevMemWriteSpeed.cpp @@ -88,9 +88,6 @@ int main(int argc, char* argv[]) { HIPCHECK(hipStreamCreate(&stream)); - - cout << "info: warm up launch for 'write_kernel' on the stream " << stream << endl; - hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval); HIPCHECK(hipMemcpy(hDst, dDst, nBytes , hipMemcpyDeviceToHost)); hipDeviceSynchronize(); @@ -106,9 +103,6 @@ int main(int argc, char* argv[]) { } } - cout << "info: data validated for warm up launch for 'write_kernel" << endl; - cout << "info: launching 'write_kernel' on the stream " << stream << " for "<< nIter << " iterations"<< endl; - auto all_start = chrono::steady_clock::now(); for(int i = 0; i < nIter; i++) { hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval); @@ -121,31 +115,8 @@ int main(int argc, char* argv[]) { double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count(); cout << "info: average write speed of " << perf << " GB/s " << "achieved for memory size of " << - nBytes / (1024 * 1024) << " MB, calculated based on host time" << endl; + nBytes / (1024 * 1024) << " MB" << endl; - // measure performance based on events time - hipEvent_t start, stop; - HIPCHECK(hipEventCreate(&start)); - HIPCHECK(hipEventCreate(&stop)); - float allEventMs = 0; - for(int i = 0; i < nIter; i++) { - HIPCHECK(hipEventRecord(start, NULL)); - - hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval); - - HIPCHECK(hipEventRecord(stop, NULL)); - HIPCHECK(hipEventSynchronize(stop)); - - float eventMs = 1.0f; - HIPCHECK(hipEventElapsedTime(&eventMs, start, stop)); - - allEventMs += eventMs; - - } - - double perfe = ((double)nBytes * nIter * (double)(1e-06)) / allEventMs; - cout << "info: average write speed of " << perfe << " GB/s " << "achieved for memory size of " << - nBytes / (1024 * 1024) << " MB, calculated based on events time" << endl; delete [] hDst; hipFree(dDst); From 8dfe3c92a7c663403b4845d1c5a4a63cbec8420e Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Thu, 13 Aug 2020 17:26:44 -0400 Subject: [PATCH 5/5] [HipPerf] relocating/renaming some hip perf tests Change-Id: Ie85d242cd68cd14a858e07ed27875a5196014688 --- .../hipPerfDispatchSpeed.cpp | 2 +- .../hipPerfBufferCopyRectSpeed.cpp | 2 +- .../hipPerfBufferCopySpeed.cpp | 2 +- ...NumaAlloc.cpp => hipPerfHostNumaAlloc.cpp} | 2 +- perftests/perfDispatch/timer.cpp | 116 ------------------ perftests/perfDispatch/timer.h | 28 ----- 6 files changed, 4 insertions(+), 148 deletions(-) rename perftests/{perfDispatch => dispatch}/hipPerfDispatchSpeed.cpp (98%) rename perftests/{perfDispatch => memory}/hipPerfBufferCopyRectSpeed.cpp (98%) rename perftests/{perfDispatch => memory}/hipPerfBufferCopySpeed.cpp (98%) rename perftests/memory/{hipHostNumaAlloc.cpp => hipPerfHostNumaAlloc.cpp} (97%) delete mode 100644 perftests/perfDispatch/timer.cpp delete mode 100644 perftests/perfDispatch/timer.h diff --git a/perftests/perfDispatch/hipPerfDispatchSpeed.cpp b/perftests/dispatch/hipPerfDispatchSpeed.cpp similarity index 98% rename from perftests/perfDispatch/hipPerfDispatchSpeed.cpp rename to perftests/dispatch/hipPerfDispatchSpeed.cpp index 84ba73c3aa..12999ebc33 100644 --- a/perftests/perfDispatch/hipPerfDispatchSpeed.cpp +++ b/perftests/dispatch/hipPerfDispatchSpeed.cpp @@ -7,7 +7,7 @@ #include "test_common.h" /* HIT_START - * BUILD: %t %s ../../src/test_common.cpp timer.cpp EXCLUDE_HIP_PLATFORM nvcc + * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ diff --git a/perftests/perfDispatch/hipPerfBufferCopyRectSpeed.cpp b/perftests/memory/hipPerfBufferCopyRectSpeed.cpp similarity index 98% rename from perftests/perfDispatch/hipPerfBufferCopyRectSpeed.cpp rename to perftests/memory/hipPerfBufferCopyRectSpeed.cpp index 5000904af9..3cb3243e80 100644 --- a/perftests/perfDispatch/hipPerfBufferCopyRectSpeed.cpp +++ b/perftests/memory/hipPerfBufferCopyRectSpeed.cpp @@ -7,7 +7,7 @@ #include "test_common.h" /* HIT_START - * BUILD: %t %s ../../src/test_common.cpp timer.cpp EXCLUDE_HIP_PLATFORM nvcc + * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ diff --git a/perftests/perfDispatch/hipPerfBufferCopySpeed.cpp b/perftests/memory/hipPerfBufferCopySpeed.cpp similarity index 98% rename from perftests/perfDispatch/hipPerfBufferCopySpeed.cpp rename to perftests/memory/hipPerfBufferCopySpeed.cpp index 6f284ae7fb..d9a2d443a2 100644 --- a/perftests/perfDispatch/hipPerfBufferCopySpeed.cpp +++ b/perftests/memory/hipPerfBufferCopySpeed.cpp @@ -7,7 +7,7 @@ #include "test_common.h" /* HIT_START - * BUILD: %t %s ../../src/test_common.cpp timer.cpp EXCLUDE_HIP_PLATFORM nvcc + * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ diff --git a/perftests/memory/hipHostNumaAlloc.cpp b/perftests/memory/hipPerfHostNumaAlloc.cpp similarity index 97% rename from perftests/memory/hipHostNumaAlloc.cpp rename to perftests/memory/hipPerfHostNumaAlloc.cpp index 38401c8046..a5e60c8549 100644 --- a/perftests/memory/hipHostNumaAlloc.cpp +++ b/perftests/memory/hipPerfHostNumaAlloc.cpp @@ -34,7 +34,7 @@ THE SOFTWARE. #include #include "hip/hip_runtime.h" /* HIT_START - * BUILD_CMD: hipHostNumaAlloc %hc -I%S/../../src %S/%s %S/../../src/test_common.cpp -lnuma -o %T/%t EXCLUDE_HIP_PLATFORM nvcc + * BUILD_CMD: hipPerfHostNumaAlloc %hc -I%S/../../src %S/%s %S/../../src/test_common.cpp -lnuma -o %T/%t EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ diff --git a/perftests/perfDispatch/timer.cpp b/perftests/perfDispatch/timer.cpp deleted file mode 100644 index ea9c6ea1d9..0000000000 --- a/perftests/perfDispatch/timer.cpp +++ /dev/null @@ -1,116 +0,0 @@ -#include "timer.h" - -#include - -#ifdef _WIN32 -#define WIN32_LEAN_AND_MEAN -#define VC_EXTRALEAN -#include -#pragma comment(lib, "user32") -#endif - -#ifdef __linux__ -#include -#define NANOSECONDS_PER_SEC 1000000000 -#endif - -CPerfCounter::CPerfCounter() : _clocks(0), _start(0) -{ - -#ifdef _WIN32 - - QueryPerformanceFrequency((LARGE_INTEGER *)&_freq); - -#endif - -#ifdef __linux__ - _freq = NANOSECONDS_PER_SEC; -#endif - -} - -CPerfCounter::~CPerfCounter() -{ - // EMPTY! -} - -void -CPerfCounter::Start(void) -{ - -#ifdef _WIN32 - - if( _start ) - { - MessageBox(NULL, "Bad Perf Counter Start", "Error", MB_OK); - exit(0); - } - QueryPerformanceCounter((LARGE_INTEGER *)&_start); - -#endif -#ifdef __linux__ - - struct timespec s; - clock_gettime(CLOCK_MONOTONIC, &s); - _start = (i64)s.tv_sec * NANOSECONDS_PER_SEC + (i64)s.tv_nsec ; - -#endif - -} - -void -CPerfCounter::Stop(void) -{ - i64 n; - -#ifdef _WIN32 - - if( !_start ) - { - MessageBox(NULL, "Bad Perf Counter Stop", "Error", MB_OK); - exit(0); - } - - QueryPerformanceCounter((LARGE_INTEGER *)&n); - -#endif -#ifdef __linux__ - - struct timespec s; - clock_gettime(CLOCK_MONOTONIC, &s); - n = (i64)s.tv_sec * NANOSECONDS_PER_SEC + (i64)s.tv_nsec ; - -#endif - - n -= _start; - _start = 0; - _clocks += n; -} - -void -CPerfCounter::Reset(void) -{ - -#ifdef _WIN32 - if( _start ) - { - MessageBox(NULL, "Bad Perf Counter Reset", "Error", MB_OK); - exit(0); - } -#endif - _clocks = 0; -} - -double -CPerfCounter::GetElapsedTime(void) -{ -#ifdef _WIN32 - if( _start ) { - MessageBox(NULL, "Trying to get time while still running.", "Error", MB_OK); - exit(0); - } -#endif - - return (double)_clocks / (double)_freq; - -} diff --git a/perftests/perfDispatch/timer.h b/perftests/perfDispatch/timer.h deleted file mode 100644 index 28bfeff74b..0000000000 --- a/perftests/perfDispatch/timer.h +++ /dev/null @@ -1,28 +0,0 @@ -#ifndef _TIMER_H_ -#define _TIMER_H_ - -#ifdef _WIN32 -typedef __int64 i64 ; -#endif -#ifdef __linux__ -typedef long long i64; -#endif - -class CPerfCounter { - -public: - CPerfCounter(); - ~CPerfCounter(); - void Start(void); - void Stop(void); - void Reset(void); - double GetElapsedTime(void); - -private: - - i64 _freq; - i64 _clocks; - i64 _start; -}; - -#endif // _TIMER_H_