From 87769017338dde5cd33b3a5025651a4856a54714 Mon Sep 17 00:00:00 2001
From: Tom Deakin <tom.deakin@bristol.ac.uk>
Date: Fri, 1 Nov 2019 15:17:05 +0000
Subject: [PATCH] [SYCL] Use the cl::sycl::id parameter in the parallel_for
 kernels

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.
---
 CHANGELOG.md   |  3 +++
 SYCLStream.cpp | 20 ++++++++------------
 2 files changed, 11 insertions(+), 12 deletions(-)

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<T>::copy()
   {
     auto ka = d_a->template get_access<access::mode::read>(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[id] = ka[id];
+      kc[idx] = ka[idx];
     });
   });
   queue->wait();
@@ -110,10 +109,9 @@ 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);
-    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[id] = scalar * kc[id];
+      kb[idx] = scalar * kc[idx];
     });
   });
   queue->wait();
@@ -127,10 +125,9 @@ 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);
-    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[id] = ka[id] + kb[id];
+      kc[idx] = ka[idx] + kb[idx];
     });
   });
   queue->wait();
@@ -145,10 +142,9 @@ 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);
-    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[id] = kb[id] + scalar * kc[id];
+      ka[idx] = kb[idx] + scalar * kc[idx];
     });
   });
   queue->wait();