diff --git a/libcudacxx/include/cuda/std/__atomic/api/common.h b/libcudacxx/include/cuda/std/__atomic/api/common.h
index 04a3135bc488..a184787636f9 100644
--- a/libcudacxx/include/cuda/std/__atomic/api/common.h
+++ b/libcudacxx/include/cuda/std/__atomic/api/common.h
@@ -163,12 +163,12 @@
   _CCCL_HOST_DEVICE_API inline _Tp fetch_add(ptrdiff_t __op, memory_order __m = memory_order_seq_cst) \
     _CONST _VOLATILE noexcept                                                                         \
   {                                                                                                   \
-    return __atomic_fetch_add_dispatch(&__a, __op, __m, __thread_scope_system_tag{});                 \
+    return __atomic_fetch_add_dispatch(&__a, __op, __m, _Sco{});                                      \
   }                                                                                                   \
   _CCCL_HOST_DEVICE_API inline _Tp fetch_sub(ptrdiff_t __op, memory_order __m = memory_order_seq_cst) \
     _CONST _VOLATILE noexcept                                                                         \
   {                                                                                                   \
-    return __atomic_fetch_sub_dispatch(&__a, __op, __m, __thread_scope_system_tag{});                 \
+    return __atomic_fetch_sub_dispatch(&__a, __op, __m, _Sco{});                                      \
   }                                                                                                   \
   _CCCL_HOST_DEVICE_API inline _Tp operator++(int) _CONST _VOLATILE noexcept                          \
   {                                                                                                   \
diff --git a/libcudacxx/test/atomic_codegen/atomic_add_non_volatile.cu b/libcudacxx/test/atomic_codegen/atomic_add_non_volatile.cu
index 4489e8942680..3f470b5a558e 100644
--- a/libcudacxx/test/atomic_codegen/atomic_add_non_volatile.cu
+++ b/libcudacxx/test/atomic_codegen/atomic_add_non_volatile.cu
@@ -6,6 +6,24 @@ __global__ void add_relaxed_device_non_volatile(int* data, int* out, int n)
   *out     = ref.fetch_add(n, cuda::std::memory_order_relaxed);
 }
 
+__global__ void add_relaxed_block_pointer_non_volatile(int** data, int** out, int n)
+{
+  auto ref = cuda::atomic_ref<int*, cuda::thread_scope_block>{*data};
+  *out     = ref.fetch_add(n, cuda::std::memory_order_relaxed);
+}
+
+__global__ void add_relaxed_device_pointer_non_volatile(int** data, int** out, int n)
+{
+  auto ref = cuda::atomic_ref<int*, cuda::thread_scope_device>{*data};
+  *out     = ref.fetch_add(n, cuda::std::memory_order_relaxed);
+}
+
+__global__ void add_relaxed_system_pointer_non_volatile(int** data, int** out, int n)
+{
+  auto ref = cuda::atomic_ref<int*, cuda::thread_scope_system>{*data};
+  *out     = ref.fetch_add(n, cuda::std::memory_order_relaxed);
+}
+
 /*
 
 ; SM8X-LABEL: .target sm_80
@@ -18,4 +36,16 @@ __global__ void add_relaxed_device_non_volatile(int* data, int* out, int n)
 ; SM8X-NEXT: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]];
 ; SM8X-NEXT: ret;
 
+; SM8X-LABEL: .visible .entry {{_.*add_relaxed_block_pointer_non_volatile.*}}(
+; SM8X: {{.*}}atom.add.relaxed.cta.u64{{.*}}
+; SM8X: ret;
+
+; SM8X-LABEL: .visible .entry {{_.*add_relaxed_device_pointer_non_volatile.*}}(
+; SM8X: {{.*}}atom.add.relaxed.gpu.u64{{.*}}
+; SM8X: ret;
+
+; SM8X-LABEL: .visible .entry {{_.*add_relaxed_system_pointer_non_volatile.*}}(
+; SM8X: {{.*}}atom.add.relaxed.sys.u64{{.*}}
+; SM8X: ret;
+
 */
diff --git a/libcudacxx/test/atomic_codegen/atomic_sub_non_volatile.cu b/libcudacxx/test/atomic_codegen/atomic_sub_non_volatile.cu
index 3b23ac14c9d5..a3986b8df81d 100644
--- a/libcudacxx/test/atomic_codegen/atomic_sub_non_volatile.cu
+++ b/libcudacxx/test/atomic_codegen/atomic_sub_non_volatile.cu
@@ -6,6 +6,24 @@ __global__ void sub_relaxed_device_non_volatile(int* data, int* out, int n)
   *out     = ref.fetch_sub(n, cuda::std::memory_order_relaxed);
 }
 
+__global__ void sub_relaxed_block_pointer_non_volatile(int** data, int** out, int n)
+{
+  auto ref = cuda::atomic_ref<int*, cuda::thread_scope_block>{*data};
+  *out     = ref.fetch_sub(n, cuda::std::memory_order_relaxed);
+}
+
+__global__ void sub_relaxed_device_pointer_non_volatile(int** data, int** out, int n)
+{
+  auto ref = cuda::atomic_ref<int*, cuda::thread_scope_device>{*data};
+  *out     = ref.fetch_sub(n, cuda::std::memory_order_relaxed);
+}
+
+__global__ void sub_relaxed_system_pointer_non_volatile(int** data, int** out, int n)
+{
+  auto ref = cuda::atomic_ref<int*, cuda::thread_scope_system>{*data};
+  *out     = ref.fetch_sub(n, cuda::std::memory_order_relaxed);
+}
+
 /*
 
 ; SM8X-LABEL: .target sm_80
@@ -19,4 +37,16 @@ __global__ void sub_relaxed_device_non_volatile(int* data, int* out, int n)
 ; SM8X-NEXT: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]];
 ; SM8X-NEXT: ret;
 
+; SM8X-LABEL: .visible .entry {{_.*sub_relaxed_block_pointer_non_volatile.*}}(
+; SM8X: {{.*}}atom.add.relaxed.cta.u64{{.*}}
+; SM8X: ret;
+
+; SM8X-LABEL: .visible .entry {{_.*sub_relaxed_device_pointer_non_volatile.*}}(
+; SM8X: {{.*}}atom.add.relaxed.gpu.u64{{.*}}
+; SM8X: ret;
+
+; SM8X-LABEL: .visible .entry {{_.*sub_relaxed_system_pointer_non_volatile.*}}(
+; SM8X: {{.*}}atom.add.relaxed.sys.u64{{.*}}
+; SM8X: ret;
+
 */