diff --git a/hipamd/tests/src/compiler/hipClassKernel.cpp b/hipamd/tests/src/compiler/hipClassKernel.cpp index b3502ca9ae..edc3808d83 100644 --- a/hipamd/tests/src/compiler/hipClassKernel.cpp +++ b/hipamd/tests/src/compiler/hipClassKernel.cpp @@ -26,6 +26,66 @@ THE SOFTWARE. */ #include "hipClassKernel.h" +#ifdef ENABLE_OVERLOAD_OVERRIDE_TESTS +__global__ void +ovrdClassKernel(bool* result_ecd){ + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testOvrD tobj1; + result_ecd[tid] = (tobj1.ovrdFunc1() == 30); +} + +void HipClassTests::TestForOverride(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(ovrdClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} + + +__global__ void +ovldClassKernel(bool* result_ecd){ + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testFuncOvld tfo1; + result_ecd[tid] = (tfo1.func1(10) == 20) + && (tfo1.func1(10,10) == 30); +} + +void HipClassTests::TestForOverload(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(ovldClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} +#endif + +#ifdef ENABLE_FRIEND_TEST +// check for friend +__global__ void +friendClassKernel(bool* result_ecd){ + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testFrndB tfb1; + result_ecd[tid] = (tfb1.showA() == 10); +} +#endif + // check sizeof empty class is 1 __global__ void emptyClassKernel(bool* result_ecd) { @@ -209,6 +269,20 @@ void HipClassTests::TestForConsrtDesrt(){ } #endif +#ifdef ENABLE_FRIEND_TEST +void HipClassTests::TestForFriend(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + hipLaunchKernelGGL(friendClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); +} +#endif + bool* HipClassTests::AllocateHostMemory(void){ bool *result_ech; HIPCHECK(hipHostMalloc(&result_ech, @@ -253,6 +327,19 @@ int main(){ test_passed(TestForClassSize); classTests.TestForPassByValue(); test_passed(TestForPassByValue); + +#ifdef ENABLE_OVERLOAD_OVERRIDE_TESTS + classTests.TestForOverload(); + test_passed(TestForOverload); + classTests.TestForOverride(); + test_passed(TestForOverride); +#endif + +#ifdef ENABLE_FRIEND_TEST + classTests.TestForFriend(); + test_passed(TestForFriend); +#endif + // classTests.TestForMallocPassByValue(); // test_passed(TestForMallocPassByValue); #this test is crashing @@ -261,8 +348,6 @@ int main(){ test_passed(TestForVirtualClassSize); #endif - - #ifdef ENABLE_DESTRUCTOR_TEST classTests.TestForConsrtDesrt(); test_passed(TestForConsrtDesrt); diff --git a/hipamd/tests/src/compiler/hipClassKernel.h b/hipamd/tests/src/compiler/hipClassKernel.h index 8538276662..09d5f22bea 100644 --- a/hipamd/tests/src/compiler/hipClassKernel.h +++ b/hipamd/tests/src/compiler/hipClassKernel.h @@ -31,10 +31,62 @@ static const int BLOCKS = 512; static const int THREADS_PER_BLOCK = 1; static const int ENABLE_DESTRUCTOR_TEST = 0; static const int ENABLE_VIRTUAL_TESTS = 0; +static const int ENABLE_FRIEND_TEST = 0; +static const int ENABLE_OVERLAD_OVERRIDE_TESTS = 0; size_t NBOOL = BLOCKS * sizeof(bool); #define test_passed(test_name) printf("%s %s PASSED!%s\n", KGRN, #test_name, KNRM); +#ifdef ENABLE_OVERLOAD_OVERRIDE_TESTS +class testFuncOvld{ + public: + int __host__ __device__ func1(int a){ + return a + 10; + } + + int __host__ __device__ func1(int a , int b){ + return a + b + 10; + } + +}; + + +class testOvrB{ + public: + int __host__ __device__ ovrdFunc1(){ + return 10; + } + + +}; + + +class testOvrD: public testOvrB{ + public: + int __host__ __device__ ovrdFunc1(){ + int x = testOvrB::ovrdFunc1(); + return x + 20; + } + +}; +#endif + +#ifdef ENABLE_FRIEND_TEST +class testFrndA{ + private: + int fa = 10; + public: + friend class testFrndB; +}; + +class testFrndB{ + public: + __host__ __device__ int showA(){ + testFrndA x; + return x.fa; + } +}; +#endif class testClassEmpty {}; @@ -177,7 +229,9 @@ class HipClassTests{ void TestForPassByValue(void); void TestForMallocPassByValue(void); void TestForConsrtDesrt(void); - + void TestForOverload(void); + void TestForOverride(void); + bool* AllocateHostMemory(void); bool* AllocateDeviceMemory(void); void VerifyResult(bool* result_ech, bool* result_ecd);