Squashed commit of the following:

commit 9548493fa754b3bf5c31cbdc2211db1e73e8c07c
Author: Jack Chung <whchung@gmail.com>
Date:   Mon May 23 11:57:23 2016 +0800

    Rename hipExternShared test to hipDynamicShared

    Change-Id: I180d9d539420fb69cfc121eceaa7db9da03483b2

commit 827081f8244a38f010789d556db0c4ff7b6422d8
Author: Jack Chung <whchung@gmail.com>
Date:   Mon May 23 11:56:27 2016 +0800

    Rename HIP_DECLARE_EXTERN_SHARED to HIP_DYNAMIC_SHARED

    Change-Id: I22362d179812ac547e0f11ba4e2bb999050e08ae

commit 4c277228ed41af187739610fa17eab1fb144c947
Author: Jack Chung <whchung@gmail.com>
Date:   Thu May 19 17:49:52 2016 +0800

    Adopt new interface to get dynamic LDS in hc.hpp

    Change-Id: I47b433b714633a4c97df87c40a0b1d3386429a00

commit 5a36117d777064113a528dc47b42e8c8413baa97
Author: Jack Chung <whchung@gmail.com>
Date:   Thu May 19 11:29:24 2016 +0800

    Add test patterns for regular expression to match "extern __shared__"

    These test patterns should better be saved as an individual test case, but I'm
    not familiar with HIP test structures so I leave them as comments in hipify as
    of now.

    Change-Id: I7fee89c89b9e73de2133357a226ec0c769733531

commit 1b26284168c7f5339f63338fd0149bed5d994656
Author: Jack Chung <whchung@gmail.com>
Date:   Thu May 19 11:25:23 2016 +0800

    Add one HIP unit test to use HIP_DECLARE_EXTERN_SHARED

    Change-Id: I4d9907815920693a74ea9d575fe26e7c67636109

commit 77b816ee5972b13d829d5bbcf06fbfd07acea2af
Author: Jack Chung <whchung@gmail.com>
Date:   Wed May 18 19:18:59 2016 +0800

    Adopt HIP_ prefix for DECLARE_EXTERN_SHARED macro

    Change-Id: I555ded16b449b67d2e20904013d86fe1ded6a2be

commit ef0997939c3578a9ae11621bf21c0416f04d2622
Author: Jack Chung <whchung@gmail.com>
Date:   Wed May 18 17:42:04 2016 +0800

    Modify hipify to support converting extern __shared__ to DECLARE_EXTERN_SHARED macro

    Added regular expression to search & replace extern __shared__ declarations to
    DECLARE_EXTERN_SHARED macro.

    Limitation:
    - Won't work if "extern __shared__" is declared at global scope

    Sample Usages:
    extern __shared__ double foo[];
    extern __shared__ unsigned int foo[];
    extern volatile __shared__ double foo[];
    extern volatile __shared__ unsigned int sdata[];
    extern __shared__ volatile unsigned int sdata[];
    extern __shared__ T s[];
    extern __shared__ T::type s[];
    extern __shared__ blah<T>::type s[];
    extern __shared__ typename mapper<Float>::type s_data[];
    extern __attribute__((used)) __shared__ typename mapper<Float>::type s_data[];

    Change-Id: I2be0b7039adeddb789f5a2b067d403a43fdc3e26

commit 93ff268724493aedfacdcd5a5aa9a100f4ebaed0
Author: Jack Chung <whchung@gmail.com>
Date:   Wed May 18 15:13:09 2016 +0800

    Introduce DECLARE_EXTERN_SHARED macro to encapsulate "extern __shared__" decls

    Change-Id: I93b2d37c763195b0ca9fd0afee78605a1e3272db

commit cff9c95412de343cc6405158b5acc4f1029267ff
Author: Jack Chung <whchung@gmail.com>
Date:   Wed May 18 12:53:54 2016 +0800

    Add __get_dynamic_groupbaseptr() to point to dynamic LDS

    Change-Id: I97b548d8a691488057617c551a8f331cad7afc77

Change-Id: I84e7875b76fa1f59e860e19c93bd4209cdd1fd2c

include/hcc_detail/hip_runtime.h

@@ -430,6 +430,8 @@ __device__ float __shfl_xor(float input, int lane_mask, int width);
 __host__ __device__ int min(int arg1, int arg2);
 __host__ __device__ int max(int arg1, int arg2);
 
+__device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr();
+
 //TODO - add a couple fast math operations here, the set here will grow :
 __device__ float __cosf(float x);
 __device__ float __expf(float x);
@@ -556,6 +558,16 @@ do {\
 
 #endif
 
+/**
+ * extern __shared__
+ */
+
+// 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(); \
+
 #endif // __HCC__
 
 

include/nvcc_detail/hip_runtime.h

@@ -95,6 +95,13 @@ kernelName<<<numblocks,numthreads,memperblock,streamId>>>(0, __VA_ARGS__);\
 #define hipGridDim_y  gridDim.y
 #define hipGridDim_z  gridDim.z
 
+/**
+ * extern __shared__
+ */
+
+#define HIP_DYNAMIC_SHARED(type, var) \
+    extern __shared__ type var[]; \
+
 #endif