diff --git a/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp b/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp
index ae210b15cf4c9..ab2c859603bf5 100644
--- a/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp
+++ b/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp
@@ -19,24 +19,27 @@ 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 +80,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");
diff --git a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
index 1fc758f4b4f8b..a9033242e9790 100644
--- a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
+++ b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp
@@ -16,18 +16,25 @@ __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 <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 _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 <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
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..0ed92dd592138
--- /dev/null
+++ b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
@@ -0,0 +1,38 @@
+// 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 enumClass { first, second };
+
+struct testIntID;
+template <int _N> struct testInt : std::integral_constant<int, _N> {
+  using type_id = testIntID;
+};
+
+struct testEnumID;
+template <enumClass _N>
+struct testEnum : std::integral_constant<enumClass, _N> {
+  using type_id = testEnumID;
+};
+
+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() {
+  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>>();
+}