fix(rccl): use MI450 cooperative 128-bit builtins and explicit vector loads/stores in LL128

projects/rccl/CHANGELOG.md

@@ -188,6 +188,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * Updated the RMA plugin interface to v15.
 * Reduced communicator host memory by allocating topology path link arrays to their actual length.
 * Widened the LL128 per-thread shared-memory slice from 8 to 32 elements on gfx1250 (MI450/MI455); other architectures are unchanged. Forced Ring/LL128 bandwidth roughly doubles at 16 MB and above. Per-block LDS rises from 36448 to 85472 bytes, and because the scratch region is shared across protocols this applies to every kernel launch, not only LL128 ones. The tuner's protocol selection is unchanged, so the gain is only visible when LL128 is selected explicitly.
+* gfx1250 LL128 comm-FIFO accesses use cooperative 128-bit atomics (`load128Fifo` / `store128Fifo`) so each 8-lane group moves one non-tearing 128-byte line at system scope. User-buffer loads and stores stay on ordinary vector operations. `-DDWORDX4_INTRINSICS=OFF` disables the cooperative builtins together with the other 128-bit intrinsics.
 * Registered-window CE AllReduce is now eligible without requiring ``NCCL_CTA_POLICY=2``
   (``CTA_POLICY_ZERO``); it is gated only by the per-arch table window and ``RCCL_CE_ALLREDUCE``.
   ``RCCL_FORCE_CE_ALLREDUCE`` now applies exclusively to the non-registered (2-shot/staged) CE

projects/rccl/src/algorithms/dda/device/CollCommon_ll128.h

@@ -182,8 +182,9 @@ __device__ __forceinline__ void ddaLL128LoadAtomic16B(const uint64_t* src, uint6
     v4i v;
     uint64_t w[2];
   } u;
-  u.v = __builtin_amdgcn_cooperative_atomic_load_8x16B(
-    (v4i_gptr)src, __ATOMIC_RELAXED, RCCL_SYSTEM_SYNCSCOPE);
+  // See rcclIdentity in rccl_ptr.h.
+  u.v = rcclIdentity(__builtin_amdgcn_cooperative_atomic_load_8x16B(
+    (v4i_gptr)src, __ATOMIC_RELAXED, RCCL_SYSTEM_SYNCSCOPE));
   lo = u.w[0];
   hi = u.w[1];
 #else

projects/rccl/src/device/op128.h

@@ -12,33 +12,29 @@
 
 #include "rccl_ptr.h"
 
+// Per-thread 128-bit access for user buffers and other callers that do not
+// execute on every lane of an 8-lane group. The gfx1250 cooperative atomics
+// live in load128Fifo/store128Fifo; using them here would issue the builtin
+// from a partial group (loadRegsBegin and storeRegs skip flag lanes).
 inline __device__ void load128(const uint64_t* ptr, uint64_t& v0, uint64_t& v1) {
+  rcclB128 u;
 #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
-  union {
-    v4u v;
-    uint64_t u64[2];
-  } u;
   u.v = __builtin_amdgcn_global_load_b128((v4u_gptr)ptr, RCCL_SYSTEM_SYNCSCOPE);
-  v0 = u.u64[0];
-  v1 = u.u64[1];
 #else
-  v0 = __builtin_nontemporal_load((u64_gptr)ptr);
-  v1 = __builtin_nontemporal_load((u64_gptr)ptr + 1);
+  u.v = __builtin_nontemporal_load((v4u_gptr)ptr);
 #endif
+  v0 = u.u64[0];
+  v1 = u.u64[1];
 }
 
 inline __device__ void store128(uint64_t* ptr, uint64_t v0, uint64_t v1) {
-#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
-  union {
-    v4u v;
-    uint64_t u64[2];
-  } u;
+  rcclB128 u;
   u.u64[0] = v0;
   u.u64[1] = v1;
+#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
   __builtin_amdgcn_global_store_b128((v4u_gptr)ptr, u.v, RCCL_SYSTEM_SYNCSCOPE);
 #else
-  *((u64_gptr)ptr) = v0;
-  *((u64_gptr)ptr + 1) = v1;
+  *((v4u_gptr)ptr) = u.v;
 #endif
 }