From 0022f53a10c1acb697d149715a5cf9357d0e039e Mon Sep 17 00:00:00 2001 From: fpadmin Date: Wed, 11 Jul 2018 16:30:33 +0530 Subject: [PATCH] Adding class size>8bytes tests --- tests/src/compiler/hipClassKernel.cpp | 270 ++++++++++++++++++++++++++ tests/src/compiler/hipClassKernel.h | 185 ++++++++++++++++++ 2 files changed, 455 insertions(+) create mode 100644 tests/src/compiler/hipClassKernel.cpp create mode 100644 tests/src/compiler/hipClassKernel.h diff --git a/tests/src/compiler/hipClassKernel.cpp b/tests/src/compiler/hipClassKernel.cpp new file mode 100644 index 0000000000..b3502ca9ae --- /dev/null +++ b/tests/src/compiler/hipClassKernel.cpp @@ -0,0 +1,270 @@ +/* +Copyright (c) 2015-Present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * HIT_END + */ +#include "hipClassKernel.h" + +// check sizeof empty class is 1 +__global__ void +emptyClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testClassEmpty ob1,ob2; + result_ecd[tid] = (sizeof(testClassEmpty) == 1) + && (&ob1 != &ob2); +} + +void HipClassTests::TestForEmptyClass(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(emptyClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} + +// tests for classes >8 bytes +__global__ void + sizeClassBKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (sizeof(testSizeB) == 12) + && (sizeof(testSizeC) == 16) + && (sizeof(testSizeP1) == 6) + && (sizeof(testSizeP2) == 13) + && (sizeof(testSizeP3) == 8); +} + +void HipClassTests::TestForClassBSize(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(sizeClassBKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} + +__global__ void +sizeClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (sizeof(testSizeA) == 16) + && (sizeof(testSizeDerived) == 24) + && (sizeof(testSizeDerived2) == 20); +} + +void HipClassTests::TestForClassSize(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(sizeClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} + +#ifdef ENABLE_VIRTUAL_TESTS +__global__ void + sizeVirtualClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (sizeof(testSizeDV) == 16) + && (sizeof(testSizeDerivedDV) == 16) + && (sizeof(testSizeVirtDerPack) == 24) + && (sizeof(testSizeVirtDer) == 24) + && (sizeof(testSizeDerMulti) == 48) ; + } + +void HipClassTests::TestForVirtualClassSize(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(sizeVirtualClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} +#endif + +// check pass by value +__global__ void +passByValueKernel(testPassByValue obj, bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (obj.exI == 10) + && (obj.exC == 'C'); +} + +void HipClassTests::TestForPassByValue(void){ + bool *result_ecd,*result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + testPassByValue exObj; + exObj.exI = 10; + exObj.exC = 'C'; + hipLaunchKernelGGL(passByValueKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + exObj, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} + + // check obj created with hipMalloc +__global__ void +mallocObjKernel(testPassByValue *obj, bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (obj->exI == 100) + && (obj->exC == 'C'); +} + +void HipClassTests::TestForMallocPassByValue(void){ + bool *result_ecd,*result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + + testPassByValue *exObjM; + HIPCHECK(hipMalloc(&exObjM, sizeof(testPassByValue))); + exObjM->exI = 100; + exObjM->exC = 'C'; + hipLaunchKernelGGL(mallocObjKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + exObjM, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); + +} + +// check if constr and destr are accessible from kernel +#ifdef ENABLE_DESTRUCTOR_TEST +__global__ void +testDeviceClassKernel() { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testDeviceClass ob1; + testDeviceClass ob2; + ob2.iVar = 10; +} + +void HipClassTests::TestForConsrtDesrt(){ + testDeviceClass tDC; + hipLaunchKernelGGL(testDeviceClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0); +} +#endif + +bool* HipClassTests::AllocateHostMemory(void){ + bool *result_ech; + HIPCHECK(hipHostMalloc(&result_ech, + NBOOL, + hipHostMallocDefault)); + return result_ech; +} + +bool* HipClassTests::AllocateDeviceMemory(void){ + bool* result_ecd; + HIPCHECK(hipMalloc(&result_ecd, + NBOOL)); + HIPCHECK(hipMemset(result_ecd, + false, + NBOOL)); + return result_ecd; +} + +void HipClassTests::VerifyResult(bool* result_ech, bool* result_ecd){ + HIPCHECK(hipMemcpy(result_ech, + result_ecd, + BLOCKS*sizeof(bool), + hipMemcpyDeviceToHost)); + // validation on host side + for (int i = 0; i < BLOCKS; i++) { + HIPASSERT(result_ech[i] == true); + } +} + +void HipClassTests::FreeMem(bool* result_ech, bool* result_ecd){ + HIPCHECK(hipHostFree(result_ech)); + HIPCHECK(hipFree(result_ecd)); +} + +int main(){ + HipClassTests classTests; + classTests.TestForEmptyClass(); + test_passed(TestForEmptyClass); + classTests.TestForClassBSize(); + test_passed(TestForClassBSize); + classTests.TestForClassSize(); + test_passed(TestForClassSize); + classTests.TestForPassByValue(); + test_passed(TestForPassByValue); +// classTests.TestForMallocPassByValue(); + // test_passed(TestForMallocPassByValue); #this test is crashing + +#ifdef ENABLE_VIRTUAL_TESTS + classTests.TestForVirtualClassSize(); + test_passed(TestForVirtualClassSize); +#endif + + + +#ifdef ENABLE_DESTRUCTOR_TEST + classTests.TestForConsrtDesrt(); + test_passed(TestForConsrtDesrt); +#endif +} diff --git a/tests/src/compiler/hipClassKernel.h b/tests/src/compiler/hipClassKernel.h new file mode 100644 index 0000000000..8538276662 --- /dev/null +++ b/tests/src/compiler/hipClassKernel.h @@ -0,0 +1,185 @@ +/* +Copyright (c) 2015-Present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#ifndef _COMPILER_HIPCLASSKERNEL_H_ +#define _COMPILER_HIPCLASSKERNEL_H_ + +#include +#include +#include "hip/hip_runtime.h" +#include "test_common.h" + +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; +size_t NBOOL = BLOCKS * sizeof(bool); + +#define test_passed(test_name) printf("%s %s PASSED!%s\n", KGRN, #test_name, KNRM); + + +class testClassEmpty {}; + +class testPassByValue{ + public: + int exI; + char exC; +}; + +class testSizeA { + public: + float xa; + int ia; + double da; + static char ca; +}; + +class testSizeDerived : testSizeA { + public: + float fd; +}; + +#pragma pack(push,4) +class testSizeDerived2 : testSizeA { + public: + float fd; +}; +#pragma pack(pop) + +class testSizeB { + public: + char ab; + int ib; + char cb; +}; + +#ifdef ENBABLE_VIRTUAL_TESTS +class testSizeVirtDer : public virtual testSizeB { + public: + int ivd; +}; + +class testSizeVirtDer1 : public virtual testSizeB { + public: + int ivd1; +}; + +class testSizeDerMulti : public testSizeVirtDer, public testSizeVirtDer1 { + public: + int ivd2; +}; + +#pragma pack(push,4) +class testSizeVirtDerPack : public virtual testSizeB { + public: + int ivd; +}; +#pragma pack(pop) +#endif + +class testSizeC { + public: + char ac; + int ic; + int bc[2]; +}; + +#ifdef ENABLE_VIRTUAL_TESTS +class testSizeDV { + public: + virtual void __host__ __device__ func1(); + private: + int iDV; + +}; + +class testSizeDerivedDV : testSizeDV { + public: + virtual void __host__ __device__ funcD1(); + private: + int iDDV; +}; +#endif + +#pragma pack(push, 1) +class testSizeP1 { + public: + char ap; + int ip; + char cp; +}; +#pragma pack(pop) + +#pragma pack(push, 1) +class testSizeP2 { + public: + char ap1; + int ip1; + int bp1[2]; +}; +#pragma pack(pop) + +#pragma pack(push, 2) +class testSizeP3 { + public: + char ap2; + int ip2; + char cp2; +}; +#pragma pack(pop) + +#ifdef ENABLE_DESTRUCTOR_TEST +class testDeviceClass { + public: + int iVar; + __host__ __device__ testDeviceClass(); + __host__ __device__ testDeviceClass(int a); + __host__ __device__ ~testDeviceClass(); +}; + +__host__ __device__ +testDeviceClass::testDeviceClass() { + iVar = 5; +} + +__host__ __device__ +testDeviceClass::testDeviceClass(int a) { + iVar = a; +} +#endif + +#endif // _HIPCLASSKERNEL_H_ + +class HipClassTests{ + public: + void TestForEmptyClass(void); + void TestForClassBSize(void); + void TestForClassSize(void); + void TestForVirtualClassSize(void); + void TestForPassByValue(void); + void TestForMallocPassByValue(void); + void TestForConsrtDesrt(void); + + bool* AllocateHostMemory(void); + bool* AllocateDeviceMemory(void); + void VerifyResult(bool* result_ech, bool* result_ecd); + void FreeMem(bool* result_ech, bool* result_ecd); +};