diff --git a/sycl/include/sycl/ext/intel/fpga_lsu.hpp b/sycl/include/sycl/ext/intel/fpga_lsu.hpp
index 4c40c58a671f2..35f99f605154e 100644
--- a/sycl/include/sycl/ext/intel/fpga_lsu.hpp
+++ b/sycl/include/sycl/ext/intel/fpga_lsu.hpp
@@ -20,27 +20,24 @@ constexpr uint8_t CACHE = 0x2;
 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> {
-  using type_id = burst_coalesce_impl_id;
+template <int32_t _N> struct burst_coalesce_impl {
+  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 cache {
+  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 prefetch_impl {
+  static constexpr int32_t value = _N;
+  static constexpr int32_t default_value = 0;
 };
 
-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 <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>;
@@ -81,21 +78,21 @@ template <class... _mem_access_params> class lsu final {
 
 private:
   static constexpr int32_t _burst_coalesce_val =
-      _GetValue<burst_coalesce_impl<0>, _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;
 
   static constexpr int32_t _cache_val =
-      _GetValue<cache<0>, _mem_access_params...>::value;
+      _GetValue<cache, _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<1>, _mem_access_params...>::value;
+      _GetValue<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<prefetch_impl<0>, _mem_access_params...>::value;
+      _GetValue<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");
diff --git a/sycl/include/sycl/ext/intel/fpga_utils.hpp b/sycl/include/sycl/ext/intel/fpga_utils.hpp
index 0baec2abfe7e5..b84293e0244a6 100644
--- a/sycl/include/sycl/ext/intel/fpga_utils.hpp
+++ b/sycl/include/sycl/ext/intel/fpga_utils.hpp
@@ -17,25 +17,18 @@ namespace sycl {
 namespace ext {
 namespace intel {
 
-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 _MatchType : std::is_same<_Type<_T::value>, _T> {};
 
-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... _T> struct _GetValue {
+  static constexpr auto value = _Type<0>::default_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;
+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;
 };
 } // namespace intel
 } // namespace ext
diff --git a/sycl/test/fpga_tests/arbitrary_template_arg.cpp b/sycl/test/fpga_tests/arbitrary_template_arg.cpp
deleted file mode 100644
index 0ed92dd592138..0000000000000
--- a/sycl/test/fpga_tests/arbitrary_template_arg.cpp
+++ /dev/null
@@ -1,38 +0,0 @@
-// 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>>();
-}