diff --git a/sycl/include/sycl/ext/intel/esimd/detail/sycl_util.hpp b/sycl/include/sycl/ext/intel/esimd/detail/sycl_util.hpp
index 891c54c11954f..2690a1363fb46 100644
--- a/sycl/include/sycl/ext/intel/esimd/detail/sycl_util.hpp
+++ b/sycl/include/sycl/ext/intel/esimd/detail/sycl_util.hpp
@@ -105,6 +105,16 @@ inline constexpr bool is_accessor_with_v =
     is_device_accessor_with_v<T, Capability> ||
     is_local_accessor_with_v<T, Capability>;
 
+template <typename T>
+inline constexpr bool is_rw_device_accessor_v =
+    is_device_accessor_with_v<T, accessor_mode_cap::can_read> &&
+    is_device_accessor_with_v<T, accessor_mode_cap::can_write>;
+
+template <typename T>
+inline constexpr bool is_rw_local_accessor_v =
+    is_local_accessor_with_v<T, accessor_mode_cap::can_read> &&
+    is_local_accessor_with_v<T, accessor_mode_cap::can_write>;
+
 template <typename T, accessor_mode_cap_val_t Capability, typename RetT>
 using EnableIfAccessor =
     std::enable_if_t<detail::is_device_accessor_with_v<T, Capability>, RetT>;
diff --git a/sycl/include/sycl/ext/intel/esimd/memory.hpp b/sycl/include/sycl/ext/intel/esimd/memory.hpp
index c8c68b6adbdd4..2d45487fe0c7b 100644
--- a/sycl/include/sycl/ext/intel/esimd/memory.hpp
+++ b/sycl/include/sycl/ext/intel/esimd/memory.hpp
@@ -3755,6 +3755,7 @@ template <atomic_op Op, typename T, int N, lsc_data_size DS, cache_hint L1H,
 __ESIMD_API std::enable_if_t<get_num_args<Op>() == 0, simd<T, N>>
 atomic_update_impl(T *p, simd<Toffset, N> offsets, simd_mask<N> pred) {
   static_assert(sizeof(T) > 1, "Unsupported data type");
+  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
   check_atomic<Op, T, N, 0, /*IsLSC*/ true>();
   check_lsc_data_size<T, DS>();
   check_cache_hint<cache_action::atomic, L1H, L2H>();
@@ -3795,6 +3796,7 @@ __ESIMD_API std::enable_if_t<get_num_args<Op>() == 1, simd<T, N>>
 atomic_update_impl(T *p, simd<Toffset, N> offsets, simd<T, N> src0,
                    simd_mask<N> pred) {
   static_assert(sizeof(T) > 1, "Unsupported data type");
+  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
   check_lsc_data_size<T, DS>();
   check_atomic<Op, T, N, 1, /*IsLSC*/ true>();
   check_cache_hint<cache_action::atomic, L1H, L2H>();
@@ -3837,6 +3839,7 @@ __ESIMD_API std::enable_if_t<get_num_args<Op>() == 2, simd<T, N>>
 atomic_update_impl(T *p, simd<Toffset, N> offsets, simd<T, N> src0,
                    simd<T, N> src1, simd_mask<N> pred) {
   static_assert(sizeof(T) > 1, "Unsupported data type");
+  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
   check_lsc_data_size<T, DS>();
   check_atomic<Op, T, N, 2, /*IsLSC*/ true>();
   check_cache_hint<cache_action::atomic, L1H, L2H>();
@@ -3880,8 +3883,10 @@ template <atomic_op Op, typename T, int N,
           typename AccessorTy, typename Toffset>
 __ESIMD_API std::enable_if_t<
     get_num_args<Op>() == 0 &&
-        sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
-        !sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
+        __ESIMD_DNS::is_device_accessor_with_v<
+            AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
+        __ESIMD_DNS::is_device_accessor_with_v<
+            AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>,
     simd<T, N>>
 atomic_update_impl(AccessorTy acc, simd<Toffset, N> byte_offsets,
                    simd_mask<N> pred) {
@@ -3933,8 +3938,10 @@ template <atomic_op Op, typename T, int N, lsc_data_size DS, cache_hint L1H,
           cache_hint L2H, typename AccessorTy, typename Toffset>
 __ESIMD_API std::enable_if_t<
     get_num_args<Op>() == 1 &&
-        sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
-        !sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
+        __ESIMD_DNS::is_device_accessor_with_v<
+            AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
+        __ESIMD_DNS::is_device_accessor_with_v<
+            AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>,
     simd<T, N>>
 atomic_update_impl(AccessorTy acc, simd<Toffset, N> byte_offset,
                    simd<T, N> src0, simd_mask<N> pred) {
diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp
index 9d1312eb953d7..e1080075b5cfc 100644
--- a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp
@@ -2854,8 +2854,11 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
           cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
           typename AccessorTy, typename Toffset>
 __ESIMD_API std::enable_if_t<
-    sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
-        !sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
+    __ESIMD_DNS::is_device_accessor_with_v<
+        AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> &&
+        (Op == __ESIMD_NS::atomic_op::load ||
+         __ESIMD_DNS::is_device_accessor_with_v<
+             AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>),
     __ESIMD_NS::simd<T, N>>
 lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
                   __ESIMD_NS::simd_mask<N> pred) {
@@ -2882,9 +2885,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
           lsc_data_size DS = lsc_data_size::default_size,
           cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
           typename AccessorTy>
-__ESIMD_API std::enable_if_t<
-    sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
-    __ESIMD_NS::simd<T, N>>
+__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>,
+                             __ESIMD_NS::simd<T, N>>
 lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
                   __ESIMD_NS::simd_mask<N> pred) {
   return lsc_slm_atomic_update<Op, T, N, DS>(
@@ -2913,10 +2915,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
           lsc_data_size DS = lsc_data_size::default_size,
           cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
           typename AccessorTy, typename Toffset>
-__ESIMD_API std::enable_if_t<
-    sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
-        !sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
-    __ESIMD_NS::simd<T, N>>
+__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>,
+                             __ESIMD_NS::simd<T, N>>
 lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
                   __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd_mask<N> pred) {
   return __ESIMD_DNS::atomic_update_impl<Op, T, N, DS, L1H, L3H>(acc, offsets,
@@ -2943,9 +2943,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
           lsc_data_size DS = lsc_data_size::default_size,
           cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
           typename AccessorTy>
-__ESIMD_API std::enable_if_t<
-    sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
-    __ESIMD_NS::simd<T, N>>
+__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>,
+                             __ESIMD_NS::simd<T, N>>
 lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
                   __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd_mask<N> pred) {
   return lsc_slm_atomic_update<Op, T, N, DS>(
@@ -2975,10 +2974,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
           lsc_data_size DS = lsc_data_size::default_size,
           cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
           typename AccessorTy, typename Toffset>
-__ESIMD_API std::enable_if_t<
-    sycl::detail::acc_properties::is_accessor_v<AccessorTy> &&
-        !sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
-    __ESIMD_NS::simd<T, N>>
+__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>,
+                             __ESIMD_NS::simd<T, N>>
 lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
                   __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
                   __ESIMD_NS::simd_mask<N> pred) {
@@ -3033,9 +3030,8 @@ template <__ESIMD_NS::atomic_op Op, typename T, int N,
           lsc_data_size DS = lsc_data_size::default_size,
           cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
           typename AccessorTy>
-__ESIMD_API std::enable_if_t<
-    sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>,
-    __ESIMD_NS::simd<T, N>>
+__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>,
+                             __ESIMD_NS::simd<T, N>>
 lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
                   __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
                   __ESIMD_NS::simd_mask<N> pred) {
@@ -3112,9 +3108,7 @@ atomic_update(T *p, simd<Toffset, N> offset, simd_mask<N> mask) {
 
 template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
           typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>>
-__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
-                                 __ESIMD_DNS::get_num_args<Op>() == 0,
-                             simd<T, N>>
+__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 0, simd<T, N>>
 atomic_update(T *p, simd_view<Toffset, RegionTy> offsets,
               simd_mask<N> mask = 1) {
   return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
@@ -3132,24 +3126,20 @@ atomic_update(T *p, Toffset offset, simd_mask<N> mask = 1) {
 
 /// LSC version of the single-argument atomic update.
 template <native::lsc::atomic_op Op, typename T, int N, typename Toffset>
-__ESIMD_API
-    __ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
-                                     __ESIMD_DNS::get_num_args<Op>() == 1,
-                                 simd<T, N>>
-    atomic_update(T *p, simd<Toffset, N> offset, simd<T, N> src0,
-                  simd_mask<N> mask) {
+__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
+                                 __ESIMD_DNS::get_num_args<Op>() == 1,
+                             simd<T, N>>
+atomic_update(T *p, simd<Toffset, N> offset, simd<T, N> src0,
+              simd_mask<N> mask) {
   return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
       p, offset, src0, mask);
 }
 
 template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
           typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>>
-__ESIMD_API
-    __ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
-                                     __ESIMD_DNS::get_num_args<Op>() == 1,
-                                 simd<T, N>>
-    atomic_update(T *p, simd_view<Toffset, RegionTy> offsets, simd<T, N> src0,
-                  simd_mask<N> mask = 1) {
+__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 1, simd<T, N>>
+atomic_update(T *p, simd_view<Toffset, RegionTy> offsets, simd<T, N> src0,
+              simd_mask<N> mask = 1) {
   return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
       p, offsets, src0, mask);
 }
@@ -3179,9 +3169,7 @@ atomic_update(T *p, simd<Toffset, N> offset, simd<T, N> src0, simd<T, N> src1,
 
 template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
           typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>>
-__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
-                                 __ESIMD_DNS::get_num_args<Op>() == 2,
-                             simd<T, N>>
+__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 2, simd<T, N>>
 atomic_update(T *p, simd_view<Toffset, RegionTy> offsets, simd<T, N> src0,
               simd<T, N> src1, simd_mask<N> mask = 1) {
   return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
@@ -3212,8 +3200,7 @@ atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd_mask<N> mask) {
 template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
           typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>,
           typename AccessorTy>
-__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
-                                 __ESIMD_DNS::get_num_args<Op>() == 0 &&
+__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 0 &&
                                  !std::is_pointer_v<AccessorTy>,
                              simd<T, N>>
 atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
@@ -3236,13 +3223,12 @@ atomic_update(AccessorTy acc, Toffset offset, simd_mask<N> mask) {
 /// LSC version of the single-argument atomic update.
 template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
           typename AccessorTy>
-__ESIMD_API
-    __ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
-                                     __ESIMD_DNS::get_num_args<Op>() == 1 &&
-                                     !std::is_pointer_v<AccessorTy>,
-                                 simd<T, N>>
-    atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<T, N> src0,
-                  simd_mask<N> mask) {
+__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
+                                 __ESIMD_DNS::get_num_args<Op>() == 1 &&
+                                 !std::is_pointer_v<AccessorTy>,
+                             simd<T, N>>
+atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<T, N> src0,
+              simd_mask<N> mask) {
   return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
       acc, offset, src0, mask);
 }
@@ -3250,13 +3236,11 @@ __ESIMD_API
 template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
           typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>,
           typename AccessorTy>
-__ESIMD_API
-    __ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
-                                     __ESIMD_DNS::get_num_args<Op>() == 1 &&
-                                     !std::is_pointer_v<AccessorTy>,
-                                 simd<T, N>>
-    atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
-                  simd<T, N> src0, simd_mask<N> mask) {
+__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 1 &&
+                                 !std::is_pointer_v<AccessorTy>,
+                             simd<T, N>>
+atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
+              simd<T, N> src0, simd_mask<N> mask) {
   return __ESIMD_ENS::lsc_atomic_update<detail::to_atomic_op<Op>(), T, N>(
       acc, offsets, src0, mask);
 }
@@ -3292,8 +3276,7 @@ atomic_update(AccessorTy acc, simd<Toffset, N> offset, simd<T, N> src0,
 template <native::lsc::atomic_op Op, typename T, int N, typename Toffset,
           typename RegionTy = __ESIMD_NS::region1d_t<Toffset, N, 1>,
           typename AccessorTy>
-__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> &&
-                                 __ESIMD_DNS::get_num_args<Op>() == 2 &&
+__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 2 &&
                                  !std::is_pointer_v<AccessorTy>,
                              simd<T, N>>
 atomic_update(AccessorTy acc, simd_view<Toffset, RegionTy> offsets,
diff --git a/sycl/test/esimd/lsc_atomic.cpp b/sycl/test/esimd/lsc_atomic.cpp
new file mode 100644
index 0000000000000..9c774f07299ff
--- /dev/null
+++ b/sycl/test/esimd/lsc_atomic.cpp
@@ -0,0 +1,131 @@
+// RUN: not %clangxx -fsycl -fsycl-device-only -fsyntax-only -Wno-unused-command-line-argument %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
+// RUN: not %clangxx %fsycl-host-only -fsyntax-only -Wno-unused-command-line-argument %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
+
+// This test checks compilation of ESIMD lsc atomic APIs.
+
+#include <limits>
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+#include <utility>
+using namespace sycl::ext::intel::esimd;
+using namespace sycl::ext::intel::experimental::esimd;
+using namespace sycl;
+
+// --- Postive tests.
+
+void kernel0(uint32_t *ptr) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  lsc_atomic_update<atomic_op::inc, uint32_t, 32>(ptr, offsets, 1);
+}
+void kernel1(uint32_t *ptr) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  simd<uint32_t, 32> v1(0, 1);
+  lsc_atomic_update<atomic_op::add, uint32_t, 32>(ptr, offsets, v1, 1);
+}
+template <class T> void kernel2(T *ptr) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  simd<T, 32> v1(0, 1);
+  lsc_atomic_update<atomic_op::cmpxchg, T, 32>(ptr, offsets, v1, v1, 1);
+}
+
+template void kernel2<uint32_t>(uint32_t *) SYCL_ESIMD_FUNCTION;
+
+void kernel3(accessor<uint32_t, 1, access::mode::read_write,
+                      access::target::device> &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+
+  lsc_atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
+}
+
+void kernel4(accessor<uint32_t, 1, access::mode::read_write,
+                      access::target::device> &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  simd<uint32_t, 32> v1(0, 1);
+
+  lsc_atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
+}
+
+void kernel5(accessor<uint32_t, 1, access::mode::read_write,
+                      access::target::device> &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  simd<uint32_t, 32> v1(0, 1);
+
+  lsc_atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
+}
+
+void kernel6(local_accessor<uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+
+  lsc_atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
+}
+
+void kernel7(local_accessor<uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  simd<uint32_t, 32> v1(0, 1);
+
+  lsc_atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
+}
+
+void kernel8(local_accessor<uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  simd<uint32_t, 32> v1(0, 1);
+
+  lsc_atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
+}
+
+// --- Negative tests.
+
+// Incompatible mode (read).
+void kernel9(accessor<uint32_t, 1, access::mode::read, access::target::device>
+                 &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+
+  // CHECK: lsc_atomic.cpp:84{{.*}}error: no matching function for call to 'lsc_atomic_update'
+  lsc_atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
+}
+
+// Incompatible mode (read).
+void kernel10(accessor<uint32_t, 1, access::mode::read, access::target::device>
+                  &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  simd<uint32_t, 32> v1(0, 1);
+
+  // CHECK: lsc_atomic.cpp:94{{.*}}error: no matching function for call to 'lsc_atomic_update'
+  lsc_atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
+}
+
+// Incompatible mode (read).
+void kernel11(accessor<uint32_t, 1, access::mode::read, access::target::device>
+                  &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  simd<uint32_t, 32> v1(0, 1);
+
+  // CHECK: lsc_atomic.cpp:104{{.*}}error: no matching function for call to 'lsc_atomic_update'
+  lsc_atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
+}
+
+// Incompatible mode (read).
+void kernel12(local_accessor<const uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+
+  // CHECK: lsc_atomic.cpp:112{{.*}}error: no matching function for call to 'lsc_atomic_update'
+  lsc_atomic_update<atomic_op::inc, uint32_t, 32>(buf, offsets, 1);
+}
+
+// Incompatible mode (read).
+void kernel13(local_accessor<const uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  simd<uint32_t, 32> v1(0, 1);
+
+  // CHECK: lsc_atomic.cpp:121{{.*}}error: no matching function for call to 'lsc_atomic_update'
+  lsc_atomic_update<atomic_op::add, uint32_t, 32>(buf, offsets, v1, 1);
+}
+
+// Incompatible mode (read).
+void kernel8(const local_accessor<const uint32_t, 1> &buf) SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  simd<uint32_t, 32> v1(0, 1);
+
+  // CHECK: lsc_atomic.cpp:130{{.*}}error: no matching function for call to 'lsc_atomic_update'
+  lsc_atomic_update<atomic_op::cmpxchg, uint32_t, 32>(buf, offsets, v1, v1, 1);
+}