From 1c44e45e1b1c5e4bcb370a56f163000f44645f16 Mon Sep 17 00:00:00 2001
From: "Niu, Shuo" <shuo.niu@intel.com>
Date: Wed, 1 Dec 2021 13:16:08 -0800
Subject: [PATCH 1/3] Create experimental FPGA latency control headers

---
 .../sycl/ext/intel/experimental/fpga_lsu.hpp  | 179 +++++++++++++++
 .../ext/intel/experimental/fpga_utils.hpp     | 105 +++++++++
 .../sycl/ext/intel/experimental/pipes.hpp     | 213 ++++++++++++++++++
 .../sycl/ext/intel/fpga_extensions.hpp        |   2 +
 4 files changed, 499 insertions(+)
 create mode 100644 sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp
 create mode 100644 sycl/include/sycl/ext/intel/experimental/fpga_utils.hpp
 create mode 100644 sycl/include/sycl/ext/intel/experimental/pipes.hpp

diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp
new file mode 100644
index 0000000000000..9819d6791db31
--- /dev/null
+++ b/sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp
@@ -0,0 +1,179 @@
+//==-------------- fpga_lsu.hpp --- SYCL FPGA LSU 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 ext {
+namespace intel {
+namespace experimental {
+
+constexpr uint8_t BURST_COALESCE = 0x1;
+constexpr uint8_t CACHE = 0x2;
+constexpr uint8_t STATICALLY_COALESCE = 0x4;
+constexpr uint8_t PREFETCH = 0x8;
+
+template <int32_t _N> struct burst_coalesce_impl {
+  static constexpr int32_t value = _N;
+  static constexpr int32_t default_value = 0;
+};
+
+template <int32_t _N> struct cache {
+  static constexpr int32_t value = _N;
+  static constexpr int32_t default_value = 0;
+};
+
+template <int32_t _N> struct prefetch_impl {
+  static constexpr int32_t value = _N;
+  static constexpr int32_t default_value = 0;
+};
+
+template <int32_t _N> struct statically_coalesce_impl {
+  static constexpr int32_t value = _N;
+  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;
+
+  template <class... _Params, typename _T, access::address_space _space>
+  static _T load(sycl::multi_ptr<_T, _space> Ptr) {
+    check_space<_space>();
+    check_load();
+#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+    static constexpr auto _anchor_id =
+        _GetValue<int32_t, latency_anchor_id, _Params...>::value;
+    static constexpr auto _constraint =
+        _GetValue3<int32_t, type, int32_t, latency_constraint,
+                   _Params...>::value;
+
+    static constexpr int32_t _target_anchor = std::get<0>(_constraint);
+    static constexpr type _control_type = std::get<1>(_constraint);
+    static constexpr int32_t _cycle = std::get<2>(_constraint);
+    int32_t _type;
+    if (_control_type == type::none) {
+      _type = 0;
+    } else if (_control_type == type::exact) {
+      _type = 1;
+    } else if (_control_type == type::max) {
+      _type = 2;
+    } else { // _control_type == type::min
+      _type = 3;
+    }
+
+    return *__latency_control_mem_wrapper((_T *)Ptr, _anchor_id, _target_anchor,
+                                          _type, _cycle);
+#else
+    return *Ptr;
+#endif
+  }
+
+  template <class... _Params, typename _T, access::address_space _space>
+  static void store(sycl::multi_ptr<_T, _space> Ptr, _T Val) {
+    check_space<_space>();
+    check_store();
+#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+    static constexpr auto _anchor_id =
+        _GetValue<int32_t, latency_anchor_id, _Params...>::value;
+    static constexpr auto _constraint =
+        _GetValue3<int32_t, type, int32_t, latency_constraint,
+                   _Params...>::value;
+
+    static constexpr int32_t _target_anchor = std::get<0>(_constraint);
+    static constexpr type _control_type = std::get<1>(_constraint);
+    static constexpr int32_t _cycle = std::get<2>(_constraint);
+    int32_t _type;
+    if (_control_type == type::none) {
+      _type = 0;
+    } else if (_control_type == type::exact) {
+      _type = 1;
+    } else if (_control_type == type::max) {
+      _type = 2;
+    } else { // _control_type == type::min
+      _type = 3;
+    }
+
+    *__latency_control_mem_wrapper((_T *)Ptr, _anchor_id, _target_anchor, _type,
+                                   _cycle) = Val;
+#else
+    *Ptr = Val;
+#endif
+  }
+
+private:
+  static constexpr int32_t _burst_coalesce_val =
+      _GetValue<int32_t, burst_coalesce_impl, _mem_access_params...>::value;
+  static constexpr uint8_t _burst_coalesce =
+      _burst_coalesce_val == 1 ? BURST_COALESCE : 0;
+
+  static constexpr int32_t _cache_val =
+      _GetValue<int32_t, cache, _mem_access_params...>::value;
+  static constexpr uint8_t _cache = (_cache_val > 0) ? CACHE : 0;
+
+  static constexpr int32_t _statically_coalesce_val =
+      _GetValue<int32_t, statically_coalesce_impl,
+                _mem_access_params...>::value;
+  static constexpr uint8_t _dont_statically_coalesce =
+      _statically_coalesce_val == 0 ? STATICALLY_COALESCE : 0;
+
+  static constexpr int32_t _prefetch_val =
+      _GetValue<int32_t, prefetch_impl, _mem_access_params...>::value;
+  static constexpr uint8_t _prefetch = _prefetch_val ? PREFETCH : 0;
+
+  static_assert(_cache_val >= 0, "cache size parameter must be non-negative");
+
+  template <access::address_space _space> static void check_space() {
+    static_assert(_space == access::address_space::global_space ||
+                      _space == access::address_space::global_device_space ||
+                      _space == access::address_space::global_host_space,
+                  "lsu controls are only supported for global_ptr, "
+                  "device_ptr, and host_ptr objects");
+  }
+
+  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.");
+  }
+
+#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+  template <typename _T>
+  static _T *__latency_control_mem_wrapper(_T *Ptr, int32_t AnchorID,
+                                           int32_t TargetAnchor, int32_t Type,
+                                           int32_t Cycle) {
+    return __builtin_intel_fpga_mem(
+        Ptr, _burst_coalesce | _cache | _dont_statically_coalesce | _prefetch,
+        _cache_val);
+  }
+#endif
+};
+
+} // namespace experimental
+} // namespace intel
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_utils.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_utils.hpp
new file mode 100644
index 0000000000000..a1c7f5e857967
--- /dev/null
+++ b/sycl/include/sycl/ext/intel/experimental/fpga_utils.hpp
@@ -0,0 +1,105 @@
+//==------------- 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/detail/stl_type_traits.hpp>
+#include <CL/sycl/stl.hpp>
+#include <tuple>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace ext {
+namespace intel {
+namespace experimental {
+
+enum class type {
+  none, // default
+  exact,
+  max,
+  min
+};
+
+template <int32_t _N> struct latency_anchor_id {
+  static constexpr int32_t value = _N;
+  static constexpr int32_t default_value = -1;
+};
+
+template <int32_t _N1, type _N2, int32_t _N3> struct latency_constraint {
+  static constexpr std::tuple<int32_t, type, int32_t> value = {_N1, _N2, _N3};
+  static constexpr std::tuple<int32_t, type, int32_t> default_value = {
+      0, type::none, 0};
+};
+
+using ignoreParam_int_t = int32_t;
+constexpr ignoreParam_int_t IgnoreParamInt{};
+using ignoreParam_enum_t = type;
+constexpr ignoreParam_enum_t IgnoreParamEnum{};
+
+template <class _VType, class _T> struct _ValueExtractorImp {
+  static constexpr auto _First = _T::value;
+  static constexpr auto _Second = IgnoreParamEnum;
+  static constexpr auto _Third = IgnoreParamInt;
+};
+
+template <class _VTypeFirst, class _VTypeSecond, class _VTypeThird, class _T>
+struct _ValueExtractorImp<
+    const std::tuple<_VTypeFirst, _VTypeSecond, _VTypeThird>, _T> {
+  static constexpr auto _First = std::get<0>(_T::value);
+  static constexpr auto _Second = std::get<1>(_T::value);
+  static constexpr auto _Third = std::get<2>(_T::value);
+};
+
+template <class _T>
+struct _ValueExtractor : _ValueExtractorImp<decltype(_T::value), _T> {};
+
+template <class _VTypeFirst, class _VTypeSecond, class _VTypeThird,
+          template <_VTypeFirst, _VTypeSecond, _VTypeThird> class _Type,
+          class _T>
+struct _MatchType
+    : std::is_same<
+          _Type<_ValueExtractor<_T>::_First, _ValueExtractor<_T>::_Second,
+                _ValueExtractor<_T>::_Third>,
+          _T> {};
+
+template <class _VTypeFirst, class _VTypeSecond, class _VTypeThird,
+          template <_VTypeFirst, _VTypeSecond, _VTypeThird> class _Type,
+          class... _T>
+struct _GetValue3 {
+  static constexpr auto value =
+      _Type<_VTypeFirst{}, _VTypeSecond{}, _VTypeThird{}>::default_value;
+};
+
+template <class _VTypeFirst, class _VTypeSecond, class _VTypeThird,
+          template <_VTypeFirst, _VTypeSecond, _VTypeThird> class _Type,
+          class _T1, class... _T>
+struct _GetValue3<_VTypeFirst, _VTypeSecond, _VTypeThird, _Type, _T1, _T...> {
+  static constexpr auto value = std::conditional<
+      _MatchType<_VTypeFirst, _VTypeSecond, _VTypeThird, _Type, _T1>::value,
+      _T1, _GetValue3<_VTypeFirst, _VTypeSecond, _VTypeThird, _Type, _T...>>::
+      type::value;
+};
+
+template <class _VType, template <_VType> class _Type, class... _T>
+struct _GetValue {
+private:
+  template <_VType _V1, ignoreParam_enum_t, ignoreParam_int_t>
+  using _Type2 = _Type<_V1>;
+
+public:
+  static constexpr auto value =
+      _GetValue3<_VType, ignoreParam_enum_t, ignoreParam_int_t, _Type2,
+                 _T...>::value;
+};
+
+} // namespace experimental
+} // namespace intel
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/sycl/ext/intel/experimental/pipes.hpp b/sycl/include/sycl/ext/intel/experimental/pipes.hpp
new file mode 100644
index 0000000000000..1e46890b4958b
--- /dev/null
+++ b/sycl/include/sycl/ext/intel/experimental/pipes.hpp
@@ -0,0 +1,213 @@
+//==---------------- pipes.hpp - SYCL pipes ------------*- C++ -*-----------==//
+//
+// 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/__spirv/spirv_ops.hpp>
+#include <CL/__spirv/spirv_types.hpp>
+#include <CL/sycl/stl.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace ext {
+namespace intel {
+namespace experimental {
+
+template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
+public:
+  // Non-blocking pipes
+  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
+  // friendly LLVM IR.
+  template <class... _Params> static _dataT read(bool &_Success) {
+#ifdef __SYCL_DEVICE_ONLY__
+    static constexpr auto _anchor_id =
+        _GetValue<int32_t, latency_anchor_id, _Params...>::value;
+    static constexpr auto _constraint =
+        _GetValue3<int32_t, type, int32_t, latency_constraint,
+                   _Params...>::value;
+
+    static constexpr int32_t _target_anchor = std::get<0>(_constraint);
+    static constexpr type _control_type = std::get<1>(_constraint);
+    static constexpr int32_t _cycle = std::get<2>(_constraint);
+    int32_t _type;
+    if (_control_type == type::none) {
+      _type = 0;
+    } else if (_control_type == type::exact) {
+      _type = 1;
+    } else if (_control_type == type::max) {
+      _type = 2;
+    } else { // _control_type == type::min
+      _type = 3;
+    }
+
+    __ocl_RPipeTy<_dataT> _RPipe =
+        __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
+    _dataT TempData;
+    _Success = !static_cast<bool>(__latency_control_nb_read_wrapper(
+        _RPipe, &TempData, _anchor_id, _target_anchor, _type, _cycle));
+    return TempData;
+#else
+    (void)_Success;
+    assert(!"Pipes are not supported on a host device!");
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
+  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
+  // friendly LLVM IR.
+  template <class... _Params>
+  static void write(const _dataT &_Data, bool &_Success) {
+#ifdef __SYCL_DEVICE_ONLY__
+    static constexpr auto _anchor_id =
+        _GetValue<int32_t, latency_anchor_id, _Params...>::value;
+    static constexpr auto _constraint =
+        _GetValue3<int32_t, type, int32_t, latency_constraint,
+                   _Params...>::value;
+
+    static constexpr int32_t _target_anchor = std::get<0>(_constraint);
+    static constexpr type _control_type = std::get<1>(_constraint);
+    static constexpr int32_t _cycle = std::get<2>(_constraint);
+    int32_t _type;
+    if (_control_type == type::none) {
+      _type = 0;
+    } else if (_control_type == type::exact) {
+      _type = 1;
+    } else if (_control_type == type::max) {
+      _type = 2;
+    } else { // _control_type == type::min
+      _type = 3;
+    }
+
+    __ocl_WPipeTy<_dataT> _WPipe =
+        __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
+    _Success = !static_cast<bool>(__latency_control_nb_write_wrapper(
+        _WPipe, &_Data, _anchor_id, _target_anchor, _type, _cycle));
+#else
+    (void)_Success;
+    (void)_Data;
+    assert(!"Pipes are not supported on a host device!");
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
+  // Blocking pipes
+  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
+  // friendly LLVM IR.
+  template <class... _Params> static _dataT read() {
+#ifdef __SYCL_DEVICE_ONLY__
+    static constexpr auto _anchor_id =
+        _GetValue<int32_t, latency_anchor_id, _Params...>::value;
+    static constexpr auto _constraint =
+        _GetValue3<int32_t, type, int32_t, latency_constraint,
+                   _Params...>::value;
+
+    static constexpr int32_t _target_anchor = std::get<0>(_constraint);
+    static constexpr type _control_type = std::get<1>(_constraint);
+    static constexpr int32_t _cycle = std::get<2>(_constraint);
+    int32_t _type;
+    if (_control_type == type::none) {
+      _type = 0;
+    } else if (_control_type == type::exact) {
+      _type = 1;
+    } else if (_control_type == type::max) {
+      _type = 2;
+    } else { // _control_type == type::min
+      _type = 3;
+    }
+
+    __ocl_RPipeTy<_dataT> _RPipe =
+        __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
+    _dataT TempData;
+    __latency_control_bl_read_wrapper(_RPipe, &TempData, _anchor_id,
+                                      _target_anchor, _type, _cycle);
+    return TempData;
+#else
+    assert(!"Pipes are not supported on a host device!");
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
+  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
+  // friendly LLVM IR.
+  template <class... _Params> static void write(const _dataT &_Data) {
+#ifdef __SYCL_DEVICE_ONLY__
+    static constexpr auto _anchor_id =
+        _GetValue<int32_t, latency_anchor_id, _Params...>::value;
+    static constexpr auto _constraint =
+        _GetValue3<int32_t, type, int32_t, latency_constraint,
+                   _Params...>::value;
+
+    static constexpr int32_t _target_anchor = std::get<0>(_constraint);
+    static constexpr type _control_type = std::get<1>(_constraint);
+    static constexpr int32_t _cycle = std::get<2>(_constraint);
+    int32_t _type;
+    if (_control_type == type::none) {
+      _type = 0;
+    } else if (_control_type == type::exact) {
+      _type = 1;
+    } else if (_control_type == type::max) {
+      _type = 2;
+    } else { // _control_type == type::min
+      _type = 3;
+    }
+
+    __ocl_WPipeTy<_dataT> _WPipe =
+        __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
+    __latency_control_bl_write_wrapper(_WPipe, &_Data, _anchor_id,
+                                       _target_anchor, _type, _cycle);
+#else
+    (void)_Data;
+    assert(!"Pipes are not supported on a host device!");
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
+private:
+  static constexpr int32_t m_Size = sizeof(_dataT);
+  static constexpr int32_t m_Alignment = alignof(_dataT);
+  static constexpr int32_t m_Capacity = _min_capacity;
+#ifdef __SYCL_DEVICE_ONLY__
+  static constexpr struct ConstantPipeStorage m_Storage = {m_Size, m_Alignment,
+                                                           m_Capacity};
+
+  template <typename _T>
+  static int32_t
+  __latency_control_nb_read_wrapper(__ocl_RPipeTy<_T> Pipe, _T *Data,
+                                    int32_t AnchorID, int32_t TargetAnchor,
+                                    int32_t Type, int32_t Cycle) {
+    return __spirv_ReadPipe(Pipe, Data, m_Size, m_Alignment);
+  }
+
+  template <typename _T>
+  static int32_t
+  __latency_control_nb_write_wrapper(__ocl_WPipeTy<_T> Pipe, const _T *Data,
+                                     int32_t AnchorID, int32_t TargetAnchor,
+                                     int32_t Type, int32_t Cycle) {
+    return __spirv_WritePipe(Pipe, Data, m_Size, m_Alignment);
+  }
+
+  template <typename _T>
+  static void __latency_control_bl_read_wrapper(__ocl_RPipeTy<_T> Pipe,
+                                                _T *Data, int32_t AnchorID,
+                                                int32_t TargetAnchor,
+                                                int32_t Type, int32_t Cycle) {
+    return __spirv_ReadPipeBlockingINTEL(Pipe, Data, m_Size, m_Alignment);
+  }
+
+  template <typename _T>
+  static void
+  __latency_control_bl_write_wrapper(__ocl_WPipeTy<_T> Pipe, const _T *Data,
+                                     int32_t AnchorID, int32_t TargetAnchor,
+                                     int32_t Type, int32_t Cycle) {
+    return __spirv_WritePipeBlockingINTEL(Pipe, Data, m_Size, m_Alignment);
+  }
+#endif // __SYCL_DEVICE_ONLY__
+};
+
+} // namespace experimental
+} // namespace intel
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/sycl/ext/intel/fpga_extensions.hpp b/sycl/include/sycl/ext/intel/fpga_extensions.hpp
index 2ff9dbbda5a83..c3b30f5558c82 100644
--- a/sycl/include/sycl/ext/intel/fpga_extensions.hpp
+++ b/sycl/include/sycl/ext/intel/fpga_extensions.hpp
@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #pragma once
+#include <sycl/ext/intel/experimental/fpga_lsu.hpp>
+#include <sycl/ext/intel/experimental/pipes.hpp>
 #include <sycl/ext/intel/fpga_device_selector.hpp>
 #include <sycl/ext/intel/fpga_dsp_control.hpp>
 #include <sycl/ext/intel/fpga_loop_fuse.hpp>

From e4e9c8ea6fb3f403b71d9e2c94d717ce24c89354 Mon Sep 17 00:00:00 2001
From: "Niu, Shuo" <shuo.niu@intel.com>
Date: Fri, 3 Dec 2021 07:11:17 -0800
Subject: [PATCH 2/3] Add constexpr and comments

---
 .../sycl/ext/intel/experimental/fpga_lsu.hpp  | 13 +++++----
 .../sycl/ext/intel/experimental/pipes.hpp     | 28 +++++++++++--------
 2 files changed, 23 insertions(+), 18 deletions(-)

diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp
index 9819d6791db31..f40147550b653 100644
--- a/sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp
@@ -65,11 +65,11 @@ template <class... _mem_access_params> class lsu final {
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
     int32_t _type;
-    if (_control_type == type::none) {
+    if constexpr (_control_type == type::none) {
       _type = 0;
-    } else if (_control_type == type::exact) {
+    } else if constexpr (_control_type == type::exact) {
       _type = 1;
-    } else if (_control_type == type::max) {
+    } else if constexpr (_control_type == type::max) {
       _type = 2;
     } else { // _control_type == type::min
       _type = 3;
@@ -97,11 +97,11 @@ template <class... _mem_access_params> class lsu final {
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
     int32_t _type;
-    if (_control_type == type::none) {
+    if constexpr (_control_type == type::none) {
       _type = 0;
-    } else if (_control_type == type::exact) {
+    } else if constexpr (_control_type == type::exact) {
       _type = 1;
-    } else if (_control_type == type::max) {
+    } else if constexpr (_control_type == type::max) {
       _type = 2;
     } else { // _control_type == type::min
       _type = 3;
@@ -161,6 +161,7 @@ template <class... _mem_access_params> class lsu final {
   }
 
 #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+  // FPGA BE will recognize this function and extract its arguments.
   template <typename _T>
   static _T *__latency_control_mem_wrapper(_T *Ptr, int32_t AnchorID,
                                            int32_t TargetAnchor, int32_t Type,
diff --git a/sycl/include/sycl/ext/intel/experimental/pipes.hpp b/sycl/include/sycl/ext/intel/experimental/pipes.hpp
index 1e46890b4958b..dfb77132578d8 100644
--- a/sycl/include/sycl/ext/intel/experimental/pipes.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/pipes.hpp
@@ -36,11 +36,11 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
     int32_t _type;
-    if (_control_type == type::none) {
+    if constexpr (_control_type == type::none) {
       _type = 0;
-    } else if (_control_type == type::exact) {
+    } else if constexpr (_control_type == type::exact) {
       _type = 1;
-    } else if (_control_type == type::max) {
+    } else if constexpr (_control_type == type::max) {
       _type = 2;
     } else { // _control_type == type::min
       _type = 3;
@@ -73,11 +73,11 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
     int32_t _type;
-    if (_control_type == type::none) {
+    if constexpr (_control_type == type::none) {
       _type = 0;
-    } else if (_control_type == type::exact) {
+    } else if constexpr (_control_type == type::exact) {
       _type = 1;
-    } else if (_control_type == type::max) {
+    } else if constexpr (_control_type == type::max) {
       _type = 2;
     } else { // _control_type == type::min
       _type = 3;
@@ -109,11 +109,11 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
     int32_t _type;
-    if (_control_type == type::none) {
+    if constexpr (_control_type == type::none) {
       _type = 0;
-    } else if (_control_type == type::exact) {
+    } else if constexpr (_control_type == type::exact) {
       _type = 1;
-    } else if (_control_type == type::max) {
+    } else if constexpr (_control_type == type::max) {
       _type = 2;
     } else { // _control_type == type::min
       _type = 3;
@@ -144,11 +144,11 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
     int32_t _type;
-    if (_control_type == type::none) {
+    if constexpr (_control_type == type::none) {
       _type = 0;
-    } else if (_control_type == type::exact) {
+    } else if constexpr (_control_type == type::exact) {
       _type = 1;
-    } else if (_control_type == type::max) {
+    } else if constexpr (_control_type == type::max) {
       _type = 2;
     } else { // _control_type == type::min
       _type = 3;
@@ -172,6 +172,7 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
   static constexpr struct ConstantPipeStorage m_Storage = {m_Size, m_Alignment,
                                                            m_Capacity};
 
+  // FPGA BE will recognize this function and extract its arguments.
   template <typename _T>
   static int32_t
   __latency_control_nb_read_wrapper(__ocl_RPipeTy<_T> Pipe, _T *Data,
@@ -180,6 +181,7 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
     return __spirv_ReadPipe(Pipe, Data, m_Size, m_Alignment);
   }
 
+  // FPGA BE will recognize this function and extract its arguments.
   template <typename _T>
   static int32_t
   __latency_control_nb_write_wrapper(__ocl_WPipeTy<_T> Pipe, const _T *Data,
@@ -188,6 +190,7 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
     return __spirv_WritePipe(Pipe, Data, m_Size, m_Alignment);
   }
 
+  // FPGA BE will recognize this function and extract its arguments.
   template <typename _T>
   static void __latency_control_bl_read_wrapper(__ocl_RPipeTy<_T> Pipe,
                                                 _T *Data, int32_t AnchorID,
@@ -196,6 +199,7 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
     return __spirv_ReadPipeBlockingINTEL(Pipe, Data, m_Size, m_Alignment);
   }
 
+  // FPGA BE will recognize this function and extract its arguments.
   template <typename _T>
   static void
   __latency_control_bl_write_wrapper(__ocl_WPipeTy<_T> Pipe, const _T *Data,

From 324ef5c5bc032c0bd0ad40d308155421ebd02d8d Mon Sep 17 00:00:00 2001
From: "Niu, Shuo" <shuo.niu@intel.com>
Date: Mon, 6 Dec 2021 14:12:47 -0800
Subject: [PATCH 3/3] Minor updates on if constexpr statements and TODO
 comments

---
 .../sycl/ext/intel/experimental/fpga_lsu.hpp  | 17 ++++-----
 .../sycl/ext/intel/experimental/pipes.hpp     | 36 +++++++++----------
 2 files changed, 23 insertions(+), 30 deletions(-)

diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp
index f40147550b653..ad3fda6729b3c 100644
--- a/sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp
@@ -64,14 +64,12 @@ template <class... _mem_access_params> class lsu final {
     static constexpr int32_t _target_anchor = std::get<0>(_constraint);
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
-    int32_t _type;
-    if constexpr (_control_type == type::none) {
-      _type = 0;
-    } else if constexpr (_control_type == type::exact) {
+    int32_t _type = 0; // Default: _control_type == type::none
+    if constexpr (_control_type == type::exact) {
       _type = 1;
     } else if constexpr (_control_type == type::max) {
       _type = 2;
-    } else { // _control_type == type::min
+    } else if constexpr (_control_type == type::min) {
       _type = 3;
     }
 
@@ -96,14 +94,12 @@ template <class... _mem_access_params> class lsu final {
     static constexpr int32_t _target_anchor = std::get<0>(_constraint);
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
-    int32_t _type;
-    if constexpr (_control_type == type::none) {
-      _type = 0;
-    } else if constexpr (_control_type == type::exact) {
+    int32_t _type = 0; // Default: _control_type == type::none
+    if constexpr (_control_type == type::exact) {
       _type = 1;
     } else if constexpr (_control_type == type::max) {
       _type = 2;
-    } else { // _control_type == type::min
+    } else if constexpr (_control_type == type::min) {
       _type = 3;
     }
 
@@ -162,6 +158,7 @@ template <class... _mem_access_params> class lsu final {
 
 #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
   // FPGA BE will recognize this function and extract its arguments.
+  // TODO: Pass latency control params via __builtin_intel_fpga_mem when ready.
   template <typename _T>
   static _T *__latency_control_mem_wrapper(_T *Ptr, int32_t AnchorID,
                                            int32_t TargetAnchor, int32_t Type,
diff --git a/sycl/include/sycl/ext/intel/experimental/pipes.hpp b/sycl/include/sycl/ext/intel/experimental/pipes.hpp
index dfb77132578d8..46494884c29ba 100644
--- a/sycl/include/sycl/ext/intel/experimental/pipes.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/pipes.hpp
@@ -35,14 +35,12 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
     static constexpr int32_t _target_anchor = std::get<0>(_constraint);
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
-    int32_t _type;
-    if constexpr (_control_type == type::none) {
-      _type = 0;
-    } else if constexpr (_control_type == type::exact) {
+    int32_t _type = 0; // Default: _control_type == type::none
+    if constexpr (_control_type == type::exact) {
       _type = 1;
     } else if constexpr (_control_type == type::max) {
       _type = 2;
-    } else { // _control_type == type::min
+    } else if constexpr (_control_type == type::min) {
       _type = 3;
     }
 
@@ -72,14 +70,12 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
     static constexpr int32_t _target_anchor = std::get<0>(_constraint);
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
-    int32_t _type;
-    if constexpr (_control_type == type::none) {
-      _type = 0;
-    } else if constexpr (_control_type == type::exact) {
+    int32_t _type = 0; // Default: _control_type == type::none
+    if constexpr (_control_type == type::exact) {
       _type = 1;
     } else if constexpr (_control_type == type::max) {
       _type = 2;
-    } else { // _control_type == type::min
+    } else if constexpr (_control_type == type::min) {
       _type = 3;
     }
 
@@ -108,14 +104,12 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
     static constexpr int32_t _target_anchor = std::get<0>(_constraint);
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
-    int32_t _type;
-    if constexpr (_control_type == type::none) {
-      _type = 0;
-    } else if constexpr (_control_type == type::exact) {
+    int32_t _type = 0; // Default: _control_type == type::none
+    if constexpr (_control_type == type::exact) {
       _type = 1;
     } else if constexpr (_control_type == type::max) {
       _type = 2;
-    } else { // _control_type == type::min
+    } else if constexpr (_control_type == type::min) {
       _type = 3;
     }
 
@@ -143,14 +137,12 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
     static constexpr int32_t _target_anchor = std::get<0>(_constraint);
     static constexpr type _control_type = std::get<1>(_constraint);
     static constexpr int32_t _cycle = std::get<2>(_constraint);
-    int32_t _type;
-    if constexpr (_control_type == type::none) {
-      _type = 0;
-    } else if constexpr (_control_type == type::exact) {
+    int32_t _type = 0; // Default: _control_type == type::none
+    if constexpr (_control_type == type::exact) {
       _type = 1;
     } else if constexpr (_control_type == type::max) {
       _type = 2;
-    } else { // _control_type == type::min
+    } else if constexpr (_control_type == type::min) {
       _type = 3;
     }
 
@@ -173,6 +165,7 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
                                                            m_Capacity};
 
   // FPGA BE will recognize this function and extract its arguments.
+  // TODO: Pass latency control parameters via the __spirv_* builtin when ready.
   template <typename _T>
   static int32_t
   __latency_control_nb_read_wrapper(__ocl_RPipeTy<_T> Pipe, _T *Data,
@@ -182,6 +175,7 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
   }
 
   // FPGA BE will recognize this function and extract its arguments.
+  // TODO: Pass latency control parameters via the __spirv_* builtin when ready.
   template <typename _T>
   static int32_t
   __latency_control_nb_write_wrapper(__ocl_WPipeTy<_T> Pipe, const _T *Data,
@@ -191,6 +185,7 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
   }
 
   // FPGA BE will recognize this function and extract its arguments.
+  // TODO: Pass latency control parameters via the __spirv_* builtin when ready.
   template <typename _T>
   static void __latency_control_bl_read_wrapper(__ocl_RPipeTy<_T> Pipe,
                                                 _T *Data, int32_t AnchorID,
@@ -200,6 +195,7 @@ template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
   }
 
   // FPGA BE will recognize this function and extract its arguments.
+  // TODO: Pass latency control parameters via the __spirv_* builtin when ready.
   template <typename _T>
   static void
   __latency_control_bl_write_wrapper(__ocl_WPipeTy<_T> Pipe, const _T *Data,