Revert "Changes to HIP to cope with Promote-free HCC"

This reverts commit 0005dd5f66.

Change-Id: I20a9bab3883ad09913b320210344d37599cb8fcd

hipamd/src/device_functions.cpp

@@ -41,8 +41,8 @@ struct holder32Bit {
   };
 } __attribute__((aligned(4)));
 
-__device__ struct holder64Bit hold64;
-__device__ struct holder32Bit hold32;
+struct holder64Bit hold64;
+struct holder32Bit hold32;
 
 __device__ float __double2float_rd(double x)
 {

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.
 */
-__device__ char gpuHeap[SIZE_OF_HEAP];
-__device__ uint32_t gpuFlags[NUM_PAGES];
+__attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP];
+__attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES];
 
 __device__ void *__hip_hc_malloc(size_t size)
 {
@@ -1083,7 +1083,7 @@ __host__ __device__ int max(int arg1, int arg2)
   return (int)(hc::precise_math::fmax((float)arg1, (float)arg2));
 }
 
-__device__ void* __get_dynamicgroupbaseptr()
+__device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr()
 {
   return hc::get_dynamic_group_segment_base_pointer();
 }

hipamd/src/hip_fp16.cpp

@@ -31,7 +31,7 @@ struct hipHalfHolder{
 
 #define HINF 65504
 
-__device__ static struct hipHalfHolder __hInfValue = {HINF};
+static struct hipHalfHolder __hInfValue = {HINF};
 
 __device__ __half __hadd(__half a, __half b) {
   return a + b;

hipamd/src/math_functions.cpp

@@ -202,8 +202,7 @@ __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);
 }