diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h index c54aef319b..eb63c7cd60 100644 --- a/projects/hip/include/hip/hcc_detail/hip_fp16.h +++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h @@ -709,6 +709,22 @@ __device__ static inline __half2 h2cos(const __half2 h) { return a; } +__device__ static inline __half2 h2exp(const __half2 h) { + __half2 factor; + factor.p[0] = 1.442694; + factor.p[1] = 1.442694; + factor.q = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.q, factor.q)); + return factor; +} + +__device__ static inline __half2 h2exp10(const __half2 h) { + __half2 factor; + factor.p[0] = 3.3219281; + factor.p[1] = 3.3219281; + factor.q = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.q, factor.q)); + return factor; +} + __device__ static inline __half2 h2exp2(const __half2 h) { __half2 a; a.q = __hip_hc_ir_h2exp2_int(h.q); @@ -721,6 +737,21 @@ __device__ static inline __half2 h2floor(const __half2 h) { return a; } +__device__ static inline __half2 h2log(const __half2 h) { + __half2 factor; + factor.p[0] = 0.693147; + factor.p[1] = 0.693147; + factor. q = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.q), factor.q); + return factor; +} + +__device__ static inline __half2 h2log10(const __half2 h) { + __half2 factor; + factor.p[0] = 0.301029; + factor.p[1] = 0.301029; + factor.q = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.q), factor.q); + return factor; +} __device__ static inline __half2 h2log2(const __half2 h) { __half2 a; a.q = __hip_hc_ir_h2log2_int(h.q); diff --git a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp index c194a42ca2..975a494b2a 100644 --- a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp +++ b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp @@ -90,12 +90,12 @@ __global__ void CheckHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __half2 Out[9] = h2div(In1[9], In2[9]); Out[10] = h2ceil(In1[10]); Out[11] = h2cos(In1[11]); -// Out[12] = h2exp(In1[12]); -// Out[13] = h2exp10(In1[13]); + Out[12] = h2exp(In1[12]); + Out[13] = h2exp10(In1[13]); Out[14] = h2exp2(In1[14]); Out[15] = h2floor(In1[15]); -// Out[16] = h2log(In1[16]); -// Out[17] = h2log10(In1[17]); + Out[16] = h2log(In1[16]); + Out[17] = h2log10(In1[17]); Out[18] = h2log2(In1[18]); Out[19] = h2rcp(In1[19]); // Out[20] = h2rint(In1[20]);