added few type reinterpret cast device functions

1. __int_as_float
2. __hiloint2double

Change-Id: Id247c196887b24a12090f0521bf91e13afeec733

hipamd/CMakeLists.txt

@@ -177,7 +177,8 @@ if(HIP_PLATFORM STREQUAL "hcc")
     set(SOURCE_FILES_DEVICE
         src/device_util.cpp
         src/hip_ldg.cpp
-        src/hip_fp16.cpp)
+        src/hip_fp16.cpp
+        src/device_functions.cpp)
 
     set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic")
     add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME})

hipamd/include/hip/device_functions.h

@@ -0,0 +1,33 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+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.
+*/
+
+#ifndef HIP_DEVICE_FUNCTIONS_H
+#define HIP_DEVICE_FUNCTIONS_H
+
+#include<hip/hip_common.h>
+
+#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
+#include <hip/hcc_detail/device_functions.h>
+#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
+#include <device_functions.h>
+#else 
+#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
+#endif 
+
+#endif

hipamd/include/hip/hcc_detail/device_functions.h

@@ -0,0 +1,32 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+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.
+*/
+
+#ifndef HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
+#define HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
+
+#include "hip_runtime.h"
+
+__device__ float __int_as_float (int x);
+
+__device__ double __hiloint2double (int hi, int lo);
+
+extern __HIP_DEVICE__ double  __longlong_as_double(long long int x);
+extern __HIP_DEVICE__ long long int __double_as_longlong(double x);
+
+#endif

hipamd/include/hip/hcc_detail/hip_vector_types.h

@@ -343,9 +343,6 @@ __HIP_DEVICE__ double2 make_double2(double, double );
 __HIP_DEVICE__ double3 make_double3(double, double, double );
 __HIP_DEVICE__ double4 make_double4(double, double, double, double );
 
-extern __HIP_DEVICE__ double  __longlong_as_double(long long int x);
-extern __HIP_DEVICE__ long long int __double_as_longlong(double x);
-
 
 /*
 ///---

hipamd/src/device_functions.cpp

@@ -0,0 +1,58 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+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/device_functions.h>
+
+extern "C" float __hip_int_as_float(int);
+
+typedef struct {
+  signed int hi;
+  signed int lo;
+} __hip_signed_2;
+
+typedef struct {
+union {
+  double d;
+  long long int lli;
+  __hip_signed_2 s2;
+} ;
+} __hip_64bit_struct;
+
+typedef struct {
+union {
+  float f;
+  unsigned int ui;
+  signed int si;
+};
+} __hip_32bit_struct;
+
+__device__ float __int_as_float (int x) {
+  __hip_32bit_struct s;
+  s.si = x;
+  return s.f;
+}
+
+__device__ double __hiloint2double (int hi, int lo) {
+  __hip_64bit_struct s;
+  s.s2.hi = hi;
+  s.s2.lo = lo;
+  return s.d;
+}
+
+