From 94e0900377e0cd49531042dd07c9134ae91999f2 Mon Sep 17 00:00:00 2001
From: James Price <j.price@bristol.ac.uk>
Date: Tue, 28 Feb 2017 13:24:45 +0000
Subject: [PATCH] Use static shared memory in dot for CUDA and HIP

---
 CUDAStream.cu | 6 ++----
 HIPStream.cpp | 5 ++---
 2 files changed, 4 insertions(+), 7 deletions(-)

diff --git a/CUDAStream.cu b/CUDAStream.cu
index 7b1e0df..9588456 100644
--- a/CUDAStream.cu
+++ b/CUDAStream.cu
@@ -182,9 +182,7 @@ void CUDAStream<T>::triad()
 template <class T>
 __global__ void dot_kernel(const T * a, const T * b, T * sum, unsigned int array_size)
 {
-
-  extern __shared__ __align__(sizeof(T)) unsigned char smem[];
-  T *tb_sum = reinterpret_cast<T*>(smem);
+  __shared__ T tb_sum[TBSIZE];
 
   int i = blockDim.x * blockIdx.x + threadIdx.x;
   const size_t local_i = threadIdx.x;
@@ -209,7 +207,7 @@ __global__ void dot_kernel(const T * a, const T * b, T * sum, unsigned int array
 template <class T>
 T CUDAStream<T>::dot()
 {
-  dot_kernel<<<DOT_NUM_BLOCKS, TBSIZE, sizeof(T)*TBSIZE>>>(d_a, d_b, d_sum, array_size);
+  dot_kernel<<<DOT_NUM_BLOCKS, TBSIZE>>>(d_a, d_b, d_sum, array_size);
   check_error();
 
   cudaMemcpy(sums, d_sum, DOT_NUM_BLOCKS*sizeof(T), cudaMemcpyDeviceToHost);
diff --git a/HIPStream.cpp b/HIPStream.cpp
index dafe2cd..7bf724a 100644
--- a/HIPStream.cpp
+++ b/HIPStream.cpp
@@ -185,8 +185,7 @@ void HIPStream<T>::triad()
 template <class T>
 __global__ void dot_kernel(hipLaunchParm lp, const T * a, const T * b, T * sum, unsigned int array_size)
 {
-
-  HIP_DYNAMIC_SHARED(T,tb_sum);
+  __shared__ T tb_sum[TBSIZE];
 
   int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
   const size_t local_i = hipThreadIdx_x;
@@ -211,7 +210,7 @@ __global__ void dot_kernel(hipLaunchParm lp, const T * a, const T * b, T * sum,
 template <class T>
 T HIPStream<T>::dot()
 {
-  hipLaunchKernel(HIP_KERNEL_NAME(dot_kernel), dim3(DOT_NUM_BLOCKS), dim3(TBSIZE), sizeof(T)*TBSIZE, 0, d_a, d_b, d_sum, array_size);
+  hipLaunchKernel(HIP_KERNEL_NAME(dot_kernel), dim3(DOT_NUM_BLOCKS), dim3(TBSIZE), 0, 0, d_a, d_b, d_sum, array_size);
   check_error();
 
   hipMemcpy(sums, d_sum, DOT_NUM_BLOCKS*sizeof(T), hipMemcpyDeviceToHost);