From ee7cd066ac50a6a270abc4669a874ba44a946383 Mon Sep 17 00:00:00 2001 From: Peter Steinbach Date: Tue, 28 Feb 2017 10:03:23 +0100 Subject: [PATCH 01/11] 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::HIPStream(const unsigned int ARRAY_SIZE, const int device_index) template HIPStream::~HIPStream() { + free(sums); + hipFree(d_a); check_error(); hipFree(d_b); check_error(); hipFree(d_c); check_error(); + hipFree(d_sum); + check_error(); } + template __global__ void init_kernel(hipLaunchParm lp, T * a, T * b, T * c, T initA, T initB, T initC) { @@ -177,22 +184,22 @@ void HIPStream::triad() check_error(); } - template __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(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 From 350a151c3b05ac251137d49adc2c18ebb3f08880 Mon Sep 17 00:00:00 2001 From: Peter Steinbach Date: Tue, 28 Feb 2017 10:04:36 +0100 Subject: [PATCH 02/11] 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 Date: Tue, 28 Feb 2017 10:07:48 +0100 Subject: [PATCH 03/11] 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 __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(smem); + HIP_DYNAMIC_SHARED(T,tb_sum); + // HIP_DYNAMIC_SHARED(unsigned char,smem); + // T *tb_sum = reinterpret_cast(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 Date: Tue, 28 Feb 2017 13:33:21 +0100 Subject: [PATCH 04/11] 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(smem); int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; const size_t local_i = hipThreadIdx_x; From e570b458a604d3c44bbb9588552626a5250b5f82 Mon Sep 17 00:00:00 2001 From: Peter Steinbach Date: Tue, 28 Feb 2017 13:43:57 +0100 Subject: [PATCH 05/11] replaced - for = so that assignment takes place --- KOKKOSStream.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/KOKKOSStream.cpp b/KOKKOSStream.cpp index 9391a13..9abcf87 100644 --- a/KOKKOSStream.cpp +++ b/KOKKOSStream.cpp @@ -42,7 +42,7 @@ void KOKKOSStream::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(); From 94e0900377e0cd49531042dd07c9134ae91999f2 Mon Sep 17 00:00:00 2001 From: James Price Date: Tue, 28 Feb 2017 13:24:45 +0000 Subject: [PATCH 06/11] 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::triad() template __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(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 T CUDAStream::dot() { - dot_kernel<<>>(d_a, d_b, d_sum, array_size); + dot_kernel<<>>(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::triad() template __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 T HIPStream::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); From ea12f2a9a13041ad114b3828ef51e2e17aa12e88 Mon Sep 17 00:00:00 2001 From: Peter Steinbach Date: Mon, 13 Mar 2017 14:41:16 +0100 Subject: [PATCH 07/11] added EXTRA_FLAGS variable to CUDA Makefile to provide the freedom to specify debug flags or gencode flags --- CUDA.make | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/CUDA.make b/CUDA.make index ef193df..bf376eb 100644 --- a/CUDA.make +++ b/CUDA.make @@ -1,6 +1,7 @@ +EXTRA_FLAGS?=-O3 cuda-stream: main.cpp CUDAStream.cu - nvcc -std=c++11 -O3 -DCUDA $^ $(EXTRA_FLAGS) -o $@ + nvcc -std=c++11 -DCUDA $^ $(EXTRA_FLAGS) -o $@ .PHONY: clean clean: From 8c7a801a8422a145ced2eb3448eaa016d6a553d4 Mon Sep 17 00:00:00 2001 From: Peter Steinbach Date: Mon, 13 Mar 2017 15:22:26 +0100 Subject: [PATCH 08/11] put -O3 into CXXFLAGS to comply with OpenMP.make --- CUDA.make | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/CUDA.make b/CUDA.make index bf376eb..50f175c 100644 --- a/CUDA.make +++ b/CUDA.make @@ -1,7 +1,7 @@ -EXTRA_FLAGS?=-O3 +CXXFLAGS?=-O3 -std=c++11 cuda-stream: main.cpp CUDAStream.cu - nvcc -std=c++11 -DCUDA $^ $(EXTRA_FLAGS) -o $@ + nvcc $(CXXFLAGS) -DCUDA $^ $(EXTRA_FLAGS) -o $@ .PHONY: clean clean: From 4f288ddc3d171ed2b1a77c238ca7219e0301709b Mon Sep 17 00:00:00 2001 From: James Price Date: Mon, 13 Mar 2017 17:15:10 +0000 Subject: [PATCH 09/11] [OpenMP] Add -qopt-streaming-stores for Intel --- OpenMP.make | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/OpenMP.make b/OpenMP.make index 28b5326..2028c46 100644 --- a/OpenMP.make +++ b/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 From 703eb945a2ebd4bfc031b818776cab3e2a216cb0 Mon Sep 17 00:00:00 2001 From: James Price Date: Mon, 13 Mar 2017 17:17:20 +0000 Subject: [PATCH 10/11] [OpenMP] Align memory (2MB by default) --- OMPStream.cpp | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/OMPStream.cpp b/OMPStream.cpp index 189cacb..06346e7 100644 --- a/OMPStream.cpp +++ b/OMPStream.cpp @@ -7,6 +7,10 @@ #include "OMPStream.h" +#ifndef ALIGNMENT +#define ALIGNMENT (2*1024*1024) // 2MB +#endif + template OMPStream::OMPStream(const unsigned int ARRAY_SIZE, T *a, T *b, T *c, int device) { @@ -22,9 +26,9 @@ OMPStream::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->b = (T*)malloc(sizeof(T)*array_size); - this->c = (T*)malloc(sizeof(T)*array_size); + this->a = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size); + this->b = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size); + this->c = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size); #endif } From d8cb7494e0c8e7beddaffadad97f0c7c0238ca71 Mon Sep 17 00:00:00 2001 From: Peter Steinbach Date: Fri, 17 Mar 2017 15:18:13 +0100 Subject: [PATCH 11/11] pulled -O3 out into CXXFLAGS, refactored CUDA compiler into CUDA_CXX make variable to cope with clang as CUDA compiler as well --- CUDA.make | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/CUDA.make b/CUDA.make index 50f175c..3edf0f5 100644 --- a/CUDA.make +++ b/CUDA.make @@ -1,7 +1,8 @@ -CXXFLAGS?=-O3 -std=c++11 +CXXFLAGS=-O3 +CUDA_CXX=nvcc cuda-stream: main.cpp CUDAStream.cu - nvcc $(CXXFLAGS) -DCUDA $^ $(EXTRA_FLAGS) -o $@ + $(CUDA_CXX) -std=c++11 $(CXXFLAGS) -DCUDA $^ $(EXTRA_FLAGS) -o $@ .PHONY: clean clean: