diff --git a/SYCL/WeakObject/weak_object_copy.cpp b/SYCL/WeakObject/weak_object_copy.cpp
new file mode 100644
index 0000000000..865ce036e4
--- /dev/null
+++ b/SYCL/WeakObject/weak_object_copy.cpp
@@ -0,0 +1,30 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// This test checks the behavior of the copy ctor and assignment operator for
+// `weak_object`.
+
+#include "weak_object_utils.hpp"
+
+template <typename SyclObjT> struct WeakObjectCheckCopy {
+  void operator()(SyclObjT Obj) {
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj{Obj};
+
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObjCopyCtor{WeakObj};
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObjCopyAssign = WeakObj;
+
+    assert(!WeakObjCopyCtor.expired());
+    assert(!WeakObjCopyAssign.expired());
+
+    assert(WeakObjCopyCtor.lock() == Obj);
+    assert(WeakObjCopyAssign.lock() == Obj);
+  }
+};
+
+int main() {
+  sycl::queue Q;
+  runTest<WeakObjectCheckCopy>(Q);
+  return 0;
+}
diff --git a/SYCL/WeakObject/weak_object_expired.cpp b/SYCL/WeakObject/weak_object_expired.cpp
new file mode 100644
index 0000000000..21b29a8078
--- /dev/null
+++ b/SYCL/WeakObject/weak_object_expired.cpp
@@ -0,0 +1,24 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// This test checks the behavior of `expired()` for `weak_object`.
+
+#include "weak_object_utils.hpp"
+
+template <typename SyclObjT> struct WeakObjectCheckExpired {
+  void operator()(SyclObjT Obj) {
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj{Obj};
+    sycl::ext::oneapi::weak_object<SyclObjT> NullWeakObj;
+
+    assert(!WeakObj.expired());
+    assert(NullWeakObj.expired());
+  }
+};
+
+int main() {
+  sycl::queue Q;
+  runTest<WeakObjectCheckExpired>(Q);
+  return 0;
+}
diff --git a/SYCL/WeakObject/weak_object_lock.cpp b/SYCL/WeakObject/weak_object_lock.cpp
new file mode 100644
index 0000000000..5991952edd
--- /dev/null
+++ b/SYCL/WeakObject/weak_object_lock.cpp
@@ -0,0 +1,32 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// This test checks the behavior of `lock()` for `weak_object`.
+
+#include "weak_object_utils.hpp"
+
+template <typename SyclObjT> struct WeakObjectCheckLock {
+  void operator()(SyclObjT Obj) {
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj{Obj};
+    sycl::ext::oneapi::weak_object<SyclObjT> NullWeakObj;
+
+    SyclObjT LObj = WeakObj.lock();
+    assert(LObj == Obj);
+
+    try {
+      SyclObjT LNull = NullWeakObj.lock();
+      assert(false && "Locking empty weak object did not throw.");
+    } catch (sycl::exception &E) {
+      assert(E.code() == sycl::make_error_code(sycl::errc::invalid) &&
+             "Unexpected thrown error code.");
+    }
+  }
+};
+
+int main() {
+  sycl::queue Q;
+  runTest<WeakObjectCheckLock>(Q);
+  return 0;
+}
diff --git a/SYCL/WeakObject/weak_object_move.cpp b/SYCL/WeakObject/weak_object_move.cpp
new file mode 100644
index 0000000000..e5f70dbf07
--- /dev/null
+++ b/SYCL/WeakObject/weak_object_move.cpp
@@ -0,0 +1,33 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// This test checks the behavior of the copy ctor and assignment operator for
+// `weak_object`.
+
+#include "weak_object_utils.hpp"
+
+template <typename SyclObjT> struct WeakObjectCheckMove {
+  void operator()(SyclObjT Obj) {
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj1{Obj};
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj2{Obj};
+
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObjMoveCtor{
+        std::move(WeakObj1)};
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObjMoveAssign =
+        std::move(WeakObj2);
+
+    assert(!WeakObjMoveCtor.expired());
+    assert(!WeakObjMoveAssign.expired());
+
+    assert(WeakObjMoveCtor.lock() == Obj);
+    assert(WeakObjMoveAssign.lock() == Obj);
+  }
+};
+
+int main() {
+  sycl::queue Q;
+  runTest<WeakObjectCheckMove>(Q);
+  return 0;
+}
diff --git a/SYCL/WeakObject/weak_object_owner_before.cpp b/SYCL/WeakObject/weak_object_owner_before.cpp
new file mode 100644
index 0000000000..28085bbf8e
--- /dev/null
+++ b/SYCL/WeakObject/weak_object_owner_before.cpp
@@ -0,0 +1,54 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// This test checks the behavior of owner_before semantics for `weak_object`.
+
+#include "weak_object_utils.hpp"
+
+template <typename SyclObjT> struct WeakObjectCheckOwnerBefore {
+  void operator()(SyclObjT Obj) {
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj{Obj};
+    sycl::ext::oneapi::weak_object<SyclObjT> NullWeakObj;
+
+    assert((WeakObj.owner_before(NullWeakObj) &&
+            !NullWeakObj.owner_before(WeakObj)) ||
+           (NullWeakObj.owner_before(WeakObj) &&
+            !WeakObj.owner_before(NullWeakObj)));
+
+    assert(!WeakObj.owner_before(Obj));
+    assert(!Obj.ext_oneapi_owner_before(WeakObj));
+
+    assert(!Obj.ext_oneapi_owner_before(Obj));
+  }
+};
+
+template <typename SyclObjT> struct WeakObjectCheckOwnerBeforeMulti {
+  void operator()(SyclObjT Obj1, SyclObjT Obj2) {
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj1{Obj1};
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj2{Obj2};
+
+    assert(
+        (WeakObj1.owner_before(WeakObj2) && !WeakObj2.owner_before(WeakObj1)) ||
+        (WeakObj2.owner_before(WeakObj1) && !WeakObj1.owner_before(WeakObj2)));
+
+    assert(!WeakObj1.owner_before(Obj1));
+    assert(!Obj1.ext_oneapi_owner_before(WeakObj1));
+
+    assert(!WeakObj2.owner_before(Obj2));
+    assert(!Obj2.ext_oneapi_owner_before(WeakObj2));
+
+    assert((Obj1.ext_oneapi_owner_before(Obj2) &&
+            !Obj2.ext_oneapi_owner_before(Obj1)) ||
+           (Obj2.ext_oneapi_owner_before(Obj1) &&
+            !Obj1.ext_oneapi_owner_before(Obj2)));
+  }
+};
+
+int main() {
+  sycl::queue Q;
+  runTest<WeakObjectCheckOwnerBefore>(Q);
+  runTestMulti<WeakObjectCheckOwnerBeforeMulti>(Q);
+  return 0;
+}
diff --git a/SYCL/WeakObject/weak_object_owner_less.cpp b/SYCL/WeakObject/weak_object_owner_less.cpp
new file mode 100644
index 0000000000..dff281be5c
--- /dev/null
+++ b/SYCL/WeakObject/weak_object_owner_less.cpp
@@ -0,0 +1,98 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// This test checks the behavior of owner_less semantics for `weak_object`.
+
+#include "weak_object_utils.hpp"
+
+#include <map>
+
+template <typename SyclObjT> struct WeakObjectCheckOwnerLess {
+  void operator()(SyclObjT Obj) {
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj{Obj};
+    sycl::ext::oneapi::weak_object<SyclObjT> NullWeakObj;
+    sycl::ext::oneapi::owner_less<SyclObjT> Comparator;
+
+    assert((Comparator(WeakObj, NullWeakObj) &&
+            !Comparator(NullWeakObj, WeakObj)) ||
+           (Comparator(NullWeakObj, WeakObj) &&
+            !Comparator(WeakObj, NullWeakObj)));
+
+    assert(!Comparator(WeakObj, Obj));
+    assert(!Comparator(Obj, WeakObj));
+  }
+};
+
+template <typename SyclObjT> struct WeakObjectCheckOwnerLessMulti {
+  void operator()(SyclObjT Obj1, SyclObjT Obj2) {
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj1{Obj1};
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj2{Obj2};
+    sycl::ext::oneapi::owner_less<SyclObjT> Comparator;
+
+    assert(
+        (Comparator(WeakObj1, WeakObj2) && !Comparator(WeakObj2, WeakObj1)) ||
+        (Comparator(WeakObj2, WeakObj1) && !Comparator(WeakObj1, WeakObj2)));
+
+    assert(!Comparator(WeakObj1, Obj1));
+    assert(!Comparator(Obj1, WeakObj1));
+
+    assert(!Comparator(WeakObj2, Obj2));
+    assert(!Comparator(Obj2, WeakObj2));
+  }
+};
+
+template <typename SyclObjT> struct WeakObjectCheckOwnerLessMap {
+  void operator()(SyclObjT Obj1, SyclObjT Obj2) {
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj1{Obj1};
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj2{Obj2};
+
+    std::map<sycl::ext::oneapi::weak_object<SyclObjT>, int,
+             sycl::ext::oneapi::owner_less<SyclObjT>>
+        Map;
+    Map[WeakObj1] = 1;
+    Map[WeakObj2] = 2;
+
+    assert(Map.size() == (size_t)2);
+    assert(Map[WeakObj1] == 1);
+    assert(Map[WeakObj2] == 2);
+    assert(Map[Obj1] == 1);
+    assert(Map[Obj2] == 2);
+
+    Map[WeakObj1] = 2;
+    Map[WeakObj2] = 3;
+
+    assert(Map.size() == (size_t)2);
+    assert(Map[WeakObj1] == 2);
+    assert(Map[WeakObj2] == 3);
+    assert(Map[Obj1] == 2);
+    assert(Map[Obj2] == 3);
+
+    Map[Obj1] = 5;
+    Map[Obj2] = 6;
+
+    assert(Map.size() == (size_t)2);
+    assert(Map[WeakObj1] == 5);
+    assert(Map[WeakObj2] == 6);
+    assert(Map[Obj1] == 5);
+    assert(Map[Obj2] == 6);
+
+    Map[sycl::ext::oneapi::weak_object<SyclObjT>{Obj1}] = 10;
+    Map[sycl::ext::oneapi::weak_object<SyclObjT>{Obj2}] = 13;
+
+    assert(Map.size() == (size_t)2);
+    assert(Map[WeakObj1] == 10);
+    assert(Map[WeakObj2] == 13);
+    assert(Map[Obj1] == 10);
+    assert(Map[Obj2] == 13);
+  }
+};
+
+int main() {
+  sycl::queue Q;
+  runTest<WeakObjectCheckOwnerLess>(Q);
+  runTestMulti<WeakObjectCheckOwnerLessMulti>(Q);
+  runTestMulti<WeakObjectCheckOwnerLessMap>(Q);
+  return 0;
+}
diff --git a/SYCL/WeakObject/weak_object_reset.cpp b/SYCL/WeakObject/weak_object_reset.cpp
new file mode 100644
index 0000000000..32bb2bb667
--- /dev/null
+++ b/SYCL/WeakObject/weak_object_reset.cpp
@@ -0,0 +1,37 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// This test checks the behavior of `reset()` for `weak_object`.
+
+#include "weak_object_utils.hpp"
+
+template <typename SyclObjT> struct WeakObjectCheckReset {
+  void operator()(SyclObjT Obj) {
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj{Obj};
+    sycl::ext::oneapi::weak_object<SyclObjT> NullWeakObj;
+
+    WeakObj.reset();
+    assert(WeakObj.expired());
+    assert(!WeakObj.owner_before(NullWeakObj));
+    assert(!NullWeakObj.owner_before(WeakObj));
+
+    std::optional<SyclObjT> TLObj = WeakObj.try_lock();
+    assert(!TLObj.has_value());
+
+    try {
+      SyclObjT LObj = WeakObj.lock();
+      assert(false && "Locking reset weak object did not throw.");
+    } catch (sycl::exception &E) {
+      assert(E.code() == sycl::make_error_code(sycl::errc::invalid) &&
+             "Unexpected thrown error code.");
+    }
+  }
+};
+
+int main() {
+  sycl::queue Q;
+  runTest<WeakObjectCheckReset>(Q);
+  return 0;
+}
diff --git a/SYCL/WeakObject/weak_object_try_lock.cpp b/SYCL/WeakObject/weak_object_try_lock.cpp
new file mode 100644
index 0000000000..a2074ae54c
--- /dev/null
+++ b/SYCL/WeakObject/weak_object_try_lock.cpp
@@ -0,0 +1,29 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// This test checks the behavior of `try_lock()` for `weak_object`.
+
+#include "weak_object_utils.hpp"
+
+template <typename SyclObjT> struct WeakObjectCheckTryLock {
+  void operator()(SyclObjT Obj) {
+    sycl::ext::oneapi::weak_object<SyclObjT> WeakObj{Obj};
+    sycl::ext::oneapi::weak_object<SyclObjT> NullWeakObj;
+
+    std::optional<SyclObjT> TLObj = WeakObj.try_lock();
+    std::optional<SyclObjT> TLNull = NullWeakObj.try_lock();
+
+    assert(TLObj.has_value());
+    assert(!TLNull.has_value());
+
+    assert(TLObj.value() == Obj);
+  }
+};
+
+int main() {
+  sycl::queue Q;
+  runTest<WeakObjectCheckTryLock>(Q);
+  return 0;
+}
diff --git a/SYCL/WeakObject/weak_object_utils.hpp b/SYCL/WeakObject/weak_object_utils.hpp
new file mode 100644
index 0000000000..64192b3624
--- /dev/null
+++ b/SYCL/WeakObject/weak_object_utils.hpp
@@ -0,0 +1,160 @@
+// Utilities for weak_object testing
+
+#include <sycl/sycl.hpp>
+
+class TestKernel1;
+class TestKernel2;
+
+void MaterializeTestKernels(sycl::queue Q) {
+  if (false) {
+    Q.single_task<TestKernel1>([]() {});
+    Q.single_task<TestKernel2>([]() {});
+  }
+}
+
+template <template <typename> typename CallableT> void runTest(sycl::queue Q) {
+  MaterializeTestKernels(Q);
+
+  sycl::context Ctx = Q.get_context();
+  sycl::device Dev = Q.get_device();
+  sycl::platform Plt = Dev.get_platform();
+  sycl::event E;
+  sycl::kernel_id KId = sycl::get_kernel_id<TestKernel1>();
+  sycl::kernel_bundle KB =
+      sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev});
+  sycl::buffer<int, 1> Buf1D{1};
+  sycl::buffer<int, 2> Buf2D{sycl::range<2>{1, 2}};
+  sycl::buffer<int, 3> Buf3D{sycl::range<3>{1, 2, 3}};
+  sycl::accessor PAcc1D{Buf1D, sycl::read_write};
+  sycl::accessor PAcc2D{Buf2D, sycl::read_write};
+  sycl::accessor PAcc3D{Buf3D, sycl::read_write};
+  sycl::accessor<int, 1, sycl::access::mode::read_write,
+                 sycl::access::target::host_buffer>
+      HAcc1D;
+  sycl::accessor<int, 2, sycl::access::mode::read_write,
+                 sycl::access::target::host_buffer>
+      HAcc2D;
+  sycl::accessor<int, 3, sycl::access::mode::read_write,
+                 sycl::access::target::host_buffer>
+      HAcc3D;
+
+  CallableT<decltype(Plt)>()(Plt);
+  CallableT<decltype(Dev)>()(Dev);
+  CallableT<decltype(Ctx)>()(Ctx);
+  CallableT<decltype(Q)>()(Q);
+  CallableT<decltype(E)>()(E);
+  CallableT<decltype(KId)>()(KId);
+  CallableT<decltype(KB)>()(KB);
+  CallableT<decltype(Buf1D)>()(Buf1D);
+  CallableT<decltype(Buf2D)>()(Buf2D);
+  CallableT<decltype(Buf3D)>()(Buf3D);
+  CallableT<decltype(PAcc1D)>()(PAcc1D);
+  CallableT<decltype(PAcc2D)>()(PAcc2D);
+  CallableT<decltype(PAcc3D)>()(PAcc3D);
+  CallableT<decltype(HAcc1D)>()(HAcc1D);
+  CallableT<decltype(HAcc2D)>()(HAcc2D);
+  CallableT<decltype(HAcc3D)>()(HAcc3D);
+
+  Q.submit([&](sycl::handler &CGH) {
+    sycl::accessor DAcc1D{Buf1D, CGH, sycl::read_only};
+    sycl::accessor DAcc2D{Buf2D, CGH, sycl::read_only};
+    sycl::accessor DAcc3D{Buf3D, CGH, sycl::read_only};
+    sycl::local_accessor<int> LAcc1D{1, CGH};
+    sycl::local_accessor<int, 2> LAcc2D{sycl::range<2>{1, 2}, CGH};
+    sycl::local_accessor<int, 3> LAcc3D{sycl::range<3>{1, 2, 3}, CGH};
+
+    CallableT<decltype(DAcc1D)>()(DAcc1D);
+    CallableT<decltype(DAcc2D)>()(DAcc2D);
+    CallableT<decltype(DAcc3D)>()(DAcc3D);
+    CallableT<decltype(LAcc1D)>()(LAcc1D);
+    CallableT<decltype(LAcc2D)>()(LAcc2D);
+    CallableT<decltype(LAcc3D)>()(LAcc3D);
+  });
+}
+
+template <template <typename> typename CallableT>
+void runTestMulti(sycl::queue Q1) {
+  MaterializeTestKernels(Q1);
+
+  sycl::context Ctx1 = Q1.get_context();
+  sycl::device Dev = Q1.get_device();
+  sycl::platform Plt = Dev.get_platform();
+
+  sycl::context Ctx2{Dev};
+  sycl::queue Q2{Dev};
+  sycl::event E1;
+  sycl::event E2;
+  sycl::kernel_id KId1 = sycl::get_kernel_id<TestKernel1>();
+  sycl::kernel_id KId2 = sycl::get_kernel_id<TestKernel2>();
+  sycl::kernel_bundle KB1 =
+      sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx1, {Dev});
+  sycl::kernel_bundle KB2 =
+      sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx2, {Dev});
+  sycl::buffer<int, 1> Buf1D1{1};
+  sycl::buffer<int, 1> Buf1D2{1};
+  sycl::buffer<int, 2> Buf2D1{sycl::range<2>{1, 2}};
+  sycl::buffer<int, 2> Buf2D2{sycl::range<2>{1, 2}};
+  sycl::buffer<int, 3> Buf3D1{sycl::range<3>{1, 2, 3}};
+  sycl::buffer<int, 3> Buf3D2{sycl::range<3>{1, 2, 3}};
+  sycl::accessor PAcc1D1{Buf1D1, sycl::read_write};
+  sycl::accessor PAcc1D2{Buf1D2, sycl::read_write};
+  sycl::accessor PAcc2D1{Buf2D1, sycl::read_write};
+  sycl::accessor PAcc2D2{Buf2D2, sycl::read_write};
+  sycl::accessor PAcc3D1{Buf3D1, sycl::read_write};
+  sycl::accessor PAcc3D2{Buf3D2, sycl::read_write};
+  sycl::accessor<int, 1, sycl::access::mode::read_write,
+                 sycl::access::target::host_buffer>
+      HAcc1D1;
+  sycl::accessor<int, 1, sycl::access::mode::read_write,
+                 sycl::access::target::host_buffer>
+      HAcc1D2;
+  sycl::accessor<int, 2, sycl::access::mode::read_write,
+                 sycl::access::target::host_buffer>
+      HAcc2D1;
+  sycl::accessor<int, 2, sycl::access::mode::read_write,
+                 sycl::access::target::host_buffer>
+      HAcc2D2;
+  sycl::accessor<int, 3, sycl::access::mode::read_write,
+                 sycl::access::target::host_buffer>
+      HAcc3D1;
+  sycl::accessor<int, 3, sycl::access::mode::read_write,
+                 sycl::access::target::host_buffer>
+      HAcc3D2;
+
+  CallableT<decltype(Ctx1)>()(Ctx1, Ctx2);
+  CallableT<decltype(Q1)>()(Q1, Q2);
+  CallableT<decltype(E1)>()(E1, E2);
+  CallableT<decltype(KId1)>()(KId1, KId2);
+  CallableT<decltype(KB1)>()(KB1, KB2);
+  CallableT<decltype(Buf1D1)>()(Buf1D1, Buf1D2);
+  CallableT<decltype(Buf2D1)>()(Buf2D1, Buf2D2);
+  CallableT<decltype(Buf3D1)>()(Buf3D1, Buf3D2);
+  CallableT<decltype(PAcc1D1)>()(PAcc1D1, PAcc1D2);
+  CallableT<decltype(PAcc2D1)>()(PAcc2D1, PAcc2D2);
+  CallableT<decltype(PAcc3D1)>()(PAcc3D1, PAcc3D2);
+  CallableT<decltype(HAcc1D1)>()(HAcc1D1, HAcc1D2);
+  CallableT<decltype(HAcc2D1)>()(HAcc2D1, HAcc2D2);
+  CallableT<decltype(HAcc3D1)>()(HAcc3D1, HAcc3D2);
+
+  Q1.submit([&](sycl::handler &CGH) {
+    sycl::accessor DAcc1D1{Buf1D1, CGH, sycl::read_only};
+    sycl::accessor DAcc1D2{Buf1D2, CGH, sycl::read_only};
+    sycl::accessor DAcc2D1{Buf2D1, CGH, sycl::read_only};
+    sycl::accessor DAcc2D2{Buf2D2, CGH, sycl::read_only};
+    sycl::accessor DAcc3D1{Buf3D1, CGH, sycl::read_only};
+    sycl::accessor DAcc3D2{Buf3D2, CGH, sycl::read_only};
+    sycl::local_accessor<int, 1> LAcc1D1{1, CGH};
+    sycl::local_accessor<int, 1> LAcc1D2{1, CGH};
+    sycl::local_accessor<int, 2> LAcc2D1{sycl::range<2>{1, 2}, CGH};
+    sycl::local_accessor<int, 2> LAcc2D2{sycl::range<2>{1, 2}, CGH};
+    sycl::local_accessor<int, 3> LAcc3D1{sycl::range<3>{1, 2, 3}, CGH};
+    sycl::local_accessor<int, 3> LAcc3D2{sycl::range<3>{1, 2, 3}, CGH};
+
+    CallableT<decltype(DAcc1D1)>()(DAcc1D1, DAcc1D2);
+    CallableT<decltype(DAcc2D1)>()(DAcc2D1, DAcc2D2);
+    CallableT<decltype(DAcc3D1)>()(DAcc3D1, DAcc3D2);
+    CallableT<decltype(LAcc1D1)>()(LAcc1D1, LAcc1D2);
+    CallableT<decltype(LAcc2D1)>()(LAcc2D1, LAcc2D2);
+    CallableT<decltype(LAcc3D1)>()(LAcc3D1, LAcc3D2);
+  });
+}