From 4e5ce57bb1c5f63e8ac18cf0012782a4fbeb9d8f Mon Sep 17 00:00:00 2001
From: Mohammad Fawaz <mohammad.fawaz@intel.com>
Date: Thu, 25 Jun 2020 20:07:41 -0400
Subject: [PATCH 1/4] [SYCL][FPGA] Adding the new header fpga_lsu.hpp that
 contains wrappers around the Clang builtin __builtin_intel_fpga_mem

---
 sycl/doc/extensions/IntelFPGA/FPGALsu.md      | 112 +++++++++++++++++
 .../include/CL/sycl/intel/fpga_extensions.hpp |   1 +
 sycl/include/CL/sycl/intel/fpga_lsu.hpp       | 116 ++++++++++++++++++
 sycl/include/CL/sycl/intel/fpga_utils.hpp     |  34 +++++
 sycl/test/fpga_tests/fpga_lsu.cpp             |  77 ++++++++++++
 5 files changed, 340 insertions(+)
 create mode 100644 sycl/doc/extensions/IntelFPGA/FPGALsu.md
 create mode 100644 sycl/include/CL/sycl/intel/fpga_lsu.hpp
 create mode 100644 sycl/include/CL/sycl/intel/fpga_utils.hpp
 create mode 100644 sycl/test/fpga_tests/fpga_lsu.cpp

diff --git a/sycl/doc/extensions/IntelFPGA/FPGALsu.md b/sycl/doc/extensions/IntelFPGA/FPGALsu.md
new file mode 100644
index 0000000000000..ed043caaa7b06
--- /dev/null
+++ b/sycl/doc/extensions/IntelFPGA/FPGALsu.md
@@ -0,0 +1,112 @@
+
+# FPGA lsu
+
+The Intel FPGA `lsu` class is implemented in `CL/sycl/intel/fpga_lsu.hpp` which
+is included in `CL/sycl/intel/fpga_extensions.hpp`.
+
+The class `cl::sycl::intel::lsu` allows users to explicitly request that the
+implementation of a global memory access is configured in a certain way. The
+class has two member functions, `load()` and `store()` which allow loading from
+and storing to a `global_ptr`, respectively, and is templated on the following
+4 optional paremeters:
+
+1.  **`cl::sycl::intel::burst_coalesce<N>`, where `N` is 0 or 1**: request, to
+the extent possible, that a dynamic burst coalescer be implemented when `load`
+or `store` are called.
+2. **`cl::sycl::intel::cache<N>`, where `N` is an integer greater or equal to
+0**: request, to the extent possible, that a read-only cache of the specified
+size be implemented when when `load` is called. It is not allowed to use that
+parameter for `store`.
+3. **`cl::sycl::intel::dont_statically_coalesce<N>`, where `N` is 0 or 1**:
+request, to the extent possible, that `load` or `store` accesses, should not be
+statically coalesced with other memory accesses at compile time.
+4. **`cl::sycl::intel::prefetch<N>`, where `N` is 0 or 1**: request, to the
+extent possible, that a prefetcher be implemented when `load` is called. It is
+not allowed to use that parameter for `store`.
+
+Currently, not every combination of parameters is allowed due to limitations in
+the backend. The following rules apply:
+1. For `store`, `cl::sycl::intel::cache` and `cl::sycl::intel::prefetch` must
+be `0`.
+2. For `load`, if `cl::sycl::intel::cache` is set to a value greater than `0`,
+then `cl::sycl::intel::burst_coalesce` must be set to `1`.
+3. For `load`, exactly one of `cl::sycl::intel::prefetch` and
+`cl::sycl::intel::burst_coalesce` is allowed to be `1`.
+4. For `load`, exactly one of `cl::sycl::intel::prefetch` and
+`cl::sycl::intel::cache` is allowed to be 1.
+
+## Implementation
+
+The implementation relies on the Clang built-in `__builtin_intel_fpga_mem` when
+parsing the SYCL device code. The built-in uses the LLVM `ptr.annotation`
+intrinsic under the hood to annotate the pointer that is being accessed.
+```c++
+template <class... mem_access_params> class lsu final {
+public:
+  lsu() = delete;
+
+  template <typename T> static T &load(sycl::global_ptr<T> Ptr) {
+  check_load();
+#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+    return *__builtin_intel_fpga_mem((T *)Ptr, _burst_coalesce | _cache |
+                                         _dont_statically_coalesce |
+                                         _prefetch, _cache_size);
+#else
+    return *Ptr;
+#endif
+  }
+
+  template <typename T> static void store(sycl::global_ptr<T> Ptr, T Val) {
+  check_store();
+#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+    *__builtin_intel_fpga_mem((T *)Ptr, _burst_coalesce | _cache |
+                                  _dont_statically_coalesce |
+                                  _prefetch, _cache_size) = Val;
+#else
+    *Ptr = Val;
+#endif
+  }
+  ...
+}
+```
+
+## Usage
+
+```c++
+#include <CL/sycl/intel/fpga_extensions.hpp>
+...
+cl::sycl::buffer<int, 1> output_buffer(output_data, 1);
+cl::sycl::buffer<int, 1> input_buffer(input_data, 1);
+
+Queue.submit([&](cl::sycl::handler &cgh) {
+  auto output_accessor = output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
+  auto input_accessor = input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
+
+  cgh.single_task<class kernel>([=]() {
+    auto input_ptr = input_accessor.get_pointer();
+    auto output_ptr = output_accessor.get_pointer();
+
+    using PrefetchingLSU =
+        cl::sycl::intel::lsu<cl::sycl::intel::prefetch<1>,
+                             cl::sycl::intel::dont_statically_coalesce<1>>;
+
+    using BurstCoalescedLSU =
+        cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
+                             cl::sycl::intel::dont_statically_coalesce<1>>;
+
+    using CachingLSU =
+        cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
+                             cl::sycl::intel::cache<1024>,
+                             cl::sycl::intel::dont_statically_coalesce<0>>;
+
+    using PipelinedLSU = cl::sycl::intel::lsu<>;
+
+    int X = PrefetchingLSU::load(input_ptr); // int X = input_ptr[0]
+    int Y = CachingLSU::load(input_ptr + 1); // int Y = input_ptr[1]
+
+    BurstCoalescedLSU::store(output_ptr, X); // output_ptr[0] = X
+    PipelinedLSU::store(output_ptr + 1, Y);  // output_ptr[1] = Y
+  });
+});
+...
+```
diff --git a/sycl/include/CL/sycl/intel/fpga_extensions.hpp b/sycl/include/CL/sycl/intel/fpga_extensions.hpp
index a9fca1e6139d2..7140421fe5189 100644
--- a/sycl/include/CL/sycl/intel/fpga_extensions.hpp
+++ b/sycl/include/CL/sycl/intel/fpga_extensions.hpp
@@ -8,5 +8,6 @@
 
 #pragma once
 #include <CL/sycl/intel/fpga_device_selector.hpp>
+#include <CL/sycl/intel/fpga_lsu.hpp>
 #include <CL/sycl/intel/fpga_reg.hpp>
 #include <CL/sycl/intel/pipes.hpp>
diff --git a/sycl/include/CL/sycl/intel/fpga_lsu.hpp b/sycl/include/CL/sycl/intel/fpga_lsu.hpp
new file mode 100644
index 0000000000000..527d9e7b790b3
--- /dev/null
+++ b/sycl/include/CL/sycl/intel/fpga_lsu.hpp
@@ -0,0 +1,116 @@
+//==-------------- fpga_lsu.hpp --- SYCL FPGA Reg Extensions ---------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#pragma once
+
+#include "fpga_utils.hpp"
+#include <CL/sycl/detail/defines.hpp>
+#include <CL/sycl/pointers.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace intel {
+constexpr unsigned BURST_COALESCE = 0x1;
+constexpr unsigned CACHE = 0x2;
+constexpr unsigned DONT_STATICALLY_COALESCE = 0x4;
+constexpr unsigned PREFETCH = 0x8;
+
+template <int N> struct burst_coalesce {
+  static constexpr int value = N;
+  static constexpr int default_value = false;
+};
+
+template <int N> struct cache {
+  static constexpr int value = N;
+  static constexpr int default_value = 0;
+};
+
+template <int N> struct prefetch {
+  static constexpr int value = N;
+  static constexpr int default_value = false;
+};
+
+template <int N> struct dont_statically_coalesce {
+  static constexpr int value = N;
+  static constexpr int default_value = 0;
+};
+
+template <class... mem_access_params> class lsu final {
+public:
+  lsu() = delete;
+
+  template <typename T> static T &load(sycl::global_ptr<T> Ptr) {
+    check_load();
+#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+    return *__builtin_intel_fpga_mem((T *)Ptr,
+                                     _burst_coalesce | _cache |
+                                         _dont_statically_coalesce | _prefetch,
+                                     _cache_val);
+#else
+    return *Ptr;
+#endif
+  }
+
+  template <typename T> static void store(sycl::global_ptr<T> Ptr, T Val) {
+    check_store();
+#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+    *__builtin_intel_fpga_mem((T *)Ptr,
+                              _burst_coalesce | _cache |
+                                  _dont_statically_coalesce | _prefetch,
+                              _cache_val) = Val;
+#else
+    *Ptr = Val;
+#endif
+  }
+
+private:
+  static constexpr int _burst_coalesce_val =
+      GetValue<burst_coalesce, mem_access_params...>::value;
+  static constexpr unsigned _burst_coalesce =
+      _burst_coalesce_val == 1 ? BURST_COALESCE : 0;
+
+  static constexpr int _cache_val =
+      GetValue<cache, mem_access_params...>::value;
+  static constexpr unsigned _cache = (_cache_val > 0) ? CACHE : 0;
+
+  static constexpr unsigned _dont_statically_coalesce_val =
+      GetValue<dont_statically_coalesce, mem_access_params...>::value;
+  static constexpr unsigned _dont_statically_coalesce =
+      _dont_statically_coalesce_val == 1 ? DONT_STATICALLY_COALESCE : 0;
+
+  static constexpr unsigned _prefetch_val =
+      GetValue<prefetch, mem_access_params...>::value;
+  static constexpr unsigned _prefetch = _prefetch_val ? PREFETCH : 0;
+
+  static_assert(_burst_coalesce_val == 0 || _burst_coalesce_val == 1,
+                "burst_coalesce parameter must be 0 or 1");
+  static_assert(_cache_val >= 0, "cache size parameter must be non-negative");
+  static_assert(_dont_statically_coalesce_val == 0 ||
+                    _dont_statically_coalesce_val == 1,
+                "dont_statically_coalesce parameter must be 0 or 1");
+  static_assert(_prefetch_val == 0 || _prefetch_val == 1,
+                "prefetch parameter must be 0 or 1");
+
+  static void check_load() {
+    static_assert(_cache == 0 || _burst_coalesce == BURST_COALESCE,
+                  "unable to implement a cache without a burst coalescer");
+    static_assert(_prefetch == 0 || _burst_coalesce == 0,
+                  "unable to implement a prefetcher and a burst coalescer "
+                  "simulataneously");
+    static_assert(
+        _prefetch == 0 || _cache == 0,
+        "unable to implement a prefetcher and a cache simulataneously");
+  }
+  static void check_store() {
+    static_assert(_cache == 0, "unable to implement a store LSU with a cache.");
+    static_assert(_prefetch == 0,
+                  "unable to implement a store LSU with a prefetcher.");
+  }
+};
+} // namespace intel
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/intel/fpga_utils.hpp b/sycl/include/CL/sycl/intel/fpga_utils.hpp
new file mode 100644
index 0000000000000..79ac4cc7b720d
--- /dev/null
+++ b/sycl/include/CL/sycl/intel/fpga_utils.hpp
@@ -0,0 +1,34 @@
+//==------------- fpga_utils.hpp --- SYCL FPGA Reg Extensions --------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/detail/defines.hpp>
+#include <CL/sycl/stl.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace intel {
+
+template <template <int> class Type, class T>
+struct MatchType : std::is_same<Type<T::value>, T> {};
+
+template <template <int> class Type, class... T> struct GetValue {
+  enum { value = Type<0>::default_value };
+};
+
+template <template <int> class Type, class T1, class... T>
+struct GetValue<Type, T1, T...> {
+  enum {
+    value = std::conditional<MatchType<Type, T1>::value, T1,
+                             GetValue<Type, T...>>::type::value
+  };
+};
+} // namespace intel
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/test/fpga_tests/fpga_lsu.cpp b/sycl/test/fpga_tests/fpga_lsu.cpp
new file mode 100644
index 0000000000000..3d34048d46918
--- /dev/null
+++ b/sycl/test/fpga_tests/fpga_lsu.cpp
@@ -0,0 +1,77 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUNx: %ACC_RUN_PLACEHOLDER %t.out
+//==--------------- fpga_lsu.cpp - SYCL FPGA pipes test --------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#include <CL/sycl.hpp>
+#include <CL/sycl/intel/fpga_extensions.hpp>
+
+// TODO: run is disabled, since no support added in FPGA backend yet. Check
+// implementation correctness from CXX and SYCL languages perspective.
+
+int test_lsu(cl::sycl::queue Queue) {
+  int output_data[2];
+  for (size_t i = 0; i < 2; i++) {
+    output_data[i] = -1;
+  }
+
+  int input_data[2];
+  for (size_t i = 0; i < 2; i++) {
+    input_data[i] = i + 1;
+  }
+
+  cl::sycl::buffer<int, 1> output_buffer(output_data, 1);
+  cl::sycl::buffer<int, 1> input_buffer(input_data, 1);
+
+  Queue.submit([&](cl::sycl::handler &cgh) {
+    auto output_accessor = output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
+    auto input_accessor = input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
+
+    cgh.single_task<class kernel>([=]() {
+      auto input_ptr = input_accessor.get_pointer();
+      auto output_ptr = output_accessor.get_pointer();
+
+      using PrefetchingLSU =
+          cl::sycl::intel::lsu<cl::sycl::intel::prefetch<1>,
+                               cl::sycl::intel::dont_statically_coalesce<1>>;
+
+      using BurstCoalescedLSU =
+          cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
+                               cl::sycl::intel::dont_statically_coalesce<1>>;
+
+      using CachingLSU =
+          cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
+                               cl::sycl::intel::cache<1024>,
+                               cl::sycl::intel::dont_statically_coalesce<0>>;
+
+      using PipelinedLSU = cl::sycl::intel::lsu<>;
+
+      int X = PrefetchingLSU::load(input_ptr); // int X = input_ptr[0]
+      int Y = CachingLSU::load(input_ptr + 1); // int Y = input_ptr[1]
+
+      BurstCoalescedLSU::store(output_ptr, X); // output_ptr[0] = X
+      PipelinedLSU::store(output_ptr + 1, Y);  // output_ptr[1] = Y
+    });
+  });
+
+  auto host_buffer = output_buffer.get_access<cl::sycl::access::mode::read>();
+  for (int i = 0; i < 2; i++) {
+    if (host_buffer[i] != input_data[i]) {
+      std::cout << "Unexpected read from output_buffer: " << host_buffer[i]
+                << ", v.s. expected " << input_data[i] << std::endl;
+
+      return -1;
+    }
+  }
+  return 0;
+}
+
+int main() {
+  cl::sycl::queue Queue{cl::sycl::intel::fpga_emulator_selector{}};
+
+  return test_lsu(Queue);
+}

From f56ab264ab89a462d191459c2e748802b38889da Mon Sep 17 00:00:00 2001
From: Mohammad Fawaz <mohammad.fawaz@intel.com>
Date: Mon, 6 Jul 2020 09:44:41 -0400
Subject: [PATCH 2/4] Changing int types and simplifying the example

---
 sycl/doc/extensions/IntelFPGA/FPGALsu.md  |  2 +-
 sycl/include/CL/sycl/intel/fpga_lsu.hpp   | 48 ++++++++++----------
 sycl/include/CL/sycl/intel/fpga_utils.hpp |  6 +--
 sycl/test/fpga_tests/fpga_lsu.cpp         | 55 ++++++++++++-----------
 4 files changed, 56 insertions(+), 55 deletions(-)

diff --git a/sycl/doc/extensions/IntelFPGA/FPGALsu.md b/sycl/doc/extensions/IntelFPGA/FPGALsu.md
index ed043caaa7b06..432a97e72a443 100644
--- a/sycl/doc/extensions/IntelFPGA/FPGALsu.md
+++ b/sycl/doc/extensions/IntelFPGA/FPGALsu.md
@@ -82,7 +82,7 @@ Queue.submit([&](cl::sycl::handler &cgh) {
   auto output_accessor = output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
   auto input_accessor = input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
 
-  cgh.single_task<class kernel>([=]() {
+  cgh.single_task<class kernel>([=] {
     auto input_ptr = input_accessor.get_pointer();
     auto output_ptr = output_accessor.get_pointer();
 
diff --git a/sycl/include/CL/sycl/intel/fpga_lsu.hpp b/sycl/include/CL/sycl/intel/fpga_lsu.hpp
index 527d9e7b790b3..4908572964a0d 100644
--- a/sycl/include/CL/sycl/intel/fpga_lsu.hpp
+++ b/sycl/include/CL/sycl/intel/fpga_lsu.hpp
@@ -14,29 +14,29 @@
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace intel {
-constexpr unsigned BURST_COALESCE = 0x1;
-constexpr unsigned CACHE = 0x2;
-constexpr unsigned DONT_STATICALLY_COALESCE = 0x4;
-constexpr unsigned PREFETCH = 0x8;
+constexpr uint8_t BURST_COALESCE = 0x1;
+constexpr uint8_t CACHE = 0x2;
+constexpr uint8_t DONT_STATICALLY_COALESCE = 0x4;
+constexpr uint8_t PREFETCH = 0x8;
 
-template <int N> struct burst_coalesce {
-  static constexpr int value = N;
-  static constexpr int default_value = false;
+template <int32_t N> struct burst_coalesce {
+  static constexpr int32_t value = N;
+  static constexpr int32_t default_value = false;
 };
 
-template <int N> struct cache {
-  static constexpr int value = N;
-  static constexpr int default_value = 0;
+template <int32_t N> struct cache {
+  static constexpr int32_t value = N;
+  static constexpr int32_t default_value = 0;
 };
 
-template <int N> struct prefetch {
-  static constexpr int value = N;
-  static constexpr int default_value = false;
+template <int32_t N> struct prefetch {
+  static constexpr int32_t value = N;
+  static constexpr int32_t default_value = false;
 };
 
-template <int N> struct dont_statically_coalesce {
-  static constexpr int value = N;
-  static constexpr int default_value = 0;
+template <int32_t N> struct dont_statically_coalesce {
+  static constexpr int32_t value = N;
+  static constexpr int32_t default_value = 0;
 };
 
 template <class... mem_access_params> class lsu final {
@@ -68,23 +68,23 @@ template <class... mem_access_params> class lsu final {
   }
 
 private:
-  static constexpr int _burst_coalesce_val =
+  static constexpr int32_t _burst_coalesce_val =
       GetValue<burst_coalesce, mem_access_params...>::value;
-  static constexpr unsigned _burst_coalesce =
+  static constexpr uint8_t _burst_coalesce =
       _burst_coalesce_val == 1 ? BURST_COALESCE : 0;
 
-  static constexpr int _cache_val =
+  static constexpr int32_t _cache_val =
       GetValue<cache, mem_access_params...>::value;
-  static constexpr unsigned _cache = (_cache_val > 0) ? CACHE : 0;
+  static constexpr uint8_t _cache = (_cache_val > 0) ? CACHE : 0;
 
-  static constexpr unsigned _dont_statically_coalesce_val =
+  static constexpr int32_t _dont_statically_coalesce_val =
       GetValue<dont_statically_coalesce, mem_access_params...>::value;
-  static constexpr unsigned _dont_statically_coalesce =
+  static constexpr uint8_t _dont_statically_coalesce =
       _dont_statically_coalesce_val == 1 ? DONT_STATICALLY_COALESCE : 0;
 
-  static constexpr unsigned _prefetch_val =
+  static constexpr int32_t _prefetch_val =
       GetValue<prefetch, mem_access_params...>::value;
-  static constexpr unsigned _prefetch = _prefetch_val ? PREFETCH : 0;
+  static constexpr uint8_t _prefetch = _prefetch_val ? PREFETCH : 0;
 
   static_assert(_burst_coalesce_val == 0 || _burst_coalesce_val == 1,
                 "burst_coalesce parameter must be 0 or 1");
diff --git a/sycl/include/CL/sycl/intel/fpga_utils.hpp b/sycl/include/CL/sycl/intel/fpga_utils.hpp
index 79ac4cc7b720d..43729970d2c6c 100644
--- a/sycl/include/CL/sycl/intel/fpga_utils.hpp
+++ b/sycl/include/CL/sycl/intel/fpga_utils.hpp
@@ -15,14 +15,14 @@ __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace intel {
 
-template <template <int> class Type, class T>
+template <template <int32_t> class Type, class T>
 struct MatchType : std::is_same<Type<T::value>, T> {};
 
-template <template <int> class Type, class... T> struct GetValue {
+template <template <int32_t> class Type, class... T> struct GetValue {
   enum { value = Type<0>::default_value };
 };
 
-template <template <int> class Type, class T1, class... T>
+template <template <int32_t> class Type, class T1, class... T>
 struct GetValue<Type, T1, T...> {
   enum {
     value = std::conditional<MatchType<Type, T1>::value, T1,
diff --git a/sycl/test/fpga_tests/fpga_lsu.cpp b/sycl/test/fpga_tests/fpga_lsu.cpp
index 3d34048d46918..8a75bba0b6db0 100644
--- a/sycl/test/fpga_tests/fpga_lsu.cpp
+++ b/sycl/test/fpga_tests/fpga_lsu.cpp
@@ -24,44 +24,45 @@ int test_lsu(cl::sycl::queue Queue) {
     input_data[i] = i + 1;
   }
 
-  cl::sycl::buffer<int, 1> output_buffer(output_data, 1);
-  cl::sycl::buffer<int, 1> input_buffer(input_data, 1);
+  {
+    cl::sycl::buffer<int, 1> output_buffer(output_data, 1);
+    cl::sycl::buffer<int, 1> input_buffer(input_data, 1);
 
-  Queue.submit([&](cl::sycl::handler &cgh) {
-    auto output_accessor = output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
-    auto input_accessor = input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
+    Queue.submit([&](cl::sycl::handler &cgh) {
+      auto output_accessor = output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
+      auto input_accessor = input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
 
-    cgh.single_task<class kernel>([=]() {
-      auto input_ptr = input_accessor.get_pointer();
-      auto output_ptr = output_accessor.get_pointer();
+      cgh.single_task<class kernel>([=] {
+        auto input_ptr = input_accessor.get_pointer();
+        auto output_ptr = output_accessor.get_pointer();
 
-      using PrefetchingLSU =
-          cl::sycl::intel::lsu<cl::sycl::intel::prefetch<1>,
-                               cl::sycl::intel::dont_statically_coalesce<1>>;
+        using PrefetchingLSU =
+            cl::sycl::intel::lsu<cl::sycl::intel::prefetch<1>,
+                                 cl::sycl::intel::dont_statically_coalesce<1>>;
 
-      using BurstCoalescedLSU =
-          cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
-                               cl::sycl::intel::dont_statically_coalesce<1>>;
+        using BurstCoalescedLSU =
+            cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
+                                 cl::sycl::intel::dont_statically_coalesce<1>>;
 
-      using CachingLSU =
-          cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
-                               cl::sycl::intel::cache<1024>,
-                               cl::sycl::intel::dont_statically_coalesce<0>>;
+        using CachingLSU =
+            cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
+                                 cl::sycl::intel::cache<1024>,
+                                 cl::sycl::intel::dont_statically_coalesce<0>>;
 
-      using PipelinedLSU = cl::sycl::intel::lsu<>;
+        using PipelinedLSU = cl::sycl::intel::lsu<>;
 
-      int X = PrefetchingLSU::load(input_ptr); // int X = input_ptr[0]
-      int Y = CachingLSU::load(input_ptr + 1); // int Y = input_ptr[1]
+        int X = PrefetchingLSU::load(input_ptr); // int X = input_ptr[0]
+        int Y = CachingLSU::load(input_ptr + 1); // int Y = input_ptr[1]
 
-      BurstCoalescedLSU::store(output_ptr, X); // output_ptr[0] = X
-      PipelinedLSU::store(output_ptr + 1, Y);  // output_ptr[1] = Y
+        BurstCoalescedLSU::store(output_ptr, X); // output_ptr[0] = X
+        PipelinedLSU::store(output_ptr + 1, Y);  // output_ptr[1] = Y
+      });
     });
-  });
+  }
 
-  auto host_buffer = output_buffer.get_access<cl::sycl::access::mode::read>();
   for (int i = 0; i < 2; i++) {
-    if (host_buffer[i] != input_data[i]) {
-      std::cout << "Unexpected read from output_buffer: " << host_buffer[i]
+    if (output_data[i] != input_data[i]) {
+      std::cout << "Unexpected read from output_data: " << output_data[i]
                 << ", v.s. expected " << input_data[i] << std::endl;
 
       return -1;

From 40f927c91855b3012d6451b23ea84071c8a4f011 Mon Sep 17 00:00:00 2001
From: Mohammad Fawaz <mohammad.fawaz@intel.com>
Date: Wed, 8 Jul 2020 10:12:47 -0400
Subject: [PATCH 3/4] Some changes to the use model: - Three of the parameters
 now accept a boolean only - Some simplifications to fpga_utils -
 dont_statically_coalesce is exposed as statically_coalesce instead.

---
 sycl/doc/extensions/IntelFPGA/FPGALsu.md  | 60 ++++++++++++-----------
 sycl/include/CL/sycl/intel/fpga_lsu.hpp   | 35 ++++++-------
 sycl/include/CL/sycl/intel/fpga_utils.hpp |  9 ++--
 sycl/test/fpga_tests/fpga_lsu.cpp         | 12 ++---
 4 files changed, 58 insertions(+), 58 deletions(-)

diff --git a/sycl/doc/extensions/IntelFPGA/FPGALsu.md b/sycl/doc/extensions/IntelFPGA/FPGALsu.md
index 432a97e72a443..761cf5869076a 100644
--- a/sycl/doc/extensions/IntelFPGA/FPGALsu.md
+++ b/sycl/doc/extensions/IntelFPGA/FPGALsu.md
@@ -10,30 +10,32 @@ class has two member functions, `load()` and `store()` which allow loading from
 and storing to a `global_ptr`, respectively, and is templated on the following
 4 optional paremeters:
 
-1.  **`cl::sycl::intel::burst_coalesce<N>`, where `N` is 0 or 1**: request, to
-the extent possible, that a dynamic burst coalescer be implemented when `load`
-or `store` are called.
+1.  **`cl::sycl::intel::burst_coalesce<B>`, where `B` is a boolean**: request,
+to the extent possible, that a dynamic burst coalescer be implemented when
+`load` or `store` are called. The default value of this parameter is `false`.
 2. **`cl::sycl::intel::cache<N>`, where `N` is an integer greater or equal to
 0**: request, to the extent possible, that a read-only cache of the specified
 size be implemented when when `load` is called. It is not allowed to use that
-parameter for `store`.
-3. **`cl::sycl::intel::dont_statically_coalesce<N>`, where `N` is 0 or 1**:
-request, to the extent possible, that `load` or `store` accesses, should not be
-statically coalesced with other memory accesses at compile time.
-4. **`cl::sycl::intel::prefetch<N>`, where `N` is 0 or 1**: request, to the
+parameter for `store`. The default value of this parameter is `0`.
+3. **`cl::sycl::intel::statically_coalesce<N>`, where `B` is a boolean**:
+request, to the extent possible, that `load` or `store` accesses, is allowed to
+be statically coalesced with other memory accesses at compile time. The default
+value of this parameter is `true`.
+4. **`cl::sycl::intel::prefetch<B>`, where `N` is a boolean**: request, to the
 extent possible, that a prefetcher be implemented when `load` is called. It is
-not allowed to use that parameter for `store`.
+not allowed to use that parameter for `store`. The default value of this
+parameter is `false`.
 
 Currently, not every combination of parameters is allowed due to limitations in
 the backend. The following rules apply:
-1. For `store`, `cl::sycl::intel::cache` and `cl::sycl::intel::prefetch` must
-be `0`.
+1. For `store`, `cl::sycl::intel::cache` must be `0` and
+`cl::sycl::intel::prefetch` must be `false`.
 2. For `load`, if `cl::sycl::intel::cache` is set to a value greater than `0`,
-then `cl::sycl::intel::burst_coalesce` must be set to `1`.
+then `cl::sycl::intel::burst_coalesce` must be set to `true`.
 3. For `load`, exactly one of `cl::sycl::intel::prefetch` and
-`cl::sycl::intel::burst_coalesce` is allowed to be `1`.
+`cl::sycl::intel::burst_coalesce` is allowed to be `true`.
 4. For `load`, exactly one of `cl::sycl::intel::prefetch` and
-`cl::sycl::intel::cache` is allowed to be 1.
+`cl::sycl::intel::cache` is allowed to be `true`.
 
 ## Implementation
 
@@ -46,22 +48,24 @@ public:
   lsu() = delete;
 
   template <typename T> static T &load(sycl::global_ptr<T> Ptr) {
-  check_load();
+    check_load();
 #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
-    return *__builtin_intel_fpga_mem((T *)Ptr, _burst_coalesce | _cache |
-                                         _dont_statically_coalesce |
-                                         _prefetch, _cache_size);
+    return *__builtin_intel_fpga_mem((T *)Ptr,
+                                     _burst_coalesce | _cache |
+                                         _dont_statically_coalesce | _prefetch,
+                                     _cache_val);
 #else
     return *Ptr;
 #endif
   }
 
   template <typename T> static void store(sycl::global_ptr<T> Ptr, T Val) {
-  check_store();
+    check_store();
 #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
-    *__builtin_intel_fpga_mem((T *)Ptr, _burst_coalesce | _cache |
-                                  _dont_statically_coalesce |
-                                  _prefetch, _cache_size) = Val;
+    *__builtin_intel_fpga_mem((T *)Ptr,
+                              _burst_coalesce | _cache |
+                                  _dont_statically_coalesce | _prefetch,
+                              _cache_val) = Val;
 #else
     *Ptr = Val;
 #endif
@@ -87,17 +91,17 @@ Queue.submit([&](cl::sycl::handler &cgh) {
     auto output_ptr = output_accessor.get_pointer();
 
     using PrefetchingLSU =
-        cl::sycl::intel::lsu<cl::sycl::intel::prefetch<1>,
-                             cl::sycl::intel::dont_statically_coalesce<1>>;
+        cl::sycl::intel::lsu<cl::sycl::intel::prefetch<true>,
+                             cl::sycl::intel::statically_coalesce<false>>;
 
     using BurstCoalescedLSU =
-        cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
-                             cl::sycl::intel::dont_statically_coalesce<1>>;
+        cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<false>,
+                             cl::sycl::intel::statically_coalesce<false>>;
 
     using CachingLSU =
-        cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
+        cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<true>,
                              cl::sycl::intel::cache<1024>,
-                             cl::sycl::intel::dont_statically_coalesce<0>>;
+                             cl::sycl::intel::statically_coalesce<true>>;
 
     using PipelinedLSU = cl::sycl::intel::lsu<>;
 
diff --git a/sycl/include/CL/sycl/intel/fpga_lsu.hpp b/sycl/include/CL/sycl/intel/fpga_lsu.hpp
index 4908572964a0d..5f8d37f802e76 100644
--- a/sycl/include/CL/sycl/intel/fpga_lsu.hpp
+++ b/sycl/include/CL/sycl/intel/fpga_lsu.hpp
@@ -16,12 +16,12 @@ namespace sycl {
 namespace intel {
 constexpr uint8_t BURST_COALESCE = 0x1;
 constexpr uint8_t CACHE = 0x2;
-constexpr uint8_t DONT_STATICALLY_COALESCE = 0x4;
+constexpr uint8_t STATICALLY_COALESCE = 0x4;
 constexpr uint8_t PREFETCH = 0x8;
 
-template <int32_t N> struct burst_coalesce {
+template <int32_t N> struct burst_coalesce_impl {
   static constexpr int32_t value = N;
-  static constexpr int32_t default_value = false;
+  static constexpr int32_t default_value = 0;
 };
 
 template <int32_t N> struct cache {
@@ -29,16 +29,20 @@ template <int32_t N> struct cache {
   static constexpr int32_t default_value = 0;
 };
 
-template <int32_t N> struct prefetch {
+template <int32_t N> struct prefetch_impl {
   static constexpr int32_t value = N;
-  static constexpr int32_t default_value = false;
+  static constexpr int32_t default_value = 0;
 };
 
-template <int32_t N> struct dont_statically_coalesce {
+template <int32_t N> struct statically_coalesce_impl {
   static constexpr int32_t value = N;
-  static constexpr int32_t default_value = 0;
+  static constexpr int32_t default_value = 1;
 };
 
+template <bool B> using burst_coalesce = burst_coalesce_impl<B>;
+template <bool B> using prefetch = prefetch_impl<B>;
+template <bool B> using statically_coalesce = statically_coalesce_impl<B>;
+
 template <class... mem_access_params> class lsu final {
 public:
   lsu() = delete;
@@ -69,7 +73,7 @@ template <class... mem_access_params> class lsu final {
 
 private:
   static constexpr int32_t _burst_coalesce_val =
-      GetValue<burst_coalesce, mem_access_params...>::value;
+      GetValue<burst_coalesce_impl, mem_access_params...>::value;
   static constexpr uint8_t _burst_coalesce =
       _burst_coalesce_val == 1 ? BURST_COALESCE : 0;
 
@@ -77,23 +81,16 @@ template <class... mem_access_params> class lsu final {
       GetValue<cache, mem_access_params...>::value;
   static constexpr uint8_t _cache = (_cache_val > 0) ? CACHE : 0;
 
-  static constexpr int32_t _dont_statically_coalesce_val =
-      GetValue<dont_statically_coalesce, mem_access_params...>::value;
+  static constexpr int32_t _statically_coalesce_val =
+      GetValue<statically_coalesce_impl, mem_access_params...>::value;
   static constexpr uint8_t _dont_statically_coalesce =
-      _dont_statically_coalesce_val == 1 ? DONT_STATICALLY_COALESCE : 0;
+      _statically_coalesce_val == 0 ? STATICALLY_COALESCE : 0;
 
   static constexpr int32_t _prefetch_val =
-      GetValue<prefetch, mem_access_params...>::value;
+      GetValue<prefetch_impl, mem_access_params...>::value;
   static constexpr uint8_t _prefetch = _prefetch_val ? PREFETCH : 0;
 
-  static_assert(_burst_coalesce_val == 0 || _burst_coalesce_val == 1,
-                "burst_coalesce parameter must be 0 or 1");
   static_assert(_cache_val >= 0, "cache size parameter must be non-negative");
-  static_assert(_dont_statically_coalesce_val == 0 ||
-                    _dont_statically_coalesce_val == 1,
-                "dont_statically_coalesce parameter must be 0 or 1");
-  static_assert(_prefetch_val == 0 || _prefetch_val == 1,
-                "prefetch parameter must be 0 or 1");
 
   static void check_load() {
     static_assert(_cache == 0 || _burst_coalesce == BURST_COALESCE,
diff --git a/sycl/include/CL/sycl/intel/fpga_utils.hpp b/sycl/include/CL/sycl/intel/fpga_utils.hpp
index 43729970d2c6c..be9bf1a6fc5af 100644
--- a/sycl/include/CL/sycl/intel/fpga_utils.hpp
+++ b/sycl/include/CL/sycl/intel/fpga_utils.hpp
@@ -19,15 +19,14 @@ template <template <int32_t> class Type, class T>
 struct MatchType : std::is_same<Type<T::value>, T> {};
 
 template <template <int32_t> class Type, class... T> struct GetValue {
-  enum { value = Type<0>::default_value };
+  static constexpr auto value = Type<0>::default_value;
 };
 
 template <template <int32_t> class Type, class T1, class... T>
 struct GetValue<Type, T1, T...> {
-  enum {
-    value = std::conditional<MatchType<Type, T1>::value, T1,
-                             GetValue<Type, T...>>::type::value
-  };
+  static constexpr auto value =
+      std::conditional<MatchType<Type, T1>::value, T1,
+                       GetValue<Type, T...>>::type::value;
 };
 } // namespace intel
 } // namespace sycl
diff --git a/sycl/test/fpga_tests/fpga_lsu.cpp b/sycl/test/fpga_tests/fpga_lsu.cpp
index 8a75bba0b6db0..8596b614dfb5c 100644
--- a/sycl/test/fpga_tests/fpga_lsu.cpp
+++ b/sycl/test/fpga_tests/fpga_lsu.cpp
@@ -37,17 +37,17 @@ int test_lsu(cl::sycl::queue Queue) {
         auto output_ptr = output_accessor.get_pointer();
 
         using PrefetchingLSU =
-            cl::sycl::intel::lsu<cl::sycl::intel::prefetch<1>,
-                                 cl::sycl::intel::dont_statically_coalesce<1>>;
+            cl::sycl::intel::lsu<cl::sycl::intel::prefetch<true>,
+                                 cl::sycl::intel::statically_coalesce<false>>;
 
         using BurstCoalescedLSU =
-            cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
-                                 cl::sycl::intel::dont_statically_coalesce<1>>;
+            cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<true>,
+                                 cl::sycl::intel::statically_coalesce<false>>;
 
         using CachingLSU =
-            cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<1>,
+            cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<true>,
                                  cl::sycl::intel::cache<1024>,
-                                 cl::sycl::intel::dont_statically_coalesce<0>>;
+                                 cl::sycl::intel::statically_coalesce<false>>;
 
         using PipelinedLSU = cl::sycl::intel::lsu<>;
 

From 4c00a2a2e3ae5f2de7c6855a28dfbeb4432087aa Mon Sep 17 00:00:00 2001
From: Mohammad Fawaz <mohammad.fawaz@intel.com>
Date: Wed, 22 Jul 2020 09:19:44 -0400
Subject: [PATCH 4/4] Removing the spec from this PR

Will create a separate PR for the spec alone.
---
 sycl/doc/extensions/IntelFPGA/FPGALsu.md | 116 -----------------------
 sycl/test/fpga_tests/fpga_lsu.cpp        |   8 +-
 2 files changed, 5 insertions(+), 119 deletions(-)
 delete mode 100644 sycl/doc/extensions/IntelFPGA/FPGALsu.md

diff --git a/sycl/doc/extensions/IntelFPGA/FPGALsu.md b/sycl/doc/extensions/IntelFPGA/FPGALsu.md
deleted file mode 100644
index 761cf5869076a..0000000000000
--- a/sycl/doc/extensions/IntelFPGA/FPGALsu.md
+++ /dev/null
@@ -1,116 +0,0 @@
-
-# FPGA lsu
-
-The Intel FPGA `lsu` class is implemented in `CL/sycl/intel/fpga_lsu.hpp` which
-is included in `CL/sycl/intel/fpga_extensions.hpp`.
-
-The class `cl::sycl::intel::lsu` allows users to explicitly request that the
-implementation of a global memory access is configured in a certain way. The
-class has two member functions, `load()` and `store()` which allow loading from
-and storing to a `global_ptr`, respectively, and is templated on the following
-4 optional paremeters:
-
-1.  **`cl::sycl::intel::burst_coalesce<B>`, where `B` is a boolean**: request,
-to the extent possible, that a dynamic burst coalescer be implemented when
-`load` or `store` are called. The default value of this parameter is `false`.
-2. **`cl::sycl::intel::cache<N>`, where `N` is an integer greater or equal to
-0**: request, to the extent possible, that a read-only cache of the specified
-size be implemented when when `load` is called. It is not allowed to use that
-parameter for `store`. The default value of this parameter is `0`.
-3. **`cl::sycl::intel::statically_coalesce<N>`, where `B` is a boolean**:
-request, to the extent possible, that `load` or `store` accesses, is allowed to
-be statically coalesced with other memory accesses at compile time. The default
-value of this parameter is `true`.
-4. **`cl::sycl::intel::prefetch<B>`, where `N` is a boolean**: request, to the
-extent possible, that a prefetcher be implemented when `load` is called. It is
-not allowed to use that parameter for `store`. The default value of this
-parameter is `false`.
-
-Currently, not every combination of parameters is allowed due to limitations in
-the backend. The following rules apply:
-1. For `store`, `cl::sycl::intel::cache` must be `0` and
-`cl::sycl::intel::prefetch` must be `false`.
-2. For `load`, if `cl::sycl::intel::cache` is set to a value greater than `0`,
-then `cl::sycl::intel::burst_coalesce` must be set to `true`.
-3. For `load`, exactly one of `cl::sycl::intel::prefetch` and
-`cl::sycl::intel::burst_coalesce` is allowed to be `true`.
-4. For `load`, exactly one of `cl::sycl::intel::prefetch` and
-`cl::sycl::intel::cache` is allowed to be `true`.
-
-## Implementation
-
-The implementation relies on the Clang built-in `__builtin_intel_fpga_mem` when
-parsing the SYCL device code. The built-in uses the LLVM `ptr.annotation`
-intrinsic under the hood to annotate the pointer that is being accessed.
-```c++
-template <class... mem_access_params> class lsu final {
-public:
-  lsu() = delete;
-
-  template <typename T> static T &load(sycl::global_ptr<T> Ptr) {
-    check_load();
-#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
-    return *__builtin_intel_fpga_mem((T *)Ptr,
-                                     _burst_coalesce | _cache |
-                                         _dont_statically_coalesce | _prefetch,
-                                     _cache_val);
-#else
-    return *Ptr;
-#endif
-  }
-
-  template <typename T> static void store(sycl::global_ptr<T> Ptr, T Val) {
-    check_store();
-#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
-    *__builtin_intel_fpga_mem((T *)Ptr,
-                              _burst_coalesce | _cache |
-                                  _dont_statically_coalesce | _prefetch,
-                              _cache_val) = Val;
-#else
-    *Ptr = Val;
-#endif
-  }
-  ...
-}
-```
-
-## Usage
-
-```c++
-#include <CL/sycl/intel/fpga_extensions.hpp>
-...
-cl::sycl::buffer<int, 1> output_buffer(output_data, 1);
-cl::sycl::buffer<int, 1> input_buffer(input_data, 1);
-
-Queue.submit([&](cl::sycl::handler &cgh) {
-  auto output_accessor = output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
-  auto input_accessor = input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
-
-  cgh.single_task<class kernel>([=] {
-    auto input_ptr = input_accessor.get_pointer();
-    auto output_ptr = output_accessor.get_pointer();
-
-    using PrefetchingLSU =
-        cl::sycl::intel::lsu<cl::sycl::intel::prefetch<true>,
-                             cl::sycl::intel::statically_coalesce<false>>;
-
-    using BurstCoalescedLSU =
-        cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<false>,
-                             cl::sycl::intel::statically_coalesce<false>>;
-
-    using CachingLSU =
-        cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<true>,
-                             cl::sycl::intel::cache<1024>,
-                             cl::sycl::intel::statically_coalesce<true>>;
-
-    using PipelinedLSU = cl::sycl::intel::lsu<>;
-
-    int X = PrefetchingLSU::load(input_ptr); // int X = input_ptr[0]
-    int Y = CachingLSU::load(input_ptr + 1); // int Y = input_ptr[1]
-
-    BurstCoalescedLSU::store(output_ptr, X); // output_ptr[0] = X
-    PipelinedLSU::store(output_ptr + 1, Y);  // output_ptr[1] = Y
-  });
-});
-...
-```
diff --git a/sycl/test/fpga_tests/fpga_lsu.cpp b/sycl/test/fpga_tests/fpga_lsu.cpp
index 8596b614dfb5c..65b35e09dbd69 100644
--- a/sycl/test/fpga_tests/fpga_lsu.cpp
+++ b/sycl/test/fpga_tests/fpga_lsu.cpp
@@ -1,6 +1,6 @@
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUNx: %ACC_RUN_PLACEHOLDER %t.out
-//==--------------- fpga_lsu.cpp - SYCL FPGA pipes test --------------------==//
+//==----------------- fpga_lsu.cpp - SYCL FPGA LSU test --------------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -29,8 +29,10 @@ int test_lsu(cl::sycl::queue Queue) {
     cl::sycl::buffer<int, 1> input_buffer(input_data, 1);
 
     Queue.submit([&](cl::sycl::handler &cgh) {
-      auto output_accessor = output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
-      auto input_accessor = input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
+      auto output_accessor =
+          output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
+      auto input_accessor =
+          input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
 
       cgh.single_task<class kernel>([=] {
         auto input_ptr = input_accessor.get_pointer();