From 77e21dc09f39f37cc00b55083cb60150568f73cb Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Tue, 14 Mar 2017 22:58:41 +0800 Subject: [PATCH] Revert "Changes to HIP to cope with Promote-free HCC" This reverts commit efb9b9e86cefa266391e0c00ac6f004d09a83145. Change-Id: I20a9bab3883ad09913b320210344d37599cb8fcd --- include/hip/hcc_detail/hip_runtime.h | 7 ++++--- include/hip/hcc_detail/host_defines.h | 2 +- src/device_functions.cpp | 4 ++-- src/device_util.cpp | 6 +++--- src/hip_fp16.cpp | 2 +- src/math_functions.cpp | 3 +-- tests/src/deviceLib/hipTestDeviceSymbol.cpp | 7 +++++++ tests/src/kernel/hipDynamicShared2.cpp | 2 +- 8 files changed, 20 insertions(+), 13 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index eb0f7bf61a..332e9bab46 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/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__ diff --git a/include/hip/hcc_detail/host_defines.h b/include/hip/hcc_detail/host_defines.h index 012d3f0346..e401cb24f3 100644 --- a/include/hip/hcc_detail/host_defines.h +++ b/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 diff --git a/src/device_functions.cpp b/src/device_functions.cpp index 10d8d3ab89..abc9db570e 100644 --- a/src/device_functions.cpp +++ b/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) { diff --git a/src/device_util.cpp b/src/device_util.cpp index 88ffa7ab4d..4b0e7efefd 100644 --- a/src/device_util.cpp +++ b/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(); } diff --git a/src/hip_fp16.cpp b/src/hip_fp16.cpp index e7f75844ff..b306a9d3de 100644 --- a/src/hip_fp16.cpp +++ b/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; diff --git a/src/math_functions.cpp b/src/math_functions.cpp index 6e919b3926..230eb2aacc 100644 --- a/src/math_functions.cpp +++ b/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); } diff --git a/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/tests/src/deviceLib/hipTestDeviceSymbol.cpp index 429c1d69ef..e58aa58877 100644 --- a/tests/src/deviceLib/hipTestDeviceSymbol.cpp +++ b/tests/src/deviceLib/hipTestDeviceSymbol.cpp @@ -31,8 +31,15 @@ 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/tests/src/kernel/hipDynamicShared2.cpp b/tests/src/kernel/hipDynamicShared2.cpp index ea24e9341f..0f6ebb4927 100644 --- a/tests/src/kernel/hipDynamicShared2.cpp +++ b/tests/src/kernel/hipDynamicShared2.cpp @@ -29,7 +29,7 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "test_common.h" -#define LEN 8*1024 +#define LEN 16*1024 #define SIZE LEN*4 __global__ void vectorAdd(hipLaunchParm lp, float *Ad, float *Bd) {