From 5f1dfc38512c1f3eef48166d7a665e2396ed73ae Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Thu, 6 Jun 2024 16:47:23 +0000 Subject: [PATCH] SWDEV-465073 - Update hipClassKernel_Virtual test Compiler has changed to match the device side struct layout with the host side and hence the test is updated to verify if struct sizes of both host and device are equal instead of checking for hardcoded struct sizes Change-Id: Idd9ea1f04b413b7aa30f03555fd8e0ee11cd0f75 --- catch/unit/compiler/hipClassKernel.cc | 24 ++++++++++++++++-------- catch/unit/compiler/hipClassKernel.h | 10 +++++++++- 2 files changed, 25 insertions(+), 9 deletions(-) diff --git a/catch/unit/compiler/hipClassKernel.cc b/catch/unit/compiler/hipClassKernel.cc index 664a7401c9..2e3d519779 100644 --- a/catch/unit/compiler/hipClassKernel.cc +++ b/catch/unit/compiler/hipClassKernel.cc @@ -164,13 +164,13 @@ TEST_CASE("Unit_hipClassKernel_Size") { } __global__ void - sizeVirtualClassKernel(bool* result_ecd) { + sizeVirtualClassKernel(bool* result_ecd, refStructSizes structSizes) { 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) ; + result_ecd[tid] = (structSizes.sizeOftestSizeDV == sizeof(testSizeDV)) + && (structSizes.sizeOftestSizeDerivedDV == sizeof(testSizeDerivedDV)) + && (structSizes.sizeOftestSizeVirtDer = sizeof(testSizeVirtDer)) + && (structSizes.sizeOftestSizeVirtDerPack = sizeof(testSizeVirtDerPack)) + && (structSizes.sizeOftestSizeDerMulti = sizeof(testSizeDerMulti)); } TEST_CASE("Unit_hipClassKernel_Virtual") { @@ -178,12 +178,20 @@ TEST_CASE("Unit_hipClassKernel_Virtual") { result_ech = AllocateHostMemory(); result_ecd = AllocateDeviceMemory(); + struct refStructSizes structSizes; + structSizes.sizeOftestSizeDV = sizeof(testSizeDV); + structSizes.sizeOftestSizeDerivedDV = sizeof(testSizeDerivedDV); + structSizes.sizeOftestSizeVirtDer = sizeof(testSizeVirtDer); + structSizes.sizeOftestSizeVirtDerPack = sizeof(testSizeVirtDerPack); + structSizes.sizeOftestSizeDerMulti = sizeof(testSizeDerMulti); + hipLaunchKernelGGL(sizeVirtualClassKernel, dim3(BLOCKS), dim3(THREADS_PER_BLOCK), 0, 0, - result_ecd); + result_ecd, + structSizes); VerifyResult(result_ech,result_ecd); FreeMem(result_ech,result_ecd); @@ -215,4 +223,4 @@ TEST_CASE("Unit_hipClassKernel_Value") { VerifyResult(result_ech,result_ecd); FreeMem(result_ech,result_ecd); -} \ No newline at end of file +} diff --git a/catch/unit/compiler/hipClassKernel.h b/catch/unit/compiler/hipClassKernel.h index 6da07f61d6..608630a6b5 100644 --- a/catch/unit/compiler/hipClassKernel.h +++ b/catch/unit/compiler/hipClassKernel.h @@ -152,6 +152,14 @@ class testSizeDerivedDV : testSizeDV { int iDDV; }; +struct refStructSizes { + size_t sizeOftestSizeDV; + size_t sizeOftestSizeDerivedDV; + size_t sizeOftestSizeVirtDer; + size_t sizeOftestSizeVirtDerPack; + size_t sizeOftestSizeDerMulti; +}; + #pragma pack(push, 1) class testSizeP1 { public: @@ -231,4 +239,4 @@ void FreeMem(bool* result_ech, bool* result_ecd){ HIPCHECK(hipFree(result_ecd)); } -#endif // _HIPCLASSKERNEL_H_ \ No newline at end of file +#endif // _HIPCLASSKERNEL_H_