diff --git a/setup.py b/setup.py
index 53f754a5d..85b7efc80 100644
--- a/setup.py
+++ b/setup.py
@@ -19,8 +19,10 @@
 print(f'Python {sys.version_info.major}.{sys.version_info.minor} detected.')
 if sys.version_info[:2] < (3, 11):
     import pkg_resources
+    # setuptools 66+ no longer supports Python 3.7, newer runtimes work with any recent version.
+    setuptools_req = 'setuptools>=45, <66' if sys.version_info[:2] < (3, 8) else 'setuptools>=45'
     try:
-        pkg_resources.require(['pip>=18', 'setuptools>=45, <66'])
+        pkg_resources.require(['pip>=18', setuptools_req])
     except (pkg_resources.VersionConflict, pkg_resources.DistributionNotFound):
         print(
             '\033[93mTry update pip/setuptools versions, for example, '
diff --git a/superbench/benchmarks/micro_benchmarks/gpu_copy_performance/gpu_copy.cu b/superbench/benchmarks/micro_benchmarks/gpu_copy_performance/gpu_copy.cu
index 9bdb93060..da24480bc 100644
--- a/superbench/benchmarks/micro_benchmarks/gpu_copy_performance/gpu_copy.cu
+++ b/superbench/benchmarks/micro_benchmarks/gpu_copy_performance/gpu_copy.cu
@@ -599,8 +599,11 @@ int DestroyEvent(BenchArgs *args) {
 // https://github.com/ROCmSoftwarePlatform/rccl/blob/5c8380ff5b5925cae4bce00b1879a5f930226e8d/src/collectives/device/common_kernel.h#L268
 inline __device__ void FetchULong2(ulong2 &v, const ulong2 *p) {
 #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
-    v.x = p->x;
-    v.y = p->y;
+    // Use a volatile access so the compiler does not cache/reorder this zero-copy read of
+    // host-mapped (or peer) memory, matching the uncached ld.volatile.global on the CUDA path.
+    const volatile ulong2 *vp = reinterpret_cast<const volatile ulong2 *>(p);
+    v.x = vp->x;
+    v.y = vp->y;
 #else
     asm volatile("ld.volatile.global.v2.u64 {%0,%1}, [%2];" : "=l"(v.x), "=l"(v.y) : "l"(p) : "memory");
 #endif
@@ -614,8 +617,11 @@ inline __device__ void FetchULong2(ulong2 &v, const ulong2 *p) {
 // https://github.com/ROCmSoftwarePlatform/rccl/blob/5c8380ff5b5925cae4bce00b1879a5f930226e8d/src/collectives/device/common_kernel.h#L276
 inline __device__ void StoreULong2(ulong2 *p, ulong2 &v) {
 #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
-    p->x = v.x;
-    p->y = v.y;
+    // Use a volatile access so the compiler does not cache/reorder this zero-copy write to
+    // host-mapped (or peer) memory, matching the uncached st.volatile.global on the CUDA path.
+    volatile ulong2 *vp = reinterpret_cast<volatile ulong2 *>(p);
+    vp->x = v.x;
+    vp->y = v.y;
 #else
     asm volatile("st.volatile.global.v2.u64 [%0], {%1,%2};" ::"l"(p), "l"(v.x), "l"(v.y) : "memory");
 #endif