diff --git a/sycl/include/sycl/atomic_ref.hpp b/sycl/include/sycl/atomic_ref.hpp
index ccd7d264d8e60..a3d6ebc5a967a 100644
--- a/sycl/include/sycl/atomic_ref.hpp
+++ b/sycl/include/sycl/atomic_ref.hpp
@@ -10,7 +10,8 @@
 
 #include <sycl/access/access.hpp> // for address_space
 #include <sycl/bit_cast.hpp>      // for bit_cast
-#include <sycl/memory_enums.hpp>  // for getStdMemoryOrder, memory_order
+#include <sycl/ext/oneapi/experimental/address_cast.hpp>
+#include <sycl/memory_enums.hpp> // for getStdMemoryOrder, memory_order
 
 #ifdef __SYCL_DEVICE_ONLY__
 #include <sycl/detail/spirv.hpp>
@@ -157,8 +158,16 @@ class atomic_ref_base {
   }
 
 #ifdef __SYCL_DEVICE_ONLY__
+#if defined(__SPIR__)
+  explicit atomic_ref_base(T &ref)
+      : ptr(ext::oneapi::experimental::static_address_cast<AddressSpace>(
+            &ref)) {}
+#else
+  // CUDA/HIP don't support `ext::oneapi::experimental::static_address_cast`
+  // yet.
   explicit atomic_ref_base(T &ref)
       : ptr(address_space_cast<AddressSpace, access::decorated::no>(&ref)) {}
+#endif
 #else
   // FIXME: This reinterpret_cast is UB, but happens to work for now
   explicit atomic_ref_base(T &ref)
diff --git a/sycl/test/check_device_code/atomic_ref.cpp b/sycl/test/check_device_code/atomic_ref.cpp
new file mode 100644
index 0000000000000..f01727a413bac
--- /dev/null
+++ b/sycl/test/check_device_code/atomic_ref.cpp
@@ -0,0 +1,20 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// RUN: %clangxx -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code  -o - %s | FileCheck %s
+
+#include <sycl/sycl.hpp>
+
+// CHECK-LABEL: define dso_local spir_func noundef i32 @_Z17atomic_ref_globalRi(
+// CHECK-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) [[I:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !srcloc [[META6:![0-9]+]] !sycl_fixed_targets [[META7:![0-9]+]] {
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z33__spirv_GenericCastToPtr_ToGlobalPvi(ptr addrspace(4) noundef [[I]], i32 noundef 5) #[[ATTR3:[0-9]+]]
+// CHECK-NEXT:    [[CALL3_I_I:%.*]] = tail call spir_func noundef i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(ptr addrspace(1) noundef [[CALL_I_I_I_I_I_I]], i32 noundef 1, i32 noundef 898) #[[ATTR4:[0-9]+]]
+// CHECK-NEXT:    ret i32 [[CALL3_I_I]]
+//
+SYCL_EXTERNAL auto atomic_ref_global(int &i) {
+  // Verify that we use _Z33__spirv_GenericCastToPtr_ToGlobalPvi that doesn't
+  // perform dynamic address space validation.
+  sycl::atomic_ref<int, sycl::memory_order::acq_rel, sycl::memory_scope::device,
+                   sycl::access::address_space::global_space>
+      a(i);
+  return a.load();
+}