diff --git a/projects/rccl/CMakeLists.txt b/projects/rccl/CMakeLists.txt index 03e458c2b0e1..36822c97e18a 100644 --- a/projects/rccl/CMakeLists.txt +++ b/projects/rccl/CMakeLists.txt @@ -110,6 +110,7 @@ option(TRACE "Enable additional tracing" option(QP_TRACKING "Enable per-device QP tracking" OFF) option(FAULT_INJECTION "Enable fault injection" ON) option(QUIET_WARNINGS "Supress compiler warnings" OFF) +option(RCCL_POISON_HIP_ATOMICS "Reject __hip_atomic_* builtins in RCCL sources" ON) option(DWORDX4_INTRINSICS "Turn off compilation with dwordx4 intrinsics" ON) if(ENABLE_TDM_SIMPLE) add_compile_definitions(ENABLE_TDM_SIMPLE=1) diff --git a/projects/rccl/README.md b/projects/rccl/README.md index 7d46e69ae054..d07ccaf61990 100644 --- a/projects/rccl/README.md +++ b/projects/rccl/README.md @@ -85,6 +85,7 @@ RCCL build & installation helper script -DENABLE_COMPRESS=OFF Disable GPU code compression (default: ON) -DENABLE_IFC=ON Enable indirect function call (default: OFF) -DFAULT_INJECTION=OFF Disable fault injection (default: ON) + -DRCCL_POISON_HIP_ATOMICS=OFF Allow __hip_atomic_* builtins in RCCL sources (default: ON) -DRCCL_ROCPROFILER_REGISTER=OFF Disable rocprofiler-register support (default: ON) -DTIMETRACE=ON Enable time-trace during compilation (default: OFF) diff --git a/projects/rccl/install.sh b/projects/rccl/install.sh index 0c87ec7c1ef9..015de3b53557 100755 --- a/projects/rccl/install.sh +++ b/projects/rccl/install.sh @@ -118,6 +118,7 @@ function display_help() echo " -DENABLE_COMPRESS=OFF Disable GPU code compression (default: ON)" echo " -DENABLE_IFC=ON Enable indirect function call (default: OFF)" echo " -DFAULT_INJECTION=OFF Disable fault injection (default: ON)" + echo " -DRCCL_POISON_HIP_ATOMICS=OFF Allow __hip_atomic_* builtins in RCCL sources (default: ON)" echo " -DRCCL_ROCPROFILER_REGISTER=OFF Disable rocprofiler-register support (default: ON)" echo " -DTIMETRACE=ON Enable time-trace during compilation (default: OFF)" echo "" diff --git a/projects/rccl/src/CMakeLists.txt b/projects/rccl/src/CMakeLists.txt index 7d130ddacb37..28e86af88c1f 100644 --- a/projects/rccl/src/CMakeLists.txt +++ b/projects/rccl/src/CMakeLists.txt @@ -1209,6 +1209,26 @@ if(RCCL_ROCPROFILER_REGISTER) target_link_libraries( rccl PRIVATE rocprofiler-register::rocprofiler-register) endif() + +# Poison the __hip_atomic_* builtins +#================================================================================================== +# RCCL spells device atomics with the Clang scoped-atomic builtins so that memory +# scope is explicit at every call site. Force-include (rather than #include) the +# poison header so it also covers hipified and generated sources, and set it here, +# before DeviceLinker.cmake snapshots COMPILE_OPTIONS, so the per-kernel device +# compiles inherit it too. See src/include/poison_hip_atomics.h for the coverage +# this does and does not give. +if(RCCL_POISON_HIP_ATOMICS) + # Joined --include= rather than separate-argument -include : as a + # single token it needs no cooperation from anything that forwards flags one + # at a time, so it survives rccl-device-compile's catch-all for unrecognized + # `-*` tokens as-is. That script also parses the separate-argument spelling + # (see the -include/-imacros branch in parse_compiler_flags), so either form + # works today; the joined one is used because it does not depend on that. + target_compile_options(rccl PRIVATE --include=${RCCL_SOURCE_DIR}/src/include/poison_hip_atomics.h) + message(STATUS "RCCL_POISON_HIP_ATOMICS enabled - __hip_atomic_* builtins rejected in RCCL sources") +endif() + # Device Linker: include pipeline and wire up dependencies # (placed after all target_compile_definitions so DeviceLinker.cmake # can read the rccl target's definitions directly) diff --git a/projects/rccl/src/algorithms/dda/device/CollCommon_ll128.h b/projects/rccl/src/algorithms/dda/device/CollCommon_ll128.h index 4b38b64d279e..27d771bba245 100644 --- a/projects/rccl/src/algorithms/dda/device/CollCommon_ll128.h +++ b/projects/rccl/src/algorithms/dda/device/CollCommon_ll128.h @@ -65,10 +65,10 @@ __host__ __device__ __forceinline__ size_t ddaLL128NumLines(size_t nWords) { // comes from gfx1250 preserving per-thread program order of system-scope stores, // not from an atomic fence. Matches the validated microbenchmark primitives. __device__ __forceinline__ void ddaLL128StoreWord(uint64_t* p, uint64_t v) { - __hip_atomic_store((u64_gptr)p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); + __scoped_atomic_store_n((u64_gptr)p, v, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM); } __device__ __forceinline__ uint64_t ddaLL128LoadWord(const uint64_t* p) { - return __hip_atomic_load((u64_gptr) const_cast(p), __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); + return __scoped_atomic_load_n((u64_gptr) const_cast(p), __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM); } // Element-wise add of the T-elements packed into two 8B payload words. An 8B diff --git a/projects/rccl/src/device/all_gather.h b/projects/rccl/src/device/all_gather.h index ab5ae9df9907..be944d8983e2 100644 --- a/projects/rccl/src/device/all_gather.h +++ b/projects/rccl/src/device/all_gather.h @@ -195,7 +195,7 @@ struct RunWorkCollpatSteps + (step % NCCL_SHMEM_PAT_STEPS); int* poll = &ps->flags; - while (__hip_atomic_load(poll, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) != 0) { + while (__scoped_atomic_load_n(poll, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) != 0) { pollCount++;// Wait for workers to be done with step 'step-NCCL_SHMEM_PAT_STEPS' } patAlgo.getNextOp(ps); @@ -224,14 +224,14 @@ struct RunWorkCollpatSteps + (step % NCCL_SHMEM_PAT_STEPS); int* poll = &ps->flags; - while (__hip_atomic_load(poll, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) == 0) { + while (__scoped_atomic_load_n(poll, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) == 0) { pollCount++; // Wait for compute thread } int last = ps->last; prims.patCopy(ps, shmem); if (tidInGroup == 0) - __hip_atomic_store(poll, 0, __ATOMIC_RELEASE, - __HIP_MEMORY_SCOPE_WORKGROUP); // Return element to compute thread + __scoped_atomic_store_n(poll, 0, __ATOMIC_RELEASE, + __MEMORY_SCOPE_WRKGRP); // Return element to compute thread if (last) break; step += nGroups; } diff --git a/projects/rccl/src/device/ce_reduce/ce_reduce_impl.h.in b/projects/rccl/src/device/ce_reduce/ce_reduce_impl.h.in index c326c9dc7b93..3394471df9f5 100644 --- a/projects/rccl/src/device/ce_reduce/ce_reduce_impl.h.in +++ b/projects/rccl/src/device/ce_reduce/ce_reduce_impl.h.in @@ -38,12 +38,12 @@ __device__ __forceinline__ void ncclCeGlobalBlockBarrier(uint32_t* d_barrierSync uint32_t* doneGen = d_barrierSync + 1; __syncthreads(); if (threadIdx.x == 0) { - uint32_t prev = __hip_atomic_fetch_add(arrival, 1, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + uint32_t prev = __scoped_atomic_fetch_add(arrival, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE); if (prev + 1 == (uint32_t)numBlocks) { - __hip_atomic_store(arrival, 0, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); - __hip_atomic_store(doneGen, gen, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_AGENT); + __scoped_atomic_store_n(arrival, 0, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE); + __scoped_atomic_store_n(doneGen, gen, __ATOMIC_RELEASE, __MEMORY_SCOPE_DEVICE); } else { - while (__hip_atomic_load(doneGen, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_AGENT) != gen) { + while (__scoped_atomic_load_n(doneGen, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_DEVICE) != gen) { __builtin_amdgcn_s_sleep(1); } } @@ -249,8 +249,8 @@ __global__ __launch_bounds__(256) void ncclCeLocalReduceKernelVec( while (true) { int ready = 0; for (int r = 0; r < nRanks; r++) { - uint32_t sv = __hip_atomic_load((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], __ATOMIC_ACQUIRE, - __HIP_MEMORY_SCOPE_SYSTEM); + uint32_t sv = __scoped_atomic_load_n((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], __ATOMIC_ACQUIRE, + __MEMORY_SCOPE_SYSTEM); if (sv == 1) { ready++; } @@ -354,8 +354,8 @@ __global__ __launch_bounds__(256) void ncclCeLocalReduceKernelVec( } if (blockIdx.x == 0) { for (int r = (int)threadIdx.x; r < nRanks; r += (int)blockDim.x) { - __hip_atomic_store((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], 0, __ATOMIC_RELEASE, - __HIP_MEMORY_SCOPE_SYSTEM); + __scoped_atomic_store_n((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], 0, __ATOMIC_RELEASE, + __MEMORY_SCOPE_SYSTEM); } } } diff --git a/projects/rccl/src/device/common.h b/projects/rccl/src/device/common.h index 676fd0fe9693..6d95d198ab30 100644 --- a/projects/rccl/src/device/common.h +++ b/projects/rccl/src/device/common.h @@ -21,8 +21,8 @@ #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS #define STORE(DST, SRC) \ { \ - __hip_atomic_store((__attribute__((address_space(1))) __typeof__(*(DST))*)(DST), (SRC), __ATOMIC_RELAXED, \ - __HIP_MEMORY_SCOPE_SYSTEM); \ + __scoped_atomic_store_n((__attribute__((address_space(1))) __typeof__(*(DST))*)(DST), (SRC), __ATOMIC_RELAXED, \ + __MEMORY_SCOPE_SYSTEM); \ } #elif defined(__GFX9__) #define STORE(DST, SRC) \ diff --git a/projects/rccl/src/device/op128.h b/projects/rccl/src/device/op128.h index 569567831aa2..5260b08af73a 100644 --- a/projects/rccl/src/device/op128.h +++ b/projects/rccl/src/device/op128.h @@ -272,8 +272,8 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value) // template<> __device__ __forceinline__ void st_relaxed_gpu_global<0>(uintptr_t addr, BytePack<0> value) {} // Used to define implementations for above prototypes. -#define DEFINE_ld_st__size_space_hip_atomic(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, \ - addr_reg_ty) \ +#define DEFINE_ld_st__size_space_scoped_atomic(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, \ + addr_reg_ty) \ template <> \ __device__ __forceinline__ BytePack ld_##space(addr_cxx_ty addr) { \ data_cxx_ty tmp; \ @@ -285,16 +285,16 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value) template <> \ __device__ __forceinline__ BytePack ld_volatile_##space(addr_cxx_ty addr) { \ data_cxx_ty tmp; \ - tmp = __hip_atomic_load((__attribute__((address_space(1))) data_cxx_ty*)addr, __ATOMIC_RELAXED, \ - __HIP_MEMORY_SCOPE_SYSTEM); \ + tmp = __scoped_atomic_load_n((__attribute__((address_space(1))) data_cxx_ty*)addr, __ATOMIC_RELAXED, \ + __MEMORY_SCOPE_SYSTEM); \ BytePack ans; \ ans.native = tmp; \ return ans; \ } \ template <> \ __device__ __forceinline__ void st_##space(addr_cxx_ty addr, BytePack value) { \ - __hip_atomic_store((__attribute__((address_space(1))) data_cxx_ty*)addr, value.native, __ATOMIC_RELAXED, \ - __HIP_MEMORY_SCOPE_SYSTEM); \ + __scoped_atomic_store_n((__attribute__((address_space(1))) data_cxx_ty*)addr, value.native, __ATOMIC_RELAXED, \ + __MEMORY_SCOPE_SYSTEM); \ } #define DEFINE_ld_st__size_space_fallback(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, \ @@ -322,7 +322,7 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value) #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS #define DEFINE_ld_st__size_space(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) \ - DEFINE_ld_st__size_space_hip_atomic(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) + DEFINE_ld_st__size_space_scoped_atomic(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) #else #define DEFINE_ld_st__size_space(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) \ DEFINE_ld_st__size_space_fallback(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) @@ -359,7 +359,7 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value) DEFINE_ld_st__size(1, uint8_t, b8, r) DEFINE_ld_st__size(2, uint16_t, b16, h) DEFINE_ld_st__size(4, uint32_t, b32, r) DEFINE_ld_st__size(8, uint64_t, b64, l) #undef DEFINE_ld_st__size_space -#undef DEFINE_ld_st__size_space_hip_atomic +#undef DEFINE_ld_st__size_space_scoped_atomic #undef DEFINE_ld_st__size_space_fallback #undef DEFINE_ld_st__size @@ -427,13 +427,13 @@ DEFINE_ld_st_16__space(global, uintptr_t, l) __device__ __forceinline__ uint64_t ld_volatile_global(uint64_t* ptr) { uint64_t ans; ans = - __hip_atomic_load((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); + __scoped_atomic_load_n((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM); return ans; } __device__ __forceinline__ uint64_t ld_relaxed_sys_global(uint64_t* ptr) { uint64_t ans; ans = - __hip_atomic_load((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); + __scoped_atomic_load_n((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM); return ans; } @@ -450,21 +450,21 @@ __device__ __forceinline__ uint64_t ld_relaxed_sys_global(uint64_t* ptr) { __device__ __forceinline__ uint64_t ld_acquire_sys_global(uint64_t* ptr) { uint64_t ans; ans = - __hip_atomic_load((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_SYSTEM); + __scoped_atomic_load_n((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_SYSTEM); return ans; } __device__ __forceinline__ void st_volatile_global(uint64_t* ptr, uint64_t val) { - __hip_atomic_store((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED, - __HIP_MEMORY_SCOPE_SYSTEM); + __scoped_atomic_store_n((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED, + __MEMORY_SCOPE_SYSTEM); } __device__ __forceinline__ void st_relaxed_sys_global(uint64_t* ptr, uint64_t val) { - __hip_atomic_store((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED, - __HIP_MEMORY_SCOPE_SYSTEM); + __scoped_atomic_store_n((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED, + __MEMORY_SCOPE_SYSTEM); } __device__ __forceinline__ void st_release_sys_global(uint64_t* ptr, uint64_t val) { - __hip_atomic_store((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELEASE, - __HIP_MEMORY_SCOPE_SYSTEM); + __scoped_atomic_store_n((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELEASE, + __MEMORY_SCOPE_SYSTEM); } __device__ __forceinline__ void fence_acq_rel_sys() { diff --git a/projects/rccl/src/device/primitives.h b/projects/rccl/src/device/primitives.h index 735eb0e88ab0..5153ce51cd5b 100644 --- a/projects/rccl/src/device/primitives.h +++ b/projects/rccl/src/device/primitives.h @@ -27,10 +27,10 @@ if (wid == 0) { \ (BARRIER_NEXT) += (NWORKERS) / WARP_SIZE; \ __THREAD_FENCE; \ - __hip_atomic_fetch_add((BARRIERS_PTR), 1, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_WORKGROUP); \ + __scoped_atomic_fetch_add((BARRIERS_PTR), 1, __ATOMIC_RELEASE, __MEMORY_SCOPE_WRKGRP); \ int spins = 0; \ int rate_limit = 50; \ - while (__hip_atomic_load((BARRIERS_PTR), __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) < (BARRIER_NEXT)) { \ + while (__scoped_atomic_load_n((BARRIERS_PTR), __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) < (BARRIER_NEXT)) { \ spins++; \ if (spins == NCCL_SPINS_BEFORE_CHECK_ABORT) { \ if (__atomic_load_n(ncclShmem.comm.abortFlag, __ATOMIC_SEQ_CST)) { \ @@ -42,7 +42,7 @@ if (spins == 0 && rate_limit > 0) { \ rate_limit--; \ traceData(__LINE__, threadIdx.x, \ - __hip_atomic_load((BARRIERS_PTR), __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP), \ + __scoped_atomic_load_n((BARRIERS_PTR), __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP), \ (BARRIER_NEXT)); \ } \ __builtin_amdgcn_s_sleep(1); \ diff --git a/projects/rccl/src/device/prims_ll.h b/projects/rccl/src/device/prims_ll.h index 36da63e34bc4..a33ad7545155 100644 --- a/projects/rccl/src/device/prims_ll.h +++ b/projects/rccl/src/device/prims_ll.h @@ -288,8 +288,8 @@ class Primitivesv + 1) = *((u64_gptr)i4.v + 1); #else #if defined(__gfx1200__) || defined(__gfx1201__) - __hip_atomic_store((u64_gptr)dst->v, i4.v[0], __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); - __hip_atomic_store((u64_gptr)dst->v + 1, i4.v[1], __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM); + __scoped_atomic_store_n((u64_gptr)dst->v, i4.v[0], __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM); + __scoped_atomic_store_n((u64_gptr)dst->v + 1, i4.v[1], __ATOMIC_RELEASE, __MEMORY_SCOPE_SYSTEM); #else *((u64_gptr)dst->v) = *((u64_gptr)i4.v); *((u64_gptr)dst->v + 1) = *((u64_gptr)i4.v + 1); @@ -321,17 +321,17 @@ class PrimitivespatSteps + (step % NCCL_SHMEM_PAT_STEPS); int* poll = &ps->flags; - while (__hip_atomic_load(poll, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) != 0) { + while (__scoped_atomic_load_n(poll, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) != 0) { // pollCount++;// unused variable - compiler warning // Wait for workers to be done with step 'step-NCCL_SHMEM_PAT_STEPS' } patAlgo.getNextOp(ps); @@ -251,14 +251,14 @@ struct RunWorkCollpatSteps + (step % NCCL_SHMEM_PAT_STEPS); int* poll = &ps->flags; - while (__hip_atomic_load(poll, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) == 0) { + while (__scoped_atomic_load_n(poll, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) == 0) { // pollCount++; // unused variable - compiler warning // Wait for compute thread } int last = ps->last; prims.patReduce(ps, shmem); if (tidInGroup == 0) - __hip_atomic_store(poll, 0, __ATOMIC_RELEASE, - __HIP_MEMORY_SCOPE_WORKGROUP); // Return element to compute thread + __scoped_atomic_store_n(poll, 0, __ATOMIC_RELEASE, + __MEMORY_SCOPE_WRKGRP); // Return element to compute thread if (last) break; step += nGroups; } diff --git a/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_ipc_copy.h b/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_ipc_copy.h index c6b3029b2cb9..40536475d749 100644 --- a/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_ipc_copy.h +++ b/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_ipc_copy.h @@ -21,8 +21,8 @@ namespace detail { // Relaxed: ordering is provided by the caller's fence (signalPeer). NCCL_DEVICE_INLINE void ipcFlatAtomicAddSys64(uint64_t* dst, uint64_t val) { - __hip_atomic_fetch_add(reinterpret_cast(dst), static_cast(val), - __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); + __scoped_atomic_fetch_add(reinterpret_cast(dst), static_cast(val), + __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM); } } // namespace detail diff --git a/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h b/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h index 5268204cf62f..c279ae72d142 100644 --- a/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h +++ b/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h @@ -94,7 +94,7 @@ NCCL_DEVICE_INLINE void markSdmaDirty(ncclGinAnvilSdmaGPUContext* rsCtx, int pee const int bitIdx = peer * numCh + effCh; if (bitIdx < 0 || bitIdx >= kSdmaDirtyBitWidth) return; uint64_t bit = 1ULL << bitIdx; - __hip_atomic_fetch_or(dirty, bit, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + __scoped_atomic_fetch_or(dirty, bit, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE); } NCCL_DEVICE_INLINE int effectiveChannel(ncclGinAnvilSdmaGPUContext* rsCtx, int blockId) { @@ -414,7 +414,7 @@ struct ncclGinApi_Flush { uint64_t* sdmaDirty = loadConst(&rsCtx->sdmaDirty); uint64_t dirty = 0; if (sdmaDirty != nullptr) { - dirty = __hip_atomic_load(sdmaDirty, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + dirty = __scoped_atomic_load_n(sdmaDirty, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE); } if (dirty != 0) { auto** handles = (::sdma_anvil::SdmaQueueDeviceHandle**)loadConst(&rsCtx->queueHandles); @@ -435,7 +435,7 @@ struct ncclGinApi_Flush { } coop.sync(); if (coop.thread_rank() == 0 && sdmaDirty != nullptr) { - __hip_atomic_store(sdmaDirty, 0, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + __scoped_atomic_store_n(sdmaDirty, 0, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE); } coop.sync(); } diff --git a/projects/rccl/src/include/nccl_device/hip_compat.h b/projects/rccl/src/include/nccl_device/hip_compat.h index 6f70f6716136..885ab8e979d3 100644 --- a/projects/rccl/src/include/nccl_device/hip_compat.h +++ b/projects/rccl/src/include/nccl_device/hip_compat.h @@ -161,26 +161,26 @@ enum thread_scope { thread_scope_system = 3 }; -// Map thread_scope to HIP memory scope -NCCL_DEVICE_INLINE constexpr int toHipMemoryScope(thread_scope scope) { +// Map thread_scope to Clang scoped-atomic memory scope +NCCL_DEVICE_INLINE constexpr int toMemoryScope(thread_scope scope) { switch (scope) { case thread_scope_thread: - return __HIP_MEMORY_SCOPE_SINGLETHREAD; + return __MEMORY_SCOPE_SINGLE; case thread_scope_block: - return __HIP_MEMORY_SCOPE_WORKGROUP; + return __MEMORY_SCOPE_WRKGRP; case thread_scope_device: - return __HIP_MEMORY_SCOPE_AGENT; + return __MEMORY_SCOPE_DEVICE; case thread_scope_system: - return __HIP_MEMORY_SCOPE_SYSTEM; + return __MEMORY_SCOPE_SYSTEM; default: - return __HIP_MEMORY_SCOPE_SYSTEM; + return __MEMORY_SCOPE_SYSTEM; } } //////////////////////////////////////////////////////////////////////////////// // atomic_ref implementation for HIP // -// Provides cuda::atomic_ref-compatible interface using HIP atomics +// Provides cuda::atomic_ref-compatible interface using scoped atomics //////////////////////////////////////////////////////////////////////////////// template @@ -191,11 +191,11 @@ struct atomic_ref { NCCL_DEVICE_INLINE void store(T val, memory_order order = memory_order_seq_cst) const { if constexpr (sizeof(T) == 4) { - __hip_atomic_store(reinterpret_cast(ptr), *reinterpret_cast(&val), order, - toHipMemoryScope(Scope)); + __scoped_atomic_store_n(reinterpret_cast(ptr), *reinterpret_cast(&val), order, + toMemoryScope(Scope)); } else if constexpr (sizeof(T) == 8) { - __hip_atomic_store(reinterpret_cast(ptr), *reinterpret_cast(&val), - order, toHipMemoryScope(Scope)); + __scoped_atomic_store_n(reinterpret_cast(ptr), *reinterpret_cast(&val), + order, toMemoryScope(Scope)); } else { __atomic_store_n(ptr, val, order); } @@ -204,11 +204,11 @@ struct atomic_ref { NCCL_DEVICE_INLINE T load(memory_order order = memory_order_seq_cst) const { T result; if constexpr (sizeof(T) == 4) { - unsigned int tmp = __hip_atomic_load(reinterpret_cast(ptr), order, toHipMemoryScope(Scope)); + unsigned int tmp = __scoped_atomic_load_n(reinterpret_cast(ptr), order, toMemoryScope(Scope)); result = *reinterpret_cast(&tmp); } else if constexpr (sizeof(T) == 8) { unsigned long long tmp = - __hip_atomic_load(reinterpret_cast(ptr), order, toHipMemoryScope(Scope)); + __scoped_atomic_load_n(reinterpret_cast(ptr), order, toMemoryScope(Scope)); result = *reinterpret_cast(&tmp); } else { result = __atomic_load_n(ptr, order); @@ -218,9 +218,9 @@ struct atomic_ref { NCCL_DEVICE_INLINE T fetch_add(T val, memory_order order = memory_order_seq_cst) const { if constexpr (sizeof(T) == 4) { - return __hip_atomic_fetch_add(ptr, val, order, toHipMemoryScope(Scope)); + return __scoped_atomic_fetch_add(ptr, val, order, toMemoryScope(Scope)); } else if constexpr (sizeof(T) == 8) { - return __hip_atomic_fetch_add(ptr, val, order, toHipMemoryScope(Scope)); + return __scoped_atomic_fetch_add(ptr, val, order, toMemoryScope(Scope)); } else { return __atomic_fetch_add(ptr, val, order); } @@ -229,15 +229,15 @@ struct atomic_ref { NCCL_DEVICE_INLINE bool compare_exchange_weak(T& expected, T desired, memory_order success, memory_order failure) const { if constexpr (sizeof(T) == 4) { - return __hip_atomic_compare_exchange_weak(reinterpret_cast(ptr), + return __scoped_atomic_compare_exchange_n(reinterpret_cast(ptr), reinterpret_cast(&expected), - *reinterpret_cast(&desired), success, failure, - toHipMemoryScope(Scope)); + *reinterpret_cast(&desired), /*weak=*/true, success, + failure, toMemoryScope(Scope)); } else if constexpr (sizeof(T) == 8) { - return __hip_atomic_compare_exchange_weak(reinterpret_cast(ptr), + return __scoped_atomic_compare_exchange_n(reinterpret_cast(ptr), reinterpret_cast(&expected), - *reinterpret_cast(&desired), success, failure, - toHipMemoryScope(Scope)); + *reinterpret_cast(&desired), /*weak=*/true, + success, failure, toMemoryScope(Scope)); } else { return __atomic_compare_exchange_n(ptr, &expected, desired, /*weak=*/true, success, failure); } @@ -251,15 +251,15 @@ struct atomic_ref { NCCL_DEVICE_INLINE bool compare_exchange_strong(T& expected, T desired, memory_order success, memory_order failure) const { if constexpr (sizeof(T) == 4) { - return __hip_atomic_compare_exchange_strong(reinterpret_cast(ptr), - reinterpret_cast(&expected), - *reinterpret_cast(&desired), success, failure, - toHipMemoryScope(Scope)); + return __scoped_atomic_compare_exchange_n(reinterpret_cast(ptr), + reinterpret_cast(&expected), + *reinterpret_cast(&desired), /*weak=*/false, success, + failure, toMemoryScope(Scope)); } else if constexpr (sizeof(T) == 8) { - return __hip_atomic_compare_exchange_strong(reinterpret_cast(ptr), - reinterpret_cast(&expected), - *reinterpret_cast(&desired), success, failure, - toHipMemoryScope(Scope)); + return __scoped_atomic_compare_exchange_n(reinterpret_cast(ptr), + reinterpret_cast(&expected), + *reinterpret_cast(&desired), /*weak=*/false, + success, failure, toMemoryScope(Scope)); } else { return __atomic_compare_exchange_n(ptr, &expected, desired, /*weak=*/false, success, failure); } diff --git a/projects/rccl/src/include/nccl_device/impl/gin__funcs.h b/projects/rccl/src/include/nccl_device/impl/gin__funcs.h index f8986fb40e4f..c4d42e66d1d7 100644 --- a/projects/rccl/src/include/nccl_device/impl/gin__funcs.h +++ b/projects/rccl/src/include/nccl_device/impl/gin__funcs.h @@ -839,7 +839,7 @@ template NCCL_DEVICE_INLINE void ncclGin_BackendMask::increaseSignalShadow(ncclGinSignal_t signal, uint64_t delta) const { #if defined(__HIP_PLATFORM_AMD__) - __hip_atomic_fetch_add(this->_signalShadows + signal, delta, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WORKGROUP); + __scoped_atomic_fetch_add(this->_signalShadows + signal, delta, __ATOMIC_RELAXED, __MEMORY_SCOPE_WRKGRP); #else asm volatile("red.relaxed.cta.add.u64 [%0],%1;" ::"l"(this->_signalShadows + signal), "l"(delta) : "memory"); #endif diff --git a/projects/rccl/src/include/poison_hip_atomics.h b/projects/rccl/src/include/poison_hip_atomics.h new file mode 100644 index 000000000000..22c92beafd18 --- /dev/null +++ b/projects/rccl/src/include/poison_hip_atomics.h @@ -0,0 +1,54 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +#ifndef NCCL_POISON_HIP_ATOMICS_H_ +#define NCCL_POISON_HIP_ATOMICS_H_ + +// Force-included (-include) ahead of every RCCL translation unit when +// RCCL_POISON_HIP_ATOMICS=ON. RCCL expresses device atomics with the Clang +// scoped-atomic builtins (__scoped_atomic_*, see device/op128.h and +// nccl_device/hip_compat.h) so that memory scope is explicit at every call +// site; the __hip_atomic_* builtins default their scope through HIP's headers +// instead. Poisoning makes the compiler reject them rather than relying on a +// source grep, which macro expansion can hide. +// +// #pragma GCC poison rejects the identifier everywhere, including inside +// system headers, and HIP's own headers spell these builtins (amd_hip_atomic.h +// and amd_hip_unsafe_atomics.h implement atomicAdd() and friends with them; +// amd_hip_fp16.h and amd_hip_bf16.h carry their own half/bfloat16 atomicAdd +// overloads). So HIP must be parsed before the poison takes effect, which is +// why the three public wrappers below are pulled in here: hip_runtime.h drags +// in the two atomics headers, and hip_fp16.h / hip_bf16.h are listed on top of +// it because hip_runtime.h alone does not reach the fp16/bf16 ones. The +// consequence is that this covers RCCL's own sources only: HIP's +// atomicAdd()/atomicMax()/atomicExch() still reach the builtins through header +// code that was already parsed, and RCCL device code still calls those. +#if defined(__cplusplus) && (defined(__HIP__) || defined(__HIPCC__)) +#include +#include +#include +#endif + +// clang-format off +// clang-format has no notion of a line-continued #pragma: it joins the whole +// directive onto one ~300-column line. Keep it off across this block so the +// list stays reviewable one name per line. +#pragma GCC poison \ + __hip_atomic_load \ + __hip_atomic_store \ + __hip_atomic_exchange \ + __hip_atomic_compare_exchange_weak \ + __hip_atomic_compare_exchange_strong \ + __hip_atomic_fetch_add \ + __hip_atomic_fetch_sub \ + __hip_atomic_fetch_and \ + __hip_atomic_fetch_or \ + __hip_atomic_fetch_xor \ + __hip_atomic_fetch_min \ + __hip_atomic_fetch_max +// clang-format on + +#endif diff --git a/projects/rccl/src/include/test_poison_hip_atomics.py b/projects/rccl/src/include/test_poison_hip_atomics.py new file mode 100644 index 000000000000..8c946090b6cc --- /dev/null +++ b/projects/rccl/src/include/test_poison_hip_atomics.py @@ -0,0 +1,289 @@ +#!/usr/bin/env python3 +"""Guard that #pragma GCC poison in poison_hip_atomics.h actually rejects +__hip_atomic_* builtins. + +Nothing in RCCL CI runs ctest, so this is invoked from the host-test pipeline +(test/host/run_host_tests.sh `guards` phase, folded into `run`) rather than +add_test(). It compiles tiny HIP probes with amdclang++; no GPU and no +librccl.so are required (-nogpulib -fsyntax-only). + +A missing HIP compiler is a failure here, not a skip: unittest exits 0 on a +fully skipped suite, so a toolchain-less run would report success without ever +compiling a probe. Set RCCL_POISON_TEST_ALLOW_SKIP=1 to downgrade that to a +skip when running on a box without ROCm. +""" + +from __future__ import annotations + +import os +import re +import shutil +import subprocess +import tempfile +import textwrap +import unittest + +HERE = os.path.dirname(os.path.abspath(__file__)) +POISON_H = os.path.join(HERE, "poison_hip_atomics.h") + +# Names the pragma in poison_hip_atomics.h must reject. Kept in lockstep with +# that file by test_pragma_matches_poisoned_tuple, which compares this tuple +# against the pragma itself -- drift in either direction fails. +POISONED = ( + "__hip_atomic_load", + "__hip_atomic_store", + "__hip_atomic_exchange", + "__hip_atomic_compare_exchange_weak", + "__hip_atomic_compare_exchange_strong", + "__hip_atomic_fetch_add", + "__hip_atomic_fetch_sub", + "__hip_atomic_fetch_and", + "__hip_atomic_fetch_or", + "__hip_atomic_fetch_xor", + "__hip_atomic_fetch_min", + "__hip_atomic_fetch_max", +) + +# Opt out of the "no compiler is a failure" rule for local runs without ROCm. +ALLOW_SKIP = os.environ.get("RCCL_POISON_TEST_ALLOW_SKIP") == "1" + +_PRAGMA_RE = re.compile(r"\s*#\s*pragma\s+GCC\s+poison\b") + + +def _pragma_poisoned_names(text): + """Identifiers listed in the `#pragma GCC poison` directive in `text`. + + Joins backslash continuations so the multi-line spelling in the header + parses the same as a single-line one. Returns () when no pragma is found, + which fails the comparison rather than silently matching an empty tuple. + """ + lines = text.splitlines() + for start, line in enumerate(lines): + if not _PRAGMA_RE.match(line): + continue + parts = [] + i = start + while i < len(lines): + stripped = lines[i].rstrip() + continued = stripped.endswith("\\") + parts.append(stripped[:-1] if continued else stripped) + i += 1 + if not continued: + break + joined = _PRAGMA_RE.sub("", " ".join(parts), count=1) + return tuple(joined.split()) + return () + + +def _find_cxx(): + for env in ("HIPCXX", "CXX"): + val = os.environ.get(env) + if val and shutil.which(val): + return val + rocm = os.environ.get("ROCM_PATH", "/opt/rocm") + for cand in ( + os.path.join(rocm, "bin", "amdclang++"), + os.path.join(rocm, "llvm", "bin", "clang++"), + os.path.join(rocm, "bin", "hipcc"), + ): + if os.path.isfile(cand) and os.access(cand, os.X_OK): + return cand + for name in ("amdclang++", "hipcc"): + found = shutil.which(name) + if found: + return found + return None + + +CXX = _find_cxx() + + +def _call_expr(name): + """A well-typed call to `name` that survives HIP host+device parsing.""" + if name in ("__hip_atomic_load",): + return f"{name}(p, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)" + if name in ("__hip_atomic_store",): + return f"{name}(p, 0u, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)" + if name in ("__hip_atomic_exchange",): + return f"{name}(p, 1u, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)" + if name in ( + "__hip_atomic_compare_exchange_weak", + "__hip_atomic_compare_exchange_strong", + ): + return ( + f"{name}(p, &expected, 1u, __ATOMIC_RELAXED, __ATOMIC_RELAXED, " + "__HIP_MEMORY_SCOPE_AGENT)" + ) + # fetch_* family + return f"{name}(p, 1u, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)" + + +def _probe_src(call): + return textwrap.dedent(f"""\ + #include + __global__ void k(unsigned int *p) {{ + unsigned int expected = 0; + (void)expected; + (void){call}; + }} + """) + + +def _compile(src, *, poison, offload): + """Return (returncode, combined stdout+stderr).""" + cmd = [CXX, "-x", "hip", "-nogpulib", "-fsyntax-only"] + if offload == "host": + cmd += ["--offload-host-only"] + elif offload == "device": + cmd += ["--offload-device-only", "--offload-arch=gfx942"] + else: + raise ValueError(offload) + if poison: + cmd += [f"--include={POISON_H}"] + with tempfile.NamedTemporaryFile("w", suffix=".cpp", delete=False) as f: + f.write(src) + path = f.name + try: + proc = subprocess.run(cmd + [path], capture_output=True, text=True) + return proc.returncode, proc.stdout + proc.stderr + finally: + os.unlink(path) + + +class PoisonHeaderTest(unittest.TestCase): + """Checks on the header text itself. These need no HIP compiler, so they + still run (and still gate) on a machine without ROCm.""" + + def test_poison_header_exists(self): + self.assertTrue(os.path.isfile(POISON_H), POISON_H) + + def test_pragma_matches_poisoned_tuple(self): + """POISONED and the pragma are two hand-maintained lists; keep them equal. + + Without this, dropping a name from POISONED just shrinks the compile + probes below and nothing fails -- the direction that silently loses + coverage. + """ + with open(POISON_H) as f: + pragma = _pragma_poisoned_names(f.read()) + self.assertEqual( + set(pragma), + set(POISONED), + "#pragma GCC poison in poison_hip_atomics.h and POISONED here have " + f"drifted.\n only in pragma: {sorted(set(pragma) - set(POISONED))}" + f"\n only in POISONED: {sorted(set(POISONED) - set(pragma))}", + ) + + +class ToolchainTest(unittest.TestCase): + """Makes a toolchain-less run red. + + Every compile probe below can only skip when no HIP compiler is found, and + a fully skipped suite exits 0. This is the test that tells "the poison + works" apart from "nothing was ever compiled". + """ + + @unittest.skipIf(ALLOW_SKIP, "RCCL_POISON_TEST_ALLOW_SKIP=1") + def test_hip_compiler_available(self): + self.assertIsNotNone( + CXX, + "no HIP compiler found: set ROCM_PATH, HIPCXX or CXX, or put " + "amdclang++/hipcc on PATH. Without one the poison probes cannot " + "run, and skipping them would let this guard pass without " + "checking anything. Set RCCL_POISON_TEST_ALLOW_SKIP=1 to skip " + "instead when running locally without ROCm.", + ) + + +@unittest.skipUnless(CXX, "amdclang++/hipcc not found (set ROCM_PATH or HIPCXX)") +class PoisonHipAtomicsTest(unittest.TestCase): + def test_hip_atomic_load_rejected_on_device(self): + rc, out = _compile( + _probe_src(_call_expr("__hip_atomic_load")), + poison=True, + offload="device", + ) + self.assertNotEqual(rc, 0, msg=out) + self.assertIn("poisoned identifier", out, msg=out) + self.assertIn("__hip_atomic_load", out, msg=out) + + def test_hip_atomic_load_rejected_on_host(self): + rc, out = _compile( + _probe_src(_call_expr("__hip_atomic_load")), + poison=True, + offload="host", + ) + self.assertNotEqual(rc, 0, msg=out) + self.assertIn("poisoned identifier", out, msg=out) + self.assertIn("__hip_atomic_load", out, msg=out) + + def test_every_probe_is_well_formed(self): + """Each _call_expr() must compile with the poison off. + + #pragma GCC poison fires in the lexer, so the rejection tests pass on + any spelling of the argument list. This is what keeps the per-name + arity table in _call_expr honest for all 12 names rather than just the + one the load tests happen to use. + """ + broken = [] + for name in POISONED: + rc, out = _compile( + _probe_src(_call_expr(name)), poison=False, offload="device" + ) + if rc != 0: + broken.append(f"{name}: rc={rc}\n{out}") + self.assertFalse( + broken, + "probe is not a well-formed call without the poison:\n" + "\n".join(broken), + ) + + def test_scoped_atomic_load_accepted_with_poison(self): + src = textwrap.dedent("""\ + #include + __global__ void k(unsigned int *p, unsigned int *out) { + *out = __scoped_atomic_load_n(p, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE); + } + """) + rc, out = _compile(src, poison=True, offload="device") + self.assertEqual(rc, 0, msg=out) + + def test_scoped_atomic_compare_exchange_accepted_with_poison(self): + """Pin the replacement compare-exchange shape for both widths. + + This is the one conversion in the scoped-atomics port whose argument + list changed rather than just its name, gaining a weak flag ahead of + the success/failure orders. Mirrors the spelling atomic_ref uses in + nccl_device/hip_compat.h: (ptr, expected, desired, weak, success, + failure, scope). It pins the builtin, not hip_compat.h -- a reshape + of that header still needs review. + """ + for cxx_ty in ("unsigned int", "unsigned long long"): + with self.subTest(type=cxx_ty): + src = textwrap.dedent(f"""\ + #include + __global__ void k({cxx_ty} *p, bool *out) {{ + {cxx_ty} expected = 0; + *out = __scoped_atomic_compare_exchange_n( + p, &expected, ({cxx_ty})1, /*weak=*/true, + __ATOMIC_ACQ_REL, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_DEVICE); + }} + """) + rc, out = _compile(src, poison=True, offload="device") + self.assertEqual(rc, 0, msg=out) + + def test_every_poisoned_builtin_is_rejected(self): + missing = [] + for name in POISONED: + rc, out = _compile( + _probe_src(_call_expr(name)), poison=True, offload="device" + ) + if rc == 0 or "poisoned identifier" not in out: + missing.append(f"{name}: rc={rc}\n{out}") + self.assertFalse( + missing, + "poison header did not reject:\n" + "\n".join(missing), + ) + + +if __name__ == "__main__": + unittest.main(verbosity=2) diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index 606930b53a71..e69763f2e912 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -35,12 +35,11 @@ if(BUILD_TESTS) "${PROJECT_SOURCE_DIR}/src/device/ce_reduce/test_generate_ce_reduce.py" ) - # The kernel-count leak guards (src/device/test_generate_kernel_counts.py and - # src/device/symmetric/test_generate_symmetric.py) are intentionally NOT - # registered as ctest cases here: nothing in RCCL CI runs `ctest`, so an - # add_test() would never gate. They run in CI via the CPU-only host-test - # pipeline instead -- see test/host/run_host_tests.sh (the `guards` phase, - # folded into `run`). + # The kernel-count leak guards and the __hip_atomic_* poison probe + # (src/include/test_poison_hip_atomics.py) are intentionally NOT registered as + # ctest cases here: nothing in RCCL CI runs `ctest`, so an add_test() would + # never gate. They run in CI via the CPU-only host-test pipeline instead -- + # see test/host/run_host_tests.sh (the `guards` phase, folded into `run`). # CPU-only net_telemetry unit test: net_telemetry.cc depends only on libc and # pthread, so this compiles it directly with its own harness (no GPU, no gtest, diff --git a/projects/rccl/test/device/GinAnvilIpcDevice_test.cpp b/projects/rccl/test/device/GinAnvilIpcDevice_test.cpp index 56f74c818f2c..22a583527327 100644 --- a/projects/rccl/test/device/GinAnvilIpcDevice_test.cpp +++ b/projects/rccl/test/device/GinAnvilIpcDevice_test.cpp @@ -185,7 +185,7 @@ __global__ void kernelDetailHelpers(const ncclGinAnvilSdmaGPUContext* ctx, int b fusedOut[0] = useSdmaFusedSignal(mut, true, true, false, ncclGinSignalInc); markSdmaDirty(mut, 1, mut->numChannels, effChOut[0]); if (mut->sdmaDirty) { - dirtyOut[0] = __hip_atomic_load(mut->sdmaDirty, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + dirtyOut[0] = __scoped_atomic_load_n(mut->sdmaDirty, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE); } } diff --git a/projects/rccl/test/device/GinAnvilSdmaTemplate_test.cpp b/projects/rccl/test/device/GinAnvilSdmaTemplate_test.cpp index e2755625420b..2f5f2d39a30d 100644 --- a/projects/rccl/test/device/GinAnvilSdmaTemplate_test.cpp +++ b/projects/rccl/test/device/GinAnvilSdmaTemplate_test.cpp @@ -156,7 +156,7 @@ __global__ void kernelPutSdmaPath(TemplateHarness* h, uint64_t* dirtyOut) { reinterpret_cast(&h->srcMh), 0, 256, sig, ncclGinSignalInc, 0, false, 0, false, nullptr, cuda::thread_scope_system, cuda::thread_scope_system); if (h->ctx.sdmaDirty) { - dirtyOut[0] = __hip_atomic_load(h->ctx.sdmaDirty, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + dirtyOut[0] = __scoped_atomic_load_n(h->ctx.sdmaDirty, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE); } } diff --git a/projects/rccl/test/host/run_host_tests.sh b/projects/rccl/test/host/run_host_tests.sh index 71921adf233b..49d8abb492b7 100755 --- a/projects/rccl/test/host/run_host_tests.sh +++ b/projects/rccl/test/host/run_host_tests.sh @@ -22,10 +22,13 @@ # prerequisite the host tests compile against # configure configure test/host # build build all host binaries (default target) +# guards kernel-count pytest plus src/include/test_poison_hip_atomics.py # run run the suite (timestamped log + JUnit XML). Always emits # llvm source-based coverage profiles (*.profraw) into # /coverage (requires the host tests to be built -# with -DHOST_TEST_COVERAGE=ON, the default) +# with -DHOST_TEST_COVERAGE=ON, the default). Also runs the +# CPU-only guards: kernel-count pytest plus the +# __hip_atomic_* poison compile probe. # coverage turn the per-binary *.profraw profiles from `run` into # reports: a per-binary text/HTML report + lcov tracefile # (clean, no hash mismatch), plus an overall line/branch union @@ -195,19 +198,43 @@ do_host_tests() { return "$rc" } +# CPU-only compile probes that #pragma GCC poison in poison_hip_atomics.h +# actually rejects __hip_atomic_* (in particular __hip_atomic_load). Needs +# amdclang++ and HIP headers from ROCM_PATH; no GPU and no librccl.so. +# A missing compiler fails here rather than skipping: the probes can only skip, +# and a fully skipped unittest run exits 0. Export RCCL_POISON_TEST_ALLOW_SKIP=1 +# to opt out when running this phase on a box without ROCm. +do_poison_hip_atomics() { + echo "==> Poison HIP atomics (src/include/test_poison_hip_atomics.py)" + python3 "$RCCL_ROOT/src/include/test_poison_hip_atomics.py" +} + # Run the kernel-count guard pytest suite (test/kernel-count) in a local venv so # the lean host-test image needs no system pytest. See that dir's README. -do_guards() { +do_kernel_count_guards() { echo "==> Kernel-count guards (pytest: test/kernel-count)" local gd="$RCCL_ROOT/test/kernel-count" local venv="$gd/venv" if [ ! -x "$venv/bin/pytest" ]; then - python3 -m venv "$venv" - "$venv/bin/pip" install -q --disable-pip-version-check -r "$gd/requirements.txt" + python3 -m venv "$venv" \ + && "$venv/bin/pip" install -q --disable-pip-version-check -r "$gd/requirements.txt" \ + || { echo "ERROR: could not provision $venv" >&2; return 1; } fi "$venv/bin/python" -m pytest "$gd/tests" -v } +# All CPU-only guards: the kernel-count pytest suite, then the __hip_atomic_* +# poison compile probe. Collected with `|| rc=1` rather than run back to back so +# that under `set -e` (line 53) a kernel-count failure still leaves the poison +# probe running and reported, instead of aborting the phase at the first one. +# Same idiom as do_host_tests above. +do_guards() { + local rc=0 + do_kernel_count_guards || rc=1 + do_poison_hip_atomics || rc=1 + return "$rc" +} + # Turn the per-binary profraw sets produced by the `run` phase into coverage # reports. Two levels of output: # @@ -366,10 +393,14 @@ do_coverage() { # (and `all` ends with it), so adding a future check here makes both CI and # local runs pick it up automatically -- no dispatch or workflow-YAML change. # do_host_tests runs first so the JUnit XML artifact is always produced before a -# later guard can gate. +# later guard can gate. Both are collected rather than chained: under `set -e` a +# gtest failure would otherwise abort the phase and drop the guards entirely, so +# one red signal would hide the other. do_run() { - do_host_tests "$@" - do_guards + local rc=0 + do_host_tests "$@" || rc=1 + do_guards || rc=1 + return "$rc" } case "$PHASE" in diff --git a/projects/rccl/tools/llx_latency_bench/llx_latency_bench.c b/projects/rccl/tools/llx_latency_bench/llx_latency_bench.c index 92d790f1dfeb..494b7d87601c 100644 --- a/projects/rccl/tools/llx_latency_bench/llx_latency_bench.c +++ b/projects/rccl/tools/llx_latency_bench/llx_latency_bench.c @@ -162,11 +162,11 @@ typedef const uint64_t __attribute__((address_space(1)))* cgptr64; __device__ __forceinline__ uint64_t sys_load64(cgptr64 p) -{ return __hip_atomic_load(p, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); } +{ return __scoped_atomic_load_n(p, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM); } __device__ __forceinline__ void sys_store64(gptr64 p, uint64_t v) -{ __hip_atomic_store(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); } +{ __scoped_atomic_store_n(p, v, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM); } // Payload hash: nonzero for any (grp, lane, flag) where flag >= 1 — keeps a // freshly-zeroed buffer from passing verify by accident. diff --git a/projects/rccl/tools/rccl-device-compile b/projects/rccl/tools/rccl-device-compile index faac990a2e97..ff028c3f6252 100755 --- a/projects/rccl/tools/rccl-device-compile +++ b/projects/rccl/tools/rccl-device-compile @@ -719,6 +719,16 @@ def parse_compiler_flags(argv): elif arg == '-I' and i + 1 < len(argv): forwarded.extend(['-I', argv[i + 1]]) i += 1 + elif arg in ('-include', '-imacros') and i + 1 < len(argv): + # Separate-argument forms: the value is a path, not a source file, + # so it must not fall through to `sources` below. RCCL's own build + # emits the joined --include= spelling (see the poison block + # in src/CMakeLists.txt), which the catch-all further down already + # forwards; this branch is here so a hand-written or third-party + # flag in the separate-argument form does not silently become a + # source file. + forwarded.extend([arg, argv[i + 1]]) + i += 1 elif arg.startswith('-'): forwarded.append(arg) else: