diff --git a/sycl/test-e2e/AtomicRef/add.h b/sycl/test-e2e/AtomicRef/add.h
index 48f24342202d2..0fea8fffb18e4 100644
--- a/sycl/test-e2e/AtomicRef/add.h
+++ b/sycl/test-e2e/AtomicRef/add.h
@@ -104,6 +104,47 @@ void add_fetch_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void add_fetch_test_usm_shared(queue q, size_t N) {
+  T *sum = malloc_shared<T>(1, q);
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  sum[0] = T(0);
+  std::fill(output_begin, output_end, T(0));
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         int gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (sum[0]);
+         output[gid] = atm.fetch_add(Difference(1), order);
+       });
+     }).wait_and_throw();
+  }
+
+  // All work-items increment by 1, so final value should be equal to N.
+  assert(sum[0] == T(N));
+
+  // Fetch returns original value: will be in [0, N-1].
+  auto min_e = std::min_element(output_begin, output_end);
+  auto max_e = std::max_element(output_begin, output_end);
+  assert(*min_e == T(0) && *max_e == T(N - 1));
+
+  // Intermediate values should be unique.
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+
+  free(sum, q);
+  free(output, q);
+}
+
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
           access::address_space space, typename T, typename Difference = T,
@@ -146,6 +187,47 @@ void add_plus_equal_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void add_plus_equal_test_usm_shared(queue q, size_t N) {
+  T *sum = malloc_shared<T>(1, q);
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  sum[0] = T(0);
+  std::fill(output_begin, output_end, T(0));
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         int gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (sum[0]);
+         output[gid] = atm += Difference(1);
+       });
+     }).wait_and_throw();
+  }
+
+  // All work-items increment by 1, so final value should be equal to N.
+  assert(sum[0] == T(N));
+
+  // += returns updated value: will be in [1, N].
+  auto min_e = std::min_element(output_begin, output_end);
+  auto max_e = std::max_element(output_begin, output_end);
+  assert(*min_e == T(1) && *max_e == T(N));
+
+  // Intermediate values should be unique.
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+
+  free(sum, q);
+  free(output, q);
+}
+
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
           access::address_space space, typename T, typename Difference = T,
@@ -188,6 +270,46 @@ void add_pre_inc_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void add_pre_inc_test_usm_shared(queue q, size_t N) {
+  T *sum = malloc_shared<T>(1, q);
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  sum[0] = T(0);
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         int gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (sum[0]);
+         output[gid] = ++atm;
+       });
+     }).wait_and_throw();
+  }
+
+  // All work-items increment by 1, so final value should be equal to N.
+  assert(sum[0] == T(N));
+
+  // Pre-increment returns updated value: will be in [1, N].
+  auto min_e = std::min_element(output_begin, output_end);
+  auto max_e = std::max_element(output_begin, output_end);
+  assert(*min_e == T(1) && *max_e == T(N));
+
+  // Intermediate values should be unique.
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+
+  free(sum, q);
+  free(output, q);
+}
+
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
           access::address_space space, typename T, typename Difference = T,
@@ -230,6 +352,46 @@ void add_post_inc_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void add_post_inc_test_usm_shared(queue q, size_t N) {
+  T *sum = malloc_shared<T>(1, q);
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  sum[0] = T(0);
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         int gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (sum[0]);
+         output[gid] = atm++;
+       });
+     }).wait_and_throw();
+  }
+
+  // All work-items increment by 1, so final value should be equal to N.
+  assert(sum[0] == T(N));
+
+  // Post-increment returns original value: will be in [0, N-1].
+  auto min_e = std::min_element(output_begin, output_end);
+  auto max_e = std::max_element(output_begin, output_end);
+  assert(*min_e == T(0) && *max_e == T(N - 1));
+
+  // Intermediate values should be unique.
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+
+  free(sum, q);
+  free(output, q);
+}
+
 template <access::address_space space, typename T, typename Difference = T,
           memory_order order = memory_order::relaxed,
           memory_scope scope = memory_scope::device>
@@ -241,6 +403,7 @@ void add_test(queue q, size_t N) {
       space == access::address_space::global_space ||
       (space == access::address_space::generic_space && !TEST_GENERIC_IN_LOCAL);
   constexpr bool do_ext_tests = space != access::address_space::generic_space;
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   if constexpr (do_local_tests) {
 #ifdef RUN_DEPRECATED
     if constexpr (do_ext_tests) {
@@ -259,11 +422,23 @@ void add_test(queue q, size_t N) {
                      order, scope>(q, N);
       add_plus_equal_test<::sycl::ext::oneapi::atomic_ref, space, T, Difference,
                           order, scope>(q, N);
+      if (do_usm_tests) {
+        add_fetch_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                  Difference, order, scope>(q, N);
+        add_plus_equal_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space,
+                                       T, Difference, order, scope>(q, N);
+      }
       if constexpr (!std::is_floating_point_v<T>) {
         add_pre_inc_test<::sycl::ext::oneapi::atomic_ref, space, T, Difference,
                          order, scope>(q, N);
         add_post_inc_test<::sycl::ext::oneapi::atomic_ref, space, T, Difference,
                           order, scope>(q, N);
+        if (do_usm_tests) {
+          add_pre_inc_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                      Difference, order, scope>(q, N);
+          add_post_inc_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space,
+                                       T, Difference, order, scope>(q, N);
+        }
       }
     }
 #else
@@ -271,11 +446,23 @@ void add_test(queue q, size_t N) {
                                                                            N);
     add_plus_equal_test<::sycl::atomic_ref, space, T, Difference, order, scope>(
         q, N);
+    if (do_usm_tests) {
+      add_fetch_test_usm_shared<::sycl::atomic_ref, space, T, Difference, order,
+                                scope>(q, N);
+      add_plus_equal_test_usm_shared<::sycl::atomic_ref, space, T, Difference,
+                                     order, scope>(q, N);
+    }
     if constexpr (!std::is_floating_point_v<T>) {
       add_pre_inc_test<::sycl::atomic_ref, space, T, Difference, order, scope>(
           q, N);
       add_post_inc_test<::sycl::atomic_ref, space, T, Difference, order, scope>(
           q, N);
+      if (do_usm_tests) {
+        add_pre_inc_test_usm_shared<::sycl::atomic_ref, space, T, Difference,
+                                    order, scope>(q, N);
+        add_post_inc_test_usm_shared<::sycl::atomic_ref, space, T, Difference,
+                                     order, scope>(q, N);
+      }
     }
 #endif
   }
diff --git a/sycl/test-e2e/AtomicRef/and.h b/sycl/test-e2e/AtomicRef/and.h
index 5762125a40b21..f31a396dcbfea 100644
--- a/sycl/test-e2e/AtomicRef/and.h
+++ b/sycl/test-e2e/AtomicRef/and.h
@@ -97,6 +97,44 @@ void and_global_test(queue q) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void and_global_test_usm_shared(queue q) {
+  const size_t N = 32;
+  const T initial = T((1ll << N) - 1);
+  T *cum = malloc_shared<T>(1, q);
+  cum[0] = initial;
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, T(0));
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         size_t gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (cum[0]);
+         output[gid] = atm.fetch_and(~T(1ll << gid), order);
+       });
+     }).wait_and_throw();
+  }
+
+  // Final value should be equal to 0.
+  assert(cum[0] == 0);
+
+  // All other values should be unique; each work-item sets one bit to 0.
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+
+  free(cum, q);
+  free(output, q);
+}
+
 template <access::address_space space, typename T,
           memory_order order = memory_order::relaxed,
           memory_scope scope = memory_scope::device>
@@ -108,6 +146,7 @@ void and_test(queue q) {
       space == access::address_space::global_space ||
       (space == access::address_space::generic_space && !TEST_GENERIC_IN_LOCAL);
   constexpr bool do_ext_tests = space != access::address_space::generic_space;
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   if constexpr (do_local_tests) {
 #ifdef RUN_DEPRECATED
     if constexpr (do_ext_tests) {
@@ -123,9 +162,16 @@ void and_test(queue q) {
     if constexpr (do_ext_tests) {
       and_global_test<::sycl::ext::oneapi::atomic_ref, space, T, order, scope>(
           q);
+      if (do_usm_tests) {
+        and_global_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                   order, scope>(q);
+      }
     }
 #else
     and_global_test<::sycl::atomic_ref, space, T, order, scope>(q);
+    if (do_usm_tests) {
+      and_global_test_usm_shared<::sycl::atomic_ref, space, T, order, scope>(q);
+    }
 #endif
   }
 }
diff --git a/sycl/test-e2e/AtomicRef/assignment.h b/sycl/test-e2e/AtomicRef/assignment.h
index 68e0e7c1c7f75..0626be60d82dc 100644
--- a/sycl/test-e2e/AtomicRef/assignment.h
+++ b/sycl/test-e2e/AtomicRef/assignment.h
@@ -40,17 +40,63 @@ void assignment_test(queue q, size_t N) {
   assert(assignment >= T(0) && assignment <= T(N - 1));
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space address_space, typename T>
+class assignment_usm_kernel;
+
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space address_space, typename T>
+void assignment_test_usm_shared(queue q, size_t N) {
+  T initial = T(N);
+  T *st = malloc_shared<T>(1, q);
+  st[0] = initial;
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for<assignment_usm_kernel<AtomicRef, address_space, T>>(
+           range<1>(N), [=](item<1> it) {
+             size_t gid = it.get_id(0);
+             auto atm = AtomicRef<T, memory_order::relaxed,
+                                  memory_scope::device, address_space>(st[0]);
+             atm = T(gid);
+           });
+     }).wait_and_throw();
+  }
+
+  // The initial value should have been overwritten by a work-item ID.
+  // Atomicity isn't tested here, but support for assignment() is.
+  assert(st[0] != initial);
+  assert(st[0] >= T(0) && st[0] <= T(N - 1));
+
+  free(st, q);
+}
+
 template <typename T> void assignment_test(queue q, size_t N) {
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
 #ifdef RUN_DEPRECATED
   assignment_test<::sycl::ext::oneapi::atomic_ref,
                   access::address_space::global_space, T>(q, N);
+  if (do_usm_tests) {
+    assignment_test_usm_shared<::sycl::ext::oneapi::atomic_ref,
+                               access::address_space::global_space, T>(q, N);
+  }
 #else
   assignment_test<::sycl::atomic_ref, access::address_space::global_space, T>(
       q, N);
+  if (do_usm_tests) {
+    assignment_test_usm_shared<::sycl::atomic_ref,
+                               access::address_space::global_space, T>(q, N);
+  }
 #endif
 }
 
 template <typename T> void assignment_generic_test(queue q, size_t N) {
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   assignment_test<::sycl::atomic_ref, access::address_space::generic_space, T>(
       q, N);
+  if (do_usm_tests) {
+    assignment_test_usm_shared<::sycl::atomic_ref,
+                               access::address_space::generic_space, T>(q, N);
+  }
 }
diff --git a/sycl/test-e2e/AtomicRef/compare_exchange.h b/sycl/test-e2e/AtomicRef/compare_exchange.h
index 01bef9d567d01..acc774bffd059 100644
--- a/sycl/test-e2e/AtomicRef/compare_exchange.h
+++ b/sycl/test-e2e/AtomicRef/compare_exchange.h
@@ -114,6 +114,50 @@ void compare_exchange_global_test(queue q, size_t N) {
   }
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void compare_exchange_global_test_usm_shared(queue q, size_t N) {
+  const T initial = T(N);
+  T *exc = malloc_shared<T>(1, q);
+  exc[0] = initial;
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, T(0));
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         size_t gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (exc[0]);
+         T result = initial; // Avoid copying pointer
+         bool success = atm.compare_exchange_strong(result, (T)gid, order);
+         if (success) {
+           output[gid] = result;
+         } else {
+           output[gid] = T(gid);
+         }
+       });
+     }).wait_and_throw();
+  }
+
+  // Only one work-item should have received the initial sentinel value.
+  assert(std::count(output_begin, output_end, initial) == 1);
+
+  // All other values should be the index itself or the sentinel value.
+  for (size_t i = 0; i < N; ++i) {
+    assert(output[i] == T(i) || output[i] == initial);
+  }
+
+  free(exc, q);
+  free(output, q);
+}
+
 template <access::address_space space, typename T,
           memory_order order = memory_order::relaxed,
           memory_scope scope = memory_scope::device>
@@ -125,6 +169,7 @@ void compare_exchange_test(queue q, size_t N) {
       space == access::address_space::global_space ||
       (space == access::address_space::generic_space && !TEST_GENERIC_IN_LOCAL);
   constexpr bool do_ext_tests = space != access::address_space::generic_space;
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   if constexpr (do_local_tests) {
 #ifdef RUN_DEPRECATED
     if constexpr (do_ext_tests) {
@@ -141,10 +186,18 @@ void compare_exchange_test(queue q, size_t N) {
     if constexpr (do_ext_tests) {
       compare_exchange_global_test<::sycl::ext::oneapi::atomic_ref, space, T,
                                    order, scope>(q, N);
+      if (do_usm_tests) {
+        compare_exchange_global_test_usm_shared<::sycl::ext::oneapi::atomic_ref,
+                                                space, T, order, scope>(q, N);
+      }
     }
 #else
     compare_exchange_global_test<::sycl::atomic_ref, space, T, order, scope>(q,
                                                                              N);
+    if (do_usm_tests) {
+      compare_exchange_global_test_usm_shared<::sycl::atomic_ref, space, T,
+                                              order, scope>(q, N);
+    }
 #endif
   }
 }
diff --git a/sycl/test-e2e/AtomicRef/exchange.h b/sycl/test-e2e/AtomicRef/exchange.h
index ae7d182c86530..2316d93a04105 100644
--- a/sycl/test-e2e/AtomicRef/exchange.h
+++ b/sycl/test-e2e/AtomicRef/exchange.h
@@ -99,6 +99,47 @@ void exchange_global_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void exchange_global_test_usm_shared(queue q, size_t N) {
+  const T initial = T(N);
+  T *exc = malloc_shared<T>(1, q);
+  exc[0] = initial;
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, T(0));
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         size_t gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (exc[0]);
+         output[gid] = atm.exchange(T(gid), order);
+       });
+     }).wait_and_throw();
+  }
+
+  // Only one work-item should have received the initial sentinel value.
+  assert(std::count(output_begin, output_end, initial) == 1);
+
+  // All other values should be unique; each work-item replaces the value it
+  // reads with its own ID
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+  for (int i = 0; i < N; ++i) {
+    assert((output[i] >= 0 && output[i] < N) || output[i] == initial);
+  }
+
+  free(exc, q);
+  free(output, q);
+}
+
 template <access::address_space space, typename T,
           memory_order order = memory_order::relaxed,
           memory_scope scope = memory_scope::device>
@@ -110,6 +151,7 @@ void exchange_test(queue q, size_t N) {
       space == access::address_space::global_space ||
       (space == access::address_space::generic_space && !TEST_GENERIC_IN_LOCAL);
   constexpr bool do_ext_tests = space != access::address_space::generic_space;
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   if constexpr (do_local_tests) {
 #ifdef RUN_DEPRECATED
     if constexpr (do_ext_tests) {
@@ -125,9 +167,17 @@ void exchange_test(queue q, size_t N) {
     if constexpr (do_ext_tests) {
       exchange_global_test<::sycl::ext::oneapi::atomic_ref, space, T, order,
                            scope>(q, N);
+      if (do_usm_tests) {
+        exchange_global_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space,
+                                        T, order, scope>(q, N);
+      }
     }
 #else
     exchange_global_test<::sycl::atomic_ref, space, T, order, scope>(q, N);
+    if (do_usm_tests) {
+      exchange_global_test_usm_shared<::sycl::atomic_ref, space, T, order,
+                                      scope>(q, N);
+    }
 #endif
   }
 }
diff --git a/sycl/test-e2e/AtomicRef/load.h b/sycl/test-e2e/AtomicRef/load.h
index f40cd60fdd42b..31eb77785bb87 100644
--- a/sycl/test-e2e/AtomicRef/load.h
+++ b/sycl/test-e2e/AtomicRef/load.h
@@ -81,6 +81,37 @@ void load_global_test(queue q, size_t N) {
                      [&](T x) { return (x == initial); }));
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void load_global_test_usm_shared(queue q, size_t N) {
+  T initial = T(42);
+  T *ld = malloc_shared<T>(1, q);
+  ld[0] = initial;
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, 0);
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         size_t gid = it.get_id(0);
+         auto atm = AtomicRef<T, memory_order::relaxed, scope, space>(ld[0]);
+         output[gid] = atm.load(order);
+       });
+     }).wait_and_throw();
+  }
+
+  // All work-items should read the same value.
+  // Atomicity isn't tested here, but support for load() is.
+  assert(std::all_of(output_begin, output_end,
+                     [&](T x) { return (x == initial); }));
+
+  free(ld, q);
+  free(output, q);
+}
+
 template <access::address_space space, typename T,
           memory_order order = memory_order::relaxed,
           memory_scope scope = memory_scope::device>
@@ -92,6 +123,7 @@ void load_test(queue q, size_t N) {
       space == access::address_space::global_space ||
       (space == access::address_space::generic_space && !TEST_GENERIC_IN_LOCAL);
   constexpr bool do_ext_tests = space != access::address_space::generic_space;
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   if constexpr (do_local_tests) {
 #ifdef RUN_DEPRECATED
     if constexpr (do_ext_tests) {
@@ -107,9 +139,17 @@ void load_test(queue q, size_t N) {
     if constexpr (do_ext_tests) {
       load_global_test<::sycl::ext::oneapi::atomic_ref, space, T, order, scope>(
           q, N);
+      if (do_usm_tests) {
+        load_global_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                    order, scope>(q, N);
+      }
     }
 #else
     load_global_test<::sycl::atomic_ref, space, T, order, scope>(q, N);
+    if (do_usm_tests) {
+      load_global_test_usm_shared<::sycl::atomic_ref, space, T, order, scope>(
+          q, N);
+    }
 #endif
   }
 }
diff --git a/sycl/test-e2e/AtomicRef/max.h b/sycl/test-e2e/AtomicRef/max.h
index 0eed88cac666c..f176b37a32681 100644
--- a/sycl/test-e2e/AtomicRef/max.h
+++ b/sycl/test-e2e/AtomicRef/max.h
@@ -108,6 +108,50 @@ void max_global_test(queue q, size_t N) {
   }
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void max_global_test_usm_shared(queue q, size_t N) {
+  T initial = std::numeric_limits<T>::lowest();
+  T *val = malloc_shared<T>(1, q);
+  val[0] = initial;
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, std::numeric_limits<T>::max());
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         int gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (val[0]);
+
+         // +max/2 to ensure correct signed/unsigned operation is applied
+         output[gid] =
+             atm.fetch_max(T(gid) + std::numeric_limits<T>::max() / 2, order);
+       });
+     }).wait_and_throw();
+  }
+
+  assert(val[0] == N - 1 + std::numeric_limits<T>::max() / 2);
+
+  // Only one work-item should have received the initial value.
+  assert(std::count(output_begin, output_end, initial) == 1);
+
+  // fetch_max returns original value.
+  // Intermediate values should all be >= initial value.
+  for (int i = 0; i < N; ++i) {
+    assert(output[i] >= initial && output[i] <= val[0]);
+  }
+
+  free(val, q);
+  free(output, q);
+}
+
 template <access::address_space space, typename T,
           memory_order order = memory_order::relaxed,
           memory_scope scope = memory_scope::device>
@@ -119,6 +163,7 @@ void max_test(queue q, size_t N) {
       space == access::address_space::global_space ||
       (space == access::address_space::generic_space && !TEST_GENERIC_IN_LOCAL);
   constexpr bool do_ext_tests = space != access::address_space::generic_space;
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   if constexpr (do_local_tests) {
 #ifdef RUN_DEPRECATED
     if constexpr (do_ext_tests) {
@@ -134,9 +179,17 @@ void max_test(queue q, size_t N) {
     if constexpr (do_ext_tests) {
       max_global_test<::sycl::ext::oneapi::atomic_ref, space, T, order, scope>(
           q, N);
+      if (do_usm_tests) {
+        max_global_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                   order, scope>(q, N);
+      }
     }
 #else
     max_global_test<::sycl::atomic_ref, space, T, order, scope>(q, N);
+    if (do_usm_tests) {
+      max_global_test_usm_shared<::sycl::atomic_ref, space, T, order, scope>(q,
+                                                                             N);
+    }
 #endif
   }
 }
diff --git a/sycl/test-e2e/AtomicRef/min.h b/sycl/test-e2e/AtomicRef/min.h
index 5b1ff078ba922..df606d3b20c98 100644
--- a/sycl/test-e2e/AtomicRef/min.h
+++ b/sycl/test-e2e/AtomicRef/min.h
@@ -106,6 +106,48 @@ void min_global_test(queue q, size_t N) {
   }
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void min_global_test_usm_shared(queue q, size_t N) {
+  T initial = std::numeric_limits<T>::max();
+  T *val = malloc_shared<T>(1, q);
+  val[0] = initial;
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, 0);
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         int gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (val[0]);
+         output[gid] = atm.fetch_min(T(gid), order);
+       });
+     }).wait_and_throw();
+  }
+
+  // Final value should be equal to 0.
+  assert(val[0] == 0);
+
+  // Only one work-item should have received the initial value.
+  assert(std::count(output_begin, output_end, initial) == 1);
+
+  // fetch_min returns original value.
+  // Intermediate values should all be <= initial value.
+  for (int i = 0; i < N; ++i) {
+    assert(output[i] <= initial && output[i] >= val[0]);
+  }
+
+  free(val, q);
+  free(output, q);
+}
+
 template <access::address_space space, typename T,
           memory_order order = memory_order::relaxed,
           memory_scope scope = memory_scope::device>
@@ -117,6 +159,7 @@ void min_test(queue q, size_t N) {
       space == access::address_space::global_space ||
       (space == access::address_space::generic_space && !TEST_GENERIC_IN_LOCAL);
   constexpr bool do_ext_tests = space != access::address_space::generic_space;
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   if constexpr (do_local_tests) {
 #ifdef RUN_DEPRECATED
     if constexpr (do_ext_tests) {
@@ -132,9 +175,17 @@ void min_test(queue q, size_t N) {
     if constexpr (do_ext_tests) {
       min_global_test<::sycl::ext::oneapi::atomic_ref, space, T, order, scope>(
           q, N);
+      if (do_usm_tests) {
+        min_global_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                   order, scope>(q, N);
+      }
     }
 #else
     min_global_test<::sycl::atomic_ref, space, T, order, scope>(q, N);
+    if (do_usm_tests) {
+      min_global_test_usm_shared<::sycl::atomic_ref, space, T, order, scope>(q,
+                                                                             N);
+    }
 #endif
   }
 }
diff --git a/sycl/test-e2e/AtomicRef/or.h b/sycl/test-e2e/AtomicRef/or.h
index 5395dba28dd78..5f600888193c3 100644
--- a/sycl/test-e2e/AtomicRef/or.h
+++ b/sycl/test-e2e/AtomicRef/or.h
@@ -97,6 +97,44 @@ void or_global_test(queue q) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void or_global_test_usm_shared(queue q) {
+  const size_t N = 32;
+  const T initial = 0;
+  T *cum = malloc_shared<T>(1, q);
+  cum[0] = initial;
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, T(0));
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         size_t gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (cum[0]);
+         output[gid] = atm.fetch_or(T(1ll << gid), order);
+       });
+     }).wait_and_throw();
+  }
+
+  // Final value should be equal to N ones.
+  assert(cum[0] == T((1ll << N) - 1));
+
+  // All other values should be unique; each work-item sets one bit to 1.
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+
+  free(cum, q);
+  free(output, q);
+}
+
 template <access::address_space space, typename T,
           memory_order order = memory_order::relaxed,
           memory_scope scope = memory_scope::device>
@@ -108,6 +146,7 @@ void or_test(queue q) {
       space == access::address_space::global_space ||
       (space == access::address_space::generic_space && !TEST_GENERIC_IN_LOCAL);
   constexpr bool do_ext_tests = space != access::address_space::generic_space;
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   if constexpr (do_local_tests) {
 #ifdef RUN_DEPRECATED
     if constexpr (do_ext_tests) {
@@ -122,9 +161,16 @@ void or_test(queue q) {
     if constexpr (do_ext_tests) {
       or_global_test<::sycl::ext::oneapi::atomic_ref, space, T, order, scope>(
           q);
+      if (do_usm_tests) {
+        or_global_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                  order, scope>(q);
+      }
     }
 #else
     or_global_test<::sycl::atomic_ref, space, T, order, scope>(q);
+    if (do_usm_tests) {
+      or_global_test_usm_shared<::sycl::atomic_ref, space, T, order, scope>(q);
+    }
 #endif
   }
 }
diff --git a/sycl/test-e2e/AtomicRef/store.h b/sycl/test-e2e/AtomicRef/store.h
index a2ea45aaab819..e3adefbe9b0cf 100644
--- a/sycl/test-e2e/AtomicRef/store.h
+++ b/sycl/test-e2e/AtomicRef/store.h
@@ -39,6 +39,33 @@ void store_global_test(queue q, size_t N) {
   assert(store >= T(0) && store <= T(N - 1));
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void store_global_test_usm_shared(queue q, size_t N) {
+  T initial = T(N);
+  T *st = malloc_shared<T>(1, q);
+  st[0] = initial;
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         size_t gid = it.get_id(0);
+         auto atm = AtomicRef<T, memory_order::relaxed, scope, space>(st[0]);
+         atm.store(T(gid), order);
+       });
+     }).wait_and_throw();
+  }
+
+  // The initial value should have been overwritten by a work-item ID.
+  // Atomicity isn't tested here, but support for store() is.
+  assert(st[0] != initial);
+  assert(st[0] >= T(0) && st[0] <= T(N - 1));
+
+  free(st, q);
+}
+
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
           access::address_space space, typename T,
@@ -80,6 +107,7 @@ void store_test(queue q, size_t N) {
       space == access::address_space::global_space ||
       (space == access::address_space::generic_space && !TEST_GENERIC_IN_LOCAL);
   constexpr bool do_ext_tests = space != access::address_space::generic_space;
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   if constexpr (do_local_tests) {
 #ifdef RUN_DEPRECATED
     if constexpr (do_ext_tests) {
@@ -95,9 +123,17 @@ void store_test(queue q, size_t N) {
     if constexpr (do_ext_tests) {
       store_global_test<::sycl::ext::oneapi::atomic_ref, space, T, order,
                         scope>(q, N);
+      if (do_usm_tests) {
+        store_global_test_shared_usm<::sycl::ext::oneapi::atomic_ref, space, T,
+                                     order, scope>(q, N);
+      }
     }
 #else
     store_global_test<::sycl::atomic_ref, space, T, order, scope>(q, N);
+    if (do_usm_tests) {
+      store_global_test_usm_shared<::sycl::atomic_ref, space, T, order, scope>(
+          q, N);
+    }
 #endif
   }
 }
diff --git a/sycl/test-e2e/AtomicRef/sub.h b/sycl/test-e2e/AtomicRef/sub.h
index 2f0cab5f433a8..2aa671c86a2e9 100644
--- a/sycl/test-e2e/AtomicRef/sub.h
+++ b/sycl/test-e2e/AtomicRef/sub.h
@@ -105,6 +105,47 @@ void sub_fetch_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void sub_fetch_test_usm_shared(queue q, size_t N) {
+  T *val = malloc_shared<T>(1, q);
+  val[0] = T(N);
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, T(0));
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         int gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (val[0]);
+         output[gid] = atm.fetch_sub(Difference(1), order);
+       });
+     }).wait_and_throw();
+  }
+
+  // All work-items decrement by 1, so final value should be equal to 0.
+  assert(val[0] == T(0));
+
+  // Fetch returns original value: will be in [1, N]
+  auto min_e = std::min_element(output_begin, output_end);
+  auto max_e = std::max_element(output_begin, output_end);
+  assert(*min_e == T(1) && *max_e == T(N));
+
+  // Intermediate values should be unique.
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+
+  free(val, q);
+  free(output, q);
+}
+
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
           access::address_space space, typename T, typename Difference = T,
@@ -147,6 +188,47 @@ void sub_minus_equal_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void sub_minus_equal_test_usm_shared(queue q, size_t N) {
+  T *val = malloc_shared<T>(1, q);
+  val[0] = T(N);
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, T(0));
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         int gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (val[0]);
+         output[gid] = atm -= Difference(1);
+       });
+     }).wait_and_throw();
+  }
+
+  // All work-items decrement by 1, so final value should be equal to 0.
+  assert(val[0] == T(0));
+
+  // -= returns updated value: will be in [0, N-1].
+  auto min_e = std::min_element(output_begin, output_end);
+  auto max_e = std::max_element(output_begin, output_end);
+  assert(*min_e == T(0) && *max_e == T(N - 1));
+
+  // Intermediate values should be unique.
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+
+  free(val, q);
+  free(output, q);
+}
+
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
           access::address_space space, typename T, typename Difference = T,
@@ -189,6 +271,47 @@ void sub_pre_dec_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void sub_pre_dec_test_usm_shared(queue q, size_t N) {
+  T *val = malloc_shared<T>(1, q);
+  val[0] = T(N);
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, T(0));
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         int gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (val[0]);
+         output[gid] = --atm;
+       });
+     }).wait_and_throw();
+  }
+
+  // All work-items decrement by 1, so final value should be equal to 0.
+  assert(val[0] == T(0));
+
+  // Pre-decrement returns updated value: will be in [0, N-1].
+  auto min_e = std::min_element(output_begin, output_end);
+  auto max_e = std::max_element(output_begin, output_end);
+  assert(*min_e == T(0) && *max_e == T(N - 1));
+
+  // Intermediate values should be unique.
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+
+  free(val, q);
+  free(output, q);
+}
+
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
           access::address_space space, typename T, typename Difference = T,
@@ -231,6 +354,47 @@ void sub_post_dec_test(queue q, size_t N) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T, typename Difference = T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void sub_post_dec_test_usm_shared(queue q, size_t N) {
+  T *val = malloc_shared<T>(1, q);
+  val[0] = T(N);
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, T(0));
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         int gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (val[0]);
+         output[gid] = atm--;
+       });
+     }).wait_and_throw();
+  }
+
+  // All work-items decrement by 1, so final value should be equal to 0.
+  assert(val[0] == T(0));
+
+  // Post-decrement returns original value: will be in [1, N]
+  auto min_e = std::min_element(output_begin, output_end);
+  auto max_e = std::max_element(output_begin, output_end);
+  assert(*min_e == T(1) && *max_e == T(N));
+
+  // Intermediate values should be unique.
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+
+  free(val, q);
+  free(output, q);
+}
+
 template <access::address_space space, typename T, typename Difference = T,
           memory_order order = memory_order::relaxed,
           memory_scope scope = memory_scope::device>
@@ -242,6 +406,7 @@ void sub_test(queue q, size_t N) {
       space == access::address_space::global_space ||
       (space == access::address_space::generic_space && !TEST_GENERIC_IN_LOCAL);
   constexpr bool do_ext_tests = space != access::address_space::generic_space;
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   if constexpr (do_local_tests) {
 #ifdef RUN_DEPRECATED
     if constexpr (do_ext_tests) {
@@ -258,13 +423,21 @@ void sub_test(queue q, size_t N) {
     if constexpr (do_ext_tests) {
       sub_fetch_test<::sycl::ext::oneapi::atomic_ref, space, T, Difference,
                      order, scope>(q, N);
+      sub_fetch_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                Difference, order, scope>(q, N);
       sub_minus_equal_test<::sycl::ext::oneapi::atomic_ref, space, T,
                            Difference, order, scope>(q, N);
+      sub_minus_equal_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                      Difference, order, scope>(q, N);
       if constexpr (!std::is_floating_point_v<T>) {
         sub_pre_dec_test<::sycl::ext::oneapi::atomic_ref, space, T, Difference,
                          order, scope>(q, N);
+        sub_pre_dec_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                    Difference, order, scope>(q, N);
         sub_post_dec_test<::sycl::ext::oneapi::atomic_ref, space, T, Difference,
                           order, scope>(q, N);
+        sub_post_dec_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                     Difference, order, scope>(q, N);
       }
     }
 #else
@@ -272,11 +445,23 @@ void sub_test(queue q, size_t N) {
                                                                            N);
     sub_minus_equal_test<::sycl::atomic_ref, space, T, Difference, order,
                          scope>(q, N);
+    if (do_usm_tests) {
+      sub_fetch_test_usm_shared<::sycl::atomic_ref, space, T, Difference, order,
+                                scope>(q, N);
+      sub_minus_equal_test_usm_shared<::sycl::atomic_ref, space, T, Difference,
+                                      order, scope>(q, N);
+    }
     if constexpr (!std::is_floating_point_v<T>) {
       sub_pre_dec_test<::sycl::atomic_ref, space, T, Difference, order, scope>(
           q, N);
       sub_post_dec_test<::sycl::atomic_ref, space, T, Difference, order, scope>(
           q, N);
+      if (do_usm_tests) {
+        sub_pre_dec_test_usm_shared<::sycl::atomic_ref, space, T, Difference,
+                                    order, scope>(q, N);
+        sub_post_dec_test_usm_shared<::sycl::atomic_ref, space, T, Difference,
+                                     order, scope>(q, N);
+      }
     }
 #endif
   }
diff --git a/sycl/test-e2e/AtomicRef/xor.h b/sycl/test-e2e/AtomicRef/xor.h
index 0d64dab8785d3..3375a1a3ab052 100644
--- a/sycl/test-e2e/AtomicRef/xor.h
+++ b/sycl/test-e2e/AtomicRef/xor.h
@@ -97,6 +97,44 @@ void xor_global_test(queue q) {
   assert(std::unique(output.begin(), output.end()) == output.end());
 }
 
+template <template <typename, memory_order, memory_scope, access::address_space>
+          class AtomicRef,
+          access::address_space space, typename T,
+          memory_order order = memory_order::relaxed,
+          memory_scope scope = memory_scope::device>
+void xor_global_test_usm_shared(queue q) {
+  const size_t N = 32;
+  const T initial = 0;
+  T *cum = malloc_shared<T>(1, q);
+  cum[0] = initial;
+  T *output = malloc_shared<T>(N, q);
+  T *output_begin = &output[0], *output_end = &output[N];
+  std::fill(output_begin, output_end, T(0));
+  {
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for(range<1>(N), [=](item<1> it) {
+         size_t gid = it.get_id(0);
+         auto atm = AtomicRef < T,
+              (order == memory_order::acquire || order == memory_order::release)
+                  ? memory_order::relaxed
+                  : order,
+              scope, space > (cum[0]);
+         output[gid] = atm.fetch_xor(T(1ll << gid), order);
+       });
+     }).wait_and_throw();
+  }
+
+  // Final value should be equal to N ones.
+  assert(cum[0] == T((1ll << N) - 1));
+
+  // All other values should be unique; each work-item sets one bit to 1.
+  std::sort(output_begin, output_end);
+  assert(std::unique(output_begin, output_end) == output_end);
+
+  free(cum, q);
+  free(output, q);
+}
+
 template <access::address_space space, typename T,
           memory_order order = memory_order::relaxed,
           memory_scope scope = memory_scope::device>
@@ -108,6 +146,7 @@ void xor_test(queue q) {
       space == access::address_space::global_space ||
       (space == access::address_space::generic_space && !TEST_GENERIC_IN_LOCAL);
   constexpr bool do_ext_tests = space != access::address_space::generic_space;
+  bool do_usm_tests = q.get_device().has(aspect::usm_shared_allocations);
   if constexpr (do_local_tests) {
 #ifdef RUN_DEPRECATED
     if constexpr (do_ext_tests) {
@@ -123,9 +162,16 @@ void xor_test(queue q) {
     if constexpr (do_ext_tests) {
       xor_global_test<::sycl::ext::oneapi::atomic_ref, space, T, order, scope>(
           q);
+      if (do_usm_tests) {
+        xor_global_test_usm_shared<::sycl::ext::oneapi::atomic_ref, space, T,
+                                   order, scope>(q);
+      }
     }
 #else
     xor_global_test<::sycl::atomic_ref, space, T, order, scope>(q);
+    if (do_usm_tests) {
+      xor_global_test_usm_shared<::sycl::atomic_ref, space, T, order, scope>(q);
+    }
 #endif
   }
 }