diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 332e9bab46..eb0f7bf61a 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/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 = \ - (__attribute__((address_space(3))) type*)__get_dynamicgroupbaseptr(); \ + type* var = (type*)__get_dynamicgroupbaseptr(); \ -#define HIP_DYNAMIC_SHARED_ATTRIBUTE __attribute__((address_space(3))) +#define HIP_DYNAMIC_SHARED_ATTRIBUTE #endif // __HCC__ diff --git a/hipamd/include/hip/hcc_detail/host_defines.h b/hipamd/include/hip/hcc_detail/host_defines.h index e401cb24f3..012d3f0346 100644 --- a/hipamd/include/hip/hcc_detail/host_defines.h +++ b/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 diff --git a/hipamd/src/device_functions.cpp b/hipamd/src/device_functions.cpp index abc9db570e..10d8d3ab89 100644 --- a/hipamd/src/device_functions.cpp +++ b/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) { diff --git a/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp index 4b0e7efefd..88ffa7ab4d 100644 --- a/hipamd/src/device_util.cpp +++ b/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]; -__attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES]; +__device__ char gpuHeap[SIZE_OF_HEAP]; +__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(); } diff --git a/hipamd/src/hip_fp16.cpp b/hipamd/src/hip_fp16.cpp index b306a9d3de..e7f75844ff 100644 --- a/hipamd/src/hip_fp16.cpp +++ b/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; diff --git a/hipamd/src/math_functions.cpp b/hipamd/src/math_functions.cpp index 230eb2aacc..6e919b3926 100644 --- a/hipamd/src/math_functions.cpp +++ 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); } diff --git a/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp index e58aa58877..429c1d69ef 100644 --- a/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp +++ b/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) { diff --git a/hipamd/tests/src/kernel/hipDynamicShared2.cpp b/hipamd/tests/src/kernel/hipDynamicShared2.cpp index 0f6ebb4927..ea24e9341f 100644 --- a/hipamd/tests/src/kernel/hipDynamicShared2.cpp +++ b/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) {