From 6bdde982ffc78f95f1f9f8e24afc274479f4efc9 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sat, 2 Dec 2017 00:01:47 +0000 Subject: [PATCH 1/2] Replace archaic use of homebrew functionality with calls to the HC maths library. This fixes a hang observed when building hipTestDeviceDouble. [ROCm/clr commit: b8c80bd0b27261fef0e66a891a2da9e0d47adac6] --- projects/clr/hipamd/src/device_util.cpp | 84 +--------------------- projects/clr/hipamd/src/math_functions.cpp | 11 +-- 2 files changed, 5 insertions(+), 90 deletions(-) diff --git a/projects/clr/hipamd/src/device_util.cpp b/projects/clr/hipamd/src/device_util.cpp index b6aebdfce0..6edad53bcb 100644 --- a/projects/clr/hipamd/src/device_util.cpp +++ b/projects/clr/hipamd/src/device_util.cpp @@ -147,91 +147,11 @@ __device__ void* __hip_hc_memset(void* dst, uint8_t val, size_t size) } __device__ float __hip_erfinvf(float x){ - float ret; - int sign; - if (x < -1 || x > 1){ - return NAN; - } - if (x == 0){ - return 0; - } - if (x > 0){ - sign = 1; - } else { - sign = -1; - x = -x; - } - if (x <= 0.7) { - float x1 = x * x; - float x2 = __hip_erfinva3 * x1 + __hip_erfinva2; - float x3 = x2 * x1 + __hip_erfinva1; - float x4 = x * (x3 * x1 + __hip_erfinva0); - - float r1 = __hip_erfinvb4 * x1 + __hip_erfinvb3; - float r2 = r1 * x1 + __hip_erfinvb2; - float r3 = r2 * x1 + __hip_erfinvb1; - ret = x4 / (r3 * x1 + __hip_erfinvb0); - } else { - float x1 = hc::precise_math::sqrtf(-hc::precise_math::logf((1 - x) / 2)); - float x2 = __hip_erfinvc3 * x1 + __hip_erfinvc2; - float x3 = x2 * x1 + __hip_erfinvc1; - float x4 = x3 * x1 + __hip_erfinvc0; - - float r1 = __hip_erfinvd2 * x1 + __hip_erfinvd1; - ret = x4 / (r1 * x1 + __hip_erfinvd0); - } - - ret = ret * sign; - x = x * sign; - - ret -= (hc::precise_math::erff(ret) - x) / (2 / HIP_SQRT_PI * hc::precise_math::expf(-ret * ret)); - ret -= (hc::precise_math::erff(ret) - x) / (2 / HIP_SQRT_PI * hc::precise_math::expf(-ret * ret)); - - return ret; + return hc::precise_math::erfinvf(x); } __device__ double __hip_erfinv(double x){ - double ret; - int sign; - if (x < -1 || x > 1){ - return NAN; - } - if (x == 0){ - return 0; - } - if (x > 0){ - sign = 1; - } else { - sign = -1; - x = -x; - } - if (x <= 0.7) { - double x1 = x * x; - double x2 = __hip_erfinva3 * x1 + __hip_erfinva2; - double x3 = x2 * x1 + __hip_erfinva1; - double x4 = x * (x3 * x1 + __hip_erfinva0); - - double r1 = __hip_erfinvb4 * x1 + __hip_erfinvb3; - double r2 = r1 * x1 + __hip_erfinvb2; - double r3 = r2 * x1 + __hip_erfinvb1; - ret = x4 / (r3 * x1 + __hip_erfinvb0); - } else { - double x1 = hc::precise_math::sqrt(-hc::precise_math::log((1 - x) / 2)); - double x2 = __hip_erfinvc3 * x1 + __hip_erfinvc2; - double x3 = x2 * x1 + __hip_erfinvc1; - double x4 = x3 * x1 + __hip_erfinvc0; - - double r1 = __hip_erfinvd2 * x1 + __hip_erfinvd1; - ret = x4 / (r1 * x1 + __hip_erfinvd0); - } - - ret = ret * sign; - x = x * sign; - - ret -= (hc::precise_math::erf(ret) - x) / (2 / HIP_SQRT_PI * hc::precise_math::exp(-ret * ret)); - ret -= (hc::precise_math::erf(ret) - x) / (2 / HIP_SQRT_PI * hc::precise_math::exp(-ret * ret)); - - return ret; + return hc::precise_math::erfinv(x); } #define __hip_j0a1 57568490574.0 diff --git a/projects/clr/hipamd/src/math_functions.cpp b/projects/clr/hipamd/src/math_functions.cpp index 80ccece1a3..9dd27a7082 100644 --- a/projects/clr/hipamd/src/math_functions.cpp +++ b/projects/clr/hipamd/src/math_functions.cpp @@ -84,7 +84,7 @@ __device__ float erfcf(float x) } __device__ float erfcinvf(float y) { - return __hip_erfinvf(1 - y); + return hc::precise_math::erfcinvf(y); } __device__ float erfcxf(float x) { @@ -96,7 +96,7 @@ __device__ float erff(float x) } __device__ float erfinvf(float y) { - return __hip_erfinvf(y); + return hc::precise_math::erfinvf(y);//__hip_erfinvf(y); } __device__ float exp10f(float x) { @@ -192,12 +192,7 @@ __device__ float ldexpf(float x, int exp) } __device__ float lgammaf(float x) { - float val = 0.0f; - float y = x - 1; - while(y > 0){ - val += logf(y--); - } - return val; + return hc::precise_math::lgammaf(x); } __device__ long long int llrintf(float x) { From 1ba0d0ca147dd5ba1c050f8ef6dd8c3772921481 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sat, 2 Dec 2017 00:03:10 +0000 Subject: [PATCH 2/2] Remove stray leftover comment. [ROCm/clr commit: 954b7dadf7b07f86737c8bec65e88bf819b44a03] --- projects/clr/hipamd/src/math_functions.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/src/math_functions.cpp b/projects/clr/hipamd/src/math_functions.cpp index 9dd27a7082..9118318469 100644 --- a/projects/clr/hipamd/src/math_functions.cpp +++ b/projects/clr/hipamd/src/math_functions.cpp @@ -96,7 +96,7 @@ __device__ float erff(float x) } __device__ float erfinvf(float y) { - return hc::precise_math::erfinvf(y);//__hip_erfinvf(y); + return hc::precise_math::erfinvf(y); } __device__ float exp10f(float x) {