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

enclosing computecpp specific code in macros, rather than removing it

Browse Source
This commit is contained in:
GeorgeWeb 2019-06-20 10:13:39 +01:00
parent a2e53d6728
commit 54737d87cb
2 changed files with 50 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

50
SYCLStream.cpp
View File

@ -16,6 +16,9 @@ using namespace cl::sycl;
bool cached = false;
std::vector<device> devices;
void getDeviceList(void);
#ifdef COMPUTECPP_CE
program * p;
#endif
template <class T>
SYCLStream<T>::SYCLStream(const unsigned int ARRAY_SIZE, const int device_index)
@ -67,6 +70,17 @@ SYCLStream<T>::SYCLStream(const unsigned int ARRAY_SIZE, const int device_index)
}
});
#ifdef COMPUTECPP_CE
/* Pre-build the kernels */
p = new program(queue->get_context());
p->build_with_kernel_type<init_kernel>();
p->build_with_kernel_type<copy_kernel>();
p->build_with_kernel_type<mul_kernel>();
p->build_with_kernel_type<add_kernel>();
p->build_with_kernel_type<triad_kernel>();
p->build_with_kernel_type<dot_kernel>();
#endif
// Create buffers
d_a = new buffer<T>(array_size);
d_b = new buffer<T>(array_size);
@ -81,8 +95,9 @@ SYCLStream<T>::~SYCLStream()
delete d_b;
delete d_c;
delete d_sum;
#ifdef COMPUTECPP_CE
delete p;
#endif
delete queue;
devices.clear();
}
@ -94,7 +109,12 @@ void SYCLStream<T>::copy()
{
auto ka = d_a->template get_access<access::mode::read>(cgh);
auto kc = d_c->template get_access<access::mode::write>(cgh);
#ifdef COMPUTECPP_CE
cgh.parallel_for<copy_kernel>(p->get_kernel<copy_kernel>(),
range<1>{array_size}, [=](item<1> item)
#else
cgh.parallel_for<copy_kernel>(range<1>{array_size}, [=](item<1> item)
#endif
{
auto id = item.get_id(0);
kc[id] = ka[id];
@ -111,7 +131,12 @@ void SYCLStream<T>::mul()
{
auto kb = d_b->template get_access<access::mode::write>(cgh);
auto kc = d_c->template get_access<access::mode::read>(cgh);
#ifdef COMPUTECPP_CE
cgh.parallel_for<mul_kernel>(p->get_kernel<mul_kernel>(),
range<1>{array_size}, [=](item<1> item)
#else
cgh.parallel_for<mul_kernel>(range<1>{array_size}, [=](item<1> item)
#endif
{
auto id = item.get_id(0);
kb[id] = scalar * kc[id];
@ -128,7 +153,12 @@ void SYCLStream<T>::add()
auto ka = d_a->template get_access<access::mode::read>(cgh);
auto kb = d_b->template get_access<access::mode::read>(cgh);
auto kc = d_c->template get_access<access::mode::write>(cgh);
#ifdef COMPUTECPP_CE
cgh.parallel_for<add_kernel>(p->get_kernel<add_kernel>(),
range<1>{array_size}, [=](item<1> item)
#else
cgh.parallel_for<add_kernel>(range<1>{array_size}, [=](item<1> item)
#endif
{
auto id = item.get_id(0);
kc[id] = ka[id] + kb[id];
@ -146,7 +176,12 @@ void SYCLStream<T>::triad()
auto ka = d_a->template get_access<access::mode::write>(cgh);
auto kb = d_b->template get_access<access::mode::read>(cgh);
auto kc = d_c->template get_access<access::mode::read>(cgh);
#ifdef COMPUTECPP_CE
cgh.parallel_for<triad_kernel>(p->get_kernel<triad_kernel>(),
range<1>{array_size}, [=](item<1> item)
#else
cgh.parallel_for<triad_kernel>(range<1>{array_size}, [=](item<1> item)
#endif
{
auto id = item.get_id(0);
ka[id] = kb[id] + scalar * kc[id];
@ -167,8 +202,12 @@ T SYCLStream<T>::dot()
auto wg_sum = accessor<T, 1, access::mode::read_write, access::target::local>(range<1>(dot_wgsize), cgh);
size_t N = array_size;
#ifdef COMPUTECPP_CE
cgh.parallel_for<dot_kernel>(p->get_kernel<dot_kernel>(),
nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item)
#else
cgh.parallel_for<dot_kernel>(nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item)
#endif
{
size_t i = item.get_global_id(0);
size_t li = item.get_local_id(0);
@ -209,7 +248,12 @@ void SYCLStream<T>::init_arrays(T initA, T initB, T initC)
auto ka = d_a->template get_access<access::mode::write>(cgh);
auto kb = d_b->template get_access<access::mode::write>(cgh);
auto kc = d_c->template get_access<access::mode::write>(cgh);
#ifdef COMPUTECPP_CE
cgh.parallel_for<init_kernel>(p->get_kernel<init_kernel>(),
range<1>{array_size}, [=](item<1> item)
#else
cgh.parallel_for<init_kernel>(range<1>{array_size}, [=](item<1> item)
#endif
{
auto id = item.get_id(0);
ka[id] = initA;

3
SYCLStream.h
View File

@ -15,6 +15,9 @@
#define IMPLEMENTATION_STRING "SYCL"
// allows a use of 'parallel_for' currently known to be supported by ComputeCpp
#define COMPUTECPP_CE
namespace sycl_kernels
{
template <class T> class init;