From c7da99b7d4160a1ad931e2689d1a9f9ed2b25a23 Mon Sep 17 00:00:00 2001
From: shuoniu-intel <81583996+shuoniu-intel@users.noreply.github.com>
Date: Fri, 18 Jun 2021 13:55:08 -0400
Subject: [PATCH 01/11] Update fpga_utils.hpp

---
 sycl/include/CL/sycl/INTEL/fpga_utils.hpp | 29 +++++++++++++++--------
 1 file changed, 19 insertions(+), 10 deletions(-)

diff --git a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
index 1fc758f4b4f8b..64e6507a67a93 100644
--- a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
+++ b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
@@ -16,19 +16,28 @@ __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace INTEL {
 
-template <template <int32_t> class _Type, class _T>
-struct _MatchType : std::is_same<_Type<_T::value>, _T> {};
+template <class _D, class _T>
+struct _MatchType : std::is_same<typename _D::type_id, typename _T::type_id> {};
 
-template <template <int32_t> class _Type, class... _T> struct _GetValue {
-  static constexpr auto value = _Type<0>::default_value;
-};
+template <class _D, class... _T>
+struct _GetValue;
+
+template <class _D>
+struct _GetValue<_D> : std::integral_constant<decltype(_D::value), _D::value> {};
+
+template <class _D, class _T1, class... _T>
+struct _GetValue<_D, _T1, _T...> {
+  template <class _D2, class _T12, class _Enable = void>
+  struct impl : 
+    std::integral_constant<decltype(_D::value), _GetValue<_D, _T...>::value> {};
 
-template <template <int32_t> class _Type, class _T1, class... _T>
-struct _GetValue<_Type, _T1, _T...> {
-  static constexpr auto value =
-      detail::conditional_t<_MatchType<_Type, _T1>::value, _T1,
-                            _GetValue<_Type, _T...>>::value;
+  template <class _D2, class _T12>
+  struct impl<_D2, _T12, std::enable_if_t<_MatchType<_D2, _T12>::value>> : 
+    std::integral_constant<decltype(_D::value), _T1::value> {};
+
+  static constexpr auto value = impl<_D, _T1>::value;
 };
+
 } // namespace INTEL
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

From 19000e8fa64a7d4a5f7ef185606e39c4bf81c131 Mon Sep 17 00:00:00 2001
From: shuoniu-intel <81583996+shuoniu-intel@users.noreply.github.com>
Date: Fri, 18 Jun 2021 14:26:25 -0400
Subject: [PATCH 02/11] Update fpga_lsu.hpp

---
 sycl/include/CL/sycl/INTEL/fpga_lsu.hpp | 32 ++++++++++++-------------
 1 file changed, 16 insertions(+), 16 deletions(-)

diff --git a/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp b/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp
index ae210b15cf4c9..c52de9543f9fc 100644
--- a/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp
+++ b/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp
@@ -19,24 +19,24 @@ 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;
+struct burst_coalesce_impl_id;
+template <int32_t _N> struct burst_coalesce_impl : std::integral_constant<int32_t, _N> {
+  using type_id = burst_coalesce_impl_id;
 };
 
-template <int32_t _N> struct cache {
-  static constexpr int32_t value = _N;
-  static constexpr int32_t default_value = 0;
+struct cache_id;
+template <int32_t _N> struct cache : std::integral_constant<int32_t, _N> {
+  using type_id = cache_id;
 };
 
-template <int32_t _N> struct prefetch_impl {
-  static constexpr int32_t value = _N;
-  static constexpr int32_t default_value = 0;
+struct prefetch_impl_id;
+template <int32_t _N> struct prefetch_impl : std::integral_constant<int32_t, _N> {
+  using type_id = prefetch_impl_id;
 };
 
-template <int32_t _N> struct statically_coalesce_impl {
-  static constexpr int32_t value = _N;
-  static constexpr int32_t default_value = 1;
+struct statically_coalesce_impl_id;
+template <int32_t _N> struct statically_coalesce_impl : std::integral_constant<int32_t, _N>{
+  using type_id = statically_coalesce_impl_id;
 };
 
 template <bool _B> using burst_coalesce = burst_coalesce_impl<_B>;
@@ -77,21 +77,21 @@ template <class... _mem_access_params> class lsu final {
 
 private:
   static constexpr int32_t _burst_coalesce_val =
-      _GetValue<burst_coalesce_impl, _mem_access_params...>::value;
+      _GetValue<burst_coalesce_impl<0>, _mem_access_params...>::value;
   static constexpr uint8_t _burst_coalesce =
       _burst_coalesce_val == 1 ? BURST_COALESCE : 0;
 
   static constexpr int32_t _cache_val =
-      _GetValue<cache, _mem_access_params...>::value;
+      _GetValue<cache<0>, _mem_access_params...>::value;
   static constexpr uint8_t _cache = (_cache_val > 0) ? CACHE : 0;
 
   static constexpr int32_t _statically_coalesce_val =
-      _GetValue<statically_coalesce_impl, _mem_access_params...>::value;
+      _GetValue<statically_coalesce_impl<1>, _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<prefetch_impl, _mem_access_params...>::value;
+      _GetValue<prefetch_impl<0>, _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");

From 7bd756db3e6218e89e1cc8d5f70e78fb5511a842 Mon Sep 17 00:00:00 2001
From: shuoniu-intel <81583996+shuoniu-intel@users.noreply.github.com>
Date: Fri, 18 Jun 2021 14:42:15 -0400
Subject: [PATCH 03/11] Clang-format fpga_lsu.hpp

---
 sycl/include/CL/sycl/INTEL/fpga_lsu.hpp | 9 ++++++---
 1 file changed, 6 insertions(+), 3 deletions(-)

diff --git a/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp b/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp
index c52de9543f9fc..ab2c859603bf5 100644
--- a/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp
+++ b/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp
@@ -20,7 +20,8 @@ constexpr uint8_t STATICALLY_COALESCE = 0x4;
 constexpr uint8_t PREFETCH = 0x8;
 
 struct burst_coalesce_impl_id;
-template <int32_t _N> struct burst_coalesce_impl : std::integral_constant<int32_t, _N> {
+template <int32_t _N>
+struct burst_coalesce_impl : std::integral_constant<int32_t, _N> {
   using type_id = burst_coalesce_impl_id;
 };
 
@@ -30,12 +31,14 @@ template <int32_t _N> struct cache : std::integral_constant<int32_t, _N> {
 };
 
 struct prefetch_impl_id;
-template <int32_t _N> struct prefetch_impl : std::integral_constant<int32_t, _N> {
+template <int32_t _N>
+struct prefetch_impl : std::integral_constant<int32_t, _N> {
   using type_id = prefetch_impl_id;
 };
 
 struct statically_coalesce_impl_id;
-template <int32_t _N> struct statically_coalesce_impl : std::integral_constant<int32_t, _N>{
+template <int32_t _N>
+struct statically_coalesce_impl : std::integral_constant<int32_t, _N> {
   using type_id = statically_coalesce_impl_id;
 };
 

From ab5a614cea7c40de500de117a86500ec8a955aac Mon Sep 17 00:00:00 2001
From: shuoniu-intel <81583996+shuoniu-intel@users.noreply.github.com>
Date: Fri, 18 Jun 2021 14:45:04 -0400
Subject: [PATCH 04/11] Clang-format fpga_utils.hpp

---
 sycl/include/CL/sycl/INTEL/fpga_utils.hpp | 14 ++++++--------
 1 file changed, 6 insertions(+), 8 deletions(-)

diff --git a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
index 64e6507a67a93..bfa03a75ee04b 100644
--- a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
+++ b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
@@ -19,21 +19,19 @@ namespace INTEL {
 template <class _D, class _T>
 struct _MatchType : std::is_same<typename _D::type_id, typename _T::type_id> {};
 
-template <class _D, class... _T>
-struct _GetValue;
+template <class _D, class... _T> struct _GetValue;
 
 template <class _D>
 struct _GetValue<_D> : std::integral_constant<decltype(_D::value), _D::value> {};
 
-template <class _D, class _T1, class... _T>
-struct _GetValue<_D, _T1, _T...> {
+template <class _D, class _T1, class... _T> struct _GetValue<_D, _T1, _T...> {
   template <class _D2, class _T12, class _Enable = void>
-  struct impl : 
-    std::integral_constant<decltype(_D::value), _GetValue<_D, _T...>::value> {};
+  struct impl : std::integral_constant<decltype(_D::value),
+                                       _GetValue<_D, _T...>::value> {};
 
   template <class _D2, class _T12>
-  struct impl<_D2, _T12, std::enable_if_t<_MatchType<_D2, _T12>::value>> : 
-    std::integral_constant<decltype(_D::value), _T1::value> {};
+  struct impl<_D2, _T12, std::enable_if_t<_MatchType<_D2, _T12>::value>>
+      : std::integral_constant<decltype(_D::value), _T1::value> {};
 
   static constexpr auto value = impl<_D, _T1>::value;
 };

From b6d2035bce0c64d96896c3d25c56c2097528e942 Mon Sep 17 00:00:00 2001
From: shuoniu-intel <81583996+shuoniu-intel@users.noreply.github.com>
Date: Fri, 18 Jun 2021 14:50:35 -0400
Subject: [PATCH 05/11] Clang-format fpga_utils.hpp

---
 sycl/include/CL/sycl/INTEL/fpga_utils.hpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
index bfa03a75ee04b..7c4b52a30b3ef 100644
--- a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
+++ b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
@@ -22,7 +22,8 @@ struct _MatchType : std::is_same<typename _D::type_id, typename _T::type_id> {};
 template <class _D, class... _T> struct _GetValue;
 
 template <class _D>
-struct _GetValue<_D> : std::integral_constant<decltype(_D::value), _D::value> {};
+struct _GetValue<_D> : std::integral_constant<decltype(_D::value), _D::value> {
+};
 
 template <class _D, class _T1, class... _T> struct _GetValue<_D, _T1, _T...> {
   template <class _D2, class _T12, class _Enable = void>

From 5498834a0005ddae3108495f9b663a7a699618ec Mon Sep 17 00:00:00 2001
From: shuoniu-intel <81583996+shuoniu-intel@users.noreply.github.com>
Date: Mon, 21 Jun 2021 08:55:10 -0400
Subject: [PATCH 06/11] Update fpga_utils.hpp

---
 sycl/include/CL/sycl/INTEL/fpga_utils.hpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
index 7c4b52a30b3ef..a9033242e9790 100644
--- a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
+++ b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
@@ -36,7 +36,6 @@ template <class _D, class _T1, class... _T> struct _GetValue<_D, _T1, _T...> {
 
   static constexpr auto value = impl<_D, _T1>::value;
 };
-
 } // namespace INTEL
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

From 2b96d496a1243db42d53e2c7667757e13526c437 Mon Sep 17 00:00:00 2001
From: shuoniu-intel <81583996+shuoniu-intel@users.noreply.github.com>
Date: Tue, 6 Jul 2021 09:34:24 -0400
Subject: [PATCH 07/11] Create arbitrary_template_arg.cpp

---
 sycl/test/fpga_tests/arbitrary_template_arg.cpp | 1 +
 1 file changed, 1 insertion(+)
 create mode 100644 sycl/test/fpga_tests/arbitrary_template_arg.cpp

diff --git a/sycl/test/fpga_tests/arbitrary_template_arg.cpp b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
new file mode 100644
index 0000000000000..8b137891791fe
--- /dev/null
+++ b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
@@ -0,0 +1 @@
+

From 17eac46f45a035970c53b27db9aaae64e89618e0 Mon Sep 17 00:00:00 2001
From: shuoniu-intel <81583996+shuoniu-intel@users.noreply.github.com>
Date: Tue, 6 Jul 2021 11:13:06 -0400
Subject: [PATCH 08/11] Update arbitrary_template_arg.cpp

---
 .../fpga_tests/arbitrary_template_arg.cpp     | 47 +++++++++++++++++++
 1 file changed, 47 insertions(+)

diff --git a/sycl/test/fpga_tests/arbitrary_template_arg.cpp b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
index 8b137891791fe..a122b0e8a3dda 100644
--- a/sycl/test/fpga_tests/arbitrary_template_arg.cpp
+++ b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
@@ -1 +1,48 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %RUN_ON_HOST %t.out
 
+// This test performs basic check of SYCL FPGA arbitrary template argument list.
+
+#include <CL/sycl.hpp>
+#include <CL/sycl/INTEL/fpga_utils.hpp>
+#include <iostream>
+
+enum class enum_class { first, second };
+
+struct test_int_id;
+template <int _N> struct test_int : std::integral_constant<int, _N> {
+  using type_id = test_int_id;
+};
+
+struct test_enum_class_id;
+template <enum_class _N>
+struct test_enum_class : std::integral_constant<enum_class, _N> {
+  using type_id = test_enum_class_id;
+};
+
+template <class... _Params> void func() {
+  static constexpr auto test_int_value =
+      _GetValue<test_int<0>, _Params...>::value;
+  static constexpr auto test_enum_class_value =
+      _GetValue<test_enum_class<enum_class::first>, _Params...>::value;
+
+  std::cout << test_int_value << " ";
+  if (test_enum_class_value == enum_class::first) {
+    std::cout << "first" << std::endl;
+  } else {
+    std::cout << "second" << std::endl;
+  }
+}
+
+int main() {
+  // CHECK: 0 first
+  func();
+  // CHECK: 1 first
+  func<test_int<1>>();
+  // CHECK: 0 second
+  func<test_enum_class<enum_class::second>>();
+  // CHECK: 1 second
+  func<test_int<1>, test_enum_class<enum_class::second>>();
+  // CHECK: 1 second
+  func<test_enum_class<enum_class::second>, test_int<1>>();
+}

From fce639180230d398f0cc55b69b3384d2b60f16ef Mon Sep 17 00:00:00 2001
From: shuoniu-intel <81583996+shuoniu-intel@users.noreply.github.com>
Date: Tue, 6 Jul 2021 11:57:42 -0400
Subject: [PATCH 09/11] Update arbitrary_template_arg.cpp

---
 sycl/test/fpga_tests/arbitrary_template_arg.cpp | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/sycl/test/fpga_tests/arbitrary_template_arg.cpp b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
index a122b0e8a3dda..2b1591d254319 100644
--- a/sycl/test/fpga_tests/arbitrary_template_arg.cpp
+++ b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
@@ -22,9 +22,10 @@ struct test_enum_class : std::integral_constant<enum_class, _N> {
 
 template <class... _Params> void func() {
   static constexpr auto test_int_value =
-      _GetValue<test_int<0>, _Params...>::value;
+      sycl::INTEL::_GetValue<test_int<0>, _Params...>::value;
   static constexpr auto test_enum_class_value =
-      _GetValue<test_enum_class<enum_class::first>, _Params...>::value;
+      sycl::INTEL::_GetValue<test_enum_class<enum_class::first>,
+                             _Params...>::value;
 
   std::cout << test_int_value << " ";
   if (test_enum_class_value == enum_class::first) {

From 4f44868cca5a0e649b0e498ff2caf57ae7513dae Mon Sep 17 00:00:00 2001
From: shuoniu-intel <81583996+shuoniu-intel@users.noreply.github.com>
Date: Wed, 7 Jul 2021 07:37:47 -0400
Subject: [PATCH 10/11] Previous tests timeout, rerun tests.

---
 sycl/test/fpga_tests/arbitrary_template_arg.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/sycl/test/fpga_tests/arbitrary_template_arg.cpp b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
index 2b1591d254319..b835cd99ca07c 100644
--- a/sycl/test/fpga_tests/arbitrary_template_arg.cpp
+++ b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
@@ -1,6 +1,5 @@
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %RUN_ON_HOST %t.out
-
 // This test performs basic check of SYCL FPGA arbitrary template argument list.
 
 #include <CL/sycl.hpp>

From 44005ae6dcf9a99bfdbeffeb814485690d19cde0 Mon Sep 17 00:00:00 2001
From: shuoniu-intel <81583996+shuoniu-intel@users.noreply.github.com>
Date: Wed, 7 Jul 2021 13:48:18 -0400
Subject: [PATCH 11/11] Update arbitrary_template_arg.cpp

Apply static_assert compile-time check.
---
 .../fpga_tests/arbitrary_template_arg.cpp     | 56 ++++++++-----------
 1 file changed, 23 insertions(+), 33 deletions(-)

diff --git a/sycl/test/fpga_tests/arbitrary_template_arg.cpp b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
index b835cd99ca07c..0ed92dd592138 100644
--- a/sycl/test/fpga_tests/arbitrary_template_arg.cpp
+++ b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
@@ -1,48 +1,38 @@
-// RUN: %clangxx -fsycl %s -o %t.out
-// RUN: %RUN_ON_HOST %t.out
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s -Xclang -verify-ignore-unexpected=note,warning
+// expected-no-diagnostics
+
 // This test performs basic check of SYCL FPGA arbitrary template argument list.
 
 #include <CL/sycl.hpp>
 #include <CL/sycl/INTEL/fpga_utils.hpp>
 #include <iostream>
 
-enum class enum_class { first, second };
+enum class enumClass { first, second };
 
-struct test_int_id;
-template <int _N> struct test_int : std::integral_constant<int, _N> {
-  using type_id = test_int_id;
+struct testIntID;
+template <int _N> struct testInt : std::integral_constant<int, _N> {
+  using type_id = testIntID;
 };
 
-struct test_enum_class_id;
-template <enum_class _N>
-struct test_enum_class : std::integral_constant<enum_class, _N> {
-  using type_id = test_enum_class_id;
+struct testEnumID;
+template <enumClass _N>
+struct testEnum : std::integral_constant<enumClass, _N> {
+  using type_id = testEnumID;
 };
 
-template <class... _Params> void func() {
-  static constexpr auto test_int_value =
-      sycl::INTEL::_GetValue<test_int<0>, _Params...>::value;
-  static constexpr auto test_enum_class_value =
-      sycl::INTEL::_GetValue<test_enum_class<enum_class::first>,
-                             _Params...>::value;
-
-  std::cout << test_int_value << " ";
-  if (test_enum_class_value == enum_class::first) {
-    std::cout << "first" << std::endl;
-  } else {
-    std::cout << "second" << std::endl;
-  }
+template <int ExpectedIntValue, enumClass ExpectedEnumValue, class... _Params>
+void func() {
+  static_assert(sycl::INTEL::_GetValue<testInt<0>, _Params...>::value ==
+                ExpectedIntValue);
+  static_assert(
+      sycl::INTEL::_GetValue<testEnum<enumClass::first>, _Params...>::value ==
+      ExpectedEnumValue);
 }
 
 int main() {
-  // CHECK: 0 first
-  func();
-  // CHECK: 1 first
-  func<test_int<1>>();
-  // CHECK: 0 second
-  func<test_enum_class<enum_class::second>>();
-  // CHECK: 1 second
-  func<test_int<1>, test_enum_class<enum_class::second>>();
-  // CHECK: 1 second
-  func<test_enum_class<enum_class::second>, test_int<1>>();
+  func<0, enumClass::first>();
+  func<1, enumClass::first, testInt<1>>();
+  func<0, enumClass::second, testEnum<enumClass::second>>();
+  func<1, enumClass::second, testInt<1>, testEnum<enumClass::second>>();
+  func<1, enumClass::second, testEnum<enumClass::second>, testInt<1>>();
 }