   const int wid = threadIdx.x / GPU_WARP_SIZE;
 
   // Shape is 2 x warp_count_in_block.
-  HIP_DYNAMIC_SHARED( unsigned char, shared_memory_)
+  extern __shared__ unsigned char shared_memory_[];
   TBuf* shared_memory = reinterpret_cast<TBuf*>(shared_memory_);
 
   if (lid == 0) {

namespace onnxruntime {
namespace rocm {

template<typename Tin, typename Tout>
__global__ void ScalarSqrtKernel(Tin* input, Tout* output) {
  *output = (Tout)_Sqrt(*input);
}

template<typename Tin, typename Tout>
void ScalarSqrt(Tin* input, Tout* output) {
  hipLaunchKernelGGL(ScalarSqrtKernel, dim3(1), dim3(1), 0, 0, input, output);
}

  }

  if (threadIdx.x == 0) {
    atomic_add(w_norm, TOutOp()(TOut(shared_memory[0])));
  }
}

