Adding class size>8bytes tests

hipamd/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 
+}

hipamd/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 <iostream>
+#include <string>
+#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);
+};