cory-master/BabelStream
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

renamed HIPStream implementation

Browse Source
This commit is contained in:
Peter Steinbach 2017-02-28 10:03:23 +01:00
parent cc90cefeeb
commit ee7cd066ac
1 changed files with 14 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

21
HIPStream.cu → HIPStream.cpp
View File

@ -1,3 +1,5 @@
// Copyright (c) 2015-16 Tom Deakin, Simon McIntosh-Smith, // Copyright (c) 2015-16 Tom Deakin, Simon McIntosh-Smith,
// University of Bristol HPC // University of Bristol HPC
// //
@ -72,14 +74,19 @@ HIPStream<T>::HIPStream(const unsigned int ARRAY_SIZE, const int device_index)
template <class T> template <class T>
HIPStream<T>::~HIPStream() HIPStream<T>::~HIPStream()
{ {
free(sums);
hipFree(d_a); hipFree(d_a);
check_error(); check_error();
hipFree(d_b); hipFree(d_b);
check_error(); check_error();
hipFree(d_c); hipFree(d_c);
check_error(); check_error();
hipFree(d_sum);
check_error();
} }
template <typename T> template <typename T>
__global__ void init_kernel(hipLaunchParm lp, T * a, T * b, T * c, T initA, T initB, T initC) __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(); check_error();
} }
template <class T> template <class T>
__global__ void dot_kernel(hipLaunchParm lp, const T * a, const T * b, T * sum, unsigned int array_size) __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); T *tb_sum = reinterpret_cast<T*>(smem);
int i = blockDim.x * blockIdx.x + threadIdx.x; int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
const size_t local_i = threadIdx.x; const size_t local_i = hipThreadIdx_x;
tb_sum[local_i] = 0.0; 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]; 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(); __syncthreads();
if (local_i < offset) 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) if (local_i == 0)
sum[blockIdx.x] = tb_sum[local_i]; sum[hipBlockIdx_x] = tb_sum[local_i];
} }
template <class T> template <class T>