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 diff --git a/HIPStream.cu b/HIPStream.cpp similarity index 94% rename from HIPStream.cu rename to HIPStream.cpp index d14fe84..dafe2cd 100644 --- a/HIPStream.cu +++ b/HIPStream.cpp @@ -72,14 +72,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 +182,20 @@ 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[]; - T *tb_sum = reinterpret_cast(smem); + HIP_DYNAMIC_SHARED(T,tb_sum); - 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 +205,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