SWDEV-271416 - Remove HIP_DYNAMIC_SHARED macro in hip

Change-Id: I12f39ea8438eb7ce76d8ffb2151b4faa93689048

samples/2_Cookbook/6_dynamic_shared/Readme.md

@@ -19,10 +19,13 @@ We will be using the Simple Matrix Transpose application from the previous tutor
 
 ## Shared Memory
 
-Shared memory is way more faster than that of global and constant memory and accessible to all the threads in the block. For  In the same sourcecode, we will use the `HIP_DYNAMIC_SHARED` keyword to declare dynamic shared memory as follows:
+Shared memory is way more faster than that of global and constant memory and accessible to all the threads in the block.
+
+Previously, it was essential to declare dynamic shared memory using the HIP_DYNAMIC_SHARED macro for accuracy, as using static shared memory in the same kernel could result in overlapping memory ranges and data-races.
+
+Now, the HIP-Clang compiler provides support for extern shared declarations, and the HIP_DYNAMIC_SHARED option is no longer required. You may use the standard extern definition:
+extern __shared__ type var[];
 
-`  HIP_DYNAMIC_SHARED(float, sharedMem)                                               `
-here the first parameter is the data type while the second one is the variable name.
 
 The other important change is:
 ```

samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp

@@ -35,8 +35,7 @@ THE SOFTWARE.
 
 // Device (Kernel) function, it must be void
 __global__ void matrixTranspose(float* out, float* in, const int width) {
-    // declare dynamic shared memory
-    HIP_DYNAMIC_SHARED(float, sharedMem);
+    extern __shared__ float sharedMem[];
 
     int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
     int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;

samples/2_Cookbook/7_streams/stream.cpp

@@ -49,8 +49,7 @@ __global__ void matrixTranspose_static_shared(float* out, float* in,
 
 __global__ void matrixTranspose_dynamic_shared(float* out, float* in,
                                                const int width) {
-    // declare dynamic shared memory
-    HIP_DYNAMIC_SHARED(float, sharedMem)
+    extern __shared__ float sharedMem[];
 
     int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
     int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;

samples/2_Cookbook/8_peer2peer/peer2peer.cpp

@@ -122,8 +122,7 @@ __global__ void matrixTranspose_static_shared(float* out, float* in,
 
 __global__ void matrixTranspose_dynamic_shared(float* out, float* in,
                                                const int width) {
-    // declare dynamic shared memory
-    HIP_DYNAMIC_SHARED(float, sharedMem)
+    extern __shared__ float sharedMem[];
 
     int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
     int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;