From f57ce61559bb37d61291d5d1910048f9c5b30d36 Mon Sep 17 00:00:00 2001 From: Hongtao Zhang Date: Sun, 26 Jul 2026 23:49:12 +0000 Subject: [PATCH 1/2] Fix gpu-copy-bw correctness: use volatile SM-copy load/store on HIP The HIP branch of FetchULong2/StoreULong2 used plain cacheable accesses for zero-copy host-mapped memory, unlike the CUDA path which uses ld/st.volatile.global. This could serve stale data and fail the CheckBuf memcmp for cpu_to_gpu*_by_sm (observed on MI300X/ROCm 6.4.4). Use volatile accesses to match the CUDA path. --- .../gpu_copy_performance/gpu_copy.cu | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) 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(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(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 From 28950ca88d5d5757fc892f87c8e60e66e9235bbd Mon Sep 17 00:00:00 2001 From: Hongtao Zhang Date: Fri, 31 Jul 2026 17:30:13 +0000 Subject: [PATCH 2/2] CI/CD - Only cap setuptools below 66 for Python 3.7 torch pulls in an unpinned setuptools during 'pip install .[test,cpuworker]', which upgrades the 65.7 pinned by the pipeline to 83.0.0 and makes the setup.py guard raise VersionConflict on the python-3.10 lint job. setuptools 66+ only drops Python 3.7, so keep the cap there only. --- setup.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) 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, '