AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'master' of https://github.com/rpathani/HIP

Browse Source 
# Conflicts:
#	tests/src/kernel/hipLaunchParmFunctor.cpp
This commit is contained in:
rohit pathania
2018-07-10 15:36:07 +05:30
parent 58b9ce3b2d 782f8613e0
commit c2108cb82a
1 changed files with 49 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/tests/src/kernel/hipLaunchParmFunctor.cpp
+49 -28
View File
@@ -18,16 +18,13 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -O3
* RUN: %t
* HIT_END
*/
#include <iostream>
#include "hip/hip_runtime.h"
#include "../test_common.h"
#include "hip/hip_runtime_api.h"
#include "../test_common.h"
#define test_passed(test_name) printf("%s %s PASSED!%s\n", KGRN, #test_name, KNRM);
@@ -59,6 +56,7 @@ class HipFunctorTests {
void TestForFunctorContainInStructObj(void);
};
static const int BLOCK_DIM_SIZE = 1024;
static const int THREADS_PER_BLOCK = 1;
@@ -70,6 +68,7 @@ class DoublerFunctor{
__device__ int operator()(int x) { return x * 2;}
};
// simple doubler functor passed to kernel
__global__ void DoublerFunctorKernel(
DoublerFunctor doubler_,
@@ -87,10 +86,12 @@ void HipFunctorTests::TestForSimpleClassFunctor(void) {
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
// initialize to false, will be set to
// true if the functor is called in device code
deviceResults[k] = false;
hostResults[k] = false;
}
hipLaunchKernelGGL(HIP_KERNEL_NAME(DoublerFunctorKernel), dim3(BLOCK_DIM_SIZE),
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
hipMemcpyHostToDevice));
hipLaunchKernelGGL(DoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
dim3(THREADS_PER_BLOCK), 0, 0, doubler, deviceResults);
// Validation part of TestForSimpleClassFunctor
@@ -102,6 +103,7 @@ void HipFunctorTests::TestForSimpleClassFunctor(void) {
HIPCHECK(hipFree(deviceResults));
}
// pointer functor passed to kernel
__global__ void PtrDoublerFunctorKernel(
DoublerFunctor *doubler_,
@@ -119,10 +121,12 @@ void HipFunctorTests::TestForClassObjPtrFunctor(void) {
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
// initialize to false, will be set to
// true if the functor is called in device code
deviceResults[k] = false;
hostResults[k] = false;
}
hipLaunchKernelGGL(HIP_KERNEL_NAME(PtrDoublerFunctorKernel), dim3(BLOCK_DIM_SIZE),
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
hipMemcpyHostToDevice));
hipLaunchKernelGGL(PtrDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
dim3(THREADS_PER_BLOCK), 0, 0, ptrdoubler, deviceResults);
// Validation part of TestForClassObjPtrFunctor
@@ -143,6 +147,7 @@ class compare {
}
};
// template functor passed to kernel
__global__ void TemplateFunctorKernel(
compare compare_,
@@ -161,10 +166,12 @@ void HipFunctorTests::TestForClassTemplateFunctor(void) {
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
// initialize to false, will be set to
// true if the functor is called in device code
deviceResults[k] = false;
hostResults[k] = false;
}
hipLaunchKernelGGL(HIP_KERNEL_NAME(TemplateFunctorKernel), dim3(BLOCK_DIM_SIZE),
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
hipMemcpyHostToDevice));
hipLaunchKernelGGL(TemplateFunctorKernel, dim3(BLOCK_DIM_SIZE),
dim3(THREADS_PER_BLOCK), 0, 0, comparefunctor, deviceResults);
// Validation part of TestForClassTemplateFunctor
@@ -178,16 +185,17 @@ void HipFunctorTests::TestForClassTemplateFunctor(void) {
// Doubler calculator
class DoublerCaluclator {
class DoublerCalculator {
public:
int a, result;
// fucntor contained in class object
DoublerFunctor doubler;
};
// doubler functor conatined in class obj passed to kernel
__global__ void DoublerCalculatorFunctorKernel(
DoublerCaluclator doubler_,
DoublerCalculator doubler_,
bool* deviceResult) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int result = doubler_.doubler(doubler_.a);
@@ -195,20 +203,22 @@ __global__ void DoublerCalculatorFunctorKernel(
}
void HipFunctorTests::TestForFunctorContainInClassObj(void) {
DoublerCaluclator Doubler;
DoublerCalculator Doubler;
bool *deviceResults, *hostResults;
HIPCHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool)));
HIPCHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool)));
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
// initialize to false, will be set to
// true if the functor is called in device code
deviceResults[k] = false;
hostResults[k] = false;
}
Doubler.a = 5;
Doubler.result = 10;
// pass comparefunctor to hipLaunchParm
hipLaunchKernelGGL(HIP_KERNEL_NAME(DoublerCalculatorFunctorKernel), dim3(BLOCK_DIM_SIZE),
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
hipMemcpyHostToDevice));
hipLaunchKernelGGL(DoublerCalculatorFunctorKernel, dim3(BLOCK_DIM_SIZE),
dim3(THREADS_PER_BLOCK), 0, 0, Doubler, deviceResults);
// Validation part of TestForStructTemplateFunctor
@@ -229,6 +239,7 @@ struct sDoublerFunctor {
__device__ int operator()(int x) { return x * 2;}
};
// simple sturct doubler functor passed to kernel
__global__ void structDoublerFunctorKernel(
sDoublerFunctor doubler_,
@@ -246,10 +257,11 @@ void HipFunctorTests::TestForSimpleStructFunctor(void) {
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
// initialize to false, will be set to
// true if the functor is called in device code
deviceResults[k] = false;
hostResults[k] = false;
}
hipLaunchKernelGGL(HIP_KERNEL_NAME(structDoublerFunctorKernel), dim3(BLOCK_DIM_SIZE),
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
hipMemcpyHostToDevice));
hipLaunchKernelGGL(structDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
dim3(THREADS_PER_BLOCK), 0, 0, doubler, deviceResults);
// Validation part of TestForSimpleStructFunctor
@@ -278,10 +290,12 @@ void HipFunctorTests::TestForStructObjPtrFunctor(void) {
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
// initialize to false, will be set to
// true if the functor is called in device code
deviceResults[k] = false;
hostResults[k] = false;
}
hipLaunchKernelGGL(HIP_KERNEL_NAME(structPtrDoublerFunctorKernel), dim3(BLOCK_DIM_SIZE),
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
hipMemcpyHostToDevice));
hipLaunchKernelGGL(structPtrDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
dim3(THREADS_PER_BLOCK), 0, 0, ptrdoubler, deviceResults);
// Validation part of TestForStructObjPtrFunctor
@@ -302,6 +316,7 @@ struct sCompare {
}
};
// template functor passed to kernel
__global__ void structTemplateFunctorKernel(
sCompare compare_,
@@ -320,11 +335,13 @@ void HipFunctorTests::TestForStructTemplateFunctor(void) {
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
// initialize to false, will be set to
// true if the functor is called in device code
deviceResults[k] = false;
hostResults[k] = false;
}
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
hipMemcpyHostToDevice));
// pass comparefunctor to hipLaunchKernelGGL
hipLaunchKernelGGL(HIP_KERNEL_NAME(structTemplateFunctorKernel), dim3(BLOCK_DIM_SIZE),
hipLaunchKernelGGL(structTemplateFunctorKernel, dim3(BLOCK_DIM_SIZE),
dim3(THREADS_PER_BLOCK), 0, 0, comparefunctor, deviceResults);
// Validation part of TestForStructTemplateFunctor
@@ -337,16 +354,17 @@ void HipFunctorTests::TestForStructTemplateFunctor(void) {
}
// Doubler calculator struct
struct sDoublerCaluclator {
struct sDoublerCalculator {
public:
int a, result;
// fucntor contained in class object
DoublerFunctor doubler;
};
// doubler functor contained in struct passed to kernel
__global__ void DoublerCalculatorFunctorKernel(
sDoublerCaluclator doubler_,
sDoublerCalculator doubler_,
bool* deviceResult) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int result = doubler_.doubler(doubler_.a);
@@ -354,20 +372,23 @@ __global__ void DoublerCalculatorFunctorKernel(
}
void HipFunctorTests::TestForFunctorContainInStructObj(void) {
sDoublerCaluclator Doubler;
sDoublerCalculator Doubler;
bool *deviceResults, *hostResults;
HIPCHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool)));
HIPCHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool)));
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
// initialize to false, will be set to
// true if the functor is called in device code
deviceResults[k] = false;
hostResults[k] = false;
}
Doubler.a = 5;
Doubler.result = 10;
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
hipMemcpyHostToDevice));
// pass comparefunctor to hipLaunchKernelGGL
hipLaunchKernelGGL(HIP_KERNEL_NAME(DoublerCalculatorFunctorKernel), dim3(BLOCK_DIM_SIZE),
hipLaunchKernelGGL(DoublerCalculatorFunctorKernel, dim3(BLOCK_DIM_SIZE),
dim3(THREADS_PER_BLOCK), 0, 0, Doubler, deviceResults);
// Validation part of TestForStructTemplateFunctor