diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 129020d9cd..95826f9b60 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -305,7 +305,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__ ADDRESS_SPACE_3 void* __get_dynamicgroupbaseptr(); +__device__ void* __get_dynamicgroupbaseptr(); /** @@ -464,10 +464,10 @@ do {\ // Macro to replace extern __shared__ declarations // to local variable definitions #define HIP_DYNAMIC_SHARED(type, var) \ - ADDRESS_SPACE_3 type* var = \ - (ADDRESS_SPACE_3 type*)__get_dynamicgroupbaseptr(); \ + type* var = \ + (type*)__get_dynamicgroupbaseptr(); \ -#define HIP_DYNAMIC_SHARED_ATTRIBUTE ADDRESS_SPACE_3 +#define HIP_DYNAMIC_SHARED_ATTRIBUTE diff --git a/src/device_util.cpp b/src/device_util.cpp index e59a44e5ba..062372f0f4 100644 --- a/src/device_util.cpp +++ b/src/device_util.cpp @@ -1101,11 +1101,13 @@ __host__ __device__ int max(int arg1, int arg2) return (int)(hc::precise_math::fmax((float)arg1, (float)arg2)); } -__device__ ADDRESS_SPACE_3 void* __get_dynamicgroupbaseptr() -{ +__device__ void* __get_dynamicgroupbaseptr() { return hc::get_dynamic_group_segment_base_pointer(); } +__host__ void* __get_dynamicgroupbaseptr() { + return nullptr; +} // Precise Math Functions __device__ float __hip_precise_cosf(float x) {