Add __device__ to needful functions for promote-free.

[ROCm/clr commit: 92fd4c8f10]

projects/clr/hipamd/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
 

projects/clr/hipamd/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)
 {

projects/clr/hipamd/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)
 {

projects/clr/hipamd/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;

projects/clr/hipamd/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);
 }