From 8b00f5657780bd7e06fd27a3f4b59d7c8075fe8f Mon Sep 17 00:00:00 2001 From: Rakesh Roy Date: Fri, 20 Jan 2023 22:39:43 +0530 Subject: [PATCH] SWDEV-375004 - Fix rocprim test failure - For !__HIPCC_RTC__ case, operator<<(std::ostream& os, const hip_bfloat16& bf16) calls itself in endless recursion - Convert hip_bfloat16 to float to fix this Change-Id: I252a656817550caf43c587cebf461474f12b8c0c --- hipamd/include/hip/amd_detail/amd_hip_bfloat16.h | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h b/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h index 8c7f78390a..deb3bfb7e2 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h +++ b/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h @@ -32,10 +32,8 @@ #include "host_defines.h" #if defined(__HIPCC_RTC__) #define __HOST_DEVICE__ __device__ - #define HIP_OSTREAM __hip_internal::ostream #else #define __HOST_DEVICE__ __host__ __device__ - #define HIP_OSTREAM std::ostream #endif #if __cplusplus < 201103L || !defined(__HIPCC__) @@ -181,12 +179,12 @@ static_assert(__hip_internal::is_trivial{}, static_assert(sizeof(hip_bfloat16) == sizeof(hip_bfloat16_public) && offsetof(hip_bfloat16, data) == offsetof(hip_bfloat16_public, data), "internal hip_bfloat16 does not match public hip_bfloat16"); -#endif -inline HIP_OSTREAM& operator<<(HIP_OSTREAM& os, const hip_bfloat16& bf16) +inline std::ostream& operator<<(std::ostream& os, const hip_bfloat16& bf16) { - return os << bf16; + return os << float(bf16); } +#endif inline __HOST_DEVICE__ hip_bfloat16 operator+(hip_bfloat16 a) {