From cc14ed0981a6386e04790036da2d8014def7cfc2 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Wed, 6 Jun 2018 11:09:00 -0400 Subject: [PATCH] Add support of extern __shared__ for hip-clang --- include/hip/hcc_detail/hip_runtime.h | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 1a6b0f7dda..865ac58759 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -657,6 +657,14 @@ __DEVICE__ void inline __assert_fail(const char * __assertion, } extern "C" __device__ __attribute__((noduplicate)) void __syncthreads(); +extern "C" __device__ void *__amdgcn_get_dynamicgroupbaseptr(); + +// Macro to replace extern __shared__ declarations +// to local variable definitions +#define HIP_DYNAMIC_SHARED(type, var) \ + type* var = (type*)__amdgcn_get_dynamicgroupbaseptr(); + +#define HIP_DYNAMIC_SHARED_ATTRIBUTE #pragma push_macro("__DEVICE__")