From b1ebf33850c8d9d51a3d5dca4faec7a9b5530867 Mon Sep 17 00:00:00 2001
From: "Chaudhary, Jatin Jaikishan" <JatinJaikishan.Chaudhary@amd.com>
Date: Tue, 6 May 2025 09:41:21 +0100
Subject: [PATCH] SWDEV-529927 - add missing operations for fp16/bf16 (#238)

---
 hipamd/include/hip/amd_detail/amd_hip_bf16.h | 15 ++++++++++++---
 hipamd/include/hip/amd_detail/amd_hip_fp16.h |  2 +-
 2 files changed, 13 insertions(+), 4 deletions(-)

diff --git a/hipamd/include/hip/amd_detail/amd_hip_bf16.h b/hipamd/include/hip/amd_detail/amd_hip_bf16.h
index 405e320bb5..09065a5247 100644
--- a/hipamd/include/hip/amd_detail/amd_hip_bf16.h
+++ b/hipamd/include/hip/amd_detail/amd_hip_bf16.h
@@ -453,7 +453,7 @@ __BF16_HOST_DEVICE_STATIC__ __hip_bfloat162 __bfloat162bfloat162(const __hip_bfl
 }
 
 /**
- * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \ingroup HIP_INTRINSIC_BFLOAT16_CONV
  * \brief Reinterprets bits in a __hip_bfloat16 as a signed short integer
  */
 __BF16_HOST_DEVICE_STATIC__ short int __bfloat16_as_short(const __hip_bfloat16 h) {
@@ -466,7 +466,7 @@ __BF16_HOST_DEVICE_STATIC__ short int __bfloat16_as_short(const __hip_bfloat16 h
 }
 
 /**
- * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \ingroup HIP_INTRINSIC_BFLOAT16_CONV
  * \brief Reinterprets bits in a __hip_bfloat16 as an unsigned signed short integer
  */
 __BF16_HOST_DEVICE_STATIC__ unsigned short int __bfloat16_as_ushort(const __hip_bfloat16 h) {
@@ -479,7 +479,7 @@ __BF16_HOST_DEVICE_STATIC__ unsigned short int __bfloat16_as_ushort(const __hip_
 }
 
 /**
- * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \ingroup HIP_INTRINSIC_BFLOAT16_CONV
  * \brief Convert double to __hip_bfloat16
  */
 __BF16_HOST_DEVICE_STATIC__ __hip_bfloat16 __double2bfloat16(const double a) {
@@ -487,6 +487,15 @@ __BF16_HOST_DEVICE_STATIC__ __hip_bfloat16 __double2bfloat16(const double a) {
   return ret;
 }
 
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_CONV
+ * \brief Convert int to __hip_bfloat16
+ */
+__BF16_HOST_DEVICE_STATIC__ __hip_bfloat16 __int2bfloat16_rn(const int a) {
+  __hip_bfloat16 ret{a};
+  return ret;
+}
+
 /**
  * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
  * \brief Convert float2 to __hip_bfloat162
diff --git a/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/hipamd/include/hip/amd_detail/amd_hip_fp16.h
index 994478d23e..e4a694a5d5 100644
--- a/hipamd/include/hip/amd_detail/amd_hip_fp16.h
+++ b/hipamd/include/hip/amd_detail/amd_hip_fp16.h
@@ -783,7 +783,7 @@ THE SOFTWARE.
 
             // int -> half
             inline
-            __device__
+            __HOST_DEVICE__
             __half __int2half_rn(int x)
             {
                 return __half_raw{static_cast<_Float16>(x)};