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

[SYCL 2020] Use unnamed lamdas

Browse Source
This commit is contained in:
Tom Deakin 2021-01-12 11:00:54 +00:00
parent 903eb40d19
commit 8c72b52f16
2 changed files with 6 additions and 24 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
SYCLStream.cpp
View File

@ -91,7 +91,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>(range<1>{array_size}, [=](id<1> idx) cgh.parallel_for(range<1>{array_size}, [=](id<1> idx)
{ {
kc[idx] = ka[idx]; kc[idx] = ka[idx];
}); });
@ -107,7 +107,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>(range<1>{array_size}, [=](id<1> idx) cgh.parallel_for(range<1>{array_size}, [=](id<1> idx)
{ {
kb[idx] = scalar * kc[idx]; kb[idx] = scalar * kc[idx];
}); });
@ -123,7 +123,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>(range<1>{array_size}, [=](id<1> idx) cgh.parallel_for(range<1>{array_size}, [=](id<1> idx)
{ {
kc[idx] = ka[idx] + kb[idx]; kc[idx] = ka[idx] + kb[idx];
}); });
@ -140,7 +140,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>(range<1>{array_size}, [=](id<1> idx) cgh.parallel_for(range<1>{array_size}, [=](id<1> idx)
{ {
ka[idx] = kb[idx] + scalar * kc[idx]; ka[idx] = kb[idx] + scalar * kc[idx];
}); });
@ -160,7 +160,7 @@ 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(nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item)
{ {
size_t i = item.get_global_id(0); size_t i = item.get_global_id(0);
size_t li = item.get_local_id(0); size_t li = item.get_local_id(0);
@ -201,7 +201,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>(range<1>{array_size}, [=](item<1> item) cgh.parallel_for(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;

18
SYCLStream.h
View File

@ -15,16 +15,6 @@
#define IMPLEMENTATION_STRING "SYCL" #define IMPLEMENTATION_STRING "SYCL"
namespace sycl_kernels
{
template <class T> class init;
template <class T> class copy;
template <class T> class mul;
template <class T> class add;
template <class T> class triad;
template <class T> class dot;
}
template <class T> template <class T>
class SYCLStream : public Stream<T> class SYCLStream : public Stream<T>
{ {
@ -39,14 +29,6 @@ class SYCLStream : public Stream<T>
cl::sycl::buffer<T> *d_c; cl::sycl::buffer<T> *d_c;
cl::sycl::buffer<T> *d_sum; cl::sycl::buffer<T> *d_sum;
// SYCL kernel names
typedef sycl_kernels::init<T> init_kernel;
typedef sycl_kernels::copy<T> copy_kernel;
typedef sycl_kernels::mul<T> mul_kernel;
typedef sycl_kernels::add<T> add_kernel;
typedef sycl_kernels::triad<T> triad_kernel;
typedef sycl_kernels::dot<T> dot_kernel;
// NDRange configuration for the dot kernel // NDRange configuration for the dot kernel
size_t dot_num_groups; size_t dot_num_groups;
size_t dot_wgsize; size_t dot_wgsize;