From 7a712aa76b44896d43898dfdc6b981784ef4d9de Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 13 Dec 2016 14:41:36 -0600 Subject: [PATCH] added few type reinterpret cast device functions 1. __int_as_float 2. __hiloint2double Change-Id: Id247c196887b24a12090f0521bf91e13afeec733 --- hipamd/CMakeLists.txt | 3 +- hipamd/include/hip/device_functions.h | 33 +++++++++++ .../include/hip/hcc_detail/device_functions.h | 32 ++++++++++ .../include/hip/hcc_detail/hip_vector_types.h | 3 - hipamd/src/device_functions.cpp | 58 +++++++++++++++++++ 5 files changed, 125 insertions(+), 4 deletions(-) create mode 100644 hipamd/include/hip/device_functions.h create mode 100644 hipamd/include/hip/hcc_detail/device_functions.h create mode 100644 hipamd/src/device_functions.cpp diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt index 6f9a819c4d..c76410f06d 100644 --- a/hipamd/CMakeLists.txt +++ b/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}) diff --git a/hipamd/include/hip/device_functions.h b/hipamd/include/hip/device_functions.h new file mode 100644 index 0000000000..838bad8f0c --- /dev/null +++ b/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 + +#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) +#include +#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__) +#include +#else +#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); +#endif + +#endif diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h new file mode 100644 index 0000000000..8fa870664f --- /dev/null +++ b/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 diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h index ffe15a27a4..7c48985996 100644 --- a/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/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); - /* ///--- diff --git a/hipamd/src/device_functions.cpp b/hipamd/src/device_functions.cpp new file mode 100644 index 0000000000..30a09e6e02 --- /dev/null +++ b/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 + +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; +} + +