Moved device code to mimic cuda header behavior

1. All fp32, fp64 math device/host functions should be in math_functions.h/.cpp
2. All fp32, fp64 fast math intrinsics for device/host functions should be in device_functions.h/.cpp
3. All the device code implementations should be in device_util.h/.cpp
4. Hence, made changes appropriately by moving code and creating new header files
5. Added math_functions.cpp/.h
6. Changed #ifndef signature to make sure no conflicts between headers with same names in hip/hip_runtime.h and hip/hcc_detail/hip_runtime.h
7. Changed tests to fit the code changes, making them to include appropriate headers
8. Added math_functions.cpp to CMakeLists.txt
9. Some of the tests are still broken, mostly host math functions will fix them in next commit
10. TODO: FIX compilation issues for host math functions

Change-Id: I7a17637d7e294a7d224ffba932c1a08668febd26


[ROCm/hip commit: b723169ee9]

projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp

@@ -19,7 +19,8 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
+#include <hip/device_functions.h>
 #include "test_common.h"
 
 #pragma GCC diagnostic ignored "-Wall"
@@ -27,18 +28,18 @@ THE SOFTWARE.
 
 __device__ void double_precision_intrinsics()
 {
-    //__dadd_rd(0.0, 1.0);
-    //__dadd_rn(0.0, 1.0);
-    //__dadd_ru(0.0, 1.0);
-    //__dadd_rz(0.0, 1.0);
-    //__ddiv_rd(0.0, 1.0);
-    //__ddiv_rn(0.0, 1.0);
-    //__ddiv_ru(0.0, 1.0);
-    //__ddiv_rz(0.0, 1.0);
-    //__dmul_rd(1.0, 2.0);
-    //__dmul_rn(1.0, 2.0);
-    //__dmul_ru(1.0, 2.0);
-    //__dmul_rz(1.0, 2.0);
+    __dadd_rd(0.0, 1.0);
+    __dadd_rn(0.0, 1.0);
+    __dadd_ru(0.0, 1.0);
+    __dadd_rz(0.0, 1.0);
+    __ddiv_rd(0.0, 1.0);
+    __ddiv_rn(0.0, 1.0);
+    __ddiv_ru(0.0, 1.0);
+    __ddiv_rz(0.0, 1.0);
+    __dmul_rd(1.0, 2.0);
+    __dmul_rn(1.0, 2.0);
+    __dmul_ru(1.0, 2.0);
+    __dmul_rz(1.0, 2.0);
     __drcp_rd(2.0);
     __drcp_rn(2.0);
     __drcp_ru(2.0);
@@ -47,10 +48,10 @@ __device__ void double_precision_intrinsics()
     __dsqrt_rn(4.0);
     __dsqrt_ru(4.0);
     __dsqrt_rz(4.0);
-    //__dsub_rd(2.0, 1.0);
-    //__dsub_rn(2.0, 1.0);
-    //__dsub_ru(2.0, 1.0);
-    //__dsub_rz(2.0, 1.0);
+    __dsub_rd(2.0, 1.0);
+    __dsub_rn(2.0, 1.0);
+    __dsub_ru(2.0, 1.0);
+    __dsub_rz(2.0, 1.0);
     __fma_rd(1.0, 2.0, 3.0);
     __fma_rn(1.0, 2.0, 3.0);
     __fma_ru(1.0, 2.0, 3.0);

projects/hip/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp

@@ -19,7 +19,8 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
+#include <hip/math_functions.h>
 #include "test_common.h"
 
 #pragma GCC diagnostic ignored "-Wall"
@@ -43,8 +44,8 @@ __device__ void double_precision_math_functions()
     cos(0.0);
     cosh(0.0);
     cospi(0.0);
-    //cyl_bessel_i0(0.0);
-    //cyl_bessel_i1(0.0);
+    cyl_bessel_i0(0.0);
+    cyl_bessel_i1(0.0);
     erf(0.0);
     erfc(0.0);
     erfcinv(2.0);
@@ -61,7 +62,7 @@ __device__ void double_precision_math_functions()
     fmax(0.0, 0.0);
     fmin(0.0, 0.0);
     fmod(0.0, 1.0);
-    //frexp(0.0, &iX);
+    frexp(0.0, &iX);
     hypot(1.0, 0.0);
     ilogb(1.0);
     isfinite(0.0);
@@ -71,7 +72,7 @@ __device__ void double_precision_math_functions()
     j1(0.0);
     jn(-1.0, 1.0);
     ldexp(0.0, 0);
-    //lgamma(1.0);
+    lgamma(1.0);
     llrint(0.0);
     llround(0.0);
     log(1.0);
@@ -81,19 +82,19 @@ __device__ void double_precision_math_functions()
     logb(1.0);
     lrint(0.0);
     lround(0.0);
-    //modf(0.0, &fX);
+    modf(0.0, &fX);
     nan("1");
     nearbyint(0.0);
-    //nextafter(0.0);
-    //fX = 1.0; norm(1, &fX);
+    nextafter(0.0, 0.0);
+    fX = 1.0; norm(1, &fX);
     norm3d(1.0, 0.0, 0.0);
     norm4d(1.0, 0.0, 0.0, 0.0);
     normcdf(0.0);
-    //normcdfinv(1.0);
+    normcdfinv(1.0);
     pow(1.0, 0.0);
     rcbrt(1.0);
     remainder(2.0, 1.0);
-    //remquo(1.0, 2.0, &iX);
+    remquo(1.0, 2.0, &iX);
     rhypot(0.0, 1.0);
     rint(1.0);
     fX = 1.0; rnorm(1, &fX);

projects/hip/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp

@@ -19,7 +19,8 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
+#include <hip/math_functions.h>
 #include "test_common.h"
 
 #pragma GCC diagnostic ignored "-Wall"
@@ -85,7 +86,7 @@ __host__ void double_precision_math_functions()
     nan("1");
     nearbyint(0.0);
     //nextafter(0.0);
-    //fX = 1.0; norm(1, &fX);
+    fX = 1.0; norm(1, &fX);
 #if defined(__HIP_PLATFORM_HCC__)
     norm3d(1.0, 0.0, 0.0);
     norm4d(1.0, 0.0, 0.0, 0.0);

projects/hip/tests/src/deviceLib/hipFloatMath.cpp

@@ -27,6 +27,7 @@ THE SOFTWARE.
  */
 
 #include "test_common.h"
+#include <hip/device_functions.h>
 
 #define LEN 512
 #define SIZE LEN<<2

projects/hip/tests/src/deviceLib/hipFloatMathPrecise.cpp

@@ -19,7 +19,8 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
+#include <hip/math_functions.h>
 #include "test_common.h"
 
 __global__ void FloatMathPrecise(hipLaunchParm lp)

projects/hip/tests/src/deviceLib/hipIntegerIntrinsics.cpp

@@ -19,8 +19,8 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
-#include "hip/device_functions.h"
+#include <hip/hip_runtime.h>
+#include <hip/device_functions.h>
 #include "test_common.h"
 
 #pragma GCC diagnostic ignored "-Wall"

projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp

@@ -19,7 +19,8 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
+#include <hip/device_functions.h>
 #include "test_common.h"
 
 #pragma GCC diagnostic ignored "-Wall"
@@ -30,44 +31,44 @@ __device__ void single_precision_intrinsics()
     float fX, fY;
 
     __cosf(0.0f);
-    //__exp10f(0.0f);
+    __exp10f(0.0f);
     __expf(0.0f);
-    //__fadd_rd(0.0f, 1.0f);
-    //__fadd_rn(0.0f, 1.0f);
-    //__fadd_ru(0.0f, 1.0f);
-    //__fadd_rz(0.0f, 1.0f);
-    //__fdiv_rd(4.0f, 2.0f);
-    //__fdiv_rn(4.0f, 2.0f);
-    //__fdiv_ru(4.0f, 2.0f);
-    //__fdiv_rz(4.0f, 2.0f);
-    //__fdividef(4.0f, 2.0f);
-    //__fmaf_rd(1.0f, 2.0f, 3.0f);
-    //__fmaf_rn(1.0f, 2.0f, 3.0f);
-    //__fmaf_ru(1.0f, 2.0f, 3.0f);
-    //__fmaf_rz(1.0f, 2.0f, 3.0f);
-    //__fmul_rd(1.0f, 2.0f);
-    //__fmul_rn(1.0f, 2.0f);
-    //__fmul_ru(1.0f, 2.0f);
-    //__fmul_rz(1.0f, 2.0f);
-    //__frcp_rd(2.0f);
-    //__frcp_rn(2.0f);
-    //__frcp_ru(2.0f);
-    //__frcp_rz(2.0f);
+    __fadd_rd(0.0f, 1.0f);
+    __fadd_rn(0.0f, 1.0f);
+    __fadd_ru(0.0f, 1.0f);
+    __fadd_rz(0.0f, 1.0f);
+    __fdiv_rd(4.0f, 2.0f);
+    __fdiv_rn(4.0f, 2.0f);
+    __fdiv_ru(4.0f, 2.0f);
+    __fdiv_rz(4.0f, 2.0f);
+    __fdividef(4.0f, 2.0f);
+    __fmaf_rd(1.0f, 2.0f, 3.0f);
+    __fmaf_rn(1.0f, 2.0f, 3.0f);
+    __fmaf_ru(1.0f, 2.0f, 3.0f);
+    __fmaf_rz(1.0f, 2.0f, 3.0f);
+    __fmul_rd(1.0f, 2.0f);
+    __fmul_rn(1.0f, 2.0f);
+    __fmul_ru(1.0f, 2.0f);
+    __fmul_rz(1.0f, 2.0f);
+    __frcp_rd(2.0f);
+    __frcp_rn(2.0f);
+    __frcp_ru(2.0f);
+    __frcp_rz(2.0f);
     __frsqrt_rn(4.0f);
     __fsqrt_rd(4.0f);
     __fsqrt_rn(4.0f);
     __fsqrt_ru(4.0f);
     __fsqrt_rz(4.0f);
-    //__fsub_rd(2.0f, 1.0f);
-    //__fsub_rn(2.0f, 1.0f);
-    //__fsub_ru(2.0f, 1.0f);
-    //__fsub_rz(2.0f, 1.0f);
+    __fsub_rd(2.0f, 1.0f);
+    __fsub_rn(2.0f, 1.0f);
+    __fsub_ru(2.0f, 1.0f);
+    __fsub_rz(2.0f, 1.0f);
     __log10f(1.0f);
     __log2f(1.0f);
     __logf(1.0f);
     __powf(1.0f, 0.0f);
-    //__saturatef(0.1f);
-    //__sincosf(0.0f, &fX, &fY);
+    __saturatef(0.1f);
+    __sincosf(0.0f, &fX, &fY);
     __sinf(0.0f);
     __tanf(0.0f);
 }

projects/hip/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp

@@ -19,7 +19,8 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
+#include <hip/math_functions.h>
 #include "test_common.h"
 
 #pragma GCC diagnostic ignored "-Wall"

projects/hip/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp

@@ -19,7 +19,8 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
+#include <hip/math_functions.h>
 #include "test_common.h"
 
 #pragma GCC diagnostic ignored "-Wall"

projects/hip/tests/src/deviceLib/hipTestDevice.cpp

@@ -24,8 +24,9 @@ THE SOFTWARE.
  */
 
 #include"test_common.h"
-#include "hip/hip_runtime.h"
-#include "hip/hip_runtime_api.h"
+#include <hip/hip_runtime.h>
+#include <hip/math_functions.h>
+#include <hip/hip_runtime_api.h>
 
 #define N 512
 #define SIZE N*sizeof(float)

projects/hip/tests/src/deviceLib/hipTestDeviceDouble.cpp

@@ -24,8 +24,9 @@ THE SOFTWARE.
  */
 
 #include"test_common.h"
-#include "hip/hip_runtime.h"
-#include "hip/hip_runtime_api.h"
+#include <hip/hip_runtime.h>
+#include <hip/hip_runtime_api.h>
+#include <hip/math_functions.h>
 
 #define N 512
 #define SIZE N*sizeof(double)

projects/hip/tests/src/deviceLib/hip_anyall.cpp

@@ -29,7 +29,8 @@ THE SOFTWARE.
 #include <stdio.h>
 #include <iostream>
 
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
+#include <hip/device_functions.h>
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 
 __global__ void

projects/hip/tests/src/deviceLib/hip_ballot.cpp

@@ -25,8 +25,8 @@ THE SOFTWARE.
 
 #include <iostream>
 
-#include "hip/hip_runtime.h"
-#include "hip/device_functions.h"
+#include <hip/hip_runtime.h>
+#include <hip/device_functions.h>
 
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 

projects/hip/tests/src/deviceLib/hip_brev.cpp

@@ -32,7 +32,7 @@ THE SOFTWARE.
 #include <stdlib.h>
 #include <iostream>
 #include "hip/hip_runtime.h"
-#include "hip/device_functions.h"
+#include <hip/device_functions.h>
 
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 

projects/hip/tests/src/deviceLib/hip_clz.cpp

@@ -32,7 +32,7 @@ THE SOFTWARE.
 #include <stdlib.h>
 #include <iostream>
 #include "hip/hip_runtime.h"
-#include "hip/device_functions.h"
+#include <hip/device_functions.h>
 
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 #define WIDTH     8

projects/hip/tests/src/deviceLib/hip_ffs.cpp

@@ -31,8 +31,8 @@ THE SOFTWARE.
 #include <algorithm>
 #include <stdlib.h>
 #include <iostream>
-#include "hip/hip_runtime.h"
-#include "hip/device_functions.h"
+#include <hip/hip_runtime.h>
+#include <hip/device_functions.h>
 
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 

projects/hip/tests/src/deviceLib/hip_popc.cpp

@@ -31,8 +31,8 @@ THE SOFTWARE.
 #include <algorithm>
 #include <stdlib.h>
 #include <iostream>
-#include "hip/hip_runtime.h"
-#include "hip/device_functions.h"
+#include <hip/hip_runtime.h>
+#include <hip/device_functions.h>
 
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 

projects/hip/tests/src/deviceLib/hip_trig.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s 
+ * BUILD: %t %s
  * RUN: %t
  * HIT_END
  */
@@ -30,6 +30,7 @@ THE SOFTWARE.
 #include<hip/hip_runtime.h>
 #include<iostream>
 #include"test_common.h"
+#include<hip/device_functions.h>
 
 #define LEN 512
 #define SIZE LEN<<2

projects/hip/tests/src/runtimeApi/stream/hipAPIStreamDisable.cpp

@@ -24,6 +24,7 @@ THE SOFTWARE.
 
 #include<iostream>
 #include"test_common.h"
+#include"hip/math_functions.h"
 
 const int NN = 1 << 21;
 
@@ -31,7 +32,7 @@ __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
 	int tid = hipThreadIdx_x;
 	if(tid < 1){
 		for(int i=0;i<n;i++){
-			x[i] = sqrt(pow(3.14159,i));
+			x[i] = sqrt(powf(3.14159,i));
 		}
 		y[tid] = y[tid] + 1.0f;
 	}

projects/hip/tests/src/runtimeApi/stream/hipAPIStreamEnable.cpp

@@ -26,6 +26,7 @@ THE SOFTWARE.
 
 #include<iostream>
 #include"test_common.h"
+#include"hip/math_functions.h"
 
 const int NN = 1 << 21;
 
@@ -33,7 +34,7 @@ __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
 	int tid = hipThreadIdx_x;
 	if(tid < 1){
 		for(int i=0;i<n;i++){
-			x[i] = sqrt(pow(3.14159,i));
+			x[i] = sqrt(powf(3.14159,i));
 		}
 		y[tid] = y[tid] + 1.0f;
 	}