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

Merge pull request #58 from GeorgeWeb/sycl-compliant

Browse Source 
Making BabelStream's SYCL code compliant
This commit is contained in:
Tom Deakin 2019-06-26 12:03:47 +01:00 committed by GitHub
commit 022793bdd6
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
1 changed files with 9 additions and 27 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

34
SYCLStream.cpp
View File

@ -11,7 +11,6 @@
using namespace cl::sycl; using namespace cl::sycl;
// Cache list of devices // Cache list of devices
bool cached = false; bool cached = false;
std::vector<device> devices; std::vector<device> devices;
@ -68,15 +67,6 @@ SYCLStream<T>::SYCLStream(const unsigned int ARRAY_SIZE, const int device_index)
} }
}); });
/* 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>();
// Create buffers // Create buffers
d_a = new buffer<T>(array_size); d_a = new buffer<T>(array_size);
d_b = new buffer<T>(array_size); d_b = new buffer<T>(array_size);
@ -91,7 +81,6 @@ SYCLStream<T>::~SYCLStream()
delete d_b; delete d_b;
delete d_c; delete d_c;
delete d_sum; delete d_sum;
delete p; delete p;
delete queue; delete queue;
devices.clear(); devices.clear();
@ -104,8 +93,7 @@ void SYCLStream<T>::copy()
{ {
auto ka = d_a->template get_access<access::mode::read>(cgh); auto ka = d_a->template get_access<access::mode::read>(cgh);
auto kc = d_c->template get_access<access::mode::write>(cgh); auto kc = d_c->template get_access<access::mode::write>(cgh);
cgh.parallel_for<copy_kernel>(p->get_kernel<copy_kernel>(), cgh.parallel_for<copy_kernel>(range<1>{array_size}, [=](item<1> item)
range<1>{array_size}, [=](item<1> item)
{ {
auto id = item.get_id(0); auto id = item.get_id(0);
kc[id] = ka[id]; kc[id] = ka[id];
@ -122,8 +110,7 @@ void SYCLStream<T>::mul()
{ {
auto kb = d_b->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::read>(cgh); auto kc = d_c->template get_access<access::mode::read>(cgh);
cgh.parallel_for<mul_kernel>(p->get_kernel<mul_kernel>(), cgh.parallel_for<mul_kernel>(range<1>{array_size}, [=](item<1> item)
range<1>{array_size}, [=](item<1> item)
{ {
auto id = item.get_id(0); auto id = item.get_id(0);
kb[id] = scalar * kc[id]; kb[id] = scalar * kc[id];
@ -140,8 +127,7 @@ void SYCLStream<T>::add()
auto ka = d_a->template get_access<access::mode::read>(cgh); auto ka = d_a->template get_access<access::mode::read>(cgh);
auto kb = d_b->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); auto kc = d_c->template get_access<access::mode::write>(cgh);
cgh.parallel_for<add_kernel>(p->get_kernel<add_kernel>(), cgh.parallel_for<add_kernel>(range<1>{array_size}, [=](item<1> item)
range<1>{array_size}, [=](item<1> item)
{ {
auto id = item.get_id(0); auto id = item.get_id(0);
kc[id] = ka[id] + kb[id]; kc[id] = ka[id] + kb[id];
@ -159,8 +145,7 @@ void SYCLStream<T>::triad()
auto ka = d_a->template get_access<access::mode::write>(cgh); auto ka = d_a->template get_access<access::mode::write>(cgh);
auto kb = d_b->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::read>(cgh); auto kc = d_c->template get_access<access::mode::read>(cgh);
cgh.parallel_for<triad_kernel>(p->get_kernel<triad_kernel>(), cgh.parallel_for<triad_kernel>(range<1>{array_size}, [=](item<1> item)
range<1>{array_size}, [=](item<1> item)
{ {
auto id = item.get_id(0); auto id = item.get_id(0);
ka[id] = kb[id] + scalar * kc[id]; ka[id] = kb[id] + scalar * kc[id];
@ -181,12 +166,10 @@ T SYCLStream<T>::dot()
auto wg_sum = accessor<T, 1, access::mode::read_write, access::target::local>(range<1>(dot_wgsize), cgh); auto wg_sum = accessor<T, 1, access::mode::read_write, access::target::local>(range<1>(dot_wgsize), cgh);
size_t N = array_size; size_t N = array_size;
cgh.parallel_for<dot_kernel>(nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item)
cgh.parallel_for<dot_kernel>(p->get_kernel<dot_kernel>(),
nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item)
{ {
size_t i = item.get_global(0); size_t i = item.get_global_id(0);
size_t li = item.get_local(0); size_t li = item.get_local_id(0);
size_t global_size = item.get_global_range()[0]; size_t global_size = item.get_global_range()[0];
wg_sum[li] = 0.0; wg_sum[li] = 0.0;
@ -224,8 +207,7 @@ void SYCLStream<T>::init_arrays(T initA, T initB, T initC)
auto ka = d_a->template get_access<access::mode::write>(cgh); auto ka = d_a->template get_access<access::mode::write>(cgh);
auto kb = d_b->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); auto kc = d_c->template get_access<access::mode::write>(cgh);
cgh.parallel_for<init_kernel>(p->get_kernel<init_kernel>(), cgh.parallel_for<init_kernel>(range<1>{array_size}, [=](item<1> item)
range<1>{array_size}, [=](item<1> item)
{ {
auto id = item.get_id(0); auto id = item.get_id(0);
ka[id] = initA; ka[id] = initA;