From 37e5903a68f550d55583f69517debdca60a2a73f Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Mon, 13 Jun 2016 11:55:12 -0500
Subject: [PATCH] added more device functions

Change-Id: I191919060b393772ee442cc19d83479217c5a4ce


[ROCm/hip commit: 9ac6e422f08dfe93c53b888a7a27f6a0c1d4f563]
---
 projects/hip/include/hcc_detail/hip_runtime.h |  2 +-
 projects/hip/src/device_util.cpp              | 63 ++++++++++++++++---
 2 files changed, 57 insertions(+), 8 deletions(-)

diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h
index fbe1ffce15..dbe75b1a9d 100644
--- a/projects/hip/include/hcc_detail/hip_runtime.h
+++ b/projects/hip/include/hcc_detail/hip_runtime.h
@@ -169,7 +169,7 @@ __device__ float normcdff(float y);
 __device__ float normcdfinvf(float y);
 __device__ float normf(int dim, const float *a);
 __device__ float powf(float x, float y);
-__device__ float rcbtrf(float x);
+__device__ 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);
diff --git a/projects/hip/src/device_util.cpp b/projects/hip/src/device_util.cpp
index 29a125dce7..b81b6341c0 100644
--- a/projects/hip/src/device_util.cpp
+++ b/projects/hip/src/device_util.cpp
@@ -136,7 +136,10 @@ __device__ float fmodf(float x, float y)
 {
     return hc::precise_math::fmodf(x, y);
 }
-__device__ float frexpf(float x, float y);
+__device__ float frexpf(float x, int *nptr)
+{
+    return hc::precise_math::frexpf(x, nptr);
+}
 __device__ float hypotf(float x, float y)
 {
     return hc::precise_math::hypotf(x, y);
@@ -164,7 +167,10 @@ __device__ float ldexpf(float x, int exp)
 {
     return hc::precise_math::ldexpf(x, exp);
 }
-__device__ float lgammaf(float x);
+__device__ float lgammaf(float x, int *sign)
+{
+    return hc::precise_math::lgammaf(x, sign);
+}
 __device__ long long int llrintf(float x)
 {
     int y = hc::precise_math::roundf(x);
@@ -207,7 +213,10 @@ __device__ long int lroundf(float x)
     long int y = hc::precise_math::roundf(x);
     return y;
 }
-__device__ float modff(float x, float *iptr);
+__device__ float modff(float x, float *iptr)
+{
+    return hc::precise_math::modff(x, iptr);
+}
 __device__ float nanf(const char* tagp)
 {
     return hc::precise_math::nanf((int)*tagp);
@@ -216,7 +225,10 @@ __device__ float nearbyintf(float x)
 {
     return hc::precise_math::nearbyintf(x);
 }
-__device__ float nextafterf(float x, float y);
+__device__ float nextafterf(float x, float y)
+{
+    return hc::precise_math::nextafter(x, y);
+}
 __device__ float norm3df(float a, float b, float c)
 {
      float x = a*a + b*b + c*c;
@@ -250,12 +262,18 @@ __device__ float powf(float x, float y)
 {
     return hc::precise_math::powf(x, y);
 }
-__device__ float rcbtrf(float x);
+__device__ float rcbrtf(float x)
+{
+    return hc::precise_math::rcbrtf(x);
+}
 __device__ float remainderf(float x, float y)
 {
     return hc::precise_math::remainderf(x, y);
 }
-__device__ float remquof(float x, float y, int *quo);
+__device__ float remquof(float x, float y, int *quo)
+{
+    return hc::precise_math::remquof(x, y, quo);
+}
 __device__ float rhypotf(float x, float y)
 {
     return 1/hc::precise_math::hypotf(x, y);
@@ -288,7 +306,10 @@ __device__ float roundf(float x)
 {
     return hc::precise_math::roundf(x);
 }
-__device__ float scalblnf(float x, long int n);
+__device__ float scalblnf(float x, long int n)
+{
+    return hc::precise_math::scalb(x, n);
+}
 __device__ float scalbnf(float x, int n)
 {
     return hc::precise_math::scalbnf(x, n);
@@ -465,6 +486,10 @@ __device__ double fmod(double x, double y)
 {
     return hc::precise_math::fmod(x, y);
 }
+__device__ double frexp(double x, int *y)
+{
+    return hc::precise_math::frexp(x, y);
+}
 __device__ double hypot(double x, double y)
 {
     return hc::precise_math::hypot(x, y);
@@ -489,6 +514,10 @@ __device__ double ldexp(double x, int exp)
 {
     return hc::precise_math::ldexp(x, exp);
 }
+__device__ double lgamma(double x, int *sign)
+{
+    return hc::precise_math::lgamma(x, sign);
+}
 __device__ double log(double x)
 {
     return hc::precise_math::log(x);
@@ -529,6 +558,10 @@ __device__ long int lround(double x)
     long int y = hc::precise_math::round(x);
     return y;
 }
+__device__ double modf(double x, double *iptr)
+{
+    return hc::precise_math::modf(x, iptr);
+}
 __device__ double nan(const char *tagp)
 {
     return hc::precise_math::nan((int)*tagp);
@@ -537,6 +570,10 @@ __device__ double nearbyint(double x)
 {
     return hc::precise_math::nearbyint(x);
 }
+__device__ double nextafter(double x, double y)
+{
+    return hc::precise_math::nextafter(x, y);
+}
 __device__ double norm3d(double a, double b, double c)
 {
     double x = a*a + b*b + c*c;
@@ -556,10 +593,18 @@ __device__ double pow(double x, double y)
 {
     return hc::precise_math::pow(x, y);
 }
+__device__ double rcbrt(double x)
+{
+    return hc::precise_math::rcbrt(x);
+}
 __device__ double remainder(double x, double y)
 {
     return hc::precise_math::remainder(x, y);
 }
+__device__ double remquo(double x, double y, int *quo)
+{
+    return hc::precise_math::remquo(x, y, quo);
+}
 __device__ double rhypot(double x, double y)
 {
     return 1/hc::precise_math::sqrt(x*x + y*y);
@@ -593,6 +638,10 @@ __device__ double rsqrt(double x)
 {
     return hc::precise_math::rsqrt(x);
 }
+__device__ double scalbln(double x, long int n)
+{
+    return hc::precise_math::scalb(x, n);
+}
 __device__ double scalbn(double x, int n)
 {
     return hc::precise_math::scalbn(x, n);