SWDEV-403471 - [catch2][dtest] Conversion of preftests to catch2
Change-Id: I68cb780a71a6094dca86718e7d427806d3a0e67d
[ROCm/hip-tests commit: b0c4a4f70f]
This commit is contained in:
mbhiutra
Rakesh Roy
@@ -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)
|
||||
|
||||
@@ -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)
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <vector>
|
||||
|
||||
#define DOT_DIM 256
|
||||
|
||||
using namespace std;
|
||||
|
||||
template <unsigned int BLOCKSIZE>
|
||||
__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 <unsigned int BLOCKSIZE>
|
||||
__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 <unsigned int BLOCKSIZE>
|
||||
__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<DOT_DIM>, blocks, threadsPerBlock, 0, 0, n, x, y,
|
||||
workspace);
|
||||
} else {
|
||||
hipLaunchKernelGGL(vectors_equal<DOT_DIM>, blocks, threadsPerBlock, 0, 0, n, x, workspace);
|
||||
}
|
||||
|
||||
// Part 2 of dot product computation
|
||||
hipLaunchKernelGGL(dot_reduction<DOT_DIM>, 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<int> 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<double> hx(size);
|
||||
vector<double> 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_cast<double>rand() / static_cast<double>RAND_MAX - 1.0;
|
||||
hy[i] = 2.0 * static_cast<double>rand() / static_cast<double>RAND_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<void**>(&dx), sizeof(double) * size));
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&dy), sizeof(double) * size));
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&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 <x,y>
|
||||
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<double> 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] <x,y> " << 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 <x,x>
|
||||
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] <x,y> " << 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));
|
||||
}
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_vector_types.h>
|
||||
#include <hip/math_functions.h>
|
||||
#include <vector>
|
||||
#include <string>
|
||||
#include <map>
|
||||
|
||||
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 <typename T>
|
||||
__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<float>(xPos + xStep*i);
|
||||
float y0 = static_cast<float>(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 <typename T>
|
||||
__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<float>(xPos + xStep*static_cast<float>(i));
|
||||
float y0 = static_cast<float>(yPos + yStep*static_cast<float>(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 <typename T>
|
||||
__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<double>(xPos + xStep*i);
|
||||
double y0 = static_cast<double>(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 <typename T>
|
||||
__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<double>(xPos + xStep*static_cast<double>(i));
|
||||
double y0 = static_cast<double>(yPos + yStep*static_cast<double>(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<std::string, std::vector<double>> 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<std::string, std::vector<double>>:: 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<float>, 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<float>, 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<double>, 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<double>, 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<void **>(&hPtr[i]), bufSize, hipHostMallocDefault));
|
||||
setData(hPtr[i], 0xdeadbeef);
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<uint **>(&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<float>(coords[coordIdx].width / (double)width_);
|
||||
float yStep = static_cast<float>(-coords[coordIdx].width / (double)width_);
|
||||
float xPos = static_cast<float>(coords[coordIdx].x - 0.5 * coords[coordIdx].width);
|
||||
float yPos = static_cast<float>(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<float>(coords[coordIdx].width / static_cast<double>(width_));
|
||||
float yStep = static_cast<float>(-coords[coordIdx].width / static_cast<double>(width_));
|
||||
float xPos = static_cast<float>(coords[coordIdx].x - 0.5 * coords[coordIdx].width);
|
||||
float yPos = static_cast<float>(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<double> all_kernel_time = all_end - all_start;
|
||||
totalTime += all_kernel_time.count();
|
||||
} else {
|
||||
double xStep = coords[coordIdx].width / static_cast<double>(width_);
|
||||
double yStep = -coords[coordIdx].width / static_cast<double>(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<double> 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<double>(totalIters*numKernels) * 7 * static_cast<double>(1e-09)) /
|
||||
(totalTime / (double)numLoops);
|
||||
|
||||
|
||||
std::vector<std::string> 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();
|
||||
}
|
||||
}
|
||||
@@ -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)
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <string.h>
|
||||
#include <complex>
|
||||
|
||||
// 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<double, std::milli>(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));
|
||||
}
|
||||
@@ -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)
|
||||
@@ -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 <hip_test_common.hh>
|
||||
|
||||
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<float>(xPos + xStep*i);
|
||||
float y0 = static_cast<float>(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<void **>(&hPtr[i]),
|
||||
bufSize, hipHostMallocDefault));
|
||||
setData(hPtr[i], 0xdeadbeef);
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<uint **>(&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<float>(coords[coordIdx].width / static_cast<double>(width_));
|
||||
yStep = static_cast<float>(-coords[coordIdx].width / static_cast<double>(width_));
|
||||
xPos = static_cast<float>(coords[coordIdx].x - 0.5 * coords[coordIdx].width);
|
||||
yPos = static_cast<float>(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<double> 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));
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_vector_types.h>
|
||||
|
||||
#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<float>(xPos + xStep*veci.x);
|
||||
x0.y = static_cast<float>(xPos + xStep*veci.y);
|
||||
x0.z = static_cast<float>(xPos + xStep*veci.z);
|
||||
x0.w = static_cast<float>(xPos + xStep*veci.w);
|
||||
float4 y0;
|
||||
y0.x = static_cast<float>(yPos + yStep*vecj.x);
|
||||
y0.y = static_cast<float>(yPos + yStep*vecj.y);
|
||||
y0.z = static_cast<float>(yPos + yStep*vecj.z);
|
||||
y0.w = static_cast<float>(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<float>(4.0f);
|
||||
stay.y = (x.y*x.y+y.y*y.y) <= static_cast<float>(4.0f);
|
||||
stay.z = (x.z*x.z+y.z*y.z) <= static_cast<float>(4.0f);
|
||||
stay.w = (x.w*x.w+y.w*y.w) <= static_cast<float>(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<float>(4.0f);
|
||||
stay.y = (x.y*x.y+y.y*y.y) <= static_cast<float>(4.0f);
|
||||
stay.z = (x.z*x.z+y.z*y.z) <= static_cast<float>(4.0f);
|
||||
stay.w = (x.w*x.w+y.w*y.w) <= static_cast<float>(4.0f);
|
||||
savx.x = static_cast<bool>(stay.x ? x.x : savx.x);
|
||||
savx.y = static_cast<bool>(stay.y ? x.y : savx.y);
|
||||
savx.z = static_cast<bool>(stay.z ? x.z : savx.z);
|
||||
savx.w = static_cast<bool>(stay.w ? x.w : savx.w);
|
||||
savy.x = static_cast<bool>(stay.x ? y.x : savy.x);
|
||||
savy.y = static_cast<bool>(stay.y ? y.y : savy.y);
|
||||
savy.z = static_cast<bool>(stay.z ? y.z : savy.z);
|
||||
savy.w = static_cast<bool>(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<uint4 *>(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_cast<float>clkFrequency / 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<void **>(&hPtr[i]),
|
||||
bufSize, hipHostMallocDefault));
|
||||
setData(hPtr[i], 0xdeadbeef);
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<void **>(&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<float>(coords[coordIdx].width / static_cast<double>(width_));
|
||||
float yStep = static_cast<float>(-coords[coordIdx].width / static_cast<double>(width_));
|
||||
float xPos = static_cast<float>(coords[coordIdx].x - 0.5 * coords[coordIdx].width);
|
||||
float yPos = static_cast<float>(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<hipDeviceptr_t>(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<double> 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<hipDeviceptr_t>(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));
|
||||
}
|
||||
}
|
||||
@@ -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 <hip_test_kernels.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
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<size_t>(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<double> diff = end - start;
|
||||
|
||||
auto time = static_cast<float>(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));
|
||||
}
|
||||
}
|
||||
Reference in New Issue
Block a user