Merge branch 'amd-master' into amd-develop

Change-Id: I01a77c34e11b1b02cb20e11b105741f6bd383543

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__ void* __get_dynamicgroupbaseptr();
+__device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr();
 
 
 /**
@@ -422,9 +422,10 @@ do {\
 // Macro to replace extern __shared__ declarations
 // to local variable definitions
 #define HIP_DYNAMIC_SHARED(type, var) \
-    type* var = (type*)__get_dynamicgroupbaseptr(); \
+    __attribute__((address_space(3))) type* var = \
+    (__attribute__((address_space(3))) type*)__get_dynamicgroupbaseptr(); \
 
-#define HIP_DYNAMIC_SHARED_ATTRIBUTE 
+#define HIP_DYNAMIC_SHARED_ATTRIBUTE __attribute__((address_space(3)))
 
 #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__((hc))
+#define __constant__   __attribute__((address_space(1)))
 
 #else
 // Non-HCC compiler