diff --git a/tests/src/kernel/hipLaunchParmFunctor.cpp b/tests/src/kernel/hipLaunchParmFunctor.cpp index 3e876ab3ea..f5c30b07a4 100755 --- a/tests/src/kernel/hipLaunchParmFunctor.cpp +++ b/tests/src/kernel/hipLaunchParmFunctor.cpp @@ -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);