From ae12867774dbc10120220ecbfda4d99e4a19eea0 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Fri, 17 Jun 2016 15:05:33 -0500
Subject: [PATCH] added tests for host math functions

Change-Id: I66a5c574a27190e32054586f07ecf20e1ff71292
---
 include/hcc_detail/hip_runtime.h |  26 +++
 src/device_util.cpp              |  33 ++-
 tests/src/hipTestHost.cpp        | 357 ++++++++++++++++++++++++++++++-
 3 files changed, 413 insertions(+), 3 deletions(-)

diff --git a/include/hcc_detail/hip_runtime.h b/include/hcc_detail/hip_runtime.h
index aa634cdf1e..b0b6ecfe6c 100644
--- a/include/hcc_detail/hip_runtime.h
+++ b/include/hcc_detail/hip_runtime.h
@@ -167,19 +167,28 @@ __device__ float nanf(const char* tagp);
 __device__ float nearbyintf(float x);
 __device__ float nextafterf(float x, float y);
 __device__ float norm3df(float a, float b, float c);
+__host__ float norm3df(float a, float b, float c);
 __device__ float norm4df(float a, float b, float c, float d);
+__host__ float norm4df(float a, float b, float c, float d);
 __device__ float normcdff(float y);
+__host__ float normcdff(float y);
 __device__ float normcdfinvf(float y);
+__host__ float normcdfinvf(float y);
 __device__ float normf(int dim, const float *a);
 __device__ float powf(float x, float y);
 __device__ float rcbrtf(float x);
+__host__ float rcbrtf(float x);
 __device__ float remainderf(float x, float y);
 __device__ float remquof(float x, float y, int *quo);
 __device__ float rhypotf(float x, float y);
+__host__ float rhypotf(float x, float y);
 __device__ float rintf(float x);
 __device__ float rnorm3df(float a, float b, float c);
+__host__ float rnorm3df(float a, float b, float c);
 __device__ float rnorm4df(float a, float b, float c, float d);
+__host__ float rnorm4df(float a, float b, float c, float d);
 __device__ float rnormf(int dim, const float* a);
+__host__ float rnormf(int dim, const float* a);
 __device__ float roundf(float x);
 __device__ float rsqrtf(float x);
 __device__ float scalblnf(float x, long int n);
@@ -187,6 +196,7 @@ __device__ float scalbnf(float x, int n);
 __host__ __device__ unsigned signbit(float a);
 __device__ void sincosf(float x, float *sptr, float *cptr);
 __device__ void sincospif(float x, float *sptr, float *cptr);
+__host__ void sincospif(float x, float *sptr, float *cptr);
 __device__ float sinf(float x);
 __device__ float sinhf(float x);
 __device__ float sinpif(float x);
@@ -230,6 +240,7 @@ __device__ double exp2(double x);
 __device__ double expm1(double x);
 __device__ double fabs(double x);
 __device__ double fdim(double x, double y);
+__device__ double fdivide(double x, double y);
 __device__ double floor(double x);
 __device__ double fma(double x, double y, double z);
 __device__ double fmax(double x, double y);
@@ -261,18 +272,27 @@ __device__ double nearbyint(double x);
 __device__ double nextafter(double x, double y);
 __device__ double norm(int dim, const double* t);
 __device__ double norm3d(double a, double b, double c);
+__host__ double norm3d(double a, double b, double c);
 __device__ double norm4d(double a, double b, double c, double d);
+__host__ double norm4d(double a, double b, double c, double d);
 __device__ double normcdf(double y);
+__host__ double normcdf(double y);
 __device__ double normcdfinv(double y);
+__host__ double normcdfinv(double y);
 __device__ double pow(double x, double y);
 __device__ double rcbrt(double x);
+__host__ double rcbrt(double x);
 __device__ double remainder(double x, double y);
 __device__ double remquo(double x, double y, int *quo);
 __device__ double rhypot(double x, double y);
+__host__ double rhypot(double x, double y);
 __device__ double rint(double x);
 __device__ double rnorm(int dim, const double* t);
+__host__ double rnorm(int dim, const double* t);
 __device__ double rnorm3d(double a, double b, double c);
+__host__ double rnorm3d(double a, double b, double c);
 __device__ double rnorm4d(double a, double b, double c, double d);
+__host__ double rnorm4d(double a, double b, double c, double d);
 __device__ double round(double x);
 __host__ __device__ double rsqrt(double x);
 __device__ double scalbln(double x, long int n);
@@ -281,6 +301,7 @@ __host__ __device__ unsigned signbit(double a);
 __device__ double sin(double a);
 __device__ void sincos(double x, double *sptr, double *cptr);
 __device__ void sincospi(double x, double *sptr, double *cptr);
+__host__ void sincospi(double x, double *sptr, double *cptr);
 __device__ double sinh(double x);
 __host__ __device__ double sinpi(double x);
 __device__ double sqrt(double x);
@@ -292,6 +313,11 @@ __device__ double y0(double x);
 __device__ double y1(double y);
 __device__ double yn(int n, double x);
 
+__host__ double erfcinv(double y);
+__host__ double erfcx(double x);
+__host__ double erfinv(double y);
+__host__ double fdivide(double x, double y);
+
 // TODO - hipify-clang - change to use the function call.
 //#define warpSize hc::__wavesize()
 extern const int warpSize;
diff --git a/src/device_util.cpp b/src/device_util.cpp
index c51f23b056..a307a1e377 100644
--- a/src/device_util.cpp
+++ b/src/device_util.cpp
@@ -1254,6 +1254,10 @@ __device__ double fdim(double x, double y)
 {
     return hc::precise_math::fdim(x, y);
 }
+__device__ double fdivide(double x, double y)
+{
+    return x/y;
+}
 __device__ double floor(double x)
 {
     return hc::precise_math::floor(x);
@@ -2884,6 +2888,11 @@ __host__ double erfinv(double x)
     return __hip_host_erfinv(x);
 }
 
+__host__ double fdivide(double x, double y)
+{
+    return x/y;
+}
+
 __host__ float normcdff(float t)
 {
      return (1 - std::erf(-t/std::sqrt(2)))/2;
@@ -2951,7 +2960,7 @@ __host__ float rnormf(int dim, const float *t)
     {
         val = val + t[i] * t[i];
     }
-    return 1 / val;
+    return 1 / std::sqrt(val);
 }
 
 __host__ double rnorm(int dim, const double *t)
@@ -2961,7 +2970,7 @@ __host__ double rnorm(int dim, const double *t)
     {
         val = val + t[i] * t[i];
     }
-    return 1 / val;
+    return 1 / std::sqrt(val);
 }
 
 __host__ float rnorm4df(float a, float b, float c, float d)
@@ -3015,3 +3024,23 @@ __host__ double nextafter(double x, double y)
 {
     return std::nextafter(x, y);
 }
+
+__host__ float norm3df(float a, float b, float c)
+{
+    return std::sqrt(a*a + b*b + c*c);
+}
+
+__host__ float norm4df(float a, float b, float c, float d)
+{
+    return std::sqrt(a*a + b*b + c*c + d*d);
+}
+
+__host__ double norm3d(double a, double b, double c)
+{
+    return std::sqrt(a*a + b*b + c*c);
+}
+
+__host__ double norm4d(double a, double b, double c, double d)
+{
+    return std::sqrt(a*a + b*b + c*c + d*d);
+}
diff --git a/tests/src/hipTestHost.cpp b/tests/src/hipTestHost.cpp
index cf912e577a..e1207f43e0 100644
--- a/tests/src/hipTestHost.cpp
+++ b/tests/src/hipTestHost.cpp
@@ -62,6 +62,350 @@ bool check_erfinvf()
     return true;
 }
 
+bool check_fdividef()
+{
+    uint32_t len = 4;
+    float Val[] = {0, -0.5, 0.9, -0.2};
+    float Out[] = {1, -0.4769, 1.1631, -0.1791};
+    for(int i=0;i<len;i++){
+        if(Val[i]/Out[i] - fdividef(Val[i], Out[i]) > 0.0001){
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_erfcinv(){
+    uint32_t len = 4;
+    double Val[] = {0.1, 1.2, 1, 0.9};
+    double Out[] = {1.16309, -0.179144, 0, 0.0889};
+    for(int i=0;i<len;i++){
+        if(Out[i] - erfcinv(Val[i]) > 0.0001)
+        {
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_erfcx(){
+    uint32_t len = 4;
+    double Val[] = {-0.5, 15, 3.2, 1};
+    double Out[] = {1.9524, 0.0375, 0.1687, 0.4276};
+    for(int i=0;i<len;i++){
+        if(Out[i] - erfcx(Val[i]) > 0.0001)
+        {
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_erfinv()
+{
+    uint32_t len = 4;
+    double Val[] = {0, -0.5, 0.9, -0.2};
+    double Out[] = {0, -0.4769, 1.1631, -0.1791};
+    for(int i=0;i<len;i++){
+        if(Out[i] - erfinv(Val[i]) > 0.0001){
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_fdivide()
+{
+    uint32_t len = 4;
+    double Val[] = {0, -0.5, 0.9, -0.2};
+    double Out[] = {1, -0.4769, 1.1631, -0.1791};
+    for(int i=0;i<len;i++){
+        if(Val[i]/Out[i] - fdivide(Val[i], Out[i]) > 0.0001){
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_modff()
+{
+    uint32_t len = 4;
+    float Val[] = {0, -0.5, 0.9, -0.2};
+    float iPtr[] = {0, 0, 0, 0};
+    float frac[] = {0, -0.5, 0.9, -0.2};
+    float Out[] = {1, 1, 1, 1};
+    for(int i=0;i<len;i++){
+        if(frac[i] - modff(Val[i], Out+i) > 0.0001 && iPtr[i] == Out[i]){
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_modf()
+{
+    uint32_t len = 4;
+    double Val[] = {0, -0.5, 0.9, -0.2};
+    double iPtr[] = {0, 0, 0, 0};
+    double frac[] = {0, -0.5, 0.9, -0.2};
+    double Out[] = {1, 1, 1, 1};
+    for(int i=0;i<len;i++){
+        if(frac[i] - modf(Val[i], Out+i) > 0.0001 && iPtr[i] == Out[i]){
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_nextafterf()
+{
+    uint32_t len = 4;
+    float Val[] = {0, -0.5, 0.9, -0.2};
+    float iPtr[] = {0, 0, 0, 0};
+    float frac[] = {0, -0.5, 0.9, -0.2};
+    float Out[] = {1, 1, 1, 1};
+    for(int i=0;i<len;i++){
+        if(nextafterf(Val[i],1) - Val[i] > 0.0001){
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_nextafter()
+{
+    uint32_t len = 4;
+    double Val[] = {0, -0.5, 0.9, -0.2};
+    double iPtr[] = {0, 0, 0, 0};
+    double frac[] = {0, -0.5, 0.9, -0.2};
+    double Out[] = {1, 1, 1, 1};
+    for(int i=0;i<len;i++){
+        if(nextafter(Val[i],1) - Val[i] > 0.0001){
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_norm3df(float *A)
+{
+    float f = norm3df(A[0], A[1], A[2]);
+    float out = sqrt(A[0]*A[0] + A[1]*A[1] + A[2]*A[2]);
+    if(f - out > 0.0001)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_norm3d(double *A)
+{
+    double f = norm3d(A[0], A[1], A[2]);
+    double out = sqrt(A[0]*A[0] + A[1]*A[1] + A[2]*A[2]);
+    if(f - out > 0.0001)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_norm4df(float *A)
+{
+    float f = norm4df(A[0], A[1], A[2], A[3]);
+    float out = sqrt(A[0]*A[0] + A[1]*A[1] + A[2]*A[2] + A[3]*A[3]);
+    if(f - out > 0.0001)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_norm4d(double *A)
+{
+    double f = norm4d(A[0], A[1], A[2], A[3]);
+    double out = sqrt(A[0]*A[0] + A[1]*A[1] + A[2]*A[2] + A[3]*A[3]);
+    if(f - out > 0.0001)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_normcdff(){
+    uint32_t len = 2;
+    float Val[] = {0,1};
+    float Out[] = {0.5, 0.8413};
+    for(int i=0;i<len;i++){
+        if(Out[i] - normcdff(Val[i]) > 0.0001)
+        {
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_normcdf(){
+    uint32_t len = 2;
+    float Val[] = {0,1};
+    float Out[] = {0.5, 0.8413};
+    for(int i=0;i<len;i++){
+        if(Out[i] - normcdf(Val[i]) > 0.0001)
+        {
+            return false;
+        }
+    }
+    return true;
+}
+
+
+bool check_normcdfinvf(){
+    uint32_t len = 2;
+    double Val[] = {0.5, 0.8413};
+    for(int i=0;i<len;i++){
+        if(Val[i] - normcdfinvf(normcdff(Val[i])) > 0.0001)
+        {
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_normcdfinv(){
+    uint32_t len = 2;
+    double Val[] = {0.5, 0.8413};
+    for(int i=0;i<len;i++){
+        if(Val[i] - normcdfinv(normcdf(Val[i])) > 0.0001)
+        {
+            return false;
+        }
+    }
+    return true;
+}
+
+bool check_rcbrtf()
+{
+    float f = 1.0f;
+    if(rcbrtf(f) != 1.0f)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_rcbrt()
+{
+    double f = 1.0;
+    if(rcbrt(f) != 1.0)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_rhypotf()
+{
+    float f = 1.0f;
+    float g = 2.0f;
+    float val = rhypotf(f, g);
+    float sq = f*f + g*g ;
+    if(1/(val*val) - sq > 0.0001)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_rhypot()
+{
+    double f = 1.0f;
+    double g = 2.0f;
+    double val = rhypot(f, g);
+    double sq = f*f + g*g;
+    if(1/(val*val) - sq > 0.0001)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_rnorm3df(float *A)
+{
+    float f = rnorm3df(A[0], A[1], A[2]);
+    float out = sqrt(A[0]*A[0] + A[1]*A[1] + A[2]*A[2]);
+    if(f - 1/out > 0.0001)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_rnorm3d(double *A)
+{
+    double f = rnorm3d(A[0], A[1], A[2]);
+    double out = sqrt(A[0]*A[0] + A[1]*A[1] + A[2]*A[2]);
+    if(f - 1/out > 0.0001)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_rnorm4df(float *A)
+{
+    float f = rnorm4df(A[0], A[1], A[2], A[3]);
+    float out = sqrt(A[0]*A[0] + A[1]*A[1] + A[2]*A[2] + A[3]*A[3]);
+    if(f - 1/out > 0.0001)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_rnorm4d(double *A)
+{
+    double f = rnorm4d(A[0], A[1], A[2], A[3]);
+    double out = sqrt(A[0]*A[0] + A[1]*A[1] + A[2]*A[2] + A[3]*A[3]);
+    if(f - 1/out > 0.0001)
+    {
+        return false;
+    }
+    return true;
+}
+
+bool check_rnormf(float *A)
+{
+    return (rnorm3df(A[0],A[1],A[2]) - rnormf(3,A) < 0.0001) && (rnorm4df(A[0],A[1],A[2],A[3]) - rnormf(4, A) < 0.0001 );
+}
+
+bool check_rnorm(double *A)
+{
+    return (rnorm3d(A[0],A[1],A[2]) - rnorm(3,A) < 0.0001) && (rnorm4d(A[0],A[1],A[2],A[3]) - rnorm(4, A) < 0.0001 );
+}
+
+bool check_sincospif()
+{
+    float s1, c1, s2, c2;
+    float in1 = 1, in2 = 0.5;
+    sincospif(in1, &s1, &c1);
+    sincospif(in2, &s2, &c2);
+    if( (s1 - 0 < 0.00001) && (s2 - 1 < 0.00001) && (c1 + 1 < 0.00001) && (c2 - 0 < 0.00001)){
+        return true;
+    }
+    return false;
+}
+
+bool check_sincospi()
+{
+    double s1, c1, s2, c2;
+    double in1 = 1, in2 = 0.5;
+    sincospi(in1, &s1, &c1);
+    sincospi(in2, &s2, &c2);
+    if( (s1 - 0 < 0.00001) && (s2 - 1 < 0.00001) && (c1 + 1 < 0.00001) && (c2 - 0 < 0.00001)){
+        return true;
+    }
+    return false;
+}
+
 int main(){
     float *Af = new float[N];
     double *A = new double[N];
@@ -69,7 +413,18 @@ int main(){
         Af[i] = i * 1.0f;
         A[i] = i * 1.0;
     }
-    if(check_erfcinvf() && check_erfcxf() && check_erfcinvf()){
+    if(check_erfcinvf() && check_erfcxf() && check_erfcinvf() &&
+        check_erfcinv() && check_erfcx() && check_erfcinv() &&
+        check_fdividef() && check_fdivide() && check_modff() &&
+        check_modf() && check_nextafterf() && check_norm3df(Af) &&
+        check_norm3d(A) && check_norm4df(Af) && check_norm4d(A) &&
+        check_normcdff() && check_normcdf() && check_normcdfinvf() &&
+        check_normcdfinv() && check_rcbrtf() && check_rcbrt() && 
+        check_rhypotf() && check_rhypot() && check_rnorm3df(Af) &&
+        check_rnorm3d(A) && check_rnorm4df(Af) && check_rnorm4d(A) &&
+        check_rnormf(Af) && check_rnorm(A) && check_sincospif() &&
+        check_sincospi()
+        ){
         passed();
     }