Adding Half Abs APIs (#1902)

hipamd/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)

hipamd/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);

hipamd/tests/src/deviceLib/hipTestHalf.cpp

@@ -96,6 +96,18 @@ void kernel_hisinf(__half* input, int* output) {
   output[tx] = __hisinf(input[tx]);
 }
 
+__global__ void testHalfAbs(float* p) {
+    auto a = __float2half(*p);
+    a = __habs(a);
+    *p = __half2float(a);
+}
+
+__global__ void testHalf2Abs(float2* p) {
+    auto a = __float22half2_rn(*p);
+    a = __habs2(a);
+    *p = __half22float2(a);
+}
+
 #endif
 
 
@@ -237,6 +249,31 @@ void checkFunctional() {
   return;
 }
 
+void checkHalfAbs() {
+    {
+        float *p;
+        hipMalloc(&p, sizeof(float));
+        float pp = -2.1f;
+        hipMemcpy(p, &pp, sizeof(float), hipMemcpyDefault);
+        hipLaunchKernelGGL(testHalfAbs, 1, 1, 0, 0, p);
+        hipMemcpy(&pp, p, sizeof(float), hipMemcpyDefault);
+        hipFree(p);
+        if(pp < 0.0f) { failed("Half Abs failed"); }
+    }
+    {
+        float2 *p;
+        hipMalloc(&p, sizeof(float2));
+        float2 pp;
+        pp.x = -2.1f;
+        pp.y = -1.1f;
+        hipMemcpy(p, &pp, sizeof(float2), hipMemcpyDefault);
+        hipLaunchKernelGGL(testHalf2Abs, 1, 1, 0, 0, p);
+        hipMemcpy(&pp, p, sizeof(float2), hipMemcpyDefault);
+        hipFree(p);
+        if(pp.x < 0.0f || pp.y < 0.0f) { failed("Half2 Abs Test Failed"); }
+    }
+}
+
 int main() {
   bool* result{nullptr};
   hipMemAllocHost((void**)&result, sizeof(result));
@@ -260,5 +297,7 @@ int main() {
   // run some functional checks
   checkFunctional();
 
+  checkHalfAbs();
+
   passed();
 }