Merge branch 'amd-master' into amd-develop

hipamd/include/hip/hcc_detail/hip_runtime.h

@@ -264,7 +264,7 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask
 __host__ __device__ int min(int arg1, int arg2);
 __host__ __device__ int max(int arg1, int arg2);
 
-__device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr();
+__device__ void* __get_dynamicgroupbaseptr();
 
 
 /**
@@ -422,10 +422,9 @@ do {\
 // Macro to replace extern __shared__ declarations
 // to local variable definitions
 #define HIP_DYNAMIC_SHARED(type, var) \
-    __attribute__((address_space(3))) type* var = \
+    type* var = (type*)__get_dynamicgroupbaseptr(); \
-    (__attribute__((address_space(3))) type*)__get_dynamicgroupbaseptr(); \
 
-#define HIP_DYNAMIC_SHARED_ATTRIBUTE __attribute__((address_space(3)))
+#define HIP_DYNAMIC_SHARED_ATTRIBUTE 
 
 #endif // __HCC__
 

hipamd/include/hip/hcc_detail/host_defines.h

@@ -47,7 +47,7 @@ THE SOFTWARE.
  */
 // _restrict is supported by the compiler
 #define __shared__     tile_static
-#define __constant__   __attribute__((address_space(1)))
+#define __constant__   __attribute__((hc))
 
 #else
 // Non-HCC compiler

hipamd/src/device_functions.cpp

@@ -41,8 +41,8 @@ struct holder32Bit {
   };
 } __attribute__((aligned(4)));
 
-struct holder64Bit hold64;
+__device__ struct holder64Bit hold64;
-struct holder32Bit hold32;
+__device__ 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.
 */
-__attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP];
+__device__ char gpuHeap[SIZE_OF_HEAP];
-__attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES];
+__device__ 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__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr()
+__device__ 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
 
-static struct hipHalfHolder __hInfValue = {HINF};
+__device__ static struct hipHalfHolder __hInfValue = {HINF};
 
 __device__ __half __hadd(__half a, __half b) {
   return a + b;

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);
 }

hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp

@@ -31,15 +31,8 @@ THE SOFTWARE.
 #define NUM 1024
 #define SIZE 1024*4
 
-#ifdef __HIP_PLATFORM_HCC__
-__attribute__((address_space(1))) int globalIn[NUM];
-__attribute__((address_space(1))) int globalOut[NUM];
-#endif
-
-#ifdef __HIP_PLATFORM_NVCC__
 __device__ int globalIn[NUM];
 __device__ int globalOut[NUM];
-#endif
 
 __global__ void Assign(hipLaunchParm lp, int* Out)
 {

hipamd/tests/src/kernel/hipDynamicShared2.cpp

@@ -29,7 +29,7 @@ THE SOFTWARE.
  #include "hip/hip_runtime.h"
  #include "test_common.h"
 
-#define LEN 16*1024
+#define LEN 8*1024
 #define SIZE LEN*4
 
 __global__ void vectorAdd(hipLaunchParm lp, float *Ad, float *Bd) {