AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-529927 - add missing operations for fp16/bf16 (#238)

Browse Source
This commit is contained in:
Chaudhary, Jatin Jaikishan
2025-05-06 09:41:21 +01:00
committed by GitHub
parent 65a0181a7c
commit b1ebf33850
2 changed files with 13 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/amd_detail/amd_hip_bf16.h
+12 -3
View File
@@ -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
hipamd/include/hip/amd_detail/amd_hip_fp16.h
+1 -1
View File
@@ -783,7 +783,7 @@ THE SOFTWARE.
// int -> half
inline
__device__
__HOST_DEVICE__
__half __int2half_rn(int x)
{
return __half_raw{static_cast<_Float16>(x)};