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

Improve performance of OpenCL dot implementation

Browse Source
This commit is contained in:
James Price 2016-10-24 21:26:09 +01:00
parent 644ebc40ef
commit d5482b74f4
2 changed files with 16 additions and 12 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

19
OCLStream.cpp
View File

@ -54,11 +54,14 @@ std::string kernels{R"CLC(
global const TYPE * restrict a, global const TYPE * restrict a,
global const TYPE * restrict b, global const TYPE * restrict b,
global TYPE * restrict sum, global TYPE * restrict sum,
local TYPE * restrict wg_sum) local TYPE * restrict wg_sum,
int array_size)
{ {
const size_t i = get_global_id(0); size_t i = get_global_id(0);
const size_t local_i = get_local_id(0); const size_t local_i = get_local_id(0);
wg_sum[local_i] = a[i] * b[i]; wg_sum[local_i] = 0.0;
for (; i < array_size; i += get_global_size(0))
wg_sum[local_i] += a[i] * b[i];
for (int offset = get_local_size(0) / 2; offset > 0; offset /= 2) for (int offset = get_local_size(0) / 2; offset > 0; offset /= 2)
{ {
@ -128,7 +131,7 @@ OCLStream<T>::OCLStream(const unsigned int ARRAY_SIZE, const int device_index)
mul_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer>(program, "mul"); mul_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer>(program, "mul");
add_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer>(program, "add"); add_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer>(program, "add");
triad_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer>(program, "triad"); triad_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer>(program, "triad");
dot_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer, cl::LocalSpaceArg>(program, "stream_dot"); dot_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer, cl::LocalSpaceArg, cl_int>(program, "stream_dot");
array_size = ARRAY_SIZE; array_size = ARRAY_SIZE;
@ -144,9 +147,9 @@ OCLStream<T>::OCLStream(const unsigned int ARRAY_SIZE, const int device_index)
d_a = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(T) * ARRAY_SIZE); d_a = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(T) * ARRAY_SIZE);
d_b = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(T) * ARRAY_SIZE); d_b = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(T) * ARRAY_SIZE);
d_c = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(T) * ARRAY_SIZE); d_c = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(T) * ARRAY_SIZE);
d_sum = cl::Buffer(context, CL_MEM_WRITE_ONLY, sizeof(T) * (ARRAY_SIZE/WGSIZE)); d_sum = cl::Buffer(context, CL_MEM_WRITE_ONLY, sizeof(T) * DOT_NUM_GROUPS);
sums = std::vector<T>(ARRAY_SIZE/WGSIZE); sums = std::vector<T>(DOT_NUM_GROUPS);
} }
template <class T> template <class T>
@ -202,8 +205,8 @@ template <class T>
T OCLStream<T>::dot() T OCLStream<T>::dot()
{ {
(*dot_kernel)( (*dot_kernel)(
cl::EnqueueArgs(queue, cl::NDRange(array_size), cl::NDRange(WGSIZE)), cl::EnqueueArgs(queue, cl::NDRange(DOT_NUM_GROUPS*DOT_WGSIZE), cl::NDRange(DOT_WGSIZE)),
d_a, d_b, d_sum, cl::Local(sizeof(T) * WGSIZE) d_a, d_b, d_sum, cl::Local(sizeof(T) * DOT_WGSIZE), array_size
); );
cl::copy(queue, d_sum, sums.begin(), sums.end()); cl::copy(queue, d_sum, sums.begin(), sums.end());

7
OCLStream.h
View File

@ -21,8 +21,9 @@
#define IMPLEMENTATION_STRING "OpenCL" #define IMPLEMENTATION_STRING "OpenCL"
// Local work-group size for dot kernel // NDRange configuration for the dot kernel
#define WGSIZE 256 #define DOT_WGSIZE 256
#define DOT_NUM_GROUPS 256
template <class T> template <class T>
class OCLStream : public Stream<T> class OCLStream : public Stream<T>
@ -49,7 +50,7 @@ class OCLStream : public Stream<T>
cl::KernelFunctor<cl::Buffer, cl::Buffer> * mul_kernel; cl::KernelFunctor<cl::Buffer, cl::Buffer> * mul_kernel;
cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer> *add_kernel; cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer> *add_kernel;
cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer> *triad_kernel; cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer> *triad_kernel;
cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer, cl::LocalSpaceArg> *dot_kernel; cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer, cl::LocalSpaceArg, cl_int> *dot_kernel;
public: public: