From bda9525b9525db5a2ddf3a370a4754403323ca69 Mon Sep 17 00:00:00 2001
From: Tom Deakin <tom.deakin@bristol.ac.uk>
Date: Tue, 2 Feb 2021 12:29:00 +0000
Subject: [PATCH] Add SYCL 1.2.1 nstream kernel

---
 SYCLStream.cpp | 17 +++++++++++++++++
 SYCLStream.h   |  3 +++
 2 files changed, 20 insertions(+)

diff --git a/SYCLStream.cpp b/SYCLStream.cpp
index 8ab642f..49ad3ac 100644
--- a/SYCLStream.cpp
+++ b/SYCLStream.cpp
@@ -148,6 +148,23 @@ void SYCLStream<T>::triad()
   queue->wait();
 }
 
+template <class T>
+void SYCLStream<T>::nstream()
+{
+  const T scalar = startScalar;
+  queue->submit([&](handler &cgh)
+  {
+    auto ka = d_a->template get_access<access::mode::read_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<nstream_kernel>(range<1>{array_size}, [=](id<1> idx)
+    {
+      ka[idx] += kb[idx] + scalar * kc[idx];
+    });
+  });
+  queue->wait();
+}
+
 template <class T>
 T SYCLStream<T>::dot()
 {
diff --git a/SYCLStream.h b/SYCLStream.h
index df10946..d3988a7 100644
--- a/SYCLStream.h
+++ b/SYCLStream.h
@@ -22,6 +22,7 @@ namespace sycl_kernels
   template <class T> class mul;
   template <class T> class add;
   template <class T> class triad;
+  template <class T> class nstream;
   template <class T> class dot;
 }
 
@@ -45,6 +46,7 @@ class SYCLStream : public Stream<T>
     typedef sycl_kernels::mul<T> mul_kernel;
     typedef sycl_kernels::add<T> add_kernel;
     typedef sycl_kernels::triad<T> triad_kernel;
+    typedef sycl_kernels::nstream<T> nstream_kernel;
     typedef sycl_kernels::dot<T> dot_kernel;
 
     // NDRange configuration for the dot kernel
@@ -60,6 +62,7 @@ class SYCLStream : public Stream<T>
     virtual void add() override;
     virtual void mul() override;
     virtual void triad() override;
+    virtual void nstream() override;
     virtual T    dot() override;
 
     virtual void init_arrays(T initA, T initB, T initC) override;