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

[SYCL] Use the cl::sycl::id parameter in the parallel_for kernels

Browse Source 
The cl::sycl::item provides extra features for extracing global/local
ids which aren't required by the kernels.
This also means the kernels don't need to extract the id from the item.
This commit is contained in:
Tom Deakin 2019-11-01 15:17:05 +00:00
parent 4bcb777100
commit 8776901733
2 changed files with 11 additions and 12 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
CHANGELOG.md
View File

@ -3,6 +3,9 @@ All notable changes to this project will be documented in this file.
## [Unreleased] ## [Unreleased]
### Changed
- Use cl::sycl::id parameters instead of cl::sycl::item.
### Removed ### Removed
- Pre-building of kernels in SYCL version to ensure compatibility with SYCL 1.2.1. - Pre-building of kernels in SYCL version to ensure compatibility with SYCL 1.2.1.
Pre-building kernels is also not required, and shows no overhead as the first iteration is not timed. Pre-building kernels is also not required, and shows no overhead as the first iteration is not timed.

20
SYCLStream.cpp
View File

@ -93,10 +93,9 @@ 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}, [=](item<1> item) cgh.parallel_for<copy_kernel>(range<1>{array_size}, [=](id<1> idx)
{ {
auto id = item.get_id(0); kc[idx] = ka[idx];
kc[id] = ka[id];
}); });
}); });
queue->wait(); queue->wait();
@ -110,10 +109,9 @@ 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}, [=](item<1> item) cgh.parallel_for<mul_kernel>(range<1>{array_size}, [=](id<1> idx)
{ {
auto id = item.get_id(0); kb[idx] = scalar * kc[idx];
kb[id] = scalar * kc[id];
}); });
}); });
queue->wait(); queue->wait();
@ -127,10 +125,9 @@ 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}, [=](item<1> item) cgh.parallel_for<add_kernel>(range<1>{array_size}, [=](id<1> idx)
{ {
auto id = item.get_id(0); kc[idx] = ka[idx] + kb[idx];
kc[id] = ka[id] + kb[id];
}); });
}); });
queue->wait(); queue->wait();
@ -145,10 +142,9 @@ 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}, [=](item<1> item) cgh.parallel_for<triad_kernel>(range<1>{array_size}, [=](id<1> idx)
{ {
auto id = item.get_id(0); ka[idx] = kb[idx] + scalar * kc[idx];
ka[id] = kb[id] + scalar * kc[id];
}); });
}); });
queue->wait(); queue->wait();