From 29232ff2832154d4919f57671d9c654a98287b62 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Fri, 17 Mar 2017 11:19:48 -0500
Subject: [PATCH] Add __device__ to needful functions for promote-free.

---
 include/hip/hcc_detail/hip_runtime.h | 2 +-
 src/device_functions.cpp             | 4 ++--
 src/device_util.cpp                  | 4 ++--
 src/hip_fp16.cpp                     | 2 +-
 src/math_functions.cpp               | 3 ++-
 5 files changed, 8 insertions(+), 7 deletions(-)

diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index 06ce7c84df..6fc68abb3a 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -443,7 +443,7 @@ do {\
 // to local variable definitions
 #define HIP_DYNAMIC_SHARED(type, var) \
     ADDRESS_SPACE_3 type* var = \
-    ADDRESS_SPACE_3 type*)__get_dynamicgroupbaseptr(); \
+    (ADDRESS_SPACE_3 type*)__get_dynamicgroupbaseptr(); \
 
 #define HIP_DYNAMIC_SHARED_ATTRIBUTE ADDRESS_SPACE_3
 
diff --git a/src/device_functions.cpp b/src/device_functions.cpp
index abc9db570e..10d8d3ab89 100644
--- a/src/device_functions.cpp
+++ b/src/device_functions.cpp
@@ -41,8 +41,8 @@ struct holder32Bit {
   };
 } __attribute__((aligned(4)));
 
-struct holder64Bit hold64;
-struct holder32Bit hold32;
+__device__ struct holder64Bit hold64;
+__device__ struct holder32Bit hold32;
 
 __device__ float __double2float_rd(double x)
 {
diff --git a/src/device_util.cpp b/src/device_util.cpp
index 6e2652449f..b0df62f43b 100644
--- a/src/device_util.cpp
+++ b/src/device_util.cpp
@@ -34,8 +34,8 @@ THE SOFTWARE.
  This is the best place to put them because the device
  global variables need to be initialized at the start.
 */
-ADDRESS_SPACE_1 char gpuHeap[SIZE_OF_HEAP];
-ADDRESS_SPACE_1 uint32_t gpuFlags[NUM_PAGES];
+__device__ ADDRESS_SPACE_1 char gpuHeap[SIZE_OF_HEAP];
+__device__ ADDRESS_SPACE_1 uint32_t gpuFlags[NUM_PAGES];
 
 __device__ void *__hip_hc_malloc(size_t size)
 {
diff --git a/src/hip_fp16.cpp b/src/hip_fp16.cpp
index b306a9d3de..e7f75844ff 100644
--- a/src/hip_fp16.cpp
+++ b/src/hip_fp16.cpp
@@ -31,7 +31,7 @@ struct hipHalfHolder{
 
 #define HINF 65504
 
-static struct hipHalfHolder __hInfValue = {HINF};
+__device__ static struct hipHalfHolder __hInfValue = {HINF};
 
 __device__ __half __hadd(__half a, __half b) {
   return a + b;
diff --git a/src/math_functions.cpp b/src/math_functions.cpp
index 230eb2aacc..6e919b3926 100644
--- a/src/math_functions.cpp
+++ b/src/math_functions.cpp
@@ -202,7 +202,8 @@ __device__ long long int llroundf(float x)
     int y = hc::precise_math::roundf(x);
     long long int z = y;
     return z;
-}__device__ float log10f(float x)
+}
+__device__ float log10f(float x)
 {
     return hc::precise_math::log10f(x);
 }