From 83a25eb5b5ee07f8fc9314bb1c5252492dceeaab Mon Sep 17 00:00:00 2001 From: Sarbojit Sarkar Date: Wed, 7 Apr 2021 04:58:51 -0700 Subject: [PATCH] SWDEV-280131 - Fix __half2 member access bug Change-Id: Ic11750d18e4ba47446a7e494c3089a19cb4df783 --- include/hip/amd_detail/hip_fp16.h | 4 ++-- include/hip/amd_detail/hip_fp16_gcc.h | 4 ++-- tests/src/deviceLib/hipTestNativeHalf.cpp | 3 +++ 3 files changed, 7 insertions(+), 4 deletions(-) diff --git a/include/hip/amd_detail/hip_fp16.h b/include/hip/amd_detail/hip_fp16.h index bd016542e4..859c3efd9e 100644 --- a/include/hip/amd_detail/hip_fp16.h +++ b/include/hip/amd_detail/hip_fp16.h @@ -329,7 +329,7 @@ THE SOFTWARE. // BEGIN STRUCT __HALF2 struct __half2 { - protected: + public: union { static_assert( sizeof(_Float16_2) == sizeof(unsigned short[2]), ""); @@ -340,7 +340,7 @@ THE SOFTWARE. unsigned short y; }; }; - public: + // CREATORS __host__ __device__ __half2() = default; diff --git a/include/hip/amd_detail/hip_fp16_gcc.h b/include/hip/amd_detail/hip_fp16_gcc.h index 480fd8120a..e76a7fff3a 100644 --- a/include/hip/amd_detail/hip_fp16_gcc.h +++ b/include/hip/amd_detail/hip_fp16_gcc.h @@ -59,10 +59,10 @@ struct __half2_raw { // BEGIN STRUCT __HALF2 struct __half2 { - protected: + public: __half x; __half y; - public: + // CREATORS __half2() = default; __half2(const __half2_raw& ix) diff --git a/tests/src/deviceLib/hipTestNativeHalf.cpp b/tests/src/deviceLib/hipTestNativeHalf.cpp index f7ff3b493a..ad06880f60 100644 --- a/tests/src/deviceLib/hipTestNativeHalf.cpp +++ b/tests/src/deviceLib/hipTestNativeHalf.cpp @@ -159,6 +159,9 @@ void __half2Test(bool* result, __half2 a) { // Dot Functions result[0] = amd_mixed_dot(a, a, 1, 1) && result[0]; #endif + + half X = a.x; + half Y = a.y; } int main() {