diff --git a/tests/src/deviceLib/hipMathFunctions.cpp b/tests/src/deviceLib/hipMathFunctions.cpp
index ba3707bc59..ff69087f82 100644
--- a/tests/src/deviceLib/hipMathFunctions.cpp
+++ b/tests/src/deviceLib/hipMathFunctions.cpp
@@ -20,12 +20,150 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
 
 #include "hip/hip_runtime.h"
 #include "test_common.h"
 
+#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__
+
+__global__ void kernel_abs_int64(hipLaunchParm lp, long long *input, long long *output) {
+    int tx = threadIdx.x;
+    output[tx] = abs(input[tx]);
+}
+
+__global__ void kernel_lgamma_double(hipLaunchParm lp, double *input, double *output) {
+    int tx = threadIdx.x;
+    output[tx] = lgamma(input[tx]);
+}
+
+#endif
+
+#define CHECK_LGAMMA_DOUBLE(IN, OUT, EXP)		 \
+  {						 \
+    if (OUT != EXP)  {				 \
+      failed("check_abs_int64 failed on %f (output = %f, expected = %fd)\n", IN, OUT, EXP); \
+    }						 \
+  }
+
+#define CHECK_ABS_INT64(IN, OUT, EXP)		 \
+  {						 \
+    if (OUT != EXP)  {				 \
+      failed("check_abs_int64 failed on %lld (output = %lld, expected = %lld)\n", IN, OUT, EXP); \
+    }						 \
+  }
+
+void check_lgamma_double() {
+
+  using datatype_t = double;
+  
+  const int NUM_INPUTS = 8;
+  auto memsize = NUM_INPUTS * sizeof(datatype_t);
+  
+  // allocate memories
+  datatype_t *inputCPU = (datatype_t *) malloc(memsize);
+  datatype_t *outputCPU = (datatype_t *) malloc(memsize);
+  datatype_t *inputGPU = nullptr; hipMalloc((void**)&inputGPU, memsize);
+  datatype_t *outputGPU = nullptr; hipMalloc((void**)&outputGPU, memsize);
+  
+  // populate input
+  inputCPU[0] = -3.5;
+  inputCPU[0] = -2.5;
+  inputCPU[0] = -1.5;
+  inputCPU[0] = -0.5;
+  inputCPU[0] =  0.5;
+  inputCPU[0] =  1.5;
+  inputCPU[0] =  2.5;
+  inputCPU[0] =  3.5;
+  
+  // copy inputs to device
+  hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice);
+
+  // launch kernel
+  hipLaunchKernel(kernel_lgamma_double, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
+
+  // copy outputs from device
+  hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost);
+
+  // check outputs
+  for (int i=0; i<NUM_INPUTS; i++) {
+    CHECK_LGAMMA_DOUBLE(inputCPU[i], outputCPU[i], lgamma(inputCPU[i]));
+  }
+  
+  // free memories
+  hipFree(inputGPU);
+  hipFree(outputGPU);
+  free(inputCPU);
+  free(outputCPU);
+  
+  // done
+  return;
+}
+
+
+void check_abs_int64() {
+
+  using datatype_t = long long;
+
+  const int NUM_INPUTS = 8;
+  auto memsize = NUM_INPUTS * sizeof(datatype_t);
+  
+  // allocate memories
+  datatype_t *inputCPU = (datatype_t *) malloc(memsize);
+  datatype_t *outputCPU = (datatype_t *) malloc(memsize);
+  datatype_t *inputGPU = nullptr; hipMalloc((void**)&inputGPU, memsize);
+  datatype_t *outputGPU = nullptr; hipMalloc((void**)&outputGPU, memsize);
+  
+  // populate input
+  inputCPU[0] = -81985529216486895ll; 
+  inputCPU[1] =  81985529216486895ll;
+  inputCPU[2] = -1250999896491ll;
+  inputCPU[3] =  1250999896491ll;
+  inputCPU[4] = -19088743ll;
+  inputCPU[5] =  19088743ll;
+  inputCPU[6] = -291ll;
+  inputCPU[7] =  291ll;
+  
+  // copy inputs to device
+  hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice);
+
+  // launch kernel
+  hipLaunchKernel(kernel_abs_int64, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
+
+  // copy outputs from device
+  hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost);
+
+  // check outputs
+  CHECK_ABS_INT64(inputCPU[0], outputCPU[0], outputCPU[1]);
+  CHECK_ABS_INT64(inputCPU[1], outputCPU[1], outputCPU[1]);
+  CHECK_ABS_INT64(inputCPU[2], outputCPU[2], outputCPU[3]);
+  CHECK_ABS_INT64(inputCPU[3], outputCPU[3], outputCPU[3]);
+  CHECK_ABS_INT64(inputCPU[4], outputCPU[4], outputCPU[5]);
+  CHECK_ABS_INT64(inputCPU[5], outputCPU[5], outputCPU[5]);
+  CHECK_ABS_INT64(inputCPU[6], outputCPU[6], outputCPU[7]);
+  CHECK_ABS_INT64(inputCPU[7], outputCPU[7], outputCPU[7]);
+  
+  // free memories
+  hipFree(inputGPU);
+  hipFree(outputGPU);
+  free(inputCPU);
+  free(outputCPU);
+  
+  // done
+  return;
+}
+  
+
 int main(int argc, char* argv[]) {
     HipTest::parseStandardArguments(argc, argv, true);
 
+    check_abs_int64();
+
+    // check_lgamma_double();
+    
     passed();
 }
diff --git a/tests/src/deviceLib/hipTestHalf.cpp b/tests/src/deviceLib/hipTestHalf.cpp
index 9e5962f38f..4a63260bf7 100644
--- a/tests/src/deviceLib/hipTestHalf.cpp
+++ b/tests/src/deviceLib/hipTestHalf.cpp
@@ -17,6 +17,12 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
 #include <iostream>
 #include <hip/hip_fp16.h>
 #include "hip/hip_runtime.h"
@@ -59,8 +65,143 @@ __global__ void __half2Math(hipLaunchParm lp, __half2* A, __half2* B, __half2* C
     c = __hmul2_sat(b, c);
 }
 
+__global__ void kernel_hisnan(hipLaunchParm lp, __half* input, int* output) {
+    int tx = threadIdx.x;
+    output[tx] = __hisnan(input[tx]);
+}
+
+__global__ void kernel_hisinf(hipLaunchParm lp, __half* input, int* output) {
+    int tx = threadIdx.x;
+    output[tx] = __hisinf(input[tx]);
+}
+
 #endif
 
+
+__half host_ushort_as_half(unsigned short s) {
+  union {__half h; unsigned short s; } converter;
+  converter.s = s;
+  return converter.h;
+}
+
+
+void check_hisnan(int NUM_INPUTS, __half* inputCPU, __half* inputGPU) {
+
+  // allocate memory
+  auto memsize = NUM_INPUTS * sizeof(int);
+  int* outputGPU = nullptr;
+  hipMalloc((void**)&outputGPU, memsize);
+
+  // launch the kernel
+  hipLaunchKernel(kernel_hisnan, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
+
+  // copy output from device
+  int* outputCPU = (int*) malloc(memsize);
+  hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost);
+
+  // check output
+  for (int i=0; i<NUM_INPUTS; i++) {
+    if ((2 <= i) && (i <= 5)) { // inputs are nan, output should be true
+      if (outputCPU[i] == 0) {
+	failed("__hisnan() returned false for %f (input idx = %d)\n", inputCPU[i], i);
+      }
+    }
+    else { // inputs are NOT nan, output should be false
+      if (outputCPU[i] != 0) {
+	failed("__hisnan() returned true for %f (input idx = %d)\n", inputCPU[i], i);
+      }
+    }
+  }
+
+  // free memory
+  free(outputCPU);
+  hipFree(outputGPU);
+
+  // done
+  return;
+}
+
+
+void check_hisinf(int NUM_INPUTS, __half* inputCPU, __half* inputGPU) {
+  // allocate memory
+  auto memsize = NUM_INPUTS * sizeof(int);
+  int* outputGPU = nullptr;
+  hipMalloc((void**)&outputGPU, memsize);
+
+  // launch the kernel
+  hipLaunchKernel(kernel_hisinf, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
+
+  // copy output from device
+  int* outputCPU = (int*) malloc(memsize);
+  hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost);
+
+  // check output
+  for (int i=0; i<NUM_INPUTS; i++) {
+    if ((0 <= i) && (i <= 1)) { // inputs are inf, output should be true
+      if (outputCPU[i] == 0) {
+	failed("__hisinf() returned false for %f (input idx = %d)\n", inputCPU[i], i);
+      }
+    }
+    else { // inputs are NOT inf, output should be false
+      if (outputCPU[i] != 0) {
+	failed("__hisinf() returned true for %f (input idx = %d)\n", inputCPU[i], i);
+      }
+    }
+  }
+
+  // free memory
+  free(outputCPU);
+  hipFree(outputGPU);
+
+  // done
+  return;
+}
+
+
+void checkFunctional() {
+
+  // allocate memory 
+  const int NUM_INPUTS = 16;
+  auto memsize = NUM_INPUTS * sizeof(__half);
+  __half* inputCPU = (__half*) malloc(memsize);
+  
+  // populate inputs
+  inputCPU[0] = host_ushort_as_half(0x7c00);  //  inf
+  inputCPU[1] = host_ushort_as_half(0xfc00);  // -inf
+  inputCPU[2] = host_ushort_as_half(0x7c01);  //  nan
+  inputCPU[3] = host_ushort_as_half(0x7e00);  //  nan
+  inputCPU[4] = host_ushort_as_half(0xfc01);  //  nan
+  inputCPU[5] = host_ushort_as_half(0xfe00);  //  nan
+  inputCPU[6] = host_ushort_as_half(0x0000);  //  0
+  inputCPU[7] = host_ushort_as_half(0x8000);  // -0
+  inputCPU[8] = host_ushort_as_half(0x7bff);  // max +ve normal
+  inputCPU[9] = host_ushort_as_half(0xfbff);  // max -ve normal
+  inputCPU[10] = host_ushort_as_half(0x0400); // min +ve normal
+  inputCPU[11] = host_ushort_as_half(0x8400); // min -ve normal
+  inputCPU[12] = host_ushort_as_half(0x03ff); // max +ve sub-normal
+  inputCPU[13] = host_ushort_as_half(0x83ff); // max -ve sub-normal
+  inputCPU[14] = host_ushort_as_half(0x0001); // min +ve sub-normal
+  inputCPU[15] = host_ushort_as_half(0x8001); // min -ve sub-normal
+
+  // copy inputs to the GPU
+  __half* inputGPU = nullptr;
+  hipMalloc((void**)&inputGPU, memsize);
+  hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice);
+
+  // run checks
+
+  check_hisnan(NUM_INPUTS, inputCPU, inputGPU);
+
+  check_hisinf(NUM_INPUTS, inputCPU, inputGPU);
+
+  // free memory
+  hipFree(inputGPU);
+  free(inputCPU);
+
+  // all done
+  return;
+}
+
 int main() {
     __half *A, *B, *C;
     hipMalloc(&A, HALF_SIZE);
@@ -78,5 +219,9 @@ int main() {
     hipFree(A2);
     hipFree(B2);
     hipFree(C2);
+
+    // run some functional checks
+    checkFunctional();
+    
     passed();
 }