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

Add nstream kernel to OpenCL

Browse Source
This commit is contained in:
Tom Deakin 2021-02-02 15:46:53 +00:00
parent fdac285110
commit 4203ccb017
2 changed files with 22 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

20
OCLStream.cpp
View File

@ -61,6 +61,14 @@ std::string kernels{R"CLC(
const size_t i = get_global_id(0); const size_t i = get_global_id(0);
a[i] = b[i] + scalar * c[i]; a[i] = b[i] + scalar * c[i];
} }
kernel void nstream(
global TYPE * restrict a,
global const TYPE * restrict b,
global const TYPE * restrict c)
{
const size_t i = get_global_id(0);
a[i] += b[i] + scalar * c[i];
}
kernel void stream_dot( kernel void stream_dot(
global const TYPE * restrict a, global const TYPE * restrict a,
@ -157,6 +165,7 @@ OCLStream<T>::OCLStream(const 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");
nstream_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer>(program, "nstream");
dot_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer, cl::LocalSpaceArg, cl_int>(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;
@ -186,6 +195,7 @@ OCLStream<T>::~OCLStream()
delete mul_kernel; delete mul_kernel;
delete add_kernel; delete add_kernel;
delete triad_kernel; delete triad_kernel;
delete nstream_kernel;
delete dot_kernel; delete dot_kernel;
devices.clear(); devices.clear();
@ -231,6 +241,16 @@ void OCLStream<T>::triad()
queue.finish(); queue.finish();
} }
template <class T>
void OCLStream<T>::nstream()
{
(*nstream_kernel)(
cl::EnqueueArgs(queue, cl::NDRange(array_size)),
d_a, d_b, d_c
);
queue.finish();
}
template <class T> template <class T>
T OCLStream<T>::dot() T OCLStream<T>::dot()
{ {

2
OCLStream.h
View File

@ -47,6 +47,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> *nstream_kernel;
cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer, cl::LocalSpaceArg, cl_int> *dot_kernel; cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer, cl::LocalSpaceArg, cl_int> *dot_kernel;
// NDRange configuration for the dot kernel // NDRange configuration for the dot kernel
@ -62,6 +63,7 @@ class OCLStream : public Stream<T>
virtual void add() override; virtual void add() override;
virtual void mul() override; virtual void mul() override;
virtual void triad() override; virtual void triad() override;
virtual void nstream() override;
virtual T dot() override; virtual T dot() override;
virtual void init_arrays(T initA, T initB, T initC) override; virtual void init_arrays(T initA, T initB, T initC) override;