Checkin to fix bugs in math functions.

This change fixes the following bugs that were discovered while debuggnig TF unit test failures (cwise_ops_test)

1. __hisinf and __hisnan routines
   Both had incorrect implementations.

2. abs
   A "long long" (64bit int) version was missing, resulting in the 32bit version being used for 64bit ints (which resulted in incorrect results, when the value passed in was outside the 32bit int range)

3. lgamma
  We seemed to have a custom version for the 'double' datatype (which was giving incorrect results). Replaced it with a call to the 'double' version of the underlying 'hc::precision_math::lgamma'

hipamd/src/hip_fp16.cpp

@@ -29,10 +29,6 @@ struct hipHalfHolder {
     };
 };
 
-#define HINF 65504
-
-__device__ static struct hipHalfHolder __hInfValue = {HINF};
-
 __device__ __half __hadd(__half a, __half b) { return a + b; }
 
 __device__ __half __hadd_sat(__half a, __half b) { return a + b; }
@@ -63,9 +59,21 @@ __device__ bool __hge(__half a, __half b) { return a >= b ? true : false; }
 
 __device__ bool __hgt(__half a, __half b) { return a > b ? true : false; }
 
-__device__ bool __hisinf(__half a) { return a == HINF ? true : false; }
+__device__ bool __hisinf(__half a) {
+  hipHalfHolder hH;
+  hH.h = a;
+  // mask with 0x7fff to drop the sign bit
+  // 0x7c00 is bit pattern for inf (exp = 11111, significand = 0)
+  return ((hH.s & 0x7fff) == 0x7c00) ? true : false;
+}
 
-__device__ bool __hisnan(__half a) { return a > HINF ? true : false; }
+__device__ bool __hisnan(__half a) {
+  hipHalfHolder hH;
+  hH.h = a;
+  // mask with 0x7fff to drop the sign bit
+  // 0x7cXX is bit pattern for inf (exp = 11111, significand = 0)
+  return ((hH.s & 0x7fff) > 0x7c00) ? true : false;
+}
 
 __device__ bool __hle(__half a, __half b) { return a <= b ? true : false; }
 
@@ -124,8 +132,8 @@ __device__ __half2 __hgt2(__half2 a, __half2 b) {
 
 __device__ __half2 __hisnan2(__half2 a) {
     __half2 c;
-    c.x = (a.x > HINF) ? (__half)1 : (__half)0;
-    c.y = (a.y > HINF) ? (__half)1 : (__half)0;
+    c.x = (__hisnan(a.x)) ? (__half)1 : (__half)0;
+    c.y = (__hisnan(a.y)) ? (__half)1 : (__half)0;
     return c;
 }