Merge remote-tracking branch 'upstream/master' into rocm_hc_support

CUDA.make

@@ -1,6 +1,8 @@
+CXXFLAGS=-O3
+CUDA_CXX=nvcc
 
 cuda-stream: main.cpp CUDAStream.cu
-	nvcc -std=c++11 -O3 -DCUDA $^ $(EXTRA_FLAGS) -o $@
+	$(CUDA_CXX) -std=c++11 $(CXXFLAGS) -DCUDA $^ $(EXTRA_FLAGS) -o $@
 
 .PHONY: clean
 clean:

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)
 {
-
+  __shared__ T tb_sum[TBSIZE];
-  extern __shared__ __align__(sizeof(T)) 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;
@@ -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);

HIP.make

@@ -3,13 +3,7 @@
 
 HIPCC = hipcc
 
-ifndef CUDA_PATH
+hip-stream: main.cpp HIPStream.cpp
-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
 	$(HIPCC) $(CXXFLAGS) -std=c++11 -DHIP $^ $(EXTRA_FLAGS) -o $@
 
 .PHONY: clean

HIPStream.cpp

@@ -72,14 +72,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 +182,19 @@ 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)
 {
+  __shared__ T tb_sum[TBSIZE];
 
-  extern __shared__ __align__(sizeof(T)) unsigned char smem[];
+  int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-  T *tb_sum = reinterpret_cast<T*>(smem);
+  const size_t local_i = hipThreadIdx_x;
-
-  int i = blockDim.x * blockIdx.x + threadIdx.x;
-  const size_t local_i = threadIdx.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,13 +204,13 @@ __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>
 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);

KOKKOSStream.cpp

@@ -42,7 +42,7 @@ void KOKKOSStream<T>::init_arrays(T initA, T initB, T initC)
   parallel_for(array_size, KOKKOS_LAMBDA (const int index)
   {
     a[index] = initA;
-    b[index] - initB;
+    b[index] = initB;
     c[index] = initC;
   });
   Kokkos::fence();

OMPStream.cpp

@@ -7,6 +7,10 @@
 
 #include "OMPStream.h"
 
+#ifndef ALIGNMENT
+#define ALIGNMENT (2*1024*1024) // 2MB
+#endif
+
 template <class T>
 OMPStream<T>::OMPStream(const unsigned int ARRAY_SIZE, T *a, T *b, T *c, int device)
 {
@@ -22,9 +26,9 @@ OMPStream<T>::OMPStream(const unsigned int ARRAY_SIZE, T *a, T *b, T *c, int dev
   {}
 #else
   // Allocate on the host
-  this->a = (T*)malloc(sizeof(T)*array_size);
+  this->a = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size);
-  this->b = (T*)malloc(sizeof(T)*array_size);
+  this->b = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size);
-  this->c = (T*)malloc(sizeof(T)*array_size);
+  this->c = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size);
 #endif
 }
 

OpenMP.make

@@ -29,7 +29,7 @@ COMPILER_XL = xlc++
 CXX = $(COMPILER_$(COMPILER))
 
 FLAGS_GNU = -O3 -std=c++11
-FLAGS_INTEL = -O3 -std=c++11 -xHOST
+FLAGS_INTEL = -O3 -std=c++11 -xHOST -qopt-streaming-stores=always
 FLAGS_CRAY = -O3 -hstd=c++11
 FLAGS_CLANG = -O3 -std=c++11
 FLAGS_XL = -O5 -qarch=pwr8 -qtune=pwr8 -std=c++11