Skip to content

Commit

Permalink
Add dot kernel to HIP implementation
Browse files Browse the repository at this point in the history
  • Loading branch information
jrprice committed Feb 23, 2017
1 parent e6ac807 commit ce4f49e
Show file tree
Hide file tree
Showing 2 changed files with 55 additions and 0 deletions.
49 changes: 49 additions & 0 deletions HIPStream.cu
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,7 @@
#include "hip/hip_runtime.h"

#define TBSIZE 1024
#define DOT_NUM_BLOCKS 256

void check_error(void)
{
Expand Down Expand Up @@ -47,6 +48,9 @@ HIPStream<T>::HIPStream(const unsigned int ARRAY_SIZE, const int device_index)

array_size = ARRAY_SIZE;

// Allocate the host array for partial sums for dot kernels
sums = (T*)malloc(sizeof(T) * DOT_NUM_BLOCKS);

// Check buffers fit on the device
hipDeviceProp_t props;
hipGetDeviceProperties(&props, 0);
Expand All @@ -60,6 +64,8 @@ HIPStream<T>::HIPStream(const unsigned int ARRAY_SIZE, const int device_index)
check_error();
hipMalloc(&d_c, ARRAY_SIZE*sizeof(T));
check_error();
hipMalloc(&d_sum, DOT_NUM_BLOCKS*sizeof(T));
check_error();
}


Expand Down Expand Up @@ -172,6 +178,49 @@ void HIPStream<T>::triad()
}


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[];
T *tb_sum = reinterpret_cast<T*>(smem);

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)
tb_sum[local_i] += a[i] * b[i];

for (int offset = blockDim.x / 2; offset > 0; offset /= 2)
{
__syncthreads();
if (local_i < offset)
{
tb_sum[local_i] += tb_sum[local_i+offset];
}
}

if (local_i == 0)
sum[blockIdx.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);
check_error();

hipMemcpy(sums, d_sum, DOT_NUM_BLOCKS*sizeof(T), hipMemcpyDeviceToHost);
check_error();

T sum = 0.0;
for (int i = 0; i < DOT_NUM_BLOCKS; i++)
sum += sums[i];

return sum;
}

void listDevices(void)
{
// Get number of devices
Expand Down
6 changes: 6 additions & 0 deletions HIPStream.h
Original file line number Diff line number Diff line change
Expand Up @@ -21,10 +21,15 @@ class HIPStream : public Stream<T>
protected:
// Size of arrays
unsigned int array_size;

// Host array for partial sums for dot kernel
T *sums;

// Device side pointers to arrays
T *d_a;
T *d_b;
T *d_c;
T *d_sum;


public:
Expand All @@ -36,6 +41,7 @@ class HIPStream : public Stream<T>
virtual void add() override;
virtual void mul() override;
virtual void triad() override;
virtual T dot() override;

virtual void init_arrays(T initA, T initB, T initC) override;
virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
Expand Down

0 comments on commit ce4f49e

Please sign in to comment.