message: special casing the declaration of GPU shared memory variables for the ROCm platform
diff: @@ -291,10 +291,15 @@ __global__ void ColumnReduceMax16ColumnsKernel(
  // This is to mimic the following, but without any constructors:
  // __shared__ storage_type<value_type> partial_sums[TF_RED_WARPSIZE *
  // (TF_RED_WARPSIZE+1)];
  +#if GOOGLE_CUDA || TENSORFLOW_COMPILER_IS_HIP_CLANG
  __shared__ __align__(alignof(value_type)) char
  partial_sums_raw[TF_RED_WARPSIZE * (TF_RED_WARPSIZE + 1) *
  sizeof(value_type)];
  value_type* partial_sums = reinterpret_cast<value_type*>(partial_sums_raw);
  +#elif TENSORFLOW_USE_ROCM
  + __shared__ storage_type<value_type>
  + partial_sums[TF_RED_WARPSIZE * (TF_RED_WARPSIZE + 1)];
  +#endif
  row += rows_per_warp * gridDim.y * blockDim.y;
  for (; row < num_rows; row += rows_per_warp * gridDim.y * blockDim.y) {
  @@ -347,10 +352,15 @@ __global__ void ColumnReduceKernel(
  // This is to mimic the following, but without constructors:
  // __shared__ storage_type<value_type> partial_sums[TF_RED_WARPSIZE *
  // (TF_RED_WARPSIZE + 1)];
  +#if GOOGLE_CUDA || TENSORFLOW_COMPILER_IS_HIP_CLANG
  __shared__ __align__(alignof(value_type)) char
  partial_sums_raw[TF_RED_WARPSIZE * (TF_RED_WARPSIZE + 1) *
  sizeof(value_type)];
  value_type* partial_sums = reinterpret_cast<value_type*>(partial_sums_raw);
  +#elif TENSORFLOW_USE_ROCM
  + __shared__ storage_type<value_type>
  + partial_sums[TF_RED_WARPSIZE * (TF_RED_WARPSIZE + 1)];
  +#endif
  row += gridDim.y * blockDim.y;
  