diff --git a/projects/hip/tests/catch/unit/CMakeLists.txt b/projects/hip/tests/catch/unit/CMakeLists.txt
index 0ce9be5ed9..b1874c23eb 100644
--- a/projects/hip/tests/catch/unit/CMakeLists.txt
+++ b/projects/hip/tests/catch/unit/CMakeLists.txt
@@ -31,6 +31,7 @@ add_subdirectory(texture)
 add_subdirectory(streamperthread)
 add_subdirectory(kernel)
 add_subdirectory(multiThread)
+add_subdirectory(compiler)
 if(HIP_PLATFORM STREQUAL "amd")
 add_subdirectory(clock)
 endif()
diff --git a/projects/hip/tests/catch/unit/compiler/CMakeLists.txt b/projects/hip/tests/catch/unit/compiler/CMakeLists.txt
new file mode 100644
index 0000000000..67eb7afcf5
--- /dev/null
+++ b/projects/hip/tests/catch/unit/compiler/CMakeLists.txt
@@ -0,0 +1,8 @@
+# Common Tests - Test independent of all platforms
+set(TEST_SRC
+  hipClassKernel.cc
+)
+
+hip_add_exe_to_target(NAME CompilerTest
+                      TEST_SRC ${TEST_SRC}
+                      TEST_TARGET_NAME build_tests)
diff --git a/projects/hip/tests/catch/unit/compiler/hipClassKernel.cc b/projects/hip/tests/catch/unit/compiler/hipClassKernel.cc
new file mode 100644
index 0000000000..b205c6bffb
--- /dev/null
+++ b/projects/hip/tests/catch/unit/compiler/hipClassKernel.cc
@@ -0,0 +1,220 @@
+/*
+Copyright (c) 2015 - 2021 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
+ * TEST: %t
+ * HIT_END
+ */
+#include "hipClassKernel.h"
+
+__global__ void
+ovrdClassKernel(bool* result_ecd){
+    int tid = threadIdx.x + blockIdx.x * blockDim.x;
+    testOvrD tobj1;
+     result_ecd[tid] = (tobj1.ovrdFunc1() == 30);
+}
+
+__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);
+}
+
+TEST_CASE("Unit_hipClassKernel_Overload_Override") {
+  bool *result_ecd, *result_ech;
+  result_ech = AllocateHostMemory();
+  result_ecd = AllocateDeviceMemory();
+
+  hipLaunchKernelGGL(ovrdClassKernel,
+                     dim3(BLOCKS),
+                     dim3(THREADS_PER_BLOCK),
+                     0,
+                     0,
+                     result_ecd);
+
+  VerifyResult(result_ech,result_ecd);
+  FreeMem(result_ech,result_ecd);
+
+  result_ech = AllocateHostMemory();
+  result_ecd = AllocateDeviceMemory();
+
+  hipLaunchKernelGGL(ovldClassKernel,
+                     dim3(BLOCKS),
+                     dim3(THREADS_PER_BLOCK),
+                     0,
+                     0,
+                     result_ecd);
+
+  VerifyResult(result_ech,result_ecd);
+  FreeMem(result_ech,result_ecd);
+}
+
+// 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);
+}
+
+TEST_CASE("Unit_hipClassKernel_Friend") {
+  bool *result_ecd;
+  result_ecd = AllocateDeviceMemory();
+  hipLaunchKernelGGL(friendClassKernel,
+                      dim3(BLOCKS),
+                      dim3(THREADS_PER_BLOCK),
+                      0,
+                      0,
+                      result_ecd);
+}
+
+// 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);
+}
+
+TEST_CASE("Unit_hipClassKernel_Empty") {
+  bool *result_ecd, *result_ech;
+  result_ech = AllocateHostMemory();
+  result_ecd = AllocateDeviceMemory();
+
+  hipLaunchKernelGGL(emptyClassKernel,
+                     dim3(BLOCKS),
+                     dim3(THREADS_PER_BLOCK),
+                     0,
+                     0,
+                     result_ecd);
+
+  VerifyResult(result_ech,result_ecd);
+  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);
+}
+
+TEST_CASE("Unit_hipClassKernel_BSize") {
+  bool *result_ecd, *result_ech;
+  result_ech = AllocateHostMemory();
+  result_ecd = AllocateDeviceMemory();
+
+   hipLaunchKernelGGL(sizeClassBKernel,
+                      dim3(BLOCKS),
+                      dim3(THREADS_PER_BLOCK),
+                      0,
+                      0,
+                      result_ecd);
+
+  VerifyResult(result_ech,result_ecd);
+  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);
+}
+
+TEST_CASE("Unit_hipClassKernel_Size") {
+    bool *result_ecd, *result_ech;
+  result_ech = AllocateHostMemory();
+  result_ecd = AllocateDeviceMemory();
+
+   hipLaunchKernelGGL(sizeClassKernel,
+                      dim3(BLOCKS),
+                      dim3(THREADS_PER_BLOCK),
+                      0,
+                      0,
+                      result_ecd);
+
+  VerifyResult(result_ech,result_ecd);
+  FreeMem(result_ech,result_ecd);
+}
+
+__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) ;
+ }
+
+TEST_CASE("Unit_hipClassKernel_Virtual") {
+  bool *result_ecd, *result_ech;
+  result_ech = AllocateHostMemory();
+  result_ecd = AllocateDeviceMemory();
+
+  hipLaunchKernelGGL(sizeVirtualClassKernel,
+                      dim3(BLOCKS),
+                      dim3(THREADS_PER_BLOCK),
+                      0,
+                      0,
+                      result_ecd);
+
+  VerifyResult(result_ech,result_ecd);
+  FreeMem(result_ech,result_ecd);
+}
+
+// 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');
+}
+
+TEST_CASE("Unit_hipClassKernel_Value") {
+  bool *result_ecd,*result_ech;
+  result_ech = AllocateHostMemory();
+  result_ecd = AllocateDeviceMemory();
+
+   testPassByValue exObj;
+   exObj.exI = 10;
+   exObj.exC = 'C';
+   hipLaunchKernelGGL(passByValueKernel,
+                      dim3(BLOCKS),
+                      dim3(THREADS_PER_BLOCK),
+                      0,
+                      0,
+                      exObj,
+                      result_ecd);
+
+  VerifyResult(result_ech,result_ecd);
+  FreeMem(result_ech,result_ecd);
+}
\ No newline at end of file
diff --git a/projects/hip/tests/catch/unit/compiler/hipClassKernel.h b/projects/hip/tests/catch/unit/compiler/hipClassKernel.h
new file mode 100644
index 0000000000..6da07f61d6
--- /dev/null
+++ b/projects/hip/tests/catch/unit/compiler/hipClassKernel.h
@@ -0,0 +1,234 @@
+/*
+Copyright (c) 2015 - 2021 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 <hip_test_common.hh>
+
+static const int BLOCKS = 512;
+static const int THREADS_PER_BLOCK = 1;
+size_t NBOOL = BLOCKS * sizeof(bool);
+
+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;
+    }
+
+};
+
+class testFrndA{
+   private:
+       int fa = 10;
+   public:
+      friend class testFrndB;
+};
+
+class testFrndB{
+    public:
+      __host__ __device__ int showA(){
+         testFrndA x;
+          return x.fa;
+     }
+};
+
+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;
+};
+
+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)
+
+class testSizeC {
+ public:
+    char ac;
+    int ic;
+    int bc[2];
+};
+
+class testSizeDV {
+ public:
+    virtual void __host__ __device__ func1();
+ private:
+    int iDV;
+
+};
+
+class testSizeDerivedDV : testSizeDV {
+ public:
+    virtual void __host__ __device__ funcD1();
+ private:
+    int iDDV;
+};
+
+#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)
+
+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;
+}
+
+bool* AllocateHostMemory(void){
+  bool *result_ech;
+  HIPCHECK(hipHostMalloc(&result_ech,
+                         NBOOL,
+                         hipHostMallocDefault));
+  return result_ech;
+}
+
+bool* AllocateDeviceMemory(void){
+  bool* result_ecd;
+  HIPCHECK(hipMalloc(&result_ecd,
+                     NBOOL));
+  HIPCHECK(hipMemset(result_ecd,
+                       false,
+                       NBOOL));
+  return result_ecd;
+}
+
+void 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 FreeMem(bool* result_ech, bool* result_ecd){
+  HIPCHECK(hipHostFree(result_ech));
+  HIPCHECK(hipFree(result_ecd));
+}
+
+#endif  // _HIPCLASSKERNEL_H_
\ No newline at end of file