SWDEV-546217 Complete hip-test Port to Catch2 Framework [Stream and Compute Folder] (#559)
* SWDEV-546498 hipPerfDeviceConcurrency * SWDEV-546500 hipPerfStreamConcurrency * SWDEV-546502 hipPerfStreamCreateCopyDestroy.c * SWDEV-546479 hipPerfDotProduct * SWDEV-546482 hipPerfMandelbrot
This commit is contained in:
Luo, Phoebe
GitHub
@@ -18,10 +18,10 @@
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipPerfDotProduct hipPerfDotProduct
|
||||
* @{
|
||||
* @ingroup perfComputeTest
|
||||
*/
|
||||
* @addtogroup hipPerfDotProduct hipPerfDotProduct
|
||||
* @{
|
||||
* @ingroup perfComputeTest
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <vector>
|
||||
@@ -31,11 +31,9 @@
|
||||
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) {
|
||||
__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;
|
||||
@@ -93,9 +91,8 @@ __global__ void vectors_not_equal(int n,
|
||||
}
|
||||
|
||||
template <unsigned int BLOCKSIZE>
|
||||
__launch_bounds__(BLOCKSIZE)
|
||||
__global__ void vectors_equal(int n, const double* __restrict__ x,
|
||||
double* __restrict__ workspace) {
|
||||
__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;
|
||||
@@ -129,7 +126,7 @@ __global__ void vectors_equal(int n, const double* __restrict__ x,
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x < 8) {
|
||||
sdata[threadIdx.x] += sdata[threadIdx.x + 8];
|
||||
sdata[threadIdx.x] += sdata[threadIdx.x + 8];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
@@ -149,12 +146,11 @@ __global__ void vectors_equal(int n, const double* __restrict__ x,
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
workspace[blockIdx.x] = sdata[0];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
template <unsigned int BLOCKSIZE>
|
||||
__launch_bounds__(BLOCKSIZE)
|
||||
__global__ void dot_reduction(double* __restrict__ workspace) {
|
||||
__launch_bounds__(BLOCKSIZE) __global__ void dot_reduction(double* __restrict__ workspace) {
|
||||
__shared__ double sdata[BLOCKSIZE];
|
||||
|
||||
sdata[threadIdx.x] = workspace[threadIdx.x];
|
||||
@@ -187,7 +183,8 @@ __global__ void dot_reduction(double* __restrict__ workspace) {
|
||||
|
||||
if (threadIdx.x < 4) {
|
||||
sdata[threadIdx.x] += sdata[threadIdx.x + 4];
|
||||
} __syncthreads();
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x < 2) {
|
||||
sdata[threadIdx.x] += sdata[threadIdx.x + 2];
|
||||
@@ -203,8 +200,7 @@ __global__ void dot_reduction(double* __restrict__ workspace) {
|
||||
}
|
||||
}
|
||||
|
||||
void computeDotProduct(int n, const double* x, const double* y, double& result,
|
||||
double* workspace) {
|
||||
void computeDotProduct(int n, const double* x, const double* y, double& result, double* workspace) {
|
||||
dim3 blocks(DOT_DIM);
|
||||
dim3 threadsPerBlock(DOT_DIM);
|
||||
|
||||
@@ -225,16 +221,16 @@ void computeDotProduct(int n, const double* x, const double* y, double& result,
|
||||
}
|
||||
|
||||
/**
|
||||
* 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 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;
|
||||
@@ -252,120 +248,120 @@ TEST_CASE("Perf_hipPerfDotProduct") {
|
||||
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 0:
|
||||
nx = vectorSize[0];
|
||||
ny = vectorSize[0];
|
||||
nz = vectorSize[0];
|
||||
break;
|
||||
|
||||
case 1:
|
||||
nx = vectorSize[1];
|
||||
ny = vectorSize[1];
|
||||
nz = vectorSize[1];
|
||||
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;
|
||||
case 2:
|
||||
nx = vectorSize[0];
|
||||
ny = vectorSize[1];
|
||||
nz = vectorSize[2];
|
||||
break;
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
int trials = 200;
|
||||
int size = nx * ny * nz;
|
||||
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;
|
||||
vector<double> hx(size);
|
||||
vector<double> hy(size);
|
||||
double hresult_xy = 0.0;
|
||||
double hresult_xx = 0.0;
|
||||
|
||||
srand(time(NULL));
|
||||
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;
|
||||
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];
|
||||
}
|
||||
hresult_xy += hx[i] * hy[i];
|
||||
hresult_xx += hx[i] * hx[i];
|
||||
}
|
||||
|
||||
double* dx;
|
||||
double* dy;
|
||||
double* workspace;
|
||||
double dresult;
|
||||
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(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));
|
||||
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) {
|
||||
// Warm up
|
||||
computeDotProduct(size, dx, dy, dresult, workspace);
|
||||
computeDotProduct(size, dx, dy, dresult, workspace);
|
||||
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();
|
||||
// Timed run for <x,y>
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
auto all_start = std::chrono::steady_clock::now();
|
||||
|
||||
time /= trials;
|
||||
for (int i = 0; i < trials; ++i) {
|
||||
computeDotProduct(size, dx, dy, dresult, workspace);
|
||||
}
|
||||
|
||||
double bw = sizeof(double) * size * 2.0 / 1e9;
|
||||
double gf = 2.0 * size / 1e9;
|
||||
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();
|
||||
|
||||
cout << "\nVector Size: " << size << "\n[ddot] <x,y> " << time << "msec ;" << bw/ (time / 1e3) << " GByte/s ;"
|
||||
<< gf/(time / 1e3) << " GFlop/s" << endl;
|
||||
time /= trials;
|
||||
|
||||
// Verify the device kernel results comparing it with the host results
|
||||
REQUIRE(std::abs(dresult - hresult_xy) < std::max(dresult * 1e-10, 1e-8));
|
||||
double bw = sizeof(double) * size * 2.0 / 1e9;
|
||||
double gf = 2.0 * size / 1e9;
|
||||
|
||||
// Warm up
|
||||
computeDotProduct(size, dx, dx, dresult, workspace);
|
||||
computeDotProduct(size, dx, dx, dresult, workspace);
|
||||
computeDotProduct(size, dx, dx, dresult, workspace);
|
||||
CONSOLE_PRINT("\nVector Size: %d\n[ddot] <x,y> %.6f msec ; %.6f GByte/s ; %.6f GFlop/s", size,
|
||||
time, bw / (time / 1e3), gf / (time / 1e3));
|
||||
|
||||
// Timed run for <x,x>
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
all_start = std::chrono::steady_clock::now();
|
||||
// Verify the device kernel results comparing it with the host results
|
||||
REQUIRE(std::abs(dresult - hresult_xy) < std::max(dresult * 1e-10, 1e-8));
|
||||
|
||||
for (int i = 0; i < trials; ++i) {
|
||||
// Warm up
|
||||
computeDotProduct(size, dx, dx, dresult, workspace);
|
||||
computeDotProduct(size, dx, dx, dresult, workspace);
|
||||
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();
|
||||
// Timed run for <x,x>
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
all_start = std::chrono::steady_clock::now();
|
||||
|
||||
time /= trials;
|
||||
bw = sizeof(double) * size / 1e9;
|
||||
for (int i = 0; i < trials; ++i) {
|
||||
computeDotProduct(size, dx, dx, dresult, workspace);
|
||||
}
|
||||
|
||||
cout << "[ddot] <x,y> " << time << "msec ;" << bw/ (time / 1e3) << " GByte/s ;"
|
||||
<< gf/(time / 1e3) << " GFlop/s" << endl;
|
||||
all_end = std::chrono::steady_clock::now();
|
||||
all_kernel_time = all_end - all_start;
|
||||
time = all_kernel_time.count();
|
||||
|
||||
// Verify the device kernel results comparing it with the host results
|
||||
REQUIRE(abs(dresult - hresult_xx) < max(dresult * 1e-10, 1e-8));
|
||||
time /= trials;
|
||||
bw = sizeof(double) * size / 1e9;
|
||||
|
||||
HIP_CHECK(hipFree(dx));
|
||||
HIP_CHECK(hipFree(dy));
|
||||
HIP_CHECK(hipFree(workspace));
|
||||
CONSOLE_PRINT("[ddot] <x,y> %.6f msec ; %.6f GByte/s ; %.6f GFlop/s", time, bw / (time / 1e3),
|
||||
gf / (time / 1e3));
|
||||
|
||||
// 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));
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* End doxygen group perfComputeTest.
|
||||
* @}
|
||||
*/
|
||||
* End doxygen group perfComputeTest.
|
||||
* @}
|
||||
*/
|
||||
|
||||
@@ -18,10 +18,10 @@
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipPerfMandelbrot hipPerfMandelbrot
|
||||
* @{
|
||||
* @ingroup perfComputeTest
|
||||
*/
|
||||
* @addtogroup hipPerfMandelbrot hipPerfMandelbrot
|
||||
* @{
|
||||
* @ingroup perfComputeTest
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_vector_types.h>
|
||||
@@ -45,36 +45,35 @@ coordRec coords[] = {
|
||||
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) {
|
||||
__global__ void float_mad_kernel(uint* out, uint width, T xPos, T yPos, T xStep, T 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 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++) {
|
||||
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);
|
||||
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) {
|
||||
|
||||
__global__ void float_mandel_unroll_kernel(uint* out, uint width, T xPos, T yPos, T xStep, T 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*static_cast<float>(i));
|
||||
float y0 = static_cast<float>(yPos + yStep*static_cast<float>(j));
|
||||
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;
|
||||
@@ -84,72 +83,71 @@ __global__ void float_mandel_unroll_kernel(uint *out, uint width, T xPos,
|
||||
float tmp;
|
||||
int stay;
|
||||
int ccount = 0;
|
||||
stay = (x*x+y*y) <= 4.0;
|
||||
stay = (x * x + y * y) <= 4.0;
|
||||
float savx = x;
|
||||
float savy = y;
|
||||
#ifdef FAST
|
||||
for (iter = 0; (iter < maxIter); iter+=16) {
|
||||
for (iter = 0; (iter < maxIter); iter += 16) {
|
||||
#else
|
||||
for (iter = 0; stay && (iter < maxIter); iter+=16) {
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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;
|
||||
stay = (x * x + y * y) <= 4.0;
|
||||
savx = (stay ? x : savx);
|
||||
savy = (stay ? y : savy);
|
||||
ccount += stay*16;
|
||||
ccount += stay * 16;
|
||||
#ifdef FAST
|
||||
if (!stay)
|
||||
break;
|
||||
if (!stay) break;
|
||||
#endif
|
||||
}
|
||||
// Handle remainder
|
||||
@@ -158,10 +156,10 @@ __global__ void float_mandel_unroll_kernel(uint *out, uint width, T xPos,
|
||||
do {
|
||||
x = savx;
|
||||
y = savy;
|
||||
stay = ((x*x+y*y) <= 4.0) && (ccount < maxIter);
|
||||
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);
|
||||
x = fma(-y, y, fma(x, x, x0));
|
||||
y = fma(2.0f * tmp, y, y0);
|
||||
ccount += stay;
|
||||
iter--;
|
||||
savx = (stay ? x : savx);
|
||||
@@ -172,36 +170,36 @@ __global__ void float_mandel_unroll_kernel(uint *out, uint width, T xPos,
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void double_mad_kernel(uint *out, uint width, T xPos, T yPos, T xStep, T yStep,
|
||||
uint maxIter) {
|
||||
__global__ void double_mad_kernel(uint* out, uint width, T xPos, T yPos, T xStep, T yStep,
|
||||
uint maxIter) {
|
||||
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 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++) {
|
||||
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);
|
||||
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) {
|
||||
__global__ void double_mandel_unroll_kernel(uint* out, uint width, T xPos, T yPos, T xStep, T yStep,
|
||||
uint maxIter) {
|
||||
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 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;
|
||||
@@ -211,13 +209,13 @@ __global__ void double_mandel_unroll_kernel(uint *out, uint width, T xPos,
|
||||
double tmp;
|
||||
int stay;
|
||||
int ccount = 0;
|
||||
stay = (x*x+y*y) <= 4.0;
|
||||
stay = (x * x + y * y) <= 4.0;
|
||||
double savx = x;
|
||||
double savy = y;
|
||||
#ifdef FAST
|
||||
for (iter = 0; (iter < maxIter); iter+=16)
|
||||
for (iter = 0; (iter < maxIter); iter += 16)
|
||||
#else
|
||||
for (iter = 0; stay && (iter < maxIter); iter+=16)
|
||||
for (iter = 0; stay && (iter < maxIter); iter += 16)
|
||||
#endif
|
||||
{
|
||||
x = savx;
|
||||
@@ -225,141 +223,131 @@ __global__ void double_mandel_unroll_kernel(uint *out, uint width, T xPos,
|
||||
|
||||
// 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);
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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;
|
||||
stay = (x * x + y * y) <= 4.0;
|
||||
savx = (stay ? x : savx);
|
||||
savy = (stay ? y : savy);
|
||||
ccount += stay*16;
|
||||
ccount += stay * 16;
|
||||
#ifdef FAST
|
||||
if (!stay)
|
||||
break;
|
||||
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;
|
||||
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;
|
||||
};
|
||||
|
||||
// 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};
|
||||
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;
|
||||
}
|
||||
void setNumKernels(unsigned int num) { numKernels = num; }
|
||||
|
||||
unsigned int getNumKernels() {
|
||||
return numKernels;
|
||||
}
|
||||
unsigned int getNumKernels() { return numKernels; }
|
||||
|
||||
void setNumStreams(unsigned int num) {
|
||||
numStreams = num;
|
||||
}
|
||||
unsigned int getNumStreams() {
|
||||
return numStreams;
|
||||
}
|
||||
void setNumStreams(unsigned int num) { numStreams = num; }
|
||||
unsigned int getNumStreams() { return numStreams; }
|
||||
|
||||
void open(int deviceID);
|
||||
bool run(unsigned int testCase);
|
||||
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);
|
||||
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_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 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_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);
|
||||
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);
|
||||
void setData(void* ptr, unsigned int value);
|
||||
void checkData(uint* ptr);
|
||||
|
||||
unsigned int numKernels;
|
||||
unsigned int numStreams;
|
||||
@@ -387,9 +375,9 @@ void hipPerfMandelBrot::open(int deviceId) {
|
||||
HIP_CHECK(hipSetDevice(deviceId));
|
||||
hipDeviceProp_t props;
|
||||
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;
|
||||
|
||||
CONSOLE_PRINT("info: running on bus 0x%x %s with %d CUs and device id: %d\n", props.pciBusID,
|
||||
props.name, props.multiProcessorCount, deviceId);
|
||||
|
||||
numCUs = props.multiProcessorCount;
|
||||
}
|
||||
@@ -397,52 +385,52 @@ void hipPerfMandelBrot::open(int deviceId) {
|
||||
void hipPerfMandelBrot::printResults() {
|
||||
int numStreams = getNumStreams();
|
||||
|
||||
std::cout << "\n" <<"Measured perf for kernels in GFLOPS on "
|
||||
<< numStreams << " streams (s)" << std::endl;
|
||||
CONSOLE_PRINT("Measured perf for kernels in GFLOPS on %d streams (s)", numStreams);
|
||||
|
||||
std::map<std::string, std::vector<double>>:: iterator itr;
|
||||
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 << " ";
|
||||
}
|
||||
}
|
||||
CONSOLE_PRINT("\n%s ", itr->first.c_str());
|
||||
for (auto i : results[itr->first]) {
|
||||
CONSOLE_PRINT("%10f ", i);
|
||||
}
|
||||
}
|
||||
results.clear();
|
||||
std::cout << std::endl;
|
||||
CONSOLE_PRINT("\n");
|
||||
}
|
||||
|
||||
// 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) {
|
||||
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);
|
||||
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) {
|
||||
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);
|
||||
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) {
|
||||
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);
|
||||
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) {
|
||||
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);
|
||||
streams[kernelCnt % streamCnt], out, width, xPos, yPos, xStep, yStep, maxIter);
|
||||
}
|
||||
|
||||
bool hipPerfMandelBrot::run(unsigned int testCase) {
|
||||
@@ -450,18 +438,18 @@ bool hipPerfMandelBrot::run(unsigned int testCase) {
|
||||
coordIdx = testCase % numCoords;
|
||||
|
||||
funPtr p[] = {&hipPerfMandelBrot::float_mad, &hipPerfMandelBrot::float_mandel_unroll,
|
||||
&hipPerfMandelBrot::double_mad, &hipPerfMandelBrot::double_mandel_unroll};
|
||||
&hipPerfMandelBrot::double_mad, &hipPerfMandelBrot::double_mandel_unroll};
|
||||
|
||||
// Maximum iteration count
|
||||
maxIter = 32768;
|
||||
|
||||
uint ** hPtr = new uint *[numKernels];
|
||||
uint ** dPtr = new uint *[numKernels];
|
||||
uint** hPtr = new uint*[numKernels];
|
||||
uint** dPtr = new uint*[numKernels];
|
||||
|
||||
// Width is divisible by 4 because the mandelbrot kernel processes 4 pixels at once.
|
||||
width_ = 256;
|
||||
|
||||
bufSize = width_ * width_ * sizeof(uint);
|
||||
bufSize = width_ * width_ * sizeof(uint);
|
||||
|
||||
// Create streams for concurrency
|
||||
for (uint i = 0; i < numStreams; i++) {
|
||||
@@ -470,15 +458,15 @@ bool hipPerfMandelBrot::run(unsigned int testCase) {
|
||||
|
||||
// Allocate memory on the host and device
|
||||
for (uint i = 0; i < numKernels; i++) {
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void **>(&hPtr[i]), bufSize, hipHostMallocDefault));
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&hPtr[i]), bufSize, hipHostMallocDefault));
|
||||
setData(hPtr[i], 0xdeadbeef);
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<uint **>(&dPtr[i]), bufSize))
|
||||
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);
|
||||
int threads = (bufSize / sizeof(uint));
|
||||
int threads_per_block = 64;
|
||||
int blocks = (threads / threads_per_block) + (threads % threads_per_block);
|
||||
|
||||
// Copy memory asynchronously and concurrently from host to device
|
||||
for (uint i = 0; i < numKernels; i++) {
|
||||
@@ -489,90 +477,88 @@ bool hipPerfMandelBrot::run(unsigned int testCase) {
|
||||
HIP_CHECK(hipStreamSynchronize(0));
|
||||
|
||||
int kernelIdx;
|
||||
if(testCase == 0 || testCase == 5 || testCase == 10) {
|
||||
if (testCase == 0 || testCase == 5 || testCase == 10) {
|
||||
kernelIdx = 0;
|
||||
} else if(testCase == 1 || testCase == 6 || testCase == 11) {
|
||||
} else if (testCase == 1 || testCase == 6 || testCase == 11) {
|
||||
kernelIdx = 1;
|
||||
} else if(testCase == 2 || testCase == 7 || testCase == 12) {
|
||||
} else if (testCase == 2 || testCase == 7 || testCase == 12) {
|
||||
kernelIdx = 2;
|
||||
} else if(testCase == 3 || testCase == 8 || testCase == 13){
|
||||
} 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);
|
||||
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();
|
||||
// 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);
|
||||
}
|
||||
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));
|
||||
// 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;
|
||||
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));
|
||||
// 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();
|
||||
}
|
||||
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++;
|
||||
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. ";
|
||||
}
|
||||
if (j == 30) {
|
||||
CONSOLE_PRINT("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);
|
||||
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"};
|
||||
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);
|
||||
}
|
||||
results[kernelName[testCase % 4]].push_back(perf);
|
||||
}
|
||||
|
||||
for(uint i = 0 ; i < numStreams; i++) {
|
||||
for (uint i = 0; i < numStreams; i++) {
|
||||
HIP_CHECK(hipStreamDestroy(streams[i]));
|
||||
}
|
||||
|
||||
@@ -581,19 +567,19 @@ bool hipPerfMandelBrot::run(unsigned int testCase) {
|
||||
HIP_CHECK(hipHostFree(hPtr[i]));
|
||||
HIP_CHECK(hipFree(dPtr[i]));
|
||||
}
|
||||
delete [] hPtr;
|
||||
delete [] dPtr;
|
||||
delete[] hPtr;
|
||||
delete[] dPtr;
|
||||
return true;
|
||||
}
|
||||
|
||||
void hipPerfMandelBrot::setData(void *ptr, unsigned int value) {
|
||||
unsigned int *ptr2 = (unsigned int *)ptr;
|
||||
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;
|
||||
ptr2[i] = value;
|
||||
}
|
||||
}
|
||||
|
||||
void hipPerfMandelBrot::checkData(uint *ptr) {
|
||||
void hipPerfMandelBrot::checkData(uint* ptr) {
|
||||
totalIters = 0;
|
||||
for (unsigned int i = 0; i < width_ * width_; i++) {
|
||||
totalIters += ptr[i];
|
||||
@@ -601,30 +587,30 @@ void hipPerfMandelBrot::checkData(uint *ptr) {
|
||||
}
|
||||
|
||||
/**
|
||||
* 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 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
|
||||
#if HT_AMD
|
||||
SECTION("warm-up kernel default stream executes serially") {
|
||||
mandelbrotCompute.setNumStreams(1);
|
||||
mandelbrotCompute.setNumKernels(1);
|
||||
REQUIRE(true == mandelbrotCompute.run(100/*Random number*/));
|
||||
REQUIRE(true == mandelbrotCompute.run(100 /*Random number*/));
|
||||
}
|
||||
#endif
|
||||
#endif
|
||||
SECTION("run all - sync") {
|
||||
int i = 0;
|
||||
do {
|
||||
@@ -632,7 +618,7 @@ TEST_CASE("Perf_hipPerfMandelbrot") {
|
||||
mandelbrotCompute.setNumKernels(1);
|
||||
REQUIRE(true == mandelbrotCompute.run(i));
|
||||
i++;
|
||||
}while(i < 12);
|
||||
} while (i < 12);
|
||||
mandelbrotCompute.printResults();
|
||||
}
|
||||
|
||||
@@ -643,12 +629,12 @@ TEST_CASE("Perf_hipPerfMandelbrot") {
|
||||
mandelbrotCompute.setNumKernels(2);
|
||||
REQUIRE(true == mandelbrotCompute.run(i));
|
||||
i++;
|
||||
}while(i < 12);
|
||||
} while (i < 12);
|
||||
mandelbrotCompute.printResults();
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* End doxygen group perfComputeTest.
|
||||
* @}
|
||||
*/
|
||||
* End doxygen group perfComputeTest.
|
||||
* @}
|
||||
*/
|
||||
|
||||
Reference in New Issue
Block a user