From ee7cd066ac50a6a270abc4669a874ba44a946383 Mon Sep 17 00:00:00 2001
From: Peter Steinbach <steinbac@mpi-cbg.de>
Date: Tue, 28 Feb 2017 10:03:23 +0100
Subject: [PATCH 1/4] renamed HIPStream implementation

---
 HIPStream.cu => HIPStream.cpp | 21 ++++++++++++++-------
 1 file changed, 14 insertions(+), 7 deletions(-)
 rename HIPStream.cu => HIPStream.cpp (93%)

diff --git a/HIPStream.cu b/HIPStream.cpp
similarity index 93%
rename from HIPStream.cu
rename to HIPStream.cpp
index d14fe84..edbcac1 100644
--- a/HIPStream.cu
+++ b/HIPStream.cpp
@@ -1,3 +1,5 @@
+
+
 // Copyright (c) 2015-16 Tom Deakin, Simon McIntosh-Smith,
 // University of Bristol HPC
 //
@@ -72,14 +74,19 @@ HIPStream<T>::HIPStream(const unsigned int ARRAY_SIZE, const int device_index)
 template <class T>
 HIPStream<T>::~HIPStream()
 {
+  free(sums);
+
   hipFree(d_a);
   check_error();
   hipFree(d_b);
   check_error();
   hipFree(d_c);
   check_error();
+  hipFree(d_sum);
+  check_error();
 }
 
+
 template <typename T>
 __global__ void init_kernel(hipLaunchParm lp, T * a, T * b, T * c, T initA, T initB, T initC)
 {
@@ -177,22 +184,22 @@ void HIPStream<T>::triad()
   check_error();
 }
 
-
 template <class T>
 __global__ void dot_kernel(hipLaunchParm lp, const T * a, const T * b, T * sum, unsigned int array_size)
 {
 
-  extern __shared__ __align__(sizeof(T)) unsigned char smem[];
+  //extern __shared__ __align__(sizeof(T)) unsigned char smem[];
+  HIP_DYNAMIC_SHARED(unsigned char,smem);
   T *tb_sum = reinterpret_cast<T*>(smem);
 
-  int i = blockDim.x * blockIdx.x + threadIdx.x;
-  const size_t local_i = threadIdx.x;
+  int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+  const size_t local_i = hipThreadIdx_x;
 
   tb_sum[local_i] = 0.0;
-  for (; i < array_size; i += blockDim.x*gridDim.x)
+  for (; i < array_size; i += hipBlockDim_x*hipGridDim_x)
     tb_sum[local_i] += a[i] * b[i];
 
-  for (int offset = blockDim.x / 2; offset > 0; offset /= 2)
+  for (int offset = hipBlockDim_x / 2; offset > 0; offset /= 2)
   {
     __syncthreads();
     if (local_i < offset)
@@ -202,7 +209,7 @@ __global__ void dot_kernel(hipLaunchParm lp, const T * a, const T * b, T * sum,
   }
 
   if (local_i == 0)
-    sum[blockIdx.x] = tb_sum[local_i];
+    sum[hipBlockIdx_x] = tb_sum[local_i];
 }
 
 template <class T>

From 350a151c3b05ac251137d49adc2c18ebb3f08880 Mon Sep 17 00:00:00 2001
From: Peter Steinbach <steinbac@mpi-cbg.de>
Date: Tue, 28 Feb 2017 10:04:36 +0100
Subject: [PATCH 2/4] removed CUDA_PATH sentinel from HIP.make

---
 HIP.make | 8 +-------
 1 file changed, 1 insertion(+), 7 deletions(-)

diff --git a/HIP.make b/HIP.make
index bbd142e..35b0a6a 100644
--- a/HIP.make
+++ b/HIP.make
@@ -3,13 +3,7 @@
 
 HIPCC = hipcc
 
-ifndef CUDA_PATH
-ifeq (,$(wildcard /usr/local/bin/nvcc))
-$(error /usr/local/bin/nvcc not found, set CUDA_PATH instead)
-endif
-endif
-
-hip-stream: main.cpp HIPStream.cu
+hip-stream: main.cpp HIPStream.cpp
 	$(HIPCC) $(CXXFLAGS) -std=c++11 -DHIP $^ $(EXTRA_FLAGS) -o $@
 
 .PHONY: clean

From ceada6922f19449af7b64d3d0ded84dd4188afa2 Mon Sep 17 00:00:00 2001
From: Peter Steinbach <steinbac@mpi-cbg.de>
Date: Tue, 28 Feb 2017 10:07:48 +0100
Subject: [PATCH 3/4] proper declaration of tb_sum with HIP_DYNAMIC_SHARED
 macro

---
 HIPStream.cpp | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/HIPStream.cpp b/HIPStream.cpp
index edbcac1..648b62a 100644
--- a/HIPStream.cpp
+++ b/HIPStream.cpp
@@ -188,9 +188,9 @@ template <class T>
 __global__ void dot_kernel(hipLaunchParm lp, const T * a, const T * b, T * sum, unsigned int array_size)
 {
 
-  //extern __shared__ __align__(sizeof(T)) unsigned char smem[];
-  HIP_DYNAMIC_SHARED(unsigned char,smem);
-  T *tb_sum = reinterpret_cast<T*>(smem);
+  HIP_DYNAMIC_SHARED(T,tb_sum);
+  // HIP_DYNAMIC_SHARED(unsigned char,smem);
+  // T *tb_sum = reinterpret_cast<T*>(smem);
 
   int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
   const size_t local_i = hipThreadIdx_x;

From 58773a79b747b21860101d22af1d53c4f1e19186 Mon Sep 17 00:00:00 2001
From: Peter Steinbach <steinbac@mpi-cbg.de>
Date: Tue, 28 Feb 2017 13:33:21 +0100
Subject: [PATCH 4/4] removed extra lines introduced by hipify, removed
 obsolete commented code

---
 HIPStream.cpp | 4 ----
 1 file changed, 4 deletions(-)

diff --git a/HIPStream.cpp b/HIPStream.cpp
index 648b62a..dafe2cd 100644
--- a/HIPStream.cpp
+++ b/HIPStream.cpp
@@ -1,5 +1,3 @@
-
-
 // Copyright (c) 2015-16 Tom Deakin, Simon McIntosh-Smith,
 // University of Bristol HPC
 //
@@ -189,8 +187,6 @@ __global__ void dot_kernel(hipLaunchParm lp, const T * a, const T * b, T * sum,
 {
 
   HIP_DYNAMIC_SHARED(T,tb_sum);
-  // HIP_DYNAMIC_SHARED(unsigned char,smem);
-  // T *tb_sum = reinterpret_cast<T*>(smem);
 
   int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
   const size_t local_i = hipThreadIdx_x;