From 0e5e2ec2f7151de2650904701195669e31fc00c4 Mon Sep 17 00:00:00 2001 From: mbhiutra Date: Mon, 30 Oct 2023 14:12:41 +0530 Subject: [PATCH] SWDEV-403471 - [catch2][dtest] Conversion of preftests to catch2 Change-Id: I68cb780a71a6094dca86718e7d427806d3a0e67d [ROCm/hip-tests commit: b0c4a4f70fe50583faea061547d228b86cc52031] --- .../hip-tests/catch/perftests/CMakeLists.txt | 3 + .../catch/perftests/compute/CMakeLists.txt | 29 + .../perftests/compute/hipPerfDotProduct.cc | 367 ++++++++++ .../perftests/compute/hipPerfMandelbrot.cc | 659 ++++++++++++++++++ .../catch/perftests/dispatch/CMakeLists.txt | 28 + .../dispatch/hipPerfDispatchSpeed.cc | 196 ++++++ .../catch/perftests/stream/CMakeLists.txt | 30 + .../stream/hipPerfDeviceConcurrency.cc | 269 +++++++ .../stream/hipPerfStreamConcurrency.cc | 415 +++++++++++ .../stream/hipPerfStreamCreateCopyDestroy.cc | 146 ++++ 10 files changed, 2142 insertions(+) create mode 100644 projects/hip-tests/catch/perftests/compute/CMakeLists.txt create mode 100644 projects/hip-tests/catch/perftests/compute/hipPerfDotProduct.cc create mode 100644 projects/hip-tests/catch/perftests/compute/hipPerfMandelbrot.cc create mode 100644 projects/hip-tests/catch/perftests/dispatch/CMakeLists.txt create mode 100644 projects/hip-tests/catch/perftests/dispatch/hipPerfDispatchSpeed.cc create mode 100644 projects/hip-tests/catch/perftests/stream/CMakeLists.txt create mode 100644 projects/hip-tests/catch/perftests/stream/hipPerfDeviceConcurrency.cc create mode 100644 projects/hip-tests/catch/perftests/stream/hipPerfStreamConcurrency.cc create mode 100644 projects/hip-tests/catch/perftests/stream/hipPerfStreamCreateCopyDestroy.cc diff --git a/projects/hip-tests/catch/perftests/CMakeLists.txt b/projects/hip-tests/catch/perftests/CMakeLists.txt index be81c63a12..941763109f 100644 --- a/projects/hip-tests/catch/perftests/CMakeLists.txt +++ b/projects/hip-tests/catch/perftests/CMakeLists.txt @@ -22,3 +22,6 @@ add_custom_target(perf_test COMMAND "${CMAKE_CTEST_COMMAND}" -R "Perf_" COMMENT "Build complete, now executing the performnce test ...") add_subdirectory(memory) +add_subdirectory(stream) +add_subdirectory(dispatch) +add_subdirectory(compute) diff --git a/projects/hip-tests/catch/perftests/compute/CMakeLists.txt b/projects/hip-tests/catch/perftests/compute/CMakeLists.txt new file mode 100644 index 0000000000..4844b92c96 --- /dev/null +++ b/projects/hip-tests/catch/perftests/compute/CMakeLists.txt @@ -0,0 +1,29 @@ +# Copyright (c) 2023 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. + +# Common Tests - Test independent of all platforms +set(TEST_SRC + hipPerfDotProduct.cc + hipPerfMandelbrot.cc +) + +hip_add_exe_to_target(NAME perfComputeTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME perf_test) diff --git a/projects/hip-tests/catch/perftests/compute/hipPerfDotProduct.cc b/projects/hip-tests/catch/perftests/compute/hipPerfDotProduct.cc new file mode 100644 index 0000000000..761adc92c4 --- /dev/null +++ b/projects/hip-tests/catch/perftests/compute/hipPerfDotProduct.cc @@ -0,0 +1,367 @@ +/* + Copyright (c) 2023 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. + */ + +/** +* @addtogroup hipPerfDotProduct hipPerfDotProduct +* @{ +* @ingroup perfComputeTest +*/ + +#include +#include + +#define DOT_DIM 256 + +using namespace std; + +template +__launch_bounds__(BLOCKSIZE) +__global__ void vectors_not_equal(int n, + const double* __restrict__ x, + const double* __restrict__ y, + double* __restrict__ workspace) { + int gid = blockIdx.x * blockDim.x + threadIdx.x; + + double sum = 0.0; + for (int idx = gid; idx < n; idx += hipGridDim_x * hipBlockDim_x) { + sum = fma(y[idx], x[idx], sum); + } + + __shared__ double sdata[BLOCKSIZE]; + sdata[threadIdx.x] = sum; + + __syncthreads(); + + if (threadIdx.x < 128) { + sdata[threadIdx.x] += sdata[threadIdx.x + 128]; + } + __syncthreads(); + + if (threadIdx.x < 64) { + sdata[threadIdx.x] += sdata[threadIdx.x + 64]; + } + __syncthreads(); + + if (threadIdx.x < 32) { + sdata[threadIdx.x] += sdata[threadIdx.x + 32]; + } + __syncthreads(); + + if (threadIdx.x < 16) { + sdata[threadIdx.x] += sdata[threadIdx.x + 16]; + } + __syncthreads(); + + if (threadIdx.x < 8) { + sdata[threadIdx.x] += sdata[threadIdx.x + 8]; + } + __syncthreads(); + + if (threadIdx.x < 4) { + sdata[threadIdx.x] += sdata[threadIdx.x + 4]; + } + __syncthreads(); + + if (threadIdx.x < 2) { + sdata[threadIdx.x] += sdata[threadIdx.x + 2]; + } + __syncthreads(); + + if (threadIdx.x < 1) { + sdata[threadIdx.x] += sdata[threadIdx.x + 1]; + } + + if (threadIdx.x == 0) { + workspace[blockIdx.x] = sdata[0]; + } +} + +template +__launch_bounds__(BLOCKSIZE) +__global__ void vectors_equal(int n, const double* __restrict__ x, + double* __restrict__ workspace) { + int gid = blockIdx.x * blockDim.x + threadIdx.x; + + double sum = 0.0; + for (int idx = gid; idx < n; idx += hipGridDim_x * blockDim.x) { + sum = fma(x[idx], x[idx], sum); + } + + __shared__ double sdata[BLOCKSIZE]; + sdata[threadIdx.x] = sum; + + __syncthreads(); + + if (threadIdx.x < 128) { + sdata[threadIdx.x] += sdata[threadIdx.x + 128]; + } + __syncthreads(); + + if (threadIdx.x < 64) { + sdata[threadIdx.x] += sdata[threadIdx.x + 64]; + } + __syncthreads(); + + if (threadIdx.x < 32) { + sdata[threadIdx.x] += sdata[threadIdx.x + 32]; + } + __syncthreads(); + + if (threadIdx.x < 16) { + sdata[threadIdx.x] += sdata[threadIdx.x + 16]; + } + __syncthreads(); + + if (threadIdx.x < 8) { + sdata[threadIdx.x] += sdata[threadIdx.x + 8]; + } + __syncthreads(); + + if (threadIdx.x < 4) { + sdata[threadIdx.x] += sdata[threadIdx.x + 4]; + } + __syncthreads(); + + if (threadIdx.x < 2) { + sdata[threadIdx.x] += sdata[threadIdx.x + 2]; + } + __syncthreads(); + + if (threadIdx.x < 1) { + sdata[threadIdx.x] += sdata[threadIdx.x + 1]; + } + + if (threadIdx.x == 0) { + workspace[blockIdx.x] = sdata[0]; + } +} + +template +__launch_bounds__(BLOCKSIZE) +__global__ void dot_reduction(double* __restrict__ workspace) { + __shared__ double sdata[BLOCKSIZE]; + + sdata[threadIdx.x] = workspace[threadIdx.x]; + __syncthreads(); + + if (threadIdx.x < 128) { + sdata[threadIdx.x] += sdata[threadIdx.x + 128]; + } + __syncthreads(); + + if (threadIdx.x < 64) { + sdata[threadIdx.x] += sdata[threadIdx.x + 64]; + } + __syncthreads(); + + if (threadIdx.x < 32) { + sdata[threadIdx.x] += sdata[threadIdx.x + 32]; + } + __syncthreads(); + + if (threadIdx.x < 16) { + sdata[threadIdx.x] += sdata[threadIdx.x + 16]; + } + __syncthreads(); + + if (threadIdx.x < 8) { + sdata[threadIdx.x] += sdata[threadIdx.x + 8]; + } + __syncthreads(); + + if (threadIdx.x < 4) { + sdata[threadIdx.x] += sdata[threadIdx.x + 4]; + } __syncthreads(); + + if (threadIdx.x < 2) { + sdata[threadIdx.x] += sdata[threadIdx.x + 2]; + } + __syncthreads(); + + if (threadIdx.x < 1) { + sdata[threadIdx.x] += sdata[threadIdx.x + 1]; + } + + if (threadIdx.x == 0) { + workspace[0] = sdata[0]; + } +} + +void computeDotProduct(int n, const double* x, const double* y, double& result, + double* workspace) { + dim3 blocks(DOT_DIM); + dim3 threadsPerBlock(DOT_DIM); + + if (x != y) { + hipLaunchKernelGGL(vectors_not_equal, blocks, threadsPerBlock, 0, 0, n, x, y, + workspace); + } else { + hipLaunchKernelGGL(vectors_equal, blocks, threadsPerBlock, 0, 0, n, x, workspace); + } + + // Part 2 of dot product computation + hipLaunchKernelGGL(dot_reduction, dim3(1), threadsPerBlock, 0, 0, workspace); + + // Copy the final dot product result back from the device + HIP_CHECK(hipMemcpy(&result, workspace, sizeof(double), hipMemcpyDeviceToHost)); + + return; +} + +/** +* Test Description +* ------------------------ +* - Verify the device kernel results comparing it with the host results. +* Test source +* ------------------------ +* - perftests/compute/hipPerfDotProduct.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfDotProduct") { + int nGpu = 0; + int p_gpuDevice = 0; + HIP_CHECK(hipGetDeviceCount(&nGpu)); + + if (nGpu < 1) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 1"); + } + hipDeviceProp_t props = {0}; + props = {0}; + HIP_CHECK(hipSetDevice(p_gpuDevice)); + HIP_CHECK(hipGetDeviceProperties(&props, p_gpuDevice)); + int nx, ny, nz; + + for (unsigned int testCase = 0; testCase < 3; testCase++) { + vector vectorSize = {200, 300, 50}; + switch (testCase) { + case 0: + nx = vectorSize[0]; + ny = vectorSize[0]; + nz = vectorSize[0]; + break; + + case 1: + nx = vectorSize[1]; + ny = vectorSize[1]; + nz = vectorSize[1]; + break; + + case 2: + nx = vectorSize[0]; + ny = vectorSize[1]; + nz = vectorSize[2]; + break; + + default: + break; + } + + int trials = 200; + int size = nx * ny * nz; + + vector hx(size); + vector hy(size); + double hresult_xy = 0.0; + double hresult_xx = 0.0; + + srand(time(NULL)); + + for (int i = 0; i < size; ++i) { + hx[i] = 2.0 * static_castrand() / static_castRAND_MAX - 1.0; + hy[i] = 2.0 * static_castrand() / static_castRAND_MAX - 1.0; + + hresult_xy += hx[i] * hy[i]; + hresult_xx += hx[i] * hx[i]; + } + + double* dx; + double* dy; + double* workspace; + double dresult; + + HIP_CHECK(hipMalloc(reinterpret_cast(&dx), sizeof(double) * size)); + HIP_CHECK(hipMalloc(reinterpret_cast(&dy), sizeof(double) * size)); + HIP_CHECK(hipMalloc(reinterpret_cast(&workspace), sizeof(double) * DOT_DIM)); + + HIP_CHECK(hipMemcpy(dx, hx.data(), sizeof(double) * size, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(dy, hy.data(), sizeof(double) * size, hipMemcpyHostToDevice)); + + // Warm up + computeDotProduct(size, dx, dy, dresult, workspace); + computeDotProduct(size, dx, dy, dresult, workspace); + computeDotProduct(size, dx, dy, dresult, workspace); + + // Timed run for + HIP_CHECK(hipDeviceSynchronize()); + auto all_start = std::chrono::steady_clock::now(); + + for (int i = 0; i < trials; ++i) { + computeDotProduct(size, dx, dy, dresult, workspace); + } + + float time = 0; + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + time = all_kernel_time.count(); + + time /= trials; + + double bw = sizeof(double) * size * 2.0 / 1e9; + double gf = 2.0 * size / 1e9; + + cout << "\nVector Size: " << size << "\n[ddot] " << time << "msec ;" << bw/ (time / 1e3) << " GByte/s ;" + << gf/(time / 1e3) << " GFlop/s" << endl; + + // Verify the device kernel results comparing it with the host results + REQUIRE(std::abs(dresult - hresult_xy) < std::max(dresult * 1e-10, 1e-8)); + + // Warm up + computeDotProduct(size, dx, dx, dresult, workspace); + computeDotProduct(size, dx, dx, dresult, workspace); + computeDotProduct(size, dx, dx, dresult, workspace); + + // Timed run for + HIP_CHECK(hipDeviceSynchronize()); + all_start = std::chrono::steady_clock::now(); + + for (int i = 0; i < trials; ++i) { + computeDotProduct(size, dx, dx, dresult, workspace); + } + + all_end = std::chrono::steady_clock::now(); + all_kernel_time = all_end - all_start; + time = all_kernel_time.count(); + + time /= trials; + bw = sizeof(double) * size / 1e9; + + cout << "[ddot] " << time << "msec ;" << bw/ (time / 1e3) << " GByte/s ;" + << gf/(time / 1e3) << " GFlop/s" << endl; + + // Verify the device kernel results comparing it with the host results + REQUIRE(abs(dresult - hresult_xx) < max(dresult * 1e-10, 1e-8)); + + HIP_CHECK(hipFree(dx)); + HIP_CHECK(hipFree(dy)); + HIP_CHECK(hipFree(workspace)); + } +} diff --git a/projects/hip-tests/catch/perftests/compute/hipPerfMandelbrot.cc b/projects/hip-tests/catch/perftests/compute/hipPerfMandelbrot.cc new file mode 100644 index 0000000000..9c4b3a7a14 --- /dev/null +++ b/projects/hip-tests/catch/perftests/compute/hipPerfMandelbrot.cc @@ -0,0 +1,659 @@ +/* + Copyright (c) 2023 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. + */ + +/** +* @addtogroup hipPerfMandelbrot hipPerfMandelbrot +* @{ +* @ingroup perfComputeTest +*/ + +#include +#include +#include +#include +#include +#include + +typedef struct { + double x; + double y; + double width; +} coordRec; + +coordRec coords[] = { + {0.0, 0.0, 4.0}, // Whole set + {0.0, 0.0, 0.00001}, // All black + {-0.0180789661868, 0.6424294066162, 0.00003824140}, // Hit detail +}; + +static unsigned int numCoords = sizeof(coords) / sizeof(coordRec); + +template +__global__ void float_mad_kernel(uint *out, uint width, T xPos, T yPos, + T xStep, T yStep, uint maxIter) { +#pragma FP_CONTRACT ON + int tid = (blockIdx.x * blockDim.x + threadIdx.x); + int i = tid % width; + int j = tid / width; + float x0 = static_cast(xPos + xStep*i); + float y0 = static_cast(yPos + yStep*j); + + float x = x0; + float y = y0; + + uint iter = 0; + float tmp; + for (iter = 0; (x*x + y*y <= 4.0f) && (iter < maxIter); iter++) { + tmp = x; + x = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*tmp, y, y0); + } + out[tid] = iter; +} + +template +__global__ void float_mandel_unroll_kernel(uint *out, uint width, T xPos, + T yPos, T xStep, T yStep, uint maxIter) { + +#pragma FP_CONTRACT ON + int tid = (blockIdx.x * blockDim.x + threadIdx.x); + int i = tid % width; + int j = tid / width; + float x0 = static_cast(xPos + xStep*static_cast(i)); + float y0 = static_cast(yPos + yStep*static_cast(j)); + + float x = x0; + float y = y0; + +#define FAST + uint iter = 0; + float tmp; + int stay; + int ccount = 0; + stay = (x*x+y*y) <= 4.0; + float savx = x; + float savy = y; +#ifdef FAST + for (iter = 0; (iter < maxIter); iter+=16) { +#else + for (iter = 0; stay && (iter < maxIter); iter+=16) { +#endif + x = savx; + y = savy; + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + stay = (x*x+y*y) <= 4.0; + savx = (stay ? x : savx); + savy = (stay ? y : savy); + ccount += stay*16; +#ifdef FAST + if (!stay) + break; +#endif + } + // Handle remainder + if (!stay) { + iter = 16; + do { + x = savx; + y = savy; + stay = ((x*x+y*y) <= 4.0) && (ccount < maxIter); + tmp = x; + x = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*tmp, y, y0); + ccount += stay; + iter--; + savx = (stay ? x : savx); + savy = (stay ? y : savy); + } while (stay && iter); + } + out[tid] = (uint)ccount; +} + +template +__global__ void double_mad_kernel(uint *out, uint width, T xPos, T yPos, T xStep, T yStep, + uint maxIter) { +#pragma FP_CONTRACT ON + int tid = (blockIdx.x * blockDim.x + threadIdx.x); + int i = tid % width; + int j = tid / width; + double x0 = static_cast(xPos + xStep*i); + double y0 = static_cast(yPos + yStep*j); + + double x = x0; + double y = y0; + + uint iter = 0; + double tmp; + for (iter = 0; (x*x + y*y <= 4.0f) && (iter < maxIter); iter++) { + tmp = x; + x = fma(-y, y,fma(x, x, x0)); + y = fma(2.0f*tmp, y, y0); + } + out[tid] = iter; +}; + +template +__global__ void double_mandel_unroll_kernel(uint *out, uint width, T xPos, + T yPos, T xStep, T yStep, uint maxIter) { +#pragma FP_CONTRACT ON + int tid = (blockIdx.x * blockDim.x + threadIdx.x); + + int i = tid % width; + int j = tid / width; + double x0 = static_cast(xPos + xStep*static_cast(i)); + double y0 = static_cast(yPos + yStep*static_cast(j)); + + double x = x0; + double y = y0; + +#define FAST + uint iter = 0; + double tmp; + int stay; + int ccount = 0; + stay = (x*x+y*y) <= 4.0; + double savx = x; + double savy = y; +#ifdef FAST + for (iter = 0; (iter < maxIter); iter+=16) +#else + for (iter = 0; stay && (iter < maxIter); iter+=16) +#endif + { + x = savx; + y = savy; + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x,y,y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + // Two iterations + tmp = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*x, y, y0); + x = fma(-y, y, fma(tmp, tmp, x0)); + y = fma(2.0f*tmp, y, y0); + + stay = (x*x+y*y) <= 4.0; + savx = (stay ? x : savx); + savy = (stay ? y : savy); + ccount += stay*16; +#ifdef FAST + if (!stay) + break; +#endif + } + // Handle remainder + if (!stay) { + iter = 16; + do { + x = savx; + y = savy; + stay = ((x*x+y*y) <= 4.0) && (ccount < maxIter); + tmp = x; + x = fma(-y,y, fma(x, x, x0)); + y = fma(2.0f*tmp,y,y0); + ccount += stay; + iter--; + savx = (stay ? x : savx); + savy = (stay ? y : savy); + } + while (stay && iter); + } + out[tid] = (uint)ccount; +}; + +static const unsigned int FMA_EXPECTEDVALUES_INDEX = 15; + +// Expected results for each kernel run at each coord +unsigned long long expectedIters[] = { + 203277748ull, 2147483648ull, 120254651ull, 203277748ull, 2147483648ull, + 120254651ull, 203277748ull, 2147483648ull, 120254651ull, 203315114ull, + 2147483648ull, 120042599ull, 203315114ull, 2147483648ull, 120042599ull, + 203280620ull, 2147483648ull, 120485704ull, 203280620ull, 2147483648ull, + 120485704ull, 203280620ull, 2147483648ull, 120485704ull, 203315114ull, + 2147483648ull, 120042599ull, 203315114ull, 2147483648ull, 120042599ull}; + +class hipPerfMandelBrot { + public: + hipPerfMandelBrot(); + ~hipPerfMandelBrot(); + + void setNumKernels(unsigned int num) { + numKernels = num; + } + + unsigned int getNumKernels() { + return numKernels; + } + + void setNumStreams(unsigned int num) { + numStreams = num; + } + unsigned int getNumStreams() { + return numStreams; + } + + void open(int deviceID); + bool run(unsigned int testCase, unsigned int deviceId); + void printResults(void); + + // array of funtion pointers + typedef void (hipPerfMandelBrot::*funPtr)(uint *out, uint width, float xPos, float yPos, + float xStep, float yStep, uint maxIter, hipStream_t* streams, int blocks, + int threads_per_block, int kernelCnt); + + // Wrappers + void float_mad(uint *out, uint width, float xPos, float yPos, + float xStep, float yStep, uint maxIter, hipStream_t* streams, + int blocks, int threads_per_block, int kernelCnt); + + void float_mandel_unroll(uint *out, uint width, float xPos, float yPos, + float xStep, float yStep, uint maxIter, hipStream_t* streams, + int blocks, int threads_per_block, int kernelCnt); + + void double_mad(uint *out, uint width, float xPos, float yPos, float xStep, + float yStep, uint maxIter, hipStream_t* streams, int blocks, + int threads_per_block, int kernelCnt); + + void double_mandel_unroll(uint *out, uint width, float xPos, float yPos, float xStep, + float yStep, uint maxIter, hipStream_t* streams, int blocks, + int threads_per_block, int kernelCnt); + + hipStream_t streams[2]; + + private: + void setData(void *ptr, unsigned int value); + void checkData(uint *ptr); + + unsigned int numKernels; + unsigned int numStreams; + + std::map> results; + unsigned int width_; + unsigned int bufSize; + unsigned int maxIter; + unsigned int coordIdx; + volatile unsigned long long totalIters = 0; + int numCUs; + static const unsigned int numLoops = 10; +}; + +hipPerfMandelBrot::hipPerfMandelBrot() {} + +hipPerfMandelBrot::~hipPerfMandelBrot() {} + +void hipPerfMandelBrot::open(int deviceId) { + int nGpu = 0; + HIP_CHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 1"); + } + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + std::cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name + << " with " << props.multiProcessorCount << " CUs" << " and device id: " << deviceId + << std::endl; + + numCUs = props.multiProcessorCount; +} + +void hipPerfMandelBrot::printResults() { + int numkernels = getNumKernels(); + int numStreams = getNumStreams(); + + std::cout << "\n" <<"Measured perf for kernels in GFLOPS on " + << numStreams << " streams (s)" << std::endl; + + std::map>:: iterator itr; + for (itr = results.begin(); itr != results.end(); itr++) { + std::cout << "\n" << std::setw(20) << itr->first << " "; + for (auto i : results[itr->first]) { + std::cout << std::setw(10) << i << " "; + } + } + results.clear(); + std::cout << std::endl; +} + +// Wrappers for the kernel launches +void hipPerfMandelBrot::float_mad(uint *out, uint width, float xPos, float yPos, float xStep, + float yStep, uint maxIter, hipStream_t* streams, + int blocks, int threads_per_block, int kernelCnt) { + int streamCnt = getNumStreams(); + hipLaunchKernelGGL(float_mad_kernel, dim3(blocks), dim3(threads_per_block), 0, + streams[kernelCnt % streamCnt], out, width_, xPos, yPos, xStep, yStep, + maxIter); +} + +void hipPerfMandelBrot::float_mandel_unroll(uint *out, uint width, float xPos, float yPos, + float xStep, float yStep, uint maxIter, hipStream_t * streams, + int blocks, int threads_per_block, int kernelCnt) { + int streamCnt = getNumStreams(); + hipLaunchKernelGGL(float_mandel_unroll_kernel, dim3(blocks), dim3(threads_per_block), 0, + streams[kernelCnt % streamCnt], out, width_, xPos, yPos, xStep, yStep, maxIter); +} + +void hipPerfMandelBrot::double_mad(uint *out, uint width, float xPos, float yPos, + float xStep, float yStep, uint maxIter, hipStream_t * streams, + int blocks, int threads_per_block, int kernelCnt) { + int streamCnt = getNumStreams(); + hipLaunchKernelGGL(double_mad_kernel, dim3(blocks), dim3(threads_per_block), 0, + streams[kernelCnt % streamCnt], out, width_, xPos, yPos, xStep, yStep, maxIter); +} + +void hipPerfMandelBrot::double_mandel_unroll(uint *out, uint width, float xPos, float yPos, + float xStep, float yStep, uint maxIter, hipStream_t * streams, + int blocks, int threads_per_block, int kernelCnt) { + int streamCnt = getNumStreams(); + hipLaunchKernelGGL(float_mandel_unroll_kernel, dim3(blocks), dim3(threads_per_block), 0, + streams[kernelCnt % streamCnt], out, width_, xPos, yPos, xStep, yStep, maxIter); +} + +bool hipPerfMandelBrot::run(unsigned int testCase,unsigned int deviceId) { + unsigned int numStreams = getNumStreams(); + coordIdx = testCase % numCoords; + + funPtr p[] = {&hipPerfMandelBrot::float_mad, &hipPerfMandelBrot::float_mandel_unroll, + &hipPerfMandelBrot::double_mad, &hipPerfMandelBrot::double_mandel_unroll}; + + // Maximum iteration count + maxIter = 32768; + + uint * hPtr[numKernels]; + uint * dPtr[numKernels]; + + // Width is divisible by 4 because the mandelbrot kernel processes 4 pixels at once. + width_ = 256; + + bufSize = width_ * width_ * sizeof(uint); + + // Create streams for concurrency + for (uint i = 0; i < numStreams; i++) { + HIP_CHECK(hipStreamCreate(&streams[i])); + } + + // Allocate memory on the host and device + for (uint i = 0; i < numKernels; i++) { + HIP_CHECK(hipHostMalloc(reinterpret_cast(&hPtr[i]), bufSize, hipHostMallocDefault)); + setData(hPtr[i], 0xdeadbeef); + HIP_CHECK(hipMalloc(reinterpret_cast(&dPtr[i]), bufSize)) + } + + // Prepare kernel launch parameters + int threads = (bufSize/sizeof(uint)); + int threads_per_block = 64; + int blocks = (threads/threads_per_block) + (threads % threads_per_block); + + float xStep = static_cast(coords[coordIdx].width / (double)width_); + float yStep = static_cast(-coords[coordIdx].width / (double)width_); + float xPos = static_cast(coords[coordIdx].x - 0.5 * coords[coordIdx].width); + float yPos = static_cast(coords[coordIdx].y + 0.5 * coords[coordIdx].width); + + // Copy memory asynchronously and concurrently from host to device + for (uint i = 0; i < numKernels; i++) { + HIP_CHECK(hipMemcpy(dPtr[i], hPtr[i], bufSize, hipMemcpyHostToDevice)); + } + + // Synchronize to make sure all the copies are completed + HIP_CHECK(hipStreamSynchronize(0)); + + int kernelIdx; + if(testCase == 0 || testCase == 5 || testCase == 10) { + kernelIdx = 0; + } else if(testCase == 1 || testCase == 6 || testCase == 11) { + kernelIdx = 1; + } else if(testCase == 2 || testCase == 7 || testCase == 12) { + kernelIdx = 2; + } else if(testCase == 3 || testCase == 8 || testCase == 13){ + kernelIdx = 3; + } + double totalTime = 0.0; + for (unsigned int k = 0; k < numLoops; k++) { + if ((testCase == 0 || testCase == 1 || testCase == 2 || + testCase == 5 || testCase == 6 || testCase == 7 || + testCase == 10 || testCase == 11 || testCase == 12)) { + float xStep = static_cast(coords[coordIdx].width / static_cast(width_)); + float yStep = static_cast(-coords[coordIdx].width / static_cast(width_)); + float xPos = static_cast(coords[coordIdx].x - 0.5 * coords[coordIdx].width); + float yPos = static_cast(coords[coordIdx].y + 0.5 * coords[coordIdx].width); + + // Time the kernel execution + auto all_start = std::chrono::steady_clock::now(); + + for (uint i = 0; i < numKernels; i++) { + (this->*p[kernelIdx])(dPtr[i], width_, xPos, yPos, xStep, yStep, maxIter, streams, blocks, + threads_per_block, i); + } + + // Synchronize all the concurrent streams to have completed execution + HIP_CHECK(hipStreamSynchronize(0)); + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + totalTime += all_kernel_time.count(); + } else { + double xStep = coords[coordIdx].width / static_cast(width_); + double yStep = -coords[coordIdx].width / static_cast(width_); + double xPos = coords[coordIdx].x - 0.5 * coords[coordIdx].width; + double yPos = coords[coordIdx].y + 0.5 * coords[coordIdx].width; + + // Time the kernel execution + auto all_start = std::chrono::steady_clock::now(); + for (uint i = 0; i < numKernels; i++) { + (this->*p[kernelIdx])(dPtr[i], width_, xPos, yPos, xStep, yStep, maxIter, streams, blocks, + threads_per_block, i); + } + // Synchronize all the concurrent streams to have completed execution + HIP_CHECK(hipStreamSynchronize(0)); + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + totalTime += all_kernel_time.count(); + } + } + + // Copy data back from device to the host + for(uint i = 0; i < numKernels; i++) { + HIP_CHECK(hipMemcpy(hPtr[i] ,dPtr[i], bufSize, hipMemcpyDeviceToHost)); + } + for(uint i = 0; i < numKernels; i++) { + checkData(hPtr[i]); + int j =0; + while((totalIters != expectedIters[j] && totalIters > expectedIters[j]) && j < 30) { + j++; + } + + if(j==30) { + std::cout << "Incorrect iteration count detected. "; + } + } + + // Compute GFLOPS. There are 7 FLOPs per iteration + double perf = (static_cast(totalIters*numKernels) * 7 * static_cast(1e-09)) / + (totalTime / (double)numLoops); + + + std::vector kernelName = {"float", "float_unroll", + "double", "double_unroll"}; + + // Print results except for Warm-up kernel + if (testCase != 100) { + results[kernelName[testCase % 4]].push_back(perf); + } + + for(uint i = 0 ; i < numStreams; i++) { + HIP_CHECK(hipStreamDestroy(streams[i])); + } + + // Free host and device memory + for (uint i = 0; i < numKernels; i++) { + HIP_CHECK(hipHostFree(hPtr[i])); + HIP_CHECK(hipFree(dPtr[i])); + } + return true; +} + +void hipPerfMandelBrot::setData(void *ptr, unsigned int value) { + unsigned int *ptr2 = (unsigned int *)ptr; + for (unsigned int i = 0; i < width_ * width_; i++) { + ptr2[i] = value; + } +} + +void hipPerfMandelBrot::checkData(uint *ptr) { + totalIters = 0; + for (unsigned int i = 0; i < width_ * width_; i++) { + totalIters += ptr[i]; + } +} + +/** +* Test Description +* ------------------------ +* - Verify the warm-up kernel default stream executes serially. +* - verify by running all kernels - sync. +* - verify by running all kernels - async. +* Test source +* ------------------------ +* - perftests/compute/hipPerfMandelbrot.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfMandelbrot") { + hipPerfMandelBrot mandelbrotCompute; + int deviceId = 0; + mandelbrotCompute.open(deviceId); + #if HT_AMD + SECTION("warm-up kernel default stream executes serially") { + mandelbrotCompute.setNumStreams(1); + mandelbrotCompute.setNumKernels(1); + REQUIRE(true == mandelbrotCompute.run(100/*Random number*/, deviceId)); + } + #endif + SECTION("run all - sync") { + int i = 0; + do { + mandelbrotCompute.setNumStreams(1); + mandelbrotCompute.setNumKernels(1); + REQUIRE(true == mandelbrotCompute.run(i, deviceId)); + i++; + }while(i < 12); + mandelbrotCompute.printResults(); + } + + SECTION("run all - async") { + int i = 0; + do { + mandelbrotCompute.setNumStreams(2); + mandelbrotCompute.setNumKernels(2); + REQUIRE(true == mandelbrotCompute.run(i, deviceId)); + i++; + }while(i < 12); + mandelbrotCompute.printResults(); + } +} diff --git a/projects/hip-tests/catch/perftests/dispatch/CMakeLists.txt b/projects/hip-tests/catch/perftests/dispatch/CMakeLists.txt new file mode 100644 index 0000000000..68455e1ecb --- /dev/null +++ b/projects/hip-tests/catch/perftests/dispatch/CMakeLists.txt @@ -0,0 +1,28 @@ +# Copyright (c) 2023 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. + +# Common Tests - Test independent of all platforms +set(TEST_SRC + hipPerfDispatchSpeed.cc +) + +hip_add_exe_to_target(NAME dispatchTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME perf_test) diff --git a/projects/hip-tests/catch/perftests/dispatch/hipPerfDispatchSpeed.cc b/projects/hip-tests/catch/perftests/dispatch/hipPerfDispatchSpeed.cc new file mode 100644 index 0000000000..a55aa27a49 --- /dev/null +++ b/projects/hip-tests/catch/perftests/dispatch/hipPerfDispatchSpeed.cc @@ -0,0 +1,196 @@ +/* + Copyright (c) 2023 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. + */ + +/** +* @addtogroup hipPerfDispatchSpeed hipPerfDispatchSpeed +* @{ +* @ingroup perfDispatchTest +*/ + +#include +#include +#include + +// Quiet pesky warnings +#ifdef WIN_OS +#define SNPRINTF sprintf_s +#else +#define SNPRINTF snprintf +#endif +#define CHAR_BUF_SIZE 512 + +typedef struct { + unsigned int iterations; + int flushEvery; +} testStruct; + +testStruct testList[] = { + { 1, -1}, + { 1, -1}, + { 10, 1}, + { 10, -1}, + { 100, 1}, + { 100, 10}, + { 100, -1}, + { 1000, 1}, + { 1000, 10}, + { 1000, 100}, + { 1000, -1}, + { 10000, 1}, + { 10000, 10}, + { 10000, 100}, + { 10000, 1000}, + { 10000, -1}, + { 100000, 1}, + { 100000, 10}, + { 100000, 100}, + { 100000, 1000}, + { 100000, 10000}, + { 100000, -1}, +}; + +unsigned int mapTestList[] = {1, 1, 10, 100, 1000, 10000, 100000}; + +__global__ void _dispatchSpeed(float *outBuf) { + int i = (blockIdx.x * blockDim.x + threadIdx.x); + if (i < 0) + outBuf[i] = 0.0f; +}; + +/** +* Test Description +* ------------------------ +* - Verify the hipPerf Dispatch speed. +* Test source +* ------------------------ +* - perftests/compute/hipPerfMandelbrot.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfDispatchSpeed") { + int p_gpuDevice = 0; + int p_tests = -1; + hipError_t err = hipSuccess; + hipDeviceProp_t props = {0}; + HIP_CHECK(hipGetDeviceProperties(&props, p_gpuDevice)); + + unsigned int testListSize = sizeof(testList) / sizeof(testStruct); + int numTests = (p_tests == -1) ? (2*2*testListSize - 1) : p_tests; + int test = (p_tests == -1) ? 0 : p_tests; + + float* srcBuffer = NULL; + unsigned int bufSize_ = 64*sizeof(float); + err = hipMalloc(&srcBuffer, bufSize_); + REQUIRE(err == hipSuccess); + + for (; test <= numTests; test++) { + int openTest = test % testListSize; + bool sleep = false; + bool doWarmup = false; + + if ((test / testListSize) % 2) { + doWarmup = true; + } + if (test >= (testListSize * 2)) { + sleep = true; + } + int threads = (bufSize_ / sizeof(float)); + int threads_per_block = 64; + int blocks = (threads/threads_per_block) + (threads % threads_per_block); + hipEvent_t start, stop; + + // NULL stream check: + err = hipEventCreate(&start); + REQUIRE(err == hipSuccess); + err = hipEventCreate(&stop); + REQUIRE(err == hipSuccess); + + if (doWarmup) { + hipLaunchKernelGGL(_dispatchSpeed, dim3(blocks), dim3(threads_per_block), + 0, hipStream_t(0), srcBuffer); + err = hipDeviceSynchronize(); + REQUIRE(err == hipSuccess); + } + auto Start = std::chrono::high_resolution_clock::now(); + for (unsigned int i = 0; i < testList[openTest].iterations; i++) { + HIP_CHECK(hipEventRecord(start, NULL)); + hipLaunchKernelGGL(_dispatchSpeed, dim3(blocks), + dim3(threads_per_block), 0, hipStream_t(0), srcBuffer); + HIP_CHECK(hipEventRecord(stop, NULL)); + if ((testList[openTest].flushEvery > 0) && + (((i + 1) % testList[openTest].flushEvery) == 0)) { + if (sleep) { + err = hipDeviceSynchronize(); + REQUIRE(err == hipSuccess); + } else { + do { + err = hipEventQuery(stop); + } while (err == hipErrorNotReady); + } + } + } + if (sleep) { + err = hipDeviceSynchronize(); + REQUIRE(err == hipSuccess); + } else { + do { + err = hipEventQuery(stop); + } while (err == hipErrorNotReady); + } + auto Stop = std::chrono::high_resolution_clock::now(); + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(stop)); + double sec = std::chrono::duration(Stop - Start).count(); + + // microseconds per launch + double perf = (1000000.f*sec/testList[openTest].iterations); + const char *waitType; + const char *extraChar; + const char *n; + const char *warmup; + if (sleep) { + waitType = "sleep"; + extraChar = ""; + n = ""; + } else { + waitType = "spin"; + n = "n"; + extraChar = " "; + } + if (doWarmup) { + warmup = "warmup"; + } else { + warmup = ""; + } + char buf[256]; + if (testList[openTest].flushEvery > 0) { + SNPRINTF(buf, sizeof(buf), "HIPPerfDispatchSpeed[%3d] %7d dispatches %s%sing every %5d %6s (us/disp) %3f", + test, testList[openTest].iterations, + waitType, n, testList[openTest].flushEvery, warmup, (float)perf); + } else { + SNPRINTF(buf, sizeof(buf), "HIPPerfDispatchSpeed[%3d] %7d dispatches (%s%s) %6s (us/disp) %3f", + test, testList[openTest].iterations, + waitType, extraChar, warmup, (float)perf); + } + printf("%s\n", buf); + } + HIP_CHECK(hipFree(srcBuffer)); +} diff --git a/projects/hip-tests/catch/perftests/stream/CMakeLists.txt b/projects/hip-tests/catch/perftests/stream/CMakeLists.txt new file mode 100644 index 0000000000..1057f154f0 --- /dev/null +++ b/projects/hip-tests/catch/perftests/stream/CMakeLists.txt @@ -0,0 +1,30 @@ +# Copyright (c) 2023 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. + +# Common Tests - Test independent of all platforms +set(TEST_SRC + hipPerfDeviceConcurrency.cc + hipPerfStreamConcurrency.cc + hipPerfStreamCreateCopyDestroy.cc +) + +hip_add_exe_to_target(NAME streamTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME perf_test) diff --git a/projects/hip-tests/catch/perftests/stream/hipPerfDeviceConcurrency.cc b/projects/hip-tests/catch/perftests/stream/hipPerfDeviceConcurrency.cc new file mode 100644 index 0000000000..6fe9ad3ffc --- /dev/null +++ b/projects/hip-tests/catch/perftests/stream/hipPerfDeviceConcurrency.cc @@ -0,0 +1,269 @@ +/* + Copyright (c) 2023 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. + */ + +/** +* @addtogroup hipPerfDeviceConcurrency hipPerfDeviceConcurrency +* @{ +* @ingroup perfStreamTest +* `hipError_t hipStreamCreate(hipStream_t* stream)` - +* Create an asynchronous stream. +*/ + +#include + +typedef struct { + double x; + double y; + double width; +} coordRec; + +static coordRec coords[] = { + {0.0, 0.0, 0.00001}, // All black +}; + +static unsigned int numCoords = sizeof(coords) / sizeof(coordRec); + +__global__ void mandelbrot(uint *out, uint width, float xPos, + float yPos, float xStep, float yStep, uint maxIter) { + int tid = (blockIdx.x * blockDim.x + threadIdx.x); + int i = tid % width; + int j = tid / width; + float x0 = static_cast(xPos + xStep*i); + float y0 = static_cast(yPos + yStep*j); + + float x = x0; + float y = y0; + + uint iter = 0; + float tmp; + for (iter = 0; (x*x + y*y <= 4.0f) && (iter < maxIter); iter++) { + tmp = x; + x = fma(-y, y, fma(x, x, x0)); + y = fma(2.0f*tmp, y, y0); + } + out[tid] = iter; +}; + +class hipPerfDeviceConcurrency { + public: + hipPerfDeviceConcurrency(); + ~hipPerfDeviceConcurrency(); + + void setNumGpus(unsigned int num) { + numDevices = num; + } + unsigned int getNumGpus() { + return numDevices; + } + + void open(void); + void close(void); + bool run(unsigned int testCase, int numGpus); + + private: + void setData(void *ptr, unsigned int value); + void checkData(uint *ptr); + + unsigned int numDevices; + unsigned int width_; + unsigned int bufSize; + unsigned int coordIdx; + unsigned long long totalIters = 0; +}; + +hipPerfDeviceConcurrency::hipPerfDeviceConcurrency() {} + +hipPerfDeviceConcurrency::~hipPerfDeviceConcurrency() {} + +void hipPerfDeviceConcurrency::open(void) { + int nGpu = 0; + HIP_CHECK(hipGetDeviceCount(&nGpu)); + setNumGpus(nGpu); + if (nGpu < 1) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 1"); + } +} + +void hipPerfDeviceConcurrency::close() { +} + +bool hipPerfDeviceConcurrency::run(unsigned int testCase, int numGpus) { + static int deviceId; + uint * hPtr[numGpus]; + uint * dPtr[numGpus]; + hipStream_t streams[numGpus]; + int numCUs[numGpus]; + unsigned int maxIter[numGpus]; + unsigned long long expectedIters[numGpus]; + + int threads, threads_per_block, blocks; + float xStep, yStep, xPos, yPos; + + for (int i = 0; i < numGpus; i++) { + if (testCase != 0) { + deviceId = i; + } + + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIP_CHECK(hipGetDeviceProperties(&props, i)); + if (testCase != 0) { + std::cout << "info: running on bus " << "0x" << props.pciBusID + << " " << props.name << " with " << props.multiProcessorCount + << " CUs" << " and device ID: " << i << std::endl; + } + + numCUs[i] = props.multiProcessorCount; + int clkFrequency = 0; + HIP_CHECK(hipDeviceGetAttribute(&clkFrequency, + hipDeviceAttributeClockRate, i)); + clkFrequency =(unsigned int)clkFrequency/1000; + + // Maximum iteration count + // maxIter = 8388608 * (engine_clock / 1000).serial execution + maxIter[i] = (unsigned int)(((8388608 * ((float)clkFrequency / 1000)) + * numCUs[i]) / 128); + maxIter[i] = (maxIter[i] + 15) & ~15; + + // Width is divisible by 4 because the mandelbrot + // kernel processes 4 pixels at once. + width_ = 256; + bufSize = width_ * width_ * sizeof(uint); + // Create streams for concurrency + HIP_CHECK(hipStreamCreate(&streams[i])); + + // Allocate memory on the host and device + HIP_CHECK(hipHostMalloc(reinterpret_cast(&hPtr[i]), + bufSize, hipHostMallocDefault)); + setData(hPtr[i], 0xdeadbeef); + HIP_CHECK(hipMalloc(reinterpret_cast(&dPtr[i]), bufSize)) + + // Prepare kernel launch parameters + threads = (bufSize/sizeof(uint)); + threads_per_block = 64; + blocks = (threads/threads_per_block) + (threads % threads_per_block); + + coordIdx = testCase % numCoords; + xStep = static_cast(coords[coordIdx].width / static_cast(width_)); + yStep = static_cast(-coords[coordIdx].width / static_cast(width_)); + xPos = static_cast(coords[coordIdx].x - 0.5 * coords[coordIdx].width); + yPos = static_cast(coords[coordIdx].y + 0.5 * coords[coordIdx].width); + + // Copy memory from host to device + HIP_CHECK(hipMemcpy(dPtr[i], hPtr[i], bufSize, hipMemcpyHostToDevice)); + } + + // Time the kernel execution + auto all_start = std::chrono::steady_clock::now(); + for (int i = 0; i < numGpus; i++) { + if (testCase != 0) { + deviceId = i; + } + + HIP_CHECK(hipSetDevice(deviceId)); + hipLaunchKernelGGL(mandelbrot, dim3(blocks), dim3(threads_per_block), 0, + streams[i], dPtr[i], width_, xPos, yPos, xStep, + yStep, maxIter[i]); + } + for (int i = 0; i < numGpus; i++) { + HIP_CHECK(hipStreamSynchronize(0)); + } + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + + for(int i = 0; i < numGpus; i++) { + if(testCase != 0) { + deviceId = i; + } + HIP_CHECK(hipSetDevice(deviceId)); + + // Copy data back from device to the host + HIP_CHECK(hipMemcpy(hPtr[i], dPtr[i], bufSize, hipMemcpyDeviceToHost)); + checkData(hPtr[i]); + expectedIters[i] = width_ * width_ * (unsigned long long) maxIter[i]; + if (testCase != 0) { + checkData(hPtr[i]); + if (totalIters != expectedIters[i]) { + std::cout << "Incorrect iteration count detected" << std::endl; + } + } + + HIP_CHECK(hipStreamDestroy(streams[i])); + // Free host and device memory + HIP_CHECK(hipHostFree(hPtr[i])); + HIP_CHECK(hipFree(dPtr[i])); + } + + if (testCase != 0) { + std::cout << '\n' << "Measured time for kernel computation on " << numGpus + << " device (s): " << all_kernel_time.count() << " (s) " + << '\n' << std::endl; + } + + if (testCase == 0) { + deviceId++; + } + return true; +} + +void hipPerfDeviceConcurrency::setData(void *ptr, unsigned int value) { + unsigned int *ptr2 = (unsigned int *)ptr; + for (unsigned int i = 0; i < width_ * width_ ; i++) { + ptr2[i] = value; + } +} + +void hipPerfDeviceConcurrency::checkData(uint *ptr) { + totalIters = 0; + for (unsigned int i = 0; i < width_ * width_; i++) { + totalIters += ptr[i]; + } +} + +/** +* Test Description +* ------------------------ +* - Verify the different levels of device concurrency. +* Test source +* ------------------------ +* - perftests/stream/hipPerfDeviceConcurrency.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfDeviceConcurrency") { + hipPerfDeviceConcurrency deviceConcurrency; + deviceConcurrency.open(); + int nGpu = deviceConcurrency.getNumGpus(); + + // testCase = 0 refers to warmup kernel run + int testCase = 0; + for (int i = 0; i < nGpu; i++) { + // Warm-up kernel on all devices + REQUIRE(true == deviceConcurrency.run(testCase, 1)); + } + + // Time for kernel on 1 device + REQUIRE(true == deviceConcurrency.run(++testCase, 1)); + + // Time for kernel on all available devices + REQUIRE(true == deviceConcurrency.run(++testCase, nGpu)); +} diff --git a/projects/hip-tests/catch/perftests/stream/hipPerfStreamConcurrency.cc b/projects/hip-tests/catch/perftests/stream/hipPerfStreamConcurrency.cc new file mode 100644 index 0000000000..e82a72341f --- /dev/null +++ b/projects/hip-tests/catch/perftests/stream/hipPerfStreamConcurrency.cc @@ -0,0 +1,415 @@ +/* + Copyright (c) 2023 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. + */ + +/** +* @addtogroup hipPerfStreamConcurrency hipPerfStreamConcurrency +* @{ +* @ingroup perfComputeTest +* `hipError_t hipStreamCreate(hipStream_t* stream)` - +* Create an asynchronous stream. +*/ + +#include +#include + +#ifdef __HIP_PLATFORM_NVIDIA__ +inline __device__ float4 operator*(float s, float4 a) { + return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); +} +inline __device__ float4 operator*(float4 a, float4 b) { + return make_float4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); +} +inline __device__ float4 operator+(float4 a, float4 b) { + return make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); +} +inline __device__ float4 operator-(float4 a, float4 b) { + return make_float4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); +} +#endif + +typedef struct { + double x; + double y; + double width; +} coordRec; + +static coordRec coords[] = { + {0.0, 0.0, 0.00001}, // All black +}; + +static unsigned int numCoords = sizeof(coords) / sizeof(coordRec); + +__global__ static void mandelbrot(uint *out, uint width, float xPos, float yPos, + float xStep, float yStep, uint maxIter) { + int tid = (blockIdx.x * blockDim.x + threadIdx.x); + int i = tid % (width/4); + int j = tid / (width/4); + int4 veci = make_int4(4*i, 4*i+1, 4*i+2, 4*i+3); + int4 vecj = make_int4(j, j, j, j); + float4 x0; + x0.x = static_cast(xPos + xStep*veci.x); + x0.y = static_cast(xPos + xStep*veci.y); + x0.z = static_cast(xPos + xStep*veci.z); + x0.w = static_cast(xPos + xStep*veci.w); + float4 y0; + y0.x = static_cast(yPos + yStep*vecj.x); + y0.y = static_cast(yPos + yStep*vecj.y); + y0.z = static_cast(yPos + yStep*vecj.z); + y0.w = static_cast(yPos + yStep*vecj.w); + float4 x = x0; + float4 y = y0; + uint iter = 0; + float4 tmp; + int4 stay; + int4 ccount = make_int4(0, 0, 0, 0); + float4 savx = x; + float4 savy = y; + stay.x = (x.x*x.x+y.x*y.x) <= static_cast(4.0f); + stay.y = (x.y*x.y+y.y*y.y) <= static_cast(4.0f); + stay.z = (x.z*x.z+y.z*y.z) <= static_cast(4.0f); + stay.w = (x.w*x.w+y.w*y.w) <= static_cast(4.0f); + for (iter = 0; (stay.x | stay.y | stay.z | stay.w) && (iter < maxIter); + iter+=16) { + x = savx; + y = savy; + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + stay.x = (x.x*x.x+y.x*y.x) <= static_cast(4.0f); + stay.y = (x.y*x.y+y.y*y.y) <= static_cast(4.0f); + stay.z = (x.z*x.z+y.z*y.z) <= static_cast(4.0f); + stay.w = (x.w*x.w+y.w*y.w) <= static_cast(4.0f); + savx.x = static_cast(stay.x ? x.x : savx.x); + savx.y = static_cast(stay.y ? x.y : savx.y); + savx.z = static_cast(stay.z ? x.z : savx.z); + savx.w = static_cast(stay.w ? x.w : savx.w); + savy.x = static_cast(stay.x ? y.x : savy.x); + savy.y = static_cast(stay.y ? y.y : savy.y); + savy.z = static_cast(stay.z ? y.z : savy.z); + savy.w = static_cast(stay.w ? y.w : savy.w); + ccount.x -= stay.x*16; + ccount.y -= stay.y*16; + ccount.z -= stay.z*16; + ccount.w -= stay.w*16; + } + // Handle remainder + if (!(stay.x & stay.y & stay.z & stay.w)) { + iter = 16; + do { + x = savx; + y = savy; + stay.x = ((x.x*x.x+y.x*y.x) <= 4.0f) && (ccount.x < maxIter); + stay.y = ((x.y*x.y+y.y*y.y) <= 4.0f) && (ccount.y < maxIter); + stay.z = ((x.z*x.z+y.z*y.z) <= 4.0f) && (ccount.z < maxIter); + stay.w = ((x.w*x.w+y.w*y.w) <= 4.0f) && (ccount.w < maxIter); + tmp = x; + x = x*x + x0 - y*y; + y = 2.0f*tmp*y + y0; + ccount.x += stay.x; + ccount.y += stay.y; + ccount.z += stay.z; + ccount.w += stay.w; + iter--; + savx.x = (stay.x ? x.x : savx.x); + savx.y = (stay.y ? x.y : savx.y); + savx.z = (stay.z ? x.z : savx.z); + savx.w = (stay.w ? x.w : savx.w); + savy.x = (stay.x ? y.x : savy.x); + savy.y = (stay.y ? y.y : savy.y); + savy.z = (stay.z ? y.z : savy.z); + savy.w = (stay.w ? y.w : savy.w); + } while ((stay.x | stay.y | stay.z | stay.w) && iter); + } + uint4 *vecOut = reinterpret_cast(out); + vecOut[tid].x = (uint)(ccount.x); + vecOut[tid].y = (uint)(ccount.y); + vecOut[tid].z = (uint)(ccount.z); + vecOut[tid].w = (uint)(ccount.w); +} + +class hipPerfStreamConcurrency { + public: + hipPerfStreamConcurrency(); + ~hipPerfStreamConcurrency(); + + void setNumKernels(unsigned int num) { + numKernels = num; + } + void setNumStreams(unsigned int num) { + numStreams = num; + } + unsigned int getNumStreams() { + return numStreams; + } + + unsigned int getNumKernels() { + return numKernels; + } + + bool open(int deviceID); + bool run(unsigned int testCase, unsigned int deviceId); + void close(void); + + private: + void setData(void *ptr, unsigned int value); + void checkData(uint *ptr); + + unsigned int numKernels; + unsigned int numStreams; + + unsigned int width_; + unsigned int bufSize; + unsigned int maxIter; + unsigned int coordIdx; + unsigned long long totalIters; + int numCUs; +}; + +hipPerfStreamConcurrency::hipPerfStreamConcurrency() {} + +hipPerfStreamConcurrency::~hipPerfStreamConcurrency() {} + +bool hipPerfStreamConcurrency::open(int deviceId) { + int nGpu = 0; + HIP_CHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 1"); + } + + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + std::cout << "info: running on bus " << "0x" << props.pciBusID + << " " << props.name << " with " << props.multiProcessorCount << " CUs" + << " and device id: " << deviceId << std::endl; + numCUs = props.multiProcessorCount; + return true; +} + +void hipPerfStreamConcurrency::close() { +} + +bool hipPerfStreamConcurrency::run(unsigned int testCase, + unsigned int deviceId) { + int clkFrequency = 0; + unsigned int numStreams = getNumStreams(); + unsigned int numKernels = getNumKernels(); + + HIP_CHECK(hipDeviceGetAttribute(&clkFrequency, + hipDeviceAttributeClockRate, deviceId)); + + clkFrequency =(unsigned int)clkFrequency/1000; + + // Maximum iteration count + // maxIter = 8388608 * (engine_clock / 1000).serial execution + maxIter = (unsigned int)(((8388608 * (static_castclkFrequency / 1000)) + * numCUs) / 128); + maxIter = (maxIter + 15) & ~15; + hipStream_t streams[numStreams]; + uint * hPtr[numKernels]; + uint * dPtr[numKernels]; + + // Width is divisible by 4 because the mandelbrot kernel + // processes 4 pixels at once. + width_ = 256; + bufSize = width_ * sizeof(uint); + // Create streams for concurrency + for (uint i = 0; i < numStreams; i++) { + HIP_CHECK(hipStreamCreate(&streams[i])); + } + + // Allocate memory on the host and device + for (uint i = 0; i < numKernels; i++) { + HIP_CHECK(hipHostMalloc(reinterpret_cast(&hPtr[i]), + bufSize, hipHostMallocDefault)); + setData(hPtr[i], 0xdeadbeef); + HIP_CHECK(hipMalloc(reinterpret_cast(&dPtr[i]), bufSize)) + } + + // Prepare kernel launch parameters + int threads = (bufSize/sizeof(uint)); + int threads_per_block = 64; + int blocks = (threads/threads_per_block) + (threads % threads_per_block); + coordIdx = testCase % numCoords; + float xStep = static_cast(coords[coordIdx].width / static_cast(width_)); + float yStep = static_cast(-coords[coordIdx].width / static_cast(width_)); + float xPos = static_cast(coords[coordIdx].x - 0.5 * coords[coordIdx].width); + float yPos = static_cast(coords[coordIdx].y + 0.5 * coords[coordIdx].width); + + // Copy memory asynchronously and concurrently from host to device + for (uint i = 0; i < numKernels; i++) { + HIP_CHECK(hipMemcpyHtoDAsync(reinterpret_cast(dPtr[i]), + hPtr[i], bufSize, streams[i % numStreams])); + } + + // Synchronize to make sure all the copies are completed + for (uint i = 0; i < numStreams; i++) { + HIP_CHECK(hipStreamSynchronize(streams[i])); + } + // Warm-up kernel with lower iteration + if (testCase == 0) { + maxIter = 256; + } + // Time the kernel execution + auto all_start = std::chrono::steady_clock::now(); + + for (uint i = 0; i < numKernels; i++) { + hipLaunchKernelGGL(mandelbrot, dim3(blocks), dim3(threads_per_block), + 0, streams[i%numStreams], dPtr[i], width_, xPos, yPos, xStep, + yStep, maxIter); + } + + // Synchronize all the concurrent streans to have completed execution + for (uint i = 0; i < numStreams; i++) { + HIP_CHECK(hipStreamSynchronize(streams[i])); + } + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + + // Copy data back from device to the host + for (uint i = 0; i < numKernels; i++) { + HIP_CHECK(hipMemcpyDtoHAsync(hPtr[i], + reinterpret_cast(dPtr[i]), bufSize, + streams[i % numStreams])); + } + + if (testCase != 0) { + std::cout <<"Measured time for " << numKernels <<" kernels (s) on " + << numStreams <<" stream (s): " << all_kernel_time.count() << std::endl; + } + + unsigned long long expected = + (unsigned long long)width_ * (unsigned long long)maxIter; + + for (uint i = 0 ; i < numStreams; i++) { + HIP_CHECK(hipStreamDestroy(streams[i])); + } + + // Free host and device memory + for (uint i = 0; i < numKernels; i++) { + HIP_CHECK(hipHostFree(hPtr[i])); + HIP_CHECK(hipFree(dPtr[i])); + } + return true; +} + +void hipPerfStreamConcurrency::setData(void *ptr, unsigned int value) { + unsigned int *ptr2 = (unsigned int *)ptr; + for (unsigned int i = 0; i < width_ ; i++) { + ptr2[i] = value; + } +} + +void hipPerfStreamConcurrency::checkData(uint *ptr) { + totalIters = 0; + for (unsigned int i = 0; i < width_; i++) { + totalIters += ptr[i]; + } +} + +/** +* Test Description +* ------------------------ +* - Verify the different levels of stream concurrency. +* Test source +* ------------------------ +* - perftests/stream/hipPerfStreamConcurrency.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfStreamConcurrency") { + hipPerfStreamConcurrency streamConcurrency; + int deviceId = 0; + REQUIRE(true == streamConcurrency.open(deviceId)); + + for (unsigned int testCase = 0; testCase < 5; testCase++) { + switch (testCase) { + case 0: + // Warm-up kernel + streamConcurrency.setNumStreams(1); + streamConcurrency.setNumKernels(1); + break; + + case 1: + // default stream executes serially + streamConcurrency.setNumStreams(1); + streamConcurrency.setNumKernels(1); + break; + + case 2: + // 2-way concurrency + streamConcurrency.setNumStreams(2); + streamConcurrency.setNumKernels(2); + break; + + case 3: + // 4-way concurrency + streamConcurrency.setNumStreams(4); + streamConcurrency.setNumKernels(4); + break; + + case 4: + streamConcurrency.setNumStreams(2); + streamConcurrency.setNumKernels(4); + break; + + case 5: + break; + + default: + break; + } + REQUIRE(true == streamConcurrency.run(testCase, deviceId)); + } +} diff --git a/projects/hip-tests/catch/perftests/stream/hipPerfStreamCreateCopyDestroy.cc b/projects/hip-tests/catch/perftests/stream/hipPerfStreamCreateCopyDestroy.cc new file mode 100644 index 0000000000..2589403ef8 --- /dev/null +++ b/projects/hip-tests/catch/perftests/stream/hipPerfStreamCreateCopyDestroy.cc @@ -0,0 +1,146 @@ +/* + Copyright (c) 2023 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. + */ + +/** +* @addtogroup hipPerfStreamCreateCopyDestroy hipPerfStreamCreateCopyDestroy +* @{ +* @ingroup perfStreamTest +* `hipError_t hipStreamCreate(hipStream_t* stream)` - +* Create an asynchronous stream. +*/ + +#include +#include +#include + +using namespace std; + +#define BufSize 0x1000 +#define Iterations 0x100 +#define TotalStreams 4 +#define TotalBufs 4 + + +class hipPerfStreamCreateCopyDestroy { + private: + unsigned int numBuffers_; + unsigned int numStreams_; + const size_t totalStreams_[TotalStreams]; + const size_t totalBuffers_[TotalBufs]; + public: + hipPerfStreamCreateCopyDestroy() : numBuffers_(0), numStreams_(0), + totalStreams_{1, 2, 4, 8}, + totalBuffers_{1, 100, 1000, 5000} {}; + ~hipPerfStreamCreateCopyDestroy() {}; + bool open(int deviceID); + bool run(unsigned int testNumber); +}; + +bool hipPerfStreamCreateCopyDestroy::open(int deviceId) { + int nGpu = 0; + HIP_CHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 1"); + } + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + std::cout << "info: running on bus " << "0x" << props.pciBusID + << " " << props.name << " with " << props.multiProcessorCount << " CUs" + << " and device id: " << deviceId << std::endl; + return true; +} + +bool hipPerfStreamCreateCopyDestroy::run(unsigned int testNumber) { + numStreams_ = totalStreams_[testNumber % TotalStreams]; + size_t iter = Iterations / (numStreams_ * (static_cast(1) + << (testNumber / TotalBufs + 1))); + hipStream_t streams[numStreams_]; + + numBuffers_ = totalBuffers_[testNumber / TotalBufs]; + float* dSrc[numBuffers_]; + size_t nBytes = BufSize * sizeof(float); + + for (size_t b = 0; b < numBuffers_; ++b) { + HIP_CHECK(hipMalloc(&dSrc[b], nBytes)); + } + + float* hSrc; + hSrc = new float[nBytes]; + HIP_CHECK(hSrc == 0 ? hipErrorOutOfMemory : hipSuccess); + for (size_t i = 0; i < BufSize; i++) { + hSrc[i] = 1.618f + i; + } + + auto start = std::chrono::steady_clock::now(); + + for (size_t i = 0; i < iter; ++i) { + for (size_t s = 0; s < numStreams_; ++s) { + HIP_CHECK(hipStreamCreate(&streams[s])); + } + + for (size_t s = 0; s < numStreams_; ++s) { + for (size_t b = 0; b < numBuffers_; ++b) { + HIP_CHECK(hipMemcpyWithStream(dSrc[b], hSrc, nBytes, + hipMemcpyHostToDevice, streams[s])); + } + } + + for (size_t s = 0; s < numStreams_; ++s) { + HIP_CHECK(hipStreamDestroy(streams[s])); + } + } + + auto end = std::chrono::steady_clock::now(); + std::chrono::duration diff = end - start; + + auto time = static_cast(diff.count() * 1000 / (iter * numStreams_)); + + cout << "Create+Copy+Destroy time for " << numStreams_ << " streams and " + << setw(4) << numBuffers_ << " buffers " << " and " << setw(4) + << iter << " iterations " << time << " (ms) " << endl; + + delete [] hSrc; + for (size_t b = 0; b < numBuffers_; ++b) { + HIP_CHECK(hipFree(dSrc[b])); + } + return true; +} + +/** +* Test Description +* ------------------------ +* - Verify the Create+Copy+Destroy time for different stream. +* Test source +* ------------------------ +* - perftests/stream/hipPerfDeviceConcurrency.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfStreamCreateCopyDestroy") { + hipPerfStreamCreateCopyDestroy streamCCD; + int deviceId = 0; + REQUIRE(true == streamCCD.open(deviceId)); + + for (auto testCase = 0; testCase < TotalStreams * TotalBufs; testCase++) { + REQUIRE(true == streamCCD.run(testCase)); + } +}