From 1e976ff1502fb86ab618920c666d2afb7f170034 Mon Sep 17 00:00:00 2001
From: James Price <j.price@bristol.ac.uk>
Date: Fri, 18 Nov 2016 00:14:46 +0000
Subject: [PATCH] [SYCL] Fix multiple template specializations

---
 SYCLStream.cpp | 39 +++++++++++++++------------------------
 SYCLStream.h   | 18 ++++++++++++++++++
 main.cpp       |  4 +---
 3 files changed, 34 insertions(+), 27 deletions(-)

diff --git a/SYCLStream.cpp b/SYCLStream.cpp
index c8a908c..6160fc1 100644
--- a/SYCLStream.cpp
+++ b/SYCLStream.cpp
@@ -18,16 +18,6 @@ std::vector<device> devices;
 void getDeviceList(void);
 program * p;
 
-/* Forward declaration of SYCL kernels */
-namespace kernels {
-  class init;
-  class copy;
-  class mul;
-  class add;
-  class triad;
-  class dot;
-}
-
 template <class T>
 SYCLStream<T>::SYCLStream(const unsigned int ARRAY_SIZE, const int device_index)
 {
@@ -61,12 +51,12 @@ SYCLStream<T>::SYCLStream(const unsigned int ARRAY_SIZE, const int device_index)
 
   /* Pre-build the kernels */
   p = new program(queue->get_context());
-  p->build_from_kernel_name<kernels::init>();
-  p->build_from_kernel_name<kernels::copy>();
-  p->build_from_kernel_name<kernels::mul>();
-  p->build_from_kernel_name<kernels::add>();
-  p->build_from_kernel_name<kernels::triad>();
-  p->build_from_kernel_name<kernels::dot>();
+  p->build_from_kernel_name<init_kernel>();
+  p->build_from_kernel_name<copy_kernel>();
+  p->build_from_kernel_name<mul_kernel>();
+  p->build_from_kernel_name<add_kernel>();
+  p->build_from_kernel_name<triad_kernel>();
+  p->build_from_kernel_name<dot_kernel>();
 
   // Create buffers
   d_a = new buffer<T>(array_size);
@@ -94,7 +84,7 @@ 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<kernels::copy>(p->get_kernel<kernels::copy>(),
+    cgh.parallel_for<copy_kernel>(p->get_kernel<copy_kernel>(),
           range<1>{array_size}, [=](item<1> item)
     {
       auto id = item.get();
@@ -112,7 +102,7 @@ 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<kernels::mul>(p->get_kernel<kernels::mul>(),
+    cgh.parallel_for<mul_kernel>(p->get_kernel<mul_kernel>(),
       range<1>{array_size}, [=](item<1> item)
     {
       auto id = item.get();
@@ -130,7 +120,7 @@ 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<kernels::add>(p->get_kernel<kernels::add>(),
+    cgh.parallel_for<add_kernel>(p->get_kernel<add_kernel>(),
       range<1>{array_size}, [=](item<1> item)
     {
       auto id = item.get();
@@ -149,7 +139,7 @@ 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<kernels::triad>(p->get_kernel<kernels::triad>(),
+    cgh.parallel_for<triad_kernel>(p->get_kernel<triad_kernel>(),
       range<1>{array_size}, [=](item<1> item)
     {
       auto id = item.get();
@@ -172,7 +162,8 @@ T SYCLStream<T>::dot()
 
     size_t N = array_size;
 
-    cgh.parallel_for<kernels::dot>(nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item)
+    cgh.parallel_for<dot_kernel>(p->get_kernel<dot_kernel>(),
+      nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item)
     {
       size_t i = item.get_global(0);
       size_t li = item.get_local(0);
@@ -210,8 +201,8 @@ void SYCLStream<T>::init_arrays(T initA, T initB, T initC)
     auto ka = d_a->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);
-    cgh.parallel_for<kernels::init>(p->get_kernel<kernels::init>(),
-          range<1>{array_size}, [=](item<1> item)
+    cgh.parallel_for<init_kernel>(p->get_kernel<init_kernel>(),
+      range<1>{array_size}, [=](item<1> item)
     {
       auto id = item.get();
       ka[id[0]] = initA;
@@ -311,5 +302,5 @@ std::string getDeviceDriver(const int device)
 
 
 // TODO: Fix kernel names to allow multiple template specializations
-//template class SYCLStream<float>;
+template class SYCLStream<float>;
 template class SYCLStream<double>;
diff --git a/SYCLStream.h b/SYCLStream.h
index f3c8d25..ab62ecd 100644
--- a/SYCLStream.h
+++ b/SYCLStream.h
@@ -15,6 +15,16 @@
 
 #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>
 class SYCLStream : public Stream<T>
 {
@@ -29,6 +39,14 @@ class SYCLStream : public Stream<T>
     cl::sycl::buffer<T> *d_c;
     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
     size_t dot_num_groups;
     size_t dot_wgsize;
diff --git a/main.cpp b/main.cpp
index 16e3241..2d80814 100644
--- a/main.cpp
+++ b/main.cpp
@@ -61,13 +61,11 @@ int main(int argc, char *argv[])
 
   parseArguments(argc, argv);
 
-  // TODO: Fix SYCL to allow multiple template specializations
-#ifndef SYCL
+  // TODO: Fix Kokkos to allow multiple template specializations
 #ifndef KOKKOS
   if (use_float)
     run<float>();
   else
-#endif
 #endif
     run<double>();