SWDEV-271416 - Remove HIP_DYNAMIC_SHARED macro in hip

Change-Id: I12f39ea8438eb7ce76d8ffb2151b4faa93689048

hipamd/include/hip/amd_detail/device_functions.h

@@ -1269,14 +1269,6 @@ unsigned __smid(void)
     return (se_id << HW_ID_CU_ID_SIZE) + cu_id;
 }
 
-// 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
-
-
 #endif //defined(__clang__) && defined(__HIP__)
 
 

hipamd/include/hip/amd_detail/hip_runtime_api.h

@@ -3065,8 +3065,8 @@ hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned
  * @param [in] blockDimX X block dimensions specified in work-items
  * @param [in] blockDimY Y grid dimension specified in work-items
  * @param [in] blockDimZ Z grid dimension specified in work-items
- * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.  The
- * kernel can access this with HIP_DYNAMIC_SHARED.
+ * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. The
+ * HIP-Clang compiler provides support for extern shared declarations.
  * @param [in] stream    Stream where the kernel should be dispatched.  May be 0, in which case th
  * default stream is used with associated synchronization rules.
  * @param [in] kernelParams
@@ -3092,8 +3092,8 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigne
  * @param [in] gridDim   Grid dimensions specified as multiple of blockDim.
  * @param [in] blockDim  Block dimensions specified in work-items
  * @param [in] kernelParams A list of kernel arguments
- * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.  The
- * kernel can access this with HIP_DYNAMIC_SHARED.
+ * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. The
+ * HIP-Clang compiler provides support for extern shared declarations.
  * @param [in] stream    Stream where the kernel should be dispatched.  May be 0, in which case th
  * default stream is used with associated synchronization rules.
  *
@@ -3294,8 +3294,8 @@ hipError_t hipProfilerStop();
  *
  * @param [in] gridDim   grid dimension specified as multiple of blockDim.
  * @param [in] blockDim  block dimensions specified in work-items
- * @param [in] sharedMem Amount of dynamic shared memory to allocate for this kernel.  The
- * kernel can access this with HIP_DYNAMIC_SHARED.
+ * @param [in] sharedMem Amount of dynamic shared memory to allocate for this kernel. The
+ * HIP-Clang compiler provides support for extern shared declarations.
  * @param [in] stream    Stream where the kernel should be dispatched.  May be 0, in which case the
  * default stream is used with associated synchronization rules.
  *
@@ -3334,8 +3334,8 @@ hipError_t hipLaunchByPtr(const void* func);
  *
  * @param [in] gridDim   grid dimension specified as multiple of blockDim.
  * @param [in] blockDim  block dimensions specified in work-items
- * @param [in] sharedMem Amount of dynamic shared memory to allocate for this kernel.  The
- * kernel can access this with HIP_DYNAMIC_SHARED.
+ * @param [in] sharedMem Amount of dynamic shared memory to allocate for this kernel. The
+ * HIP-Clang compiler provides support for extern shared declarations.
  * @param [in] stream    Stream where the kernel should be dispatched.  May be 0, in which case the
  * default stream is used with associated synchronization rules.
  *
@@ -3354,7 +3354,7 @@ hipError_t __hipPushCallConfiguration(dim3 gridDim,
  * @param [out] gridDim   grid dimension specified as multiple of blockDim.
  * @param [out] blockDim  block dimensions specified in work-items
  * @param [out] sharedMem Amount of dynamic shared memory to allocate for this kernel.  The
- * kernel can access this with HIP_DYNAMIC_SHARED.
+ * HIP-Clang compiler provides support for extern shared declarations.
  * @param [out] stream    Stream where the kernel should be dispatched.  May be 0, in which case the
  * default stream is used with associated synchronization rules.
  *
@@ -3373,8 +3373,8 @@ hipError_t __hipPopCallConfiguration(dim3 *gridDim,
  * @param [in] numBlocks - number of blocks
  * @param [in] dimBlocks - dimension of a block
  * @param [in] args - kernel arguments
- * @param [in] sharedMemBytes - Amount of dynamic shared memory to allocate for this kernel.  The
- *  Kernel can access this with HIP_DYNAMIC_SHARED.
+ * @param [in] sharedMemBytes - Amount of dynamic shared memory to allocate for this kernel. The
+ * HIP-Clang compiler provides support for extern shared declarations.
  * @param [in] stream - Stream where the kernel should be dispatched.  May be 0, in which case th
  *  default stream is used with associated synchronization rules.
  *

hipamd/include/hip/hip_ext.h

@@ -43,7 +43,7 @@ THE SOFTWARE.
  * @param [in] blockDimY Y grid dimension specified in work-items
  * @param [in] blockDimZ Z grid dimension specified in work-items
  * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.  The
- kernel can access this with HIP_DYNAMIC_SHARED.
+ HIP-Clang compiler provides support for extern shared declarations
  * @param [in] stream Stream where the kernel should be dispatched.  May be 0, in which case th
  default stream is used with associated synchronization rules.
  * @param [in] kernelParams

hipamd/include/hip/nvidia_detail/hip_runtime.h

@@ -95,14 +95,6 @@ typedef int hipLaunchParm;
 
 #define HIP_SYMBOL(X) &X
 
-/**
- * extern __shared__
- */
-
-#define HIP_DYNAMIC_SHARED(type, var) extern __shared__ type var[];
-
-#define HIP_DYNAMIC_SHARED_ATTRIBUTE
-
 #ifdef __HIP_DEVICE_COMPILE__
 #define abort_()                                                                                    \
     { asm("trap;"); }