From f4256cb57c5206f0f34b54efc76ceafaea71e096 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Fri, 17 Mar 2017 11:04:39 -0500
Subject: [PATCH] Add USE_PROMOTE_FREE_HCC for smooth transition to new HCC
 caps.

ADDRESS_SPACE_1 defines
---
 hipamd/include/hip/hcc_detail/hip_runtime.h   | 19 ++++++++++++++-----
 hipamd/include/hip/hcc_detail/host_defines.h  |  6 +++++-
 hipamd/src/device_util.cpp                    |  6 +++---
 .../src/deviceLib/hipTestDeviceSymbol.cpp     |  5 +++--
 4 files changed, 25 insertions(+), 11 deletions(-)

diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h
index 673463635f..06ce7c84df 100644
--- a/hipamd/include/hip/hcc_detail/hip_runtime.h
+++ b/hipamd/include/hip/hcc_detail/hip_runtime.h
@@ -47,7 +47,16 @@ THE SOFTWARE.
 
 #include <hip/hip_runtime_api.h>
 
-//#include "hip/hcc_detail/hip_hcc.h"
+#define USE_PROMOTE_FREE_HCC 0
+
+#if USE_PROMOTE_FREE_HCC == 1
+#define ADDRESS_SPACE_1 
+#define ADDRESS_SPACE_3 
+#else
+#define ADDRESS_SPACE_1 __attribute__((address_space(1)))
+#define ADDRESS_SPACE_3 __attribute__((address_space(3)))
+#endif
+
 //---
 // Remainder of this file only compiles with HCC
 #ifdef __HCC__
@@ -275,7 +284,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__ ADDRESS_SPACE_3 void* __get_dynamicgroupbaseptr();
 
 
 /**
@@ -433,10 +442,10 @@ 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(); \
+    ADDRESS_SPACE_3 type* var = \
+    ADDRESS_SPACE_3 type*)__get_dynamicgroupbaseptr(); \
 
-#define HIP_DYNAMIC_SHARED_ATTRIBUTE __attribute__((address_space(3)))
+#define HIP_DYNAMIC_SHARED_ATTRIBUTE ADDRESS_SPACE_3
 
 #endif // __HCC__
 
diff --git a/hipamd/include/hip/hcc_detail/host_defines.h b/hipamd/include/hip/hcc_detail/host_defines.h
index e401cb24f3..0e10d0075c 100644
--- a/hipamd/include/hip/hcc_detail/host_defines.h
+++ b/hipamd/include/hip/hcc_detail/host_defines.h
@@ -47,7 +47,11 @@ THE SOFTWARE.
  */
 // _restrict is supported by the compiler
 #define __shared__     tile_static
-#define __constant__   __attribute__((address_space(1)))
+#if USE_PROMOTE_FREE_HCC==1
+#define __constant__   __attribute__((hc))
+#else
+#define __constant__   ADDRESS_SPACE_1
+#endif
 
 #else
 // Non-HCC compiler
diff --git a/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp
index 4b0e7efefd..6e2652449f 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];
+ADDRESS_SPACE_1 char gpuHeap[SIZE_OF_HEAP];
+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__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr()
+__device__ ADDRESS_SPACE_3 void* __get_dynamicgroupbaseptr()
 {
   return hc::get_dynamic_group_segment_base_pointer();
 }
diff --git a/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
index 9e188e9f17..6ffaedf659 100644
--- a/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
@@ -31,9 +31,10 @@ THE SOFTWARE.
 #define NUM 1024
 #define SIZE 1024*4
 
+// TODO - collapse:
 #ifdef __HIP_PLATFORM_HCC__
-__attribute__((address_space(1))) int globalIn[NUM];
-__attribute__((address_space(1))) int globalOut[NUM];
+__device__ ADDRESS_SPACE_1 int globalIn[NUM];
+__device__ ADDRESS_SPACE_1 int globalOut[NUM];
 #endif
 
 #ifdef __HIP_PLATFORM_NVCC__