Merge branch 'master' of https://github.com/rpathani/HIP
# Conflicts: # tests/src/kernel/hipLaunchParmFunctor.cpp
This commit is contained in:
@@ -57,6 +57,8 @@ class HipFunctorTests {
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
static const int BLOCK_DIM_SIZE = 1024;
|
||||
static const int THREADS_PER_BLOCK = 1;
|
||||
|
||||
@@ -69,6 +71,8 @@ class DoublerFunctor{
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
// simple doubler functor passed to kernel
|
||||
__global__ void DoublerFunctorKernel(
|
||||
DoublerFunctor doubler_,
|
||||
@@ -89,6 +93,7 @@ void HipFunctorTests::TestForSimpleClassFunctor(void) {
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(DoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
@@ -104,6 +109,8 @@ void HipFunctorTests::TestForSimpleClassFunctor(void) {
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
// pointer functor passed to kernel
|
||||
__global__ void PtrDoublerFunctorKernel(
|
||||
DoublerFunctor *doubler_,
|
||||
@@ -124,6 +131,7 @@ void HipFunctorTests::TestForClassObjPtrFunctor(void) {
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(PtrDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
@@ -148,6 +156,8 @@ class compare {
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
// template functor passed to kernel
|
||||
__global__ void TemplateFunctorKernel(
|
||||
compare compare_,
|
||||
@@ -193,6 +203,7 @@ class DoublerCalculator {
|
||||
};
|
||||
|
||||
|
||||
|
||||
// doubler functor conatined in class obj passed to kernel
|
||||
__global__ void DoublerCalculatorFunctorKernel(
|
||||
DoublerCalculator doubler_,
|
||||
@@ -216,6 +227,7 @@ void HipFunctorTests::TestForFunctorContainInClassObj(void) {
|
||||
Doubler.a = 5;
|
||||
Doubler.result = 10;
|
||||
// pass comparefunctor to hipLaunchParm
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(DoublerCalculatorFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
@@ -240,6 +252,8 @@ struct sDoublerFunctor {
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
// simple sturct doubler functor passed to kernel
|
||||
__global__ void structDoublerFunctorKernel(
|
||||
sDoublerFunctor doubler_,
|
||||
@@ -259,6 +273,7 @@ void HipFunctorTests::TestForSimpleStructFunctor(void) {
|
||||
// true if the functor is called in device code
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(structDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
@@ -293,6 +308,7 @@ void HipFunctorTests::TestForStructObjPtrFunctor(void) {
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(structPtrDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
@@ -317,6 +333,8 @@ struct sCompare {
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
// template functor passed to kernel
|
||||
__global__ void structTemplateFunctorKernel(
|
||||
sCompare compare_,
|
||||
@@ -340,6 +358,7 @@ void HipFunctorTests::TestForStructTemplateFunctor(void) {
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
// pass comparefunctor to hipLaunchKernelGGL
|
||||
hipLaunchKernelGGL(structTemplateFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, comparefunctor, deviceResults);
|
||||
@@ -362,6 +381,7 @@ struct sDoublerCalculator {
|
||||
};
|
||||
|
||||
|
||||
|
||||
// doubler functor contained in struct passed to kernel
|
||||
__global__ void DoublerCalculatorFunctorKernel(
|
||||
sDoublerCalculator doubler_,
|
||||
@@ -387,6 +407,7 @@ void HipFunctorTests::TestForFunctorContainInStructObj(void) {
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
|
||||
// pass comparefunctor to hipLaunchKernelGGL
|
||||
hipLaunchKernelGGL(DoublerCalculatorFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, Doubler, deviceResults);
|
||||
|
||||
Referens i nytt ärende
Block a user