Adding Half Abs APIs (#1902)

include/hip/hcc_detail/hip_fp16.h

@@ -1268,6 +1268,13 @@ THE SOFTWARE.
                     static_cast<__half_raw>(x).data +
                     static_cast<__half_raw>(y).data};
             }
+	    inline
+	    __device__
+	    __half __habs(__half x)
+	    {
+	        return __half_raw{
+		    __ocml_fabs_f16(static_cast<__half_raw>(x).data)};
+	    }
             inline
             __device__
             __half __hsub(__half x, __half y)
@@ -1334,6 +1341,13 @@ THE SOFTWARE.
                     static_cast<__half2_raw>(x).data +
                     static_cast<__half2_raw>(y).data};
             }
+	    inline
+	    __device__
+	    __half2 __habs2(__half2 x)
+	    {
+	        return __half2_raw{
+		    __ocml_fabs_2f16(static_cast<__half2_raw>(x).data)};
+	    }
             inline
             __device__
             __half2 __hsub2(__half2 x, __half2 y)

include/hip/hcc_detail/hip_fp16_math_fwd.h

@@ -38,6 +38,7 @@ extern "C"
     __device__ __attribute__((const)) _Float16 __ocml_floor_f16(_Float16);
     __device__ __attribute__((const))
     _Float16 __ocml_fma_f16(_Float16, _Float16, _Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_fabs_f16(_Float16);
     __device__ __attribute__((const)) int __ocml_isinf_f16(_Float16);
     __device__ __attribute__((const)) int __ocml_isnan_f16(_Float16);
     __device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
@@ -58,6 +59,7 @@ extern "C"
     #endif
 
     __device__ __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_fabs_2f16(__2f16);
     __device__ __2f16 __ocml_cos_2f16(__2f16);
     __device__ __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
     __device__ __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16);