From ce2dff449d4de1732163be8d04bf28665c351e8c Mon Sep 17 00:00:00 2001 From: Julia Jiang Date: Thu, 4 Feb 2021 16:22:01 -0500 Subject: [PATCH] SWDEV-271416 - Remove HIP_DYNAMIC_SHARED macro in hip Change-Id: I12f39ea8438eb7ce76d8ffb2151b4faa93689048 --- samples/2_Cookbook/6_dynamic_shared/Readme.md | 9 ++++++--- samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp | 3 +-- samples/2_Cookbook/7_streams/stream.cpp | 3 +-- samples/2_Cookbook/8_peer2peer/peer2peer.cpp | 3 +-- 4 files changed, 9 insertions(+), 9 deletions(-) diff --git a/samples/2_Cookbook/6_dynamic_shared/Readme.md b/samples/2_Cookbook/6_dynamic_shared/Readme.md index 68782807bf..02a8cb3da2 100644 --- a/samples/2_Cookbook/6_dynamic_shared/Readme.md +++ b/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: ``` diff --git a/samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp b/samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp index 743da8f63c..65922afc86 100644 --- a/samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp +++ b/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; diff --git a/samples/2_Cookbook/7_streams/stream.cpp b/samples/2_Cookbook/7_streams/stream.cpp index dd87e4363e..f2bf289c9d 100644 --- a/samples/2_Cookbook/7_streams/stream.cpp +++ b/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; diff --git a/samples/2_Cookbook/8_peer2peer/peer2peer.cpp b/samples/2_Cookbook/8_peer2peer/peer2peer.cpp index 8e011b1281..1f61e65876 100644 --- a/samples/2_Cookbook/8_peer2peer/peer2peer.cpp +++ b/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;