diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp
index 1723fe40535db..c8987fd2e57da 100644
--- a/clang/lib/CodeGen/CodeGenTypes.cpp
+++ b/clang/lib/CodeGen/CodeGenTypes.cpp
@@ -711,11 +711,6 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) {
                              "__spv::__spirv_CooperativeMatrixKHR") {
           ResultType = ConvertSPVCooperativeMatrixType(RD);
           break;
-        } else if (RD && RD->getQualifiedNameAsString() ==
-                             "__spv::__spirv_TaskSequenceINTEL") {
-          ResultType = llvm::TargetExtType::get(getLLVMContext(),
-                                                "spirv.TaskSequenceINTEL");
-          break;
         }
       }
     }
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 651dd5f578234..60f3c1b4d9d5b 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -10531,7 +10531,7 @@ static void getTripleBasedSPIRVTransOpts(Compilation &C,
   std::string DefaultExtArg =
       ",+SPV_EXT_shader_atomic_float_add,+SPV_EXT_shader_atomic_float_min_max"
       ",+SPV_KHR_no_integer_wrap_decoration,+SPV_KHR_float_controls"
-      ",+SPV_KHR_expect_assume,+SPV_KHR_linkonce_odr";
+      ",+SPV_KHR_expect_assume,+SPV_KHR_linkonce_odr,+SPV_KHR_bit_instructions";
   std::string INTELExtArg =
       ",+SPV_INTEL_subgroups,+SPV_INTEL_media_block_io"
       ",+SPV_INTEL_device_side_avc_motion_estimation"
@@ -10553,7 +10553,6 @@ static void getTripleBasedSPIRVTransOpts(Compilation &C,
       ",+SPV_INTEL_fpga_argument_interfaces"
       ",+SPV_INTEL_fpga_invocation_pipelining_attributes"
       ",+SPV_INTEL_fpga_latency_control"
-      ",+SPV_INTEL_task_sequence"
       ",+SPV_KHR_shader_clock"
       ",+SPV_INTEL_bindless_images";
   ExtArg = ExtArg + DefaultExtArg + INTELExtArg;
diff --git a/clang/test/CodeGenSYCL/intel-task-sequence.cpp b/clang/test/CodeGenSYCL/intel-task-sequence.cpp
deleted file mode 100644
index 6dc5ece84ac53..0000000000000
--- a/clang/test/CodeGenSYCL/intel-task-sequence.cpp
+++ /dev/null
@@ -1,11 +0,0 @@
-// RUN: %clang_cc1 -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
-
-// Test that SPIR-V codegen generates the expected LLVM struct name for the
-// TaskSequenceINTEL type
-
-namespace __spv {
-  struct __spirv_TaskSequenceINTEL;
-} // namespace __spv
-
-// CHECK: @_Z4func{{.*}}(target("spirv.TaskSequenceINTEL")
-void func(__spv::__spirv_TaskSequenceINTEL *task_seq) {}
\ No newline at end of file
diff --git a/clang/test/Driver/sycl-spirv-ext.c b/clang/test/Driver/sycl-spirv-ext.c
index eb4d24197b1af..b0e39f67e6132 100644
--- a/clang/test/Driver/sycl-spirv-ext.c
+++ b/clang/test/Driver/sycl-spirv-ext.c
@@ -47,7 +47,6 @@
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_fpga_argument_interfaces
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_fpga_invocation_pipelining_attributes
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_fpga_latency_control
-// CHECK-DEFAULT-SAME:,+SPV_INTEL_task_sequence
 // CHECK-DEFAULT-SAME:,+SPV_KHR_shader_clock
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_bindless_images
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_token_type
@@ -82,7 +81,6 @@
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_fpga_buffer_location
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_fpga_argument_interfaces
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_fpga_latency_control
-// CHECK-FPGA-HW-SAME:,+SPV_INTEL_task_sequence
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_usm_storage_classes
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_runtime_aligned
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_fpga_cluster_attributes,+SPV_INTEL_loop_fuse
@@ -114,7 +112,6 @@
 // CHECK-CPU-SAME:,+SPV_INTEL_fpga_argument_interfaces
 // CHECK-CPU-SAME:,+SPV_INTEL_fpga_invocation_pipelining_attributes
 // CHECK-CPU-SAME:,+SPV_INTEL_fpga_latency_control
-// CHECK-CPU-SAME:,+SPV_INTEL_task_sequence
 // CHECK-CPU-SAME:,+SPV_INTEL_token_type
 // CHECK-CPU-SAME:,+SPV_INTEL_bfloat16_conversion
 // CHECK-CPU-SAME:,+SPV_INTEL_joint_matrix
diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td
index 4210a0d150abf..d0df9098ab32b 100644
--- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td
+++ b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td
@@ -77,7 +77,6 @@ def AspectExt_intel_matrix : Aspect<"ext_intel_matrix">;
 def AspectExt_oneapi_is_composite : Aspect<"ext_oneapi_is_composite">;
 def AspectExt_oneapi_is_component : Aspect<"ext_oneapi_is_component">;
 def AspectExt_oneapi_graph : Aspect<"ext_oneapi_graph">;
-def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">;
 def AspectExt_oneapi_limited_graph : Aspect<"ext_oneapi_limited_graph">;
 def AspectExt_oneapi_private_alloca : Aspect<"ext_oneapi_private_alloca">;
 def AspectExt_oneapi_queue_profiling_tag : Aspect<"ext_oneapi_queue_profiling_tag">;
@@ -138,8 +137,8 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_oneapi_bindless_sampled_image_fetch_3d_usm, AspectExt_oneapi_bindless_sampled_image_fetch_3d,
     AspectExt_intel_esimd,
     AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group,
-    AspectExt_oneapi_tangle_group, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component,
-    AspectExt_oneapi_graph, AspectExt_intel_fpga_task_sequence, AspectExt_oneapi_limited_graph,
+    AspectExt_oneapi_tangle_group, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component, 
+    AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph,
     AspectExt_oneapi_private_alloca, AspectExt_oneapi_queue_profiling_tag, AspectExt_oneapi_virtual_mem, AspectExt_oneapi_cuda_cluster_group],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
index 36adf1e52ff56..63065e76faf6d 100644
--- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
+++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -416,7 +416,6 @@ attributeToExecModeMetadata(const Attribute &Attr, Function &F) {
 
   if (AttrKindStr == "sycl-streaming-interface") {
     // generate either:
-    //   !ip_interface !N
     //   !N = !{!"streaming"} or
     //   !N = !{!"streaming", !"stall_free_return"}
     SmallVector<Metadata *, 2> MD;
@@ -429,7 +428,6 @@ attributeToExecModeMetadata(const Attribute &Attr, Function &F) {
 
   if (AttrKindStr == "sycl-register-map-interface") {
     // generate either:
-    //   !ip_interface !N
     //   !N = !{!"csr"} or
     //   !N = !{!"csr", !"wait_for_done_write"}
     SmallVector<Metadata *, 2> MD;
@@ -440,20 +438,6 @@ attributeToExecModeMetadata(const Attribute &Attr, Function &F) {
                                             MDNode::get(Ctx, MD));
   }
 
-  if (AttrKindStr == "sycl-fpga-cluster") {
-    // generate either:
-    //   !stall_free !N
-    //   !N = !{i32 1} or
-    //   !stall_enable !N
-    //   !N = !{i32 1}
-    std::string ClusterType =
-        getAttributeAsInteger<uint32_t>(Attr) ? "stall_enable" : "stall_free";
-    Metadata *ClusterMDArgs[] = {
-        ConstantAsMetadata::get(ConstantInt::get(Type::getInt32Ty(Ctx), 1))};
-    return std::pair<std::string, MDNode *>(ClusterType,
-                                            MDNode::get(Ctx, ClusterMDArgs));
-  }
-
   if ((AttrKindStr == SYCL_REGISTER_ALLOC_MODE_ATTR ||
        AttrKindStr == SYCL_GRF_SIZE_ATTR) &&
       !llvm::esimd::isESIMD(F)) {
diff --git a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/fpga-cluster.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/fpga-cluster.ll
deleted file mode 100644
index cad7c4398e929..0000000000000
--- a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/fpga-cluster.ll
+++ /dev/null
@@ -1,33 +0,0 @@
-; Check conversion of sycl-fpga-cluster attribute
-; RUN: opt -passes="compile-time-properties" %s -S -o - | FileCheck %s --check-prefix CHECK-IR
-
-; CHECK-IR-DAG: @stallFree() #0 {{.*}}!stall_free [[MD_TRUE:![0-9]+]] {
-; Function Attrs: convergent norecurse
-define weak_odr dso_local spir_kernel void @stallFree() #0 {
-entry:
-  ret void
-}
-
-; CHECK-IR-DAG: @stallEnable() #1 {{.*}}!stall_enable [[MD_TRUE:![0-9]+]] {
-; Function Attrs: convergent norecurse
-define weak_odr dso_local spir_kernel void @stallEnable() #1 {
-entry:
-  ret void
-}
-
-attributes #0 = { convergent norecurse "frame-pointer"="all" "sycl-fpga-cluster"="0" }
-attributes #1 = { convergent norecurse "frame-pointer"="all" "sycl-fpga-cluster"="1" }
-
-!opencl.spir.version = !{!0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0}
-!spirv.Source = !{!1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1}
-!llvm.ident = !{!2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2}
-!llvm.module.flags = !{!3, !4}
-
-!0 = !{i32 1, i32 2}
-!1 = !{i32 4, i32 100000}
-!2 = !{!"clang version 13.0.0 (https://github.com/intel/llvm)"}
-!3 = !{i32 1, !"wchar_size", i32 4}
-!4 = !{i32 7, !"frame-pointer", i32 2}
-
-; Confirm the decorations for the functions
-; CHECK-IR-DAG: [[MD_TRUE]] = !{i32 1}
diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp
index f4504a72b4d36..695db502dd818 100644
--- a/sycl/include/CL/__spirv/spirv_ops.hpp
+++ b/sycl/include/CL/__spirv/spirv_ops.hpp
@@ -1300,25 +1300,6 @@ extern __DPCPP_SYCL_EXTERNAL
     std::enable_if_t<std::is_integral_v<to> && std::is_unsigned_v<to>, to>
     __spirv_ConvertPtrToU(from val) noexcept;
 
-template <typename RetT, typename... ArgsT>
-extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_TaskSequenceINTEL *
-__spirv_TaskSequenceCreateINTEL(RetT (*f)(ArgsT...), int Pipelined = -1,
-                                int ClusterMode = -1,
-                                unsigned int ResponseCapacity = 0,
-                                unsigned int InvocationCapacity = 0) noexcept;
-
-template <typename... ArgsT>
-extern __DPCPP_SYCL_EXTERNAL void
-__spirv_TaskSequenceAsyncINTEL(__spv::__spirv_TaskSequenceINTEL *TaskSequence,
-                               ArgsT... Args) noexcept;
-
-template <typename RetT>
-extern __DPCPP_SYCL_EXTERNAL RetT __spirv_TaskSequenceGetINTEL(
-    __spv::__spirv_TaskSequenceINTEL *TaskSequence) noexcept;
-
-extern __DPCPP_SYCL_EXTERNAL void __spirv_TaskSequenceReleaseINTEL(
-    __spv::__spirv_TaskSequenceINTEL *TaskSequence) noexcept;
-
 #else  // if !__SYCL_DEVICE_ONLY__
 
 template <typename dataT>
diff --git a/sycl/include/CL/__spirv/spirv_types.hpp b/sycl/include/CL/__spirv/spirv_types.hpp
index 013c2e1acc564..3246e41f41190 100644
--- a/sycl/include/CL/__spirv/spirv_types.hpp
+++ b/sycl/include/CL/__spirv/spirv_types.hpp
@@ -123,8 +123,6 @@ template <typename T, std::size_t R, std::size_t C, MatrixLayout L,
           MatrixUse U = MatrixUse::MatrixA>
 struct __spirv_JointMatrixINTEL;
 
-struct __spirv_TaskSequenceINTEL;
-
 } // namespace __spv
 
 #ifdef __SYCL_DEVICE_ONLY__
diff --git a/sycl/include/sycl/device_aspect_macros.hpp b/sycl/include/sycl/device_aspect_macros.hpp
index 55ed162863079..ffb93f9da491f 100644
--- a/sycl/include/sycl/device_aspect_macros.hpp
+++ b/sycl/include/sycl/device_aspect_macros.hpp
@@ -313,11 +313,6 @@
 #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_graph__ 0
 #endif
 
-#ifndef __SYCL_ALL_DEVICES_HAVE_ext_intel_fpga_task_sequence__
-// __SYCL_ASPECT(ext_intel_fpga_task_sequence, 62)
-#define __SYCL_ALL_DEVICES_HAVE_ext_intel_fpga_task_sequence__ 0
-#endif
-
 #ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_limited_graph__
 // __SYCL_ASPECT(ext_oneapi_limited_graph, 63)
 #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_limited_graph__ 0
@@ -691,11 +686,6 @@
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_graph__ 0
 #endif
 
-#ifndef __SYCL_ANY_DEVICE_HAS_ext_intel_fpga_task_sequence__
-// __SYCL_ASPECT(ext_intel_fpga_task_sequence__, 62)
-#define __SYCL_ANY_DEVICE_HAS_ext_intel_fpga_task_sequence__ 0
-#endif
-
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_limited_graph__
 // __SYCL_ASPECT(ext_oneapi_limited_graph, 63)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_limited_graph__ 0
diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
index 957de80500a5e..0db22e9407a5e 100644
--- a/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
@@ -1,5 +1,5 @@
-//===--------------------- fpga_kernel_properties.hpp ---------------------===//
-// SYCL properties associated with FPGA kernel properties
+//==----- fpga_kernel_properties.hpp - SYCL properties associated with FPGA
+// kernel properties ---==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -12,15 +12,11 @@
 #include <sycl/ext/oneapi/properties/property.hpp>
 #include <sycl/ext/oneapi/properties/property_value.hpp>
 
-#include <cstdint>     // for uint16_t
-#include <type_traits> // for true_type
-
 namespace sycl {
 inline namespace _V1 {
 namespace ext::intel::experimental {
 
 template <typename T, typename PropertyListT> class fpga_kernel_attribute;
-template <auto &f, typename PropertyListT> class task_sequence;
 
 enum class streaming_interface_options_enum : std::uint16_t {
   accept_downstream_stall,
@@ -32,11 +28,6 @@ enum class register_map_interface_options_enum : std::uint16_t {
   wait_for_done_write,
 };
 
-enum class fpga_cluster_options_enum : std::uint16_t {
-  stall_free,
-  stall_enable
-};
-
 struct streaming_interface_key
     : oneapi::experimental::detail::compile_time_property_key<
           oneapi::experimental::detail::PropKind::StreamingInterface> {
@@ -63,15 +54,6 @@ struct pipelined_key : oneapi::experimental::detail::compile_time_property_key<
       std::integral_constant<int, pipeline_directive_or_initiation_interval>>;
 };
 
-struct fpga_cluster_key
-    : oneapi::experimental::detail::compile_time_property_key<
-          oneapi::experimental::detail::PropKind::FPGACluster> {
-  template <fpga_cluster_options_enum option>
-  using value_t = ext::oneapi::experimental::property_value<
-      fpga_cluster_key,
-      std::integral_constant<fpga_cluster_options_enum, option>>;
-};
-
 template <streaming_interface_options_enum option =
               streaming_interface_options_enum::accept_downstream_stall>
 inline constexpr streaming_interface_key::value_t<option> streaming_interface;
@@ -102,18 +84,6 @@ inline constexpr pipelined_key::value_t<
     pipeline_directive_or_initiation_interval>
     pipelined;
 
-template <fpga_cluster_options_enum option =
-              fpga_cluster_options_enum::stall_free>
-inline constexpr fpga_cluster_key::value_t<option> fpga_cluster;
-
-inline constexpr fpga_cluster_key::value_t<
-    fpga_cluster_options_enum::stall_free>
-    stall_free_clusters;
-
-inline constexpr fpga_cluster_key::value_t<
-    fpga_cluster_options_enum::stall_enable>
-    stall_enable_clusters;
-
 } // namespace ext::intel::experimental
 
 namespace ext::oneapi::experimental {
@@ -133,22 +103,6 @@ struct is_property_key_of<
     intel::experimental::fpga_kernel_attribute<T, PropertyListT>>
     : std::true_type {};
 
-template <typename T, typename PropertyListT>
-struct is_property_key_of<
-    intel::experimental::fpga_cluster_key,
-    intel::experimental::fpga_kernel_attribute<T, PropertyListT>>
-    : std::true_type {};
-
-template <auto &f, typename PropertyListT>
-struct is_property_key_of<intel::experimental::pipelined_key,
-                          intel::experimental::task_sequence<f, PropertyListT>>
-    : std::true_type {};
-
-template <auto &f, typename PropertyListT>
-struct is_property_key_of<intel::experimental::fpga_cluster_key,
-                          intel::experimental::task_sequence<f, PropertyListT>>
-    : std::true_type {};
-
 namespace detail {
 template <intel::experimental::streaming_interface_options_enum Stall_Free>
 struct PropertyMetaInfo<
@@ -170,14 +124,6 @@ struct PropertyMetaInfo<intel::experimental::pipelined_key::value_t<Value>> {
   static constexpr int value = Value;
 };
 
-template <intel::experimental::fpga_cluster_options_enum ClusterType>
-struct PropertyMetaInfo<
-    intel::experimental::fpga_cluster_key::value_t<ClusterType>> {
-  static constexpr const char *name = "sycl-fpga-cluster";
-  static constexpr intel::experimental::fpga_cluster_options_enum value =
-      ClusterType;
-};
-
 } // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1
diff --git a/sycl/include/sycl/ext/intel/experimental/task_sequence.hpp b/sycl/include/sycl/ext/intel/experimental/task_sequence.hpp
deleted file mode 100644
index e4d39b1f73c68..0000000000000
--- a/sycl/include/sycl/ext/intel/experimental/task_sequence.hpp
+++ /dev/null
@@ -1,133 +0,0 @@
-//==--------------------------- task_sequence.hpp --------------------------==//
-//
-// 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/__spirv/spirv_ops.hpp>
-#include <CL/__spirv/spirv_types.hpp>
-#include <sycl/aspects.hpp>
-#include <sycl/exception.hpp>
-#include <sycl/ext/intel/experimental/fpga_kernel_properties.hpp>
-#include <sycl/ext/intel/experimental/task_sequence_properties.hpp>
-#include <sycl/ext/oneapi/properties/properties.hpp>
-
-#include <type_traits>
-
-namespace sycl {
-inline namespace _V1 {
-namespace ext::intel::experimental {
-
-template <auto &f> struct task_sequence_checker : std::false_type {};
-
-template <typename ReturnT, typename... ArgsT, ReturnT (&f)(ArgsT...)>
-struct task_sequence_checker<f> : std::true_type {};
-
-template <auto &f,
-          typename PropertyListT = oneapi::experimental::empty_properties_t>
-class task_sequence {
-
-  static_assert(task_sequence_checker<f>::value);
-
-  static_assert(oneapi::experimental::is_property_list<PropertyListT>::value,
-                "Property list is invalid.");
-};
-
-template <typename ReturnT, typename... ArgsT, ReturnT (&f)(ArgsT...),
-          typename... Props>
-#if defined(__SYCL_DEVICE_ONLY__)
-class
-    [[__sycl_detail__::__uses_aspects__(aspect::ext_intel_fpga_task_sequence)]] task_sequence<
-#else
-class task_sequence<
-#endif
-        f, oneapi::experimental::detail::properties_t<Props...>> {
-  using property_list_t = oneapi::experimental::detail::properties_t<Props...>;
-
-public:
-  task_sequence(const task_sequence &) = delete;
-  task_sequence &operator=(const task_sequence &) = delete;
-  task_sequence(task_sequence &&) = delete;
-  task_sequence &operator=(task_sequence &&) = delete;
-
-  task_sequence() {
-#if defined(__SYCL_DEVICE_ONLY__)
-    taskSequence = __spirv_TaskSequenceCreateINTEL(
-        &f, pipelined, fpga_cluster, response_capacity, invocation_capacity);
-#else
-    throw exception{make_error_code(errc::feature_not_supported),
-                    "task_sequence is not supported on the host"};
-#endif
-  }
-
-  void async([[maybe_unused]] ArgsT... Args) {
-#if defined(__SYCL_DEVICE_ONLY__)
-    ++outstanding;
-    __spirv_TaskSequenceAsyncINTEL(taskSequence, Args...);
-#else
-    throw exception{make_error_code(errc::feature_not_supported),
-                    "task_sequence is not supported on the host"};
-#endif
-  }
-
-  ReturnT get() {
-#if defined(__SYCL_DEVICE_ONLY__)
-    --outstanding;
-    return __spirv_TaskSequenceGetINTEL<ReturnT>(taskSequence);
-#else
-    throw exception{make_error_code(errc::feature_not_supported),
-                    "task_sequence is not supported on the host"};
-#endif
-  }
-
-  ~task_sequence() {
-#if defined(__SYCL_DEVICE_ONLY__)
-    if constexpr (!has_property<balanced_key>()) {
-      while (outstanding > 0)
-        get();
-    }
-    __spirv_TaskSequenceReleaseINTEL(taskSequence);
-#endif
-    // Destructor shouldn't throw exception.
-  }
-
-  template <typename propertyT> static constexpr bool has_property() {
-    return property_list_t::template has_property<propertyT>();
-  }
-
-  template <typename propertyT> static constexpr auto get_property() {
-    return property_list_t::template get_property<propertyT>();
-  }
-
-private:
-#if defined(__SYCL_DEVICE_ONLY__)
-  unsigned outstanding = 0;
-  __spv::__spirv_TaskSequenceINTEL *taskSequence;
-#endif
-  static constexpr int32_t pipelined =
-      oneapi::experimental::detail::ValueOrDefault<
-          property_list_t, pipelined_key>::template get<int32_t>(-1);
-  static constexpr int32_t fpga_cluster =
-      has_property<fpga_cluster_key>()
-          ? static_cast<
-                typename std::underlying_type<fpga_cluster_options_enum>::type>(
-                oneapi::experimental::detail::ValueOrDefault<property_list_t,
-                                                             fpga_cluster_key>::
-                    template get<fpga_cluster_options_enum>(
-                        fpga_cluster_options_enum::stall_free))
-          : -1;
-  static constexpr uint32_t response_capacity =
-      oneapi::experimental::detail::ValueOrDefault<
-          property_list_t, response_capacity_key>::template get<uint32_t>(0);
-  static constexpr uint32_t invocation_capacity =
-      oneapi::experimental::detail::ValueOrDefault<
-          property_list_t, invocation_capacity_key>::template get<uint32_t>(0);
-};
-
-} // namespace ext::intel::experimental
-} // namespace _V1
-} // namespace sycl
diff --git a/sycl/include/sycl/ext/intel/experimental/task_sequence_properties.hpp b/sycl/include/sycl/ext/intel/experimental/task_sequence_properties.hpp
deleted file mode 100644
index f2faecc0926f3..0000000000000
--- a/sycl/include/sycl/ext/intel/experimental/task_sequence_properties.hpp
+++ /dev/null
@@ -1,71 +0,0 @@
-//==-- task_sequence_properties.hpp - Specific properties of task_sequence -==//
-//
-// 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 <sycl/ext/oneapi/properties/property.hpp>       // for PropKind
-#include <sycl/ext/oneapi/properties/property_value.hpp> // for property_value
-
-#include <type_traits> // for true_type
-
-namespace sycl {
-inline namespace _V1 {
-namespace ext::intel::experimental {
-// Forward declare a class that these properties can be applied to
-template <auto &f, typename PropertyListT> class task_sequence;
-
-// Make sure that we are using the right namespace
-template <typename PropertyT, typename... Ts>
-using property_value =
-    sycl::ext::oneapi::experimental::property_value<PropertyT, Ts...>;
-
-struct balanced_key : oneapi::experimental::detail::compile_time_property_key<
-          oneapi::experimental::detail::PropKind::Balanced> {
-  using value_t = property_value<balanced_key>;
-};
-
-struct invocation_capacity_key
-    : oneapi::experimental::detail::compile_time_property_key<
-          oneapi::experimental::detail::PropKind::InvocationCapacity> {
-  template <unsigned int Size>
-  using value_t = property_value<invocation_capacity_key,
-                                 std::integral_constant<unsigned int, Size>>;
-};
-
-struct response_capacity_key
-    : oneapi::experimental::detail::compile_time_property_key<
-          oneapi::experimental::detail::PropKind::ResponseCapacity> {
-  template <unsigned int Size>
-  using value_t = property_value<response_capacity_key,
-                                 std::integral_constant<unsigned int, Size>>;
-};
-
-inline constexpr balanced_key::value_t balanced;
-template <unsigned int Size>
-inline constexpr invocation_capacity_key::value_t<Size> invocation_capacity;
-template <unsigned int Size>
-inline constexpr response_capacity_key::value_t<Size> response_capacity;
-
-} // namespace ext::intel::experimental
-
-namespace ext::oneapi::experimental {
-template <auto &f, typename PropertyListT>
-struct is_property_key_of<intel::experimental::balanced_key,
-                          intel::experimental::task_sequence<f, PropertyListT>>
-    : std::true_type {};
-template <auto &f, typename PropertyListT>
-struct is_property_key_of<intel::experimental::invocation_capacity_key,
-                          intel::experimental::task_sequence<f, PropertyListT>>
-    : std::true_type {};
-template <auto &f, typename PropertyListT>
-struct is_property_key_of<intel::experimental::response_capacity_key,
-                          intel::experimental::task_sequence<f, PropertyListT>>
-    : std::true_type {};
-} // namespace ext::oneapi::experimental
-} // namespace _V1
-} // namespace sycl
diff --git a/sycl/include/sycl/ext/intel/fpga_extensions.hpp b/sycl/include/sycl/ext/intel/fpga_extensions.hpp
index 5f82d0bde6b9f..fb2d814c67df5 100644
--- a/sycl/include/sycl/ext/intel/fpga_extensions.hpp
+++ b/sycl/include/sycl/ext/intel/fpga_extensions.hpp
@@ -16,8 +16,6 @@
 #include <sycl/ext/intel/experimental/fpga_mem/fpga_mem.hpp>
 #include <sycl/ext/intel/experimental/fpga_mem/properties.hpp>
 #include <sycl/ext/intel/experimental/pipes.hpp>
-#include <sycl/ext/intel/experimental/task_sequence.hpp>
-#include <sycl/ext/intel/experimental/task_sequence_properties.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>
diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp
index e225928c4cd68..f8c365240a706 100644
--- a/sycl/include/sycl/ext/oneapi/properties/property.hpp
+++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp
@@ -192,28 +192,24 @@ enum PropKind : uint32_t {
   BuildOptions = 51,
   BuildLog = 52,
   FloatingPointControls = 53,
-  FPGACluster = 54,
-  Balanced = 55,
-  InvocationCapacity = 56,
-  ResponseCapacity = 57,
-  DataPlacement = 58,
-  ContiguousMemory = 59,
-  FullGroup = 60,
-  Naive = 61,
-  WorkGroupProgress = 62,
-  SubGroupProgress = 63,
-  WorkItemProgress = 64,
-  NDRangeKernel = 65,
-  SingleTaskKernel = 66,
-  IndirectlyCallable = 67,
-  CallsIndirectly = 68,
-  InputDataPlacement = 69,
-  OutputDataPlacement = 70,
-  IncludeFiles = 71,
-  RegisteredKernelNames = 72,
-  ClusterLaunch = 73,
+  DataPlacement = 54,
+  ContiguousMemory = 55,
+  FullGroup = 56,
+  Naive = 57,
+  WorkGroupProgress = 58,
+  SubGroupProgress = 59,
+  WorkItemProgress = 60,
+  NDRangeKernel = 61,
+  SingleTaskKernel = 62,
+  IndirectlyCallable = 63,
+  CallsIndirectly = 64,
+  InputDataPlacement = 65,
+  OutputDataPlacement = 66,
+  IncludeFiles = 67,
+  RegisteredKernelNames = 68,
+  ClusterLaunch = 69,
   // PropKindSize must always be the last value.
-  PropKindSize = 74,
+  PropKindSize = 70,
 };
 
 struct property_key_base_tag {};
diff --git a/sycl/include/sycl/info/aspects.def b/sycl/include/sycl/info/aspects.def
index f9a14e40dcd3e..824ab3ae632f8 100644
--- a/sycl/include/sycl/info/aspects.def
+++ b/sycl/include/sycl/info/aspects.def
@@ -55,17 +55,16 @@ __SYCL_ASPECT(ext_intel_matrix, 58)
 __SYCL_ASPECT(ext_oneapi_is_composite, 59)
 __SYCL_ASPECT(ext_oneapi_is_component, 60)
 __SYCL_ASPECT(ext_oneapi_graph, 61)
-__SYCL_ASPECT(ext_intel_fpga_task_sequence, 62)
-__SYCL_ASPECT(ext_oneapi_limited_graph, 63)
-__SYCL_ASPECT(ext_oneapi_private_alloca, 64)
-__SYCL_ASPECT(ext_oneapi_cubemap, 65)
-__SYCL_ASPECT(ext_oneapi_cubemap_seamless_filtering, 66)
-__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_1d_usm, 67)
-__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_1d, 68)
-__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_2d_usm, 69)
-__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_2d, 70)
-__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d_usm, 71)
-__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d, 72)
-__SYCL_ASPECT(ext_oneapi_queue_profiling_tag, 73)
-__SYCL_ASPECT(ext_oneapi_virtual_mem, 74)
-__SYCL_ASPECT(ext_oneapi_cuda_cluster_group, 75)
+__SYCL_ASPECT(ext_oneapi_limited_graph, 62)
+__SYCL_ASPECT(ext_oneapi_private_alloca, 63)
+__SYCL_ASPECT(ext_oneapi_cubemap, 64)
+__SYCL_ASPECT(ext_oneapi_cubemap_seamless_filtering, 65)
+__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_1d_usm, 66)
+__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_1d, 67)
+__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_2d_usm, 68)
+__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_2d, 69)
+__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d_usm, 70)
+__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d, 71)
+__SYCL_ASPECT(ext_oneapi_queue_profiling_tag, 72)
+__SYCL_ASPECT(ext_oneapi_virtual_mem, 73)
+__SYCL_ASPECT(ext_oneapi_cuda_cluster_group, 74)
diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp
index 1bdb816efcb39..fa5041ff0ece2 100644
--- a/sycl/include/sycl/sycl.hpp
+++ b/sycl/include/sycl/sycl.hpp
@@ -70,8 +70,6 @@
 #include <sycl/ext/intel/experimental/fpga_mem/properties.hpp>
 #include <sycl/ext/intel/experimental/pipe_properties.hpp>
 #include <sycl/ext/intel/experimental/pipes.hpp>
-#include <sycl/ext/intel/experimental/task_sequence.hpp>
-#include <sycl/ext/intel/experimental/task_sequence_properties.hpp>
 #include <sycl/ext/intel/experimental/usm_properties.hpp>
 #include <sycl/ext/intel/usm_pointers.hpp>
 #include <sycl/ext/oneapi/bindless_images.hpp>
diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp
index 00a9425049aa4..9f37762dd69f9 100644
--- a/sycl/source/detail/device_impl.cpp
+++ b/sycl/source/detail/device_impl.cpp
@@ -699,9 +699,6 @@ bool device_impl::has(aspect Aspect) const {
 
     return SupportsCommandBuffers;
   }
-  case aspect::ext_intel_fpga_task_sequence: {
-    return is_accelerator();
-  }
   case aspect::ext_oneapi_private_alloca: {
     // Extension only supported on SPIR-V targets.
     backend be = getBackend();
diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in
index 55ca1bf52d5eb..61e90cd72b959 100644
--- a/sycl/source/feature_test.hpp.in
+++ b/sycl/source/feature_test.hpp.in
@@ -103,7 +103,6 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS 1
 #define SYCL_EXT_ONEAPI_IN_ORDER_QUEUE_EVENTS 1
 #define SYCL_EXT_INTEL_MATRIX 1
-#define SYCL_EXT_INTEL_FPGA_TASK_SEQUENCE 1
 #define SYCL_EXT_ONEAPI_PRIVATE_ALLOCA 1
 #define SYCL_EXT_ONEAPI_FORWARD_PROGRESS 1
 #define SYCL_EXT_ONEAPI_FREE_FUNCTION_KERNELS 1
diff --git a/sycl/test-e2e/TaskSequence/common.hpp b/sycl/test-e2e/TaskSequence/common.hpp
deleted file mode 100644
index 2b781141c0fff..0000000000000
--- a/sycl/test-e2e/TaskSequence/common.hpp
+++ /dev/null
@@ -1,13 +0,0 @@
-#pragma once
-
-#include <sycl/builtins.hpp>
-#include <sycl/detail/core.hpp>
-#include <sycl/ext/intel/experimental/pipe_properties.hpp>
-#include <sycl/ext/intel/experimental/pipes.hpp>
-#include <sycl/ext/intel/experimental/task_sequence.hpp>
-#include <sycl/ext/intel/experimental/task_sequence_properties.hpp>
-#include <sycl/pipes.hpp>
-
-using namespace sycl;
-using namespace sycl::ext::oneapi::experimental;
-using namespace sycl::ext::intel::experimental;
diff --git a/sycl/test-e2e/TaskSequence/concurrent-loops.cpp b/sycl/test-e2e/TaskSequence/concurrent-loops.cpp
deleted file mode 100644
index 8cfb6335182cc..0000000000000
--- a/sycl/test-e2e/TaskSequence/concurrent-loops.cpp
+++ /dev/null
@@ -1,94 +0,0 @@
-//==-------------- concurrent-loops.cpp - DPC++ task_sequence --------------==//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-
-// REQUIRES: aspect-ext_intel_fpga_task_sequence
-// RUN: %clangxx -fsycl -fintelfpga %s -o %t.out
-// RUN: %{run} %t.out
-
-#include "common.hpp"
-
-constexpr int kPrecision = 20;
-constexpr double kX = 0.9;
-constexpr double kA = 2.5;
-constexpr double kB = 3.7;
-
-using intertask_pipe = ext::intel::pipe<class p, double>;
-
-template <typename OutPipe> double task_a(double x, double a, int precision) {
-  double res = 1;
-  double term = 1;
-  double multiplier = a;
-
-  // compute the non-constant common term of the Maclaurin series of exp(x)
-  // feed reuseable result through intertask pipe
-  // compute remaining terms of Maclaurin series of exp(ax) and accummulate
-  for (int iter = 1; iter <= precision; ++iter) {
-    term *= x / iter;
-    // reuse the shared value for both exponential series
-    OutPipe::write(term);
-    res += term * multiplier;
-    multiplier *= a;
-  }
-
-  return res;
-}
-
-template <typename InPipe> double task_b(double b, int precision) {
-  double res = 1;
-  double multiplier = b;
-
-  // reuse the non-constant common term from task_a
-  for (int iter = 1; iter <= precision; ++iter) {
-    // compute terms of Maclaurin series of exp(bx) and accumulate
-    res += InPipe::read() * multiplier;
-    multiplier *= b;
-  }
-
-  return res;
-};
-
-class ConcurrentLoop;
-
-int main() {
-
-  double golden = std::exp(kA * kX) / std::exp(kB * kX);
-
-  // initialize inputs
-  std::vector<double> in = {kX, kA, kB};
-  std::vector<double> out = {0.0};
-
-  {
-    sycl::queue q;
-
-    buffer in_buf(in);
-    buffer out_buf(out);
-
-    q.submit([&](handler &h) {
-      accessor in_acc(in_buf, h, read_only);
-      accessor out_acc(out_buf, h, write_only);
-      h.single_task<ConcurrentLoop>([=]() {
-        task_sequence<task_a<intertask_pipe>> ts_a;
-        task_sequence<task_b<intertask_pipe>> ts_b;
-
-        double x = in_acc[0];
-        double a = in_acc[1];
-        double b = in_acc[2];
-
-        // approximate e^(ax)/e^(bx) using Maclaurin series
-        ts_a.async(x, a, kPrecision);
-        ts_b.async(b, kPrecision);
-
-        out_acc[0] = ts_a.get() / ts_b.get();
-      });
-    });
-    q.wait();
-  }
-
-  assert(std::abs(out[0] - golden) < 0.001);
-  return 0;
-}
diff --git a/sycl/test-e2e/TaskSequence/in-order-async-get.cpp b/sycl/test-e2e/TaskSequence/in-order-async-get.cpp
deleted file mode 100644
index 8e9853d5fa15a..0000000000000
--- a/sycl/test-e2e/TaskSequence/in-order-async-get.cpp
+++ /dev/null
@@ -1,69 +0,0 @@
-//==------------- in-order-async-get.cpp - DPC++ task_sequence -------------==//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-
-// REQUIRES: aspect-ext_intel_fpga_task_sequence
-// RUN: %clangxx -fsycl -fintelfpga %s -o %t.out
-// RUN: %{run} %t.out
-
-#include "common.hpp"
-
-constexpr int kNumInputs = 16;
-
-class InOrderTasks;
-
-// task that computes the sigmoid function
-double sigmoid(double in) {
-  double Ex = sycl::exp(in);
-  return (Ex / (Ex + 1));
-}
-
-int main() {
-  std::vector<double> values(kNumInputs);
-  std::vector<double> golden(kNumInputs);
-  std::vector<double> results(kNumInputs, 0);
-
-  // test data
-  for (int i = 0; i < kNumInputs; ++i) {
-    values[i] = (static_cast<double>(-kNumInputs / 2.0 + i) / 1000);
-    golden[i] = 1 / (1 + std::exp(-values[i]));
-  }
-
-  {
-    queue q;
-
-    buffer values_buf(values);
-    buffer results_buf(values);
-
-    q.submit([&](handler &h) {
-      accessor values_acc(values_buf, h, read_only);
-      accessor results_acc(results_buf, h, write_only);
-      h.single_task<InOrderTasks>([=]() {
-        task_sequence<sigmoid,
-                      decltype(properties{invocation_capacity<kNumInputs>})>
-            ts;
-
-        for (int i = 0; i < kNumInputs; ++i) {
-          ts.async(values_acc[i]);
-        }
-
-        // multiple async calls before first get call
-        for (int i = 0; i < kNumInputs; ++i) {
-          results_acc[i] = ts.get();
-        }
-      });
-    });
-    q.wait();
-  }
-
-  // verification
-  for (int i = 0; i < kNumInputs; ++i) {
-    assert(std::abs(results[i] - golden[i]) < 0.001);
-  }
-
-  return 0;
-}
diff --git a/sycl/test-e2e/TaskSequence/mult-and-add.cpp b/sycl/test-e2e/TaskSequence/mult-and-add.cpp
deleted file mode 100644
index 3b8119a44223b..0000000000000
--- a/sycl/test-e2e/TaskSequence/mult-and-add.cpp
+++ /dev/null
@@ -1,41 +0,0 @@
-//==---------------- mult-and-add.cpp - DPC++ task_sequence ----------------==//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-
-// REQUIRES: aspect-ext_intel_fpga_task_sequence
-// RUN: %clangxx -fsycl -fintelfpga %s -o %t.out
-// RUN: %{run} %t.out
-
-#include "common.hpp"
-
-int mult(int a, int b) { return a * b; }
-
-int mult_and_add(int a1, int b1, int a2, int b2) {
-  task_sequence<mult, decltype(properties{invocation_capacity<1>,
-                                          response_capacity<1>})>
-      task1, task2;
-
-  task1.async(a1, b1);
-  task2.async(a2, b2);
-  return task1.get() + task2.get();
-}
-
-int main() {
-  queue myQueue;
-
-  int result = 0;
-  buffer<int, 1> res_buf(&result, range<1>(1));
-
-  myQueue.submit([&](handler &cgh) {
-    auto res_acc = res_buf.get_access<access::mode::write>(cgh);
-    cgh.single_task(
-        [=](kernel_handler kh) { res_acc[0] = mult_and_add(1, 2, 3, 4); });
-  });
-  myQueue.wait();
-
-  assert(result == (1 * 2 + 3 * 4));
-}
\ No newline at end of file
diff --git a/sycl/test-e2e/TaskSequence/multi-kernel-task-function-reuse.cpp b/sycl/test-e2e/TaskSequence/multi-kernel-task-function-reuse.cpp
deleted file mode 100644
index cb3abf85a1ead..0000000000000
--- a/sycl/test-e2e/TaskSequence/multi-kernel-task-function-reuse.cpp
+++ /dev/null
@@ -1,80 +0,0 @@
-//==------ multi-kernel-task-function-reuse.cpp - DPC++ task_sequence ------==//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-
-// REQUIRES: aspect-ext_intel_fpga_task_sequence
-// RUN: %clangxx -fsycl -fintelfpga %s -o %t.out
-// RUN: %{run} %t.out
-
-#include "common.hpp"
-
-constexpr int kCount = 1024;
-
-int vectorSum(const int *v, size_t s, size_t sz) {
-  int result = 0;
-  for (size_t i = s; i < s + sz; i++) {
-    result += v[i];
-  }
-
-  return result;
-}
-
-// Kernel identifiers
-class SequentialTask;
-class ParallelTask;
-
-int main() {
-  queue q;
-
-  // create input and golden output data
-  std::vector<int> in(kCount), out(2);
-  for (int i = 0; i < kCount; i++) {
-    in[i] = i;
-  }
-
-  int golden = vectorSum(in.data(), 0, kCount);
-
-  {
-    buffer in_buf(in);
-    buffer out_buf(out);
-
-    q.submit([&](handler &h) {
-      accessor in_acc(in_buf, h, read_only);
-      accessor out_acc(out_buf, h, write_only);
-      h.single_task<SequentialTask>([=]() {
-        task_sequence<vectorSum> whole;
-        whole.async(in_acc.get_pointer(), 0, kCount);
-        out_acc[0] = whole.get();
-      });
-    });
-
-    q.submit([&](sycl::handler &h) {
-      accessor in_acc(in_buf, h, read_only);
-      accessor out_acc(out_buf, h, write_only);
-      h.single_task<ParallelTask>([=]() {
-        task_sequence<vectorSum> firstQuarter;
-        task_sequence<vectorSum> secondQuarter;
-        task_sequence<vectorSum> thirdQuarter;
-        task_sequence<vectorSum> fourthQuarter;
-        constexpr int quarterCount = kCount / 4;
-        firstQuarter.async(in_acc.get_pointer(), 0, quarterCount);
-        secondQuarter.async(in_acc.get_pointer(), quarterCount, quarterCount);
-        thirdQuarter.async(in_acc.get_pointer(), 2 * quarterCount,
-                           quarterCount);
-        fourthQuarter.async(in_acc.get_pointer(), 3 * quarterCount,
-                            quarterCount);
-        out_acc[1] = firstQuarter.get() + secondQuarter.get() +
-                     thirdQuarter.get() + fourthQuarter.get();
-      });
-    });
-    q.wait();
-  }
-
-  assert(out[0] == golden);
-  assert(out[1] == golden);
-  return 0;
-}
diff --git a/sycl/test-e2e/TaskSequence/producer-consumer.cpp b/sycl/test-e2e/TaskSequence/producer-consumer.cpp
deleted file mode 100644
index 2b5e46ba0e386..0000000000000
--- a/sycl/test-e2e/TaskSequence/producer-consumer.cpp
+++ /dev/null
@@ -1,61 +0,0 @@
-//==-------------- producer-consumer.cpp - DPC++ task_sequence -------------==//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-
-// REQUIRES: aspect-ext_intel_fpga_task_sequence
-// RUN: %clangxx -fsycl -fintelfpga %s -o %t.out
-// RUN: %{run} %t.out
-
-#include "common.hpp"
-
-constexpr int kSize = 128;
-
-using intertask_pipe = ext::intel::pipe<class p, int>;
-
-template <typename OutPipe> void producer(int *a, int *b, int N) {
-  for (int i = 0; i < N; ++i) {
-    OutPipe::write(a[i] * b[i]);
-  }
-}
-
-template <typename InPipe> int consumer(int N) {
-  int total = 0;
-  for (int i = 0; i < N; ++i) {
-    total += (InPipe::read() + 42);
-  }
-  return total;
-}
-
-int main() {
-  queue myQueue;
-
-  int result = 0;
-  buffer<int, 1> res_buf(&result, range<1>(1));
-
-  myQueue.submit([&](handler &cgh) {
-    auto res_acc = res_buf.get_access<access::mode::write>(cgh);
-    cgh.single_task([=](kernel_handler kh) {
-      int d1[kSize], d2[kSize];
-      for (int i = 0; i < kSize; ++i)
-        d1[i] = d2[i] = i;
-      task_sequence<producer<intertask_pipe>> producerTask;
-      task_sequence<consumer<intertask_pipe>> consumerTask;
-
-      producerTask.async(d1, d2, kSize);
-      consumerTask.async(kSize);
-      res_acc[0] = consumerTask.get();
-    });
-  });
-  myQueue.wait();
-
-  // Check result:
-  int sum = 0;
-  for (int i = 0; i < kSize; ++i)
-    sum += i * i + 42;
-  assert(result == sum);
-  return 0;
-}
\ No newline at end of file
diff --git a/sycl/test-e2e/TaskSequence/struct-array-args-and-return.cpp b/sycl/test-e2e/TaskSequence/struct-array-args-and-return.cpp
deleted file mode 100644
index fe802c679fbf2..0000000000000
--- a/sycl/test-e2e/TaskSequence/struct-array-args-and-return.cpp
+++ /dev/null
@@ -1,116 +0,0 @@
-//==-------- struct-array-args-and-return.cpp - DPC++ task_sequence --------==//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-
-// REQUIRES: aspect-ext_intel_fpga_task_sequence
-// RUN: %clangxx -fsycl -fintelfpga %s -o %t.out
-// RUN: %{run} %t.out
-#include "common.hpp"
-
-class TaskSequenceTest;
-
-
-template <typename T, int COUNT> class FixedVect {
-  T data[COUNT];
-
-public:
-  T &operator[](int i) { return data[i]; }
-};
-
-struct DataStruct {
-  float val;
-  bool isValid;
-
-  DataStruct() : val(-1.0f), isValid(false) {}
-  DataStruct(float val_, bool isValid_) : val(val_), isValid(isValid_) {}
-};
-
-using DataArray = FixedVect<float, 2>;
-using intertask_pipe1 = ext::intel::pipe<class p1, float, 1>;
-using intertask_pipe2 = ext::intel::pipe<class p2, float, 1>;
-
-template <typename OutPipe> void argStruct(DataStruct data) {
-  if (data.isValid) {
-    OutPipe::write(data.val * data.val);
-  }
-}
-
-template <typename OutPipe> void argArray(DataArray data) {
-  for (int i = 0; i < 2; i++) {
-    float a = data[i];
-    OutPipe::write(a * a);
-  }
-}
-
-template <typename InPipe> DataStruct returnStruct() {
-  float a = InPipe::read();
-  DataStruct res{sycl::sqrt(a), true};
-  return res;
-}
-
-template <typename InPipe> DataArray returnArray() {
-  DataArray res;
-  for (int i = 0; i < 2; i++) {
-    float a = InPipe::read();
-    res[i] = sycl::sqrt(a);
-  }
-
-  return res;
-}
-
-int main() {
-  std::vector<DataStruct> vec_in_struct(1);
-  std::vector<DataArray> vec_in_array(1);
-  std::vector<DataStruct> res_struct(1);
-  std::vector<DataArray> res_array(1);
-
-  vec_in_struct[0] = DataStruct{5.0, true};
-
-  DataArray d;
-  d[0] = 7.0;
-  d[1] = 9.0;
-  vec_in_array[0] = d;
-
-  {
-    queue q;
-
-    buffer buffer_in_struct(vec_in_struct);
-    buffer buffer_in_array(vec_in_array);
-    buffer buffer_out_struct(res_struct);
-    buffer buffer_out_array(res_array);
-
-    q.submit([&](handler &h) {
-      accessor in_acc_struct(buffer_in_struct, h, read_only);
-      accessor in_acc_array(buffer_in_array, h, read_only);
-      accessor out_acc_struct(buffer_out_struct, h, write_only, no_init);
-      accessor out_acc_array(buffer_out_array, h, write_only, no_init);
-      h.single_task<TaskSequenceTest>([=]() {
-        task_sequence<argStruct<intertask_pipe1>> ts_func1Struct;
-        task_sequence<argArray<intertask_pipe2>> ts_func1Array;
-        task_sequence<returnStruct<intertask_pipe1>> ts_func2Struct;
-        task_sequence<returnArray<intertask_pipe2>> ts_func2Array;
-
-        // launch struct task functions
-        ts_func1Struct.async(in_acc_struct[0]);
-        ts_func2Struct.async();
-
-        // launch array task functions
-        ts_func1Array.async(in_acc_array[0]);
-        ts_func2Array.async();
-
-        out_acc_struct[0] = ts_func2Struct.get();
-        out_acc_array[0] = ts_func2Array.get();
-      });
-    });
-    q.wait();
-  }
-  assert((std::abs(res_struct[0].val - vec_in_struct[0].val) < 0.001) &&
-         res_struct[0].isValid);
-  assert((std::abs(res_array[0][0] - vec_in_array[0][0]) < 0.001) &&
-         (std::abs(res_array[0][1] - vec_in_array[0][1]) < 0.001));
-  return 0;
-}
diff --git a/sycl/test/check_device_code/task_sequence_intel_balanced.cpp b/sycl/test/check_device_code/task_sequence_intel_balanced.cpp
deleted file mode 100644
index a06dc2b92755c..0000000000000
--- a/sycl/test/check_device_code/task_sequence_intel_balanced.cpp
+++ /dev/null
@@ -1,39 +0,0 @@
-// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s
-
-#include <sycl/sycl.hpp>
-
-using namespace sycl::ext::intel::experimental;
-using namespace sycl::ext::oneapi::experimental;
-
-constexpr int kSize = 128;
-
-int arrayAdd(int *data1, int *data2, int N) {
-  int ret = 0;
-  for (int i = 0; i < N; ++i) {
-    ret += data1[i] + data2[i];
-  }
-
-  return ret;
-}
-
-// CHECK: [[TASK_SEQUENCE:%.*]] ={{.*}} call spir_func target("spirv.TaskSequenceINTEL") @_Z31__spirv_TaskSequenceCreateINTEL{{.*}}(ptr{{.*}}@_Z8arrayAdd{{.*}}, i32 -1, i32 -1, i32 0, i32 128)
-// CHECK: call spir_func void @_Z30__spirv_TaskSequenceAsyncINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]], ptr addrspace(4) {{.*}}, ptr addrspace(4) {{.*}}, i32 128)
-// CHECK: call spir_func void @_Z32__spirv_TaskSequenceReleaseINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
-// CHECK-COUNT-1: call spir_func i32 @_Z28__spirv_TaskSequenceGetINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
-SYCL_EXTERNAL void task_sequence_intel_balanced(
-    sycl::accessor<
-        sycl::vec<int, 1>, 1, sycl::access::mode::write,
-        sycl::access::target::device, sycl::access::placeholder::false_t,
-        sycl::ext::oneapi::accessor_property_list<sycl::property::no_init>>
-        results_acc) {
-  int d1[kSize], d2[kSize];
-  task_sequence<arrayAdd,
-                decltype(properties{balanced, invocation_capacity<kSize>})>
-      arrayAddTask;
-  for (int i = 0; i < kSize; i++) {
-    arrayAddTask.async(d1, d2, kSize);
-  }
-  for (int i = 0; i < kSize; i++) {
-    results_acc[i] = arrayAddTask.get();
-  }
-}
\ No newline at end of file
diff --git a/sycl/test/check_device_code/task_sequence_intel_explicit_get.cpp b/sycl/test/check_device_code/task_sequence_intel_explicit_get.cpp
deleted file mode 100644
index 7d8f45334e99f..0000000000000
--- a/sycl/test/check_device_code/task_sequence_intel_explicit_get.cpp
+++ /dev/null
@@ -1,40 +0,0 @@
-// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s
-
-#include <sycl/sycl.hpp>
-
-using namespace sycl::ext::intel::experimental;
-using namespace sycl::ext::oneapi::experimental;
-
-constexpr int kSize = 128;
-
-int arrayAdd(int *data1, int *data2, int N) {
-  int ret = 0;
-  for (int i = 0; i < N; ++i) {
-    ret += data1[i] + data2[i];
-  }
-
-  return ret;
-}
-
-// CHECK: [[TASK_SEQUENCE:%.*]] ={{.*}} call spir_func target("spirv.TaskSequenceINTEL") @_Z31__spirv_TaskSequenceCreateINTEL{{.*}}(ptr{{.*}}@_Z8arrayAdd{{.*}}, i32 -1, i32 -1, i32 0, i32 128)
-// CHECK: call spir_func void @_Z30__spirv_TaskSequenceAsyncINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]], ptr addrspace(4) {{.*}}, ptr addrspace(4) {{.*}}, i32 128)
-// CHECK: call spir_func i32 @_Z28__spirv_TaskSequenceGetINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
-// CHECK: call spir_func void @_Z32__spirv_TaskSequenceReleaseINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
-// CHECK: call spir_func i32 @_Z28__spirv_TaskSequenceGetINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
-SYCL_EXTERNAL void task_sequence_intel_explicit_get(
-    sycl::accessor<
-        sycl::vec<int, 1>, 1, sycl::access::mode::write,
-        sycl::access::target::device, sycl::access::placeholder::false_t,
-        sycl::ext::oneapi::accessor_property_list<sycl::property::no_init>>
-        results_acc) {
-  int d1[kSize], d2[kSize];
-  task_sequence<arrayAdd, decltype(properties{invocation_capacity<kSize>})>
-      arrayAddTask;
-  for (int i = 0; i < kSize; i++) {
-    arrayAddTask.async(d1, d2, kSize);
-  }
-
-  for (int i = 0; i < kSize; i++) {
-    results_acc[i] = arrayAddTask.get();
-  }
-}
\ No newline at end of file
diff --git a/sycl/test/check_device_code/task_sequence_intel_no_explicit_get.cpp b/sycl/test/check_device_code/task_sequence_intel_no_explicit_get.cpp
deleted file mode 100644
index bd208f9e0e0e1..0000000000000
--- a/sycl/test/check_device_code/task_sequence_intel_no_explicit_get.cpp
+++ /dev/null
@@ -1,31 +0,0 @@
-// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s
-
-// CHECK: [[TASK_SEQUENCE:%.*]] ={{.*}} call spir_func target("spirv.TaskSequenceINTEL") @_Z31__spirv_TaskSequenceCreateINTEL{{.*}}(ptr{{.*}}@_Z8arrayAdd{{.*}}, i32 0, i32 1, i32 1, i32 1)
-// CHECK: call spir_func void @_Z30__spirv_TaskSequenceAsyncINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]], ptr addrspace(4) {{.*}}, ptr addrspace(4) {{.*}}, i32 128)
-// CHECK-COUNT-1: call spir_func i32 @_Z28__spirv_TaskSequenceGetINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
-// CHECK: call spir_func void @_Z32__spirv_TaskSequenceReleaseINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
-
-#include <sycl/sycl.hpp>
-
-using namespace sycl::ext::intel::experimental;
-using namespace sycl::ext::oneapi::experimental;
-
-constexpr int kSize = 128;
-
-int arrayAdd(int *data1, int *data2, int N) {
-  int ret = 0;
-  for (int i = 0; i < N; ++i) {
-    ret += data1[i] + data2[i];
-  }
-
-  return ret;
-}
-
-SYCL_EXTERNAL void task_sequence_no_explicit_get() {
-  int d1[kSize], d2[kSize];
-  task_sequence<arrayAdd, decltype(properties{
-                              pipelined<0>, stall_enable_clusters,
-                              invocation_capacity<1>, response_capacity<1>})>
-      arrayAddTask;
-  arrayAddTask.async(d1, d2, kSize);
-}
\ No newline at end of file
diff --git a/sycl/test/extensions/properties/properties_kernel_fpga.cpp b/sycl/test/extensions/properties/properties_kernel_fpga.cpp
index f329fa3ddcd3e..2acd767c386ae 100644
--- a/sycl/test/extensions/properties/properties_kernel_fpga.cpp
+++ b/sycl/test/extensions/properties/properties_kernel_fpga.cpp
@@ -15,18 +15,21 @@ int main() {
                 intel::experimental::register_map_interface_key>::value);
   static_assert(oneapi::experimental::is_property_key<
                 intel::experimental::pipelined_key>::value);
-  static_assert(oneapi::experimental::is_property_key<
-                intel::experimental::fpga_cluster_key>::value);
 
   // Check that oneapi::experimental::is_property_value is correctly specialized
   static_assert(oneapi::experimental::is_property_value<
                 decltype(intel::experimental::streaming_interface<
                          intel::experimental::streaming_interface_options_enum::
                              accept_downstream_stall>)>::value);
-  static_assert(oneapi::experimental::is_property_value<
-                decltype(intel::experimental::streaming_interface<
-                         intel::experimental::streaming_interface_options_enum::
-                             remove_downstream_stall>)>::value);
+  static_assert(
+      oneapi::experimental::is_property_value<
+          decltype(intel::experimental::
+                       streaming_interface_accept_downstream_stall)>::value);
+  static_assert(
+      oneapi::experimental::is_property_value<
+          decltype(intel::experimental::
+                       streaming_interface_remove_downstream_stall)>::value);
+
   static_assert(
       oneapi::experimental::is_property_value<
           decltype(intel::experimental::register_map_interface<
@@ -34,103 +37,44 @@ int main() {
                        wait_for_done_write>)>::value);
   static_assert(
       oneapi::experimental::is_property_value<
-          decltype(intel::experimental::register_map_interface<
-                   intel::experimental::register_map_interface_options_enum::
-                       do_not_wait_for_done_write>)>::value);
+          decltype(intel::experimental::
+                       register_map_interface_wait_for_done_write)>::value);
+  static_assert(
+      oneapi::experimental::is_property_value<
+          decltype(intel::experimental::
+                       register_map_interface_do_not_wait_for_done_write)>::
+          value);
+
   static_assert(oneapi::experimental::is_property_value<
                 decltype(intel::experimental::pipelined<-1>)>::value);
   static_assert(oneapi::experimental::is_property_value<
                 decltype(intel::experimental::pipelined<0>)>::value);
   static_assert(oneapi::experimental::is_property_value<
                 decltype(intel::experimental::pipelined<4>)>::value);
-  static_assert(oneapi::experimental::is_property_value<
-                decltype(intel::experimental::fpga_cluster<
-                         intel::experimental::fpga_cluster_options_enum::
-                             stall_enable>)>::value);
-  static_assert(oneapi::experimental::is_property_value<
-                decltype(intel::experimental::fpga_cluster<
-                         intel::experimental::fpga_cluster_options_enum::
-                             stall_free>)>::value);
-
-  // Checks that fully specialized properties are the same as the templated
-  // variants.
-  static_assert(std::is_same_v<
-                decltype(intel::experimental::
-                             streaming_interface_accept_downstream_stall),
-                decltype(intel::experimental::streaming_interface<
-                         intel::experimental::streaming_interface_options_enum::
-                             accept_downstream_stall>)>);
-  static_assert(std::is_same_v<
-                decltype(intel::experimental::
-                             streaming_interface_remove_downstream_stall),
-                decltype(intel::experimental::streaming_interface<
-                         intel::experimental::streaming_interface_options_enum::
-                             remove_downstream_stall>)>);
-  static_assert(
-      std::is_same_v<
-          decltype(intel::experimental::
-                       register_map_interface_wait_for_done_write),
-          decltype(intel::experimental::register_map_interface<
-                   intel::experimental::register_map_interface_options_enum::
-                       wait_for_done_write>)>);
-  static_assert(
-      std::is_same_v<
-          decltype(intel::experimental::
-                       register_map_interface_do_not_wait_for_done_write),
-          decltype(intel::experimental::register_map_interface<
-                   intel::experimental::register_map_interface_options_enum::
-                       do_not_wait_for_done_write>)>);
-  static_assert(
-      std::is_same_v<decltype(intel::experimental::stall_enable_clusters),
-                     decltype(intel::experimental::fpga_cluster<
-                              intel::experimental::fpga_cluster_options_enum::
-                                  stall_enable>)>);
-  static_assert(
-      std::is_same_v<decltype(intel::experimental::stall_free_clusters),
-                     decltype(intel::experimental::fpga_cluster<
-                              intel::experimental::fpga_cluster_options_enum::
-                                  stall_free>)>);
 
   // Check that property lists will accept the new properties
   using PS = decltype(oneapi::experimental::properties(
-      intel::experimental::streaming_interface_remove_downstream_stall,
-      intel::experimental::pipelined<1>,
-      intel::experimental::stall_enable_clusters));
+      intel::experimental::streaming_interface_remove_downstream_stall));
   static_assert(oneapi::experimental::is_property_list_v<PS>);
   static_assert(
       PS::has_property<intel::experimental::streaming_interface_key>());
-  static_assert(PS::has_property<intel::experimental::pipelined_key>());
-  static_assert(PS::has_property<intel::experimental::fpga_cluster_key>());
   static_assert(
       PS::get_property<intel::experimental::streaming_interface_key>() ==
-      intel::experimental::streaming_interface<
-          intel::experimental::streaming_interface_options_enum::
-              remove_downstream_stall>);
-  static_assert(PS::get_property<intel::experimental::pipelined_key>() ==
-                intel::experimental::pipelined<1>);
-  static_assert(
-      PS::get_property<intel::experimental::fpga_cluster_key>() ==
-      intel::experimental::fpga_cluster<
-          intel::experimental::fpga_cluster_options_enum::stall_enable>);
+      intel::experimental::streaming_interface_remove_downstream_stall);
 
   using PR = decltype(oneapi::experimental::properties(
-      intel::experimental::register_map_interface_wait_for_done_write,
-      intel::experimental::pipelined<0>,
-      intel::experimental::stall_free_clusters));
+      intel::experimental::register_map_interface_wait_for_done_write));
   static_assert(oneapi::experimental::is_property_list_v<PR>);
   static_assert(
       PR::has_property<intel::experimental::register_map_interface_key>());
-  static_assert(PR::has_property<intel::experimental::pipelined_key>());
-  static_assert(PR::has_property<intel::experimental::fpga_cluster_key>());
   static_assert(
       PR::get_property<intel::experimental::register_map_interface_key>() ==
-      intel::experimental::register_map_interface<
-          intel::experimental::register_map_interface_options_enum::
-              wait_for_done_write>);
-  static_assert(PR::get_property<intel::experimental::pipelined_key>() ==
-                intel::experimental::pipelined<0>);
-  static_assert(
-      PR::get_property<intel::experimental::fpga_cluster_key>() ==
-      intel::experimental::fpga_cluster<
-          intel::experimental::fpga_cluster_options_enum::stall_free>);
+      intel::experimental::register_map_interface_wait_for_done_write);
+
+  using PP = decltype(oneapi::experimental::properties(
+      intel::experimental::pipelined<-1>));
+  static_assert(oneapi::experimental::is_property_list_v<PP>);
+  static_assert(PP::has_property<intel::experimental::pipelined_key>());
+  static_assert(PP::get_property<intel::experimental::pipelined_key>() ==
+                intel::experimental::pipelined<-1>);
 }
diff --git a/sycl/test/extensions/properties/properties_task_sequence.cpp b/sycl/test/extensions/properties/properties_task_sequence.cpp
deleted file mode 100644
index 75d2894b098df..0000000000000
--- a/sycl/test/extensions/properties/properties_task_sequence.cpp
+++ /dev/null
@@ -1,52 +0,0 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
-// expected-no-diagnostics
-
-#include <sycl/sycl.hpp>
-
-#include <sycl/ext/intel/experimental/task_sequence_properties.hpp>
-
-using namespace sycl::ext;
-
-constexpr uint32_t TestResponseCapacity = 7;
-constexpr uint32_t TestInvocationCapacity = 5;
-int main() {
-  // Check that is_property_key is correctly specialized
-  static_assert(oneapi::experimental::is_property_key<
-                intel::experimental::balanced_key>::value);
-  static_assert(oneapi::experimental::is_property_key<
-                intel::experimental::response_capacity_key>::value);
-  static_assert(oneapi::experimental::is_property_key<
-                intel::experimental::invocation_capacity_key>::value);
-
-  // Check that is_property_value is correctly specialized
-  static_assert(oneapi::experimental::is_property_value<
-                decltype(intel::experimental::balanced)>::value);
-  static_assert(oneapi::experimental::is_property_value<
-                decltype(intel::experimental::response_capacity<0>)>::value);
-  static_assert(oneapi::experimental::is_property_value<
-                decltype(intel::experimental::response_capacity<
-                         TestResponseCapacity>)>::value);
-  static_assert(oneapi::experimental::is_property_value<
-                decltype(intel::experimental::invocation_capacity<0>)>::value);
-  static_assert(oneapi::experimental::is_property_value<
-                decltype(intel::experimental::invocation_capacity<
-                         TestInvocationCapacity>)>::value);
-
-  // Check that property lists will accept the new properties.
-  using P = decltype(oneapi::experimental::properties(
-      intel::experimental::balanced,
-      intel::experimental::response_capacity<TestResponseCapacity>,
-      intel::experimental::invocation_capacity<TestInvocationCapacity>));
-  static_assert(oneapi::experimental::is_property_list_v<P>);
-  static_assert(P::has_property<intel::experimental::balanced_key>());
-  static_assert(P::has_property<intel::experimental::response_capacity_key>());
-  static_assert(
-      P::has_property<intel::experimental::invocation_capacity_key>());
-  static_assert(P::get_property<intel::experimental::balanced_key>() ==
-                intel::experimental::balanced);
-  static_assert(P::get_property<intel::experimental::response_capacity_key>() ==
-                intel::experimental::response_capacity<TestResponseCapacity>);
-  static_assert(
-      P::get_property<intel::experimental::invocation_capacity_key>() ==
-      intel::experimental::invocation_capacity<TestInvocationCapacity>);
-}
\ No newline at end of file
diff --git a/sycl/test/extensions/task_sequence/task_sequence_properties.cpp b/sycl/test/extensions/task_sequence/task_sequence_properties.cpp
deleted file mode 100644
index d5ca162bc4080..0000000000000
--- a/sycl/test/extensions/task_sequence/task_sequence_properties.cpp
+++ /dev/null
@@ -1,96 +0,0 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
-// expected-no-diagnostics
-
-#include <sycl/ext/intel/fpga_extensions.hpp>
-#include <sycl/sycl.hpp>
-
-using namespace sycl::ext::oneapi::experimental;
-using namespace sycl::ext::intel::experimental;
-
-void dummyTask() {}
-
-static task_sequence<dummyTask,
-                     decltype(properties(balanced, response_capacity<0>,
-                                         invocation_capacity<0>, pipelined<-1>,
-                                         stall_free_clusters))>
-    TaskSequence1;
-static task_sequence<
-    dummyTask, decltype(properties(response_capacity<2>, invocation_capacity<3>,
-                                   pipelined<1>, stall_enable_clusters))>
-    TaskSequence2;
-static task_sequence<dummyTask, decltype(properties(pipelined<0>))>
-    TaskSequence3;
-
-// Checks is_property_key_of and is_property_value_of for T.
-template <typename T> void checkIsPropertyOf() {
-  static_assert(is_property_key_of<balanced_key, T>::value);
-  static_assert(is_property_key_of<response_capacity_key, T>::value);
-  static_assert(is_property_key_of<invocation_capacity_key, T>::value);
-  static_assert(is_property_key_of<pipelined_key, T>::value);
-  static_assert(is_property_key_of<fpga_cluster_key, T>::value);
-
-  static_assert(is_property_value_of<decltype(balanced), T>::value);
-  static_assert(is_property_value_of<decltype(response_capacity<1>), T>::value);
-  static_assert(
-      is_property_value_of<decltype(invocation_capacity<1>), T>::value);
-  static_assert(is_property_value_of<decltype(pipelined<1>), T>::value);
-  static_assert(is_property_value_of<
-                decltype(fpga_cluster<fpga_cluster_options_enum::stall_enable>),
-                T>::value);
-}
-
-int main() {
-  static_assert(is_property_key<balanced_key>::value);
-  static_assert(is_property_key<response_capacity_key>::value);
-  static_assert(is_property_key<invocation_capacity_key>::value);
-  static_assert(is_property_key<pipelined_key>::value);
-  static_assert(is_property_key<fpga_cluster_key>::value);
-
-  static_assert(is_property_value<decltype(balanced)>::value);
-  static_assert(is_property_value<decltype(response_capacity<1>)>::value);
-  static_assert(is_property_value<decltype(invocation_capacity<1>)>::value);
-  static_assert(is_property_value<decltype(pipelined<-1>)>::value);
-  static_assert(is_property_value<decltype(pipelined<0>)>::value);
-  static_assert(is_property_value<decltype(pipelined<1>)>::value);
-  static_assert(is_property_value<decltype(stall_enable_clusters)>::value);
-  static_assert(is_property_value<decltype(stall_free_clusters)>::value);
-
-  checkIsPropertyOf<decltype(TaskSequence1)>();
-  static_assert(TaskSequence1.has_property<balanced_key>());
-  static_assert(TaskSequence1.has_property<response_capacity_key>());
-  static_assert(TaskSequence1.has_property<invocation_capacity_key>());
-  static_assert(TaskSequence1.has_property<pipelined_key>());
-  static_assert(TaskSequence1.has_property<fpga_cluster_key>());
-  static_assert(TaskSequence1.get_property<balanced_key>() == balanced);
-  static_assert(TaskSequence1.get_property<response_capacity_key>() ==
-                response_capacity<0>);
-  static_assert(TaskSequence1.get_property<invocation_capacity_key>() ==
-                invocation_capacity<0>);
-  static_assert(TaskSequence1.get_property<pipelined_key>() == pipelined<-1>);
-  static_assert(TaskSequence1.get_property<fpga_cluster_key>() ==
-                stall_free_clusters);
-
-  checkIsPropertyOf<decltype(TaskSequence2)>();
-  static_assert(!TaskSequence2.has_property<balanced_key>());
-  static_assert(TaskSequence2.has_property<response_capacity_key>());
-  static_assert(TaskSequence2.has_property<invocation_capacity_key>());
-  static_assert(TaskSequence2.has_property<pipelined_key>());
-  static_assert(TaskSequence2.has_property<fpga_cluster_key>());
-  static_assert(TaskSequence2.get_property<response_capacity_key>() ==
-                response_capacity<2>);
-  static_assert(TaskSequence2.get_property<invocation_capacity_key>() ==
-                invocation_capacity<3>);
-  static_assert(TaskSequence2.get_property<pipelined_key>() == pipelined<1>);
-  static_assert(TaskSequence2.get_property<fpga_cluster_key>() ==
-                stall_enable_clusters);
-
-  checkIsPropertyOf<decltype(TaskSequence3)>();
-  static_assert(!TaskSequence3.has_property<balanced_key>());
-  static_assert(!TaskSequence3.has_property<response_capacity_key>());
-  static_assert(!TaskSequence3.has_property<invocation_capacity_key>());
-  static_assert(TaskSequence3.has_property<pipelined_key>());
-  static_assert(!TaskSequence3.has_property<fpga_cluster_key>());
-  static_assert(TaskSequence3.get_property<pipelined_key>() == pipelined<0>);
-
-  return 0;
-}