diff --git a/CHANGELOG.md b/CHANGELOG.md index ca64c6a..9378d09 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,9 @@ All notable changes to this project will be documented in this file. ## [Unreleased] +### Changed +- Use cl::sycl::id parameters instead of cl::sycl::item. + ### Removed - 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. diff --git a/SYCLStream.cpp b/SYCLStream.cpp index ebe5f63..8e33588 100644 --- a/SYCLStream.cpp +++ b/SYCLStream.cpp @@ -93,10 +93,9 @@ void SYCLStream::copy() { auto ka = d_a->template get_access(cgh); auto kc = d_c->template get_access(cgh); - cgh.parallel_for(range<1>{array_size}, [=](item<1> item) + cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) { - auto id = item.get_id(0); - kc[id] = ka[id]; + kc[idx] = ka[idx]; }); }); queue->wait(); @@ -110,10 +109,9 @@ void SYCLStream::mul() { auto kb = d_b->template get_access(cgh); auto kc = d_c->template get_access(cgh); - cgh.parallel_for(range<1>{array_size}, [=](item<1> item) + cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) { - auto id = item.get_id(0); - kb[id] = scalar * kc[id]; + kb[idx] = scalar * kc[idx]; }); }); queue->wait(); @@ -127,10 +125,9 @@ void SYCLStream::add() auto ka = d_a->template get_access(cgh); auto kb = d_b->template get_access(cgh); auto kc = d_c->template get_access(cgh); - cgh.parallel_for(range<1>{array_size}, [=](item<1> item) + cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) { - auto id = item.get_id(0); - kc[id] = ka[id] + kb[id]; + kc[idx] = ka[idx] + kb[idx]; }); }); queue->wait(); @@ -145,10 +142,9 @@ void SYCLStream::triad() auto ka = d_a->template get_access(cgh); auto kb = d_b->template get_access(cgh); auto kc = d_c->template get_access(cgh); - cgh.parallel_for(range<1>{array_size}, [=](item<1> item) + cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) { - auto id = item.get_id(0); - ka[id] = kb[id] + scalar * kc[id]; + ka[idx] = kb[idx] + scalar * kc[idx]; }); }); queue->wait();