fix(device-ir): remove dead epoch accessors, fix sync and guard issues

Author: MC952-arch

bindings/ir/flagcx_device_scalar_ir.h

@@ -8,7 +8,6 @@
  *
  * Design:
  *   - CoopAny/Team replaced by flagcxCoopKind_t / flagcxTeamKind_t enums
- *   - Barrier split arrive/wait use stampEpoch internally
  *   - Net (transport) obtained via flagcxDevNetGetFromCommS (pre-allocated)
  *
  * When compiled to LLVM bitcode (clang -x cuda --cuda-device-only),

bindings/ir/public_symbols.txt

@@ -102,10 +102,12 @@ flagcxGetTeamInter
 flagcxGetTeamIntra
 flagcxGetTeamWorld
 flagcxInterBarrierArriveS
+flagcxInterBarrierSessionInit
 flagcxInterBarrierSessionSync
 flagcxInterBarrierSyncS
 flagcxInterBarrierWaitS
 flagcxIntraBarrierArriveS
+flagcxIntraBarrierSessionInit
 flagcxIntraBarrierSessionArrive
 flagcxIntraBarrierSessionSync
 flagcxIntraBarrierSessionWait
@@ -116,6 +118,7 @@ flagcxTeamRankToIntraS
 flagcxTeamRankToWorldC
 flagcxTeamRankToWorldS
 flagcxWorldBarrierArriveS
+flagcxWorldBarrierSessionInit
 flagcxWorldBarrierSessionSync
 flagcxWorldBarrierSyncS
 flagcxWorldBarrierWaitS

flagcx/adaptor/flagcx_device.cc

@@ -1324,8 +1324,8 @@ extern "C" flagcxResult_t flagcxDevCommGetDevicePtr(flagcxDevComm_t devComm,
         deviceAdaptor->deviceMemcpy(netCtxField, &netDevPtr, sizeof(void *),
                                     flagcxMemcpyHostToDevice, NULL, NULL),
         res, fail);
-    // Ensure net construction is complete before pointer escapes
-    FLAGCXCHECKGOTO(deviceAdaptor->streamSynchronize(NULL), res, fail);
+    // Ensure net construction + pointer patch are visible to all streams
+    FLAGCXCHECKGOTO(deviceAdaptor->deviceSynchronize(), res, fail);
   }
 
   devComm->cachedDevicePtr = dPtr;

flagcx/adaptor/include/device_api/default_comm_traits.h

@@ -854,10 +854,6 @@ struct Barrier<Default<P>, flagcxTeamTagIntra, Coop> {
     arrive(order);
     wait(order);
   }
-
-  // S-API epoch accessors
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getEpoch() const { return _epoch; }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setEpoch(uint64_t e) { _epoch = e; }
 };
 
 // ---- Barrier<Default<P>, flagcxTeamTagInter, Coop> ----
@@ -937,10 +933,6 @@ struct Barrier<Default<P>, flagcxTeamTagInter, Coop> {
     arrive(order);
     wait(order);
   }
-
-  // S-API epoch accessors
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getEpoch() const { return _epoch; }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setEpoch(uint64_t e) { _epoch = e; }
 };
 
 // ---- Barrier<Default<P>, flagcxTeamTagWorld, Coop> ----
@@ -1026,20 +1018,6 @@ struct Barrier<Default<P>, flagcxTeamTagWorld, Coop> {
       _intra.wait(order);
     }
   }
-
-  // S-API epoch accessors (world = intra + inter)
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getIntraEpoch() const {
-    return _intra.getEpoch();
-  }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setIntraEpoch(uint64_t e) {
-    _intra.setEpoch(e);
-  }
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getInterEpoch() const {
-    return _inter.getEpoch();
-  }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setInterEpoch(uint64_t e) {
-    _inter.setEpoch(e);
-  }
 };
 
 #endif // FLAGCX_FALLBACK_DEVICE_TRAITS_H_

flagcx/adaptor/include/device_api/flagcx_device_core.h

@@ -534,14 +534,6 @@ struct flagcxDevBarrier<flagcxTeamTagIntra, Coop> {
   sync(flagcxDeviceMemoryOrder_t order = flagcxDeviceMemoryOrderAcqRel) {
     _impl.sync(order);
   }
-
-  // S-API epoch accessors
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getEpoch() const {
-    return _impl.getEpoch();
-  }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setEpoch(uint64_t e) {
-    _impl.setEpoch(e);
-  }
 };
 
 // ============================================================
@@ -918,14 +910,6 @@ struct flagcxDevBarrier<flagcxTeamTagInter, Coop> {
        flagcxDevNetFenceLevel fence = flagcxDevNetFenceLevel::Relaxed) {
     _impl.sync(order, fence);
   }
-
-  // S-API epoch accessors
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getEpoch() const {
-    return _impl.getEpoch();
-  }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setEpoch(uint64_t e) {
-    _impl.setEpoch(e);
-  }
 };
 
 // ---- World ----
@@ -970,20 +954,6 @@ struct flagcxDevBarrier<flagcxTeamTagWorld, Coop> {
        flagcxDevNetFenceLevel fence = flagcxDevNetFenceLevel::Relaxed) {
     _impl.sync(order, fence);
   }
-
-  // S-API epoch accessors (world = intra + inter)
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getIntraEpoch() const {
-    return _impl.getIntraEpoch();
-  }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setIntraEpoch(uint64_t e) {
-    _impl.setIntraEpoch(e);
-  }
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getInterEpoch() const {
-    return _impl.getInterEpoch();
-  }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setInterEpoch(uint64_t e) {
-    _impl.setInterEpoch(e);
-  }
 };
 
 // Backward-compatible aliases

flagcx/adaptor/include/device_api/nvidia_comm_traits.h

@@ -500,10 +500,6 @@ struct Barrier<NvidiaVendor, flagcxTeamTagIntra, Coop> {
     reinterpret_cast<ncclLsaBarrierSession<ncclCoopCta> *>(_implStorage)
         ->sync(ncclCoopCta(), Atomic::toNativeOrder(order));
   }
-
-  // S-API epoch accessors — NCCL manages epoch internally; no-op
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getEpoch() const { return 0; }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setEpoch(uint64_t) {}
 };
 
 // ---- Barrier<NvidiaVendor, flagcxTeamTagInter, Coop> ----
@@ -559,10 +555,6 @@ struct Barrier<NvidiaVendor, flagcxTeamTagInter, Coop> {
                  flagcxDevNetFenceLevelMap[static_cast<int>(fence)]);
     }
   }
-
-  // S-API epoch accessors — NCCL manages epoch internally; no-op
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getEpoch() const { return 0; }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setEpoch(uint64_t) {}
 };
 
 // ---- Barrier<NvidiaVendor, flagcxTeamTagWorld, Coop> ----
@@ -661,12 +653,6 @@ struct Barrier<NvidiaVendor, flagcxTeamTagWorld, Coop> {
                  flagcxDevNetFenceLevelMap[static_cast<int>(fence)]);
     }
   }
-
-  // S-API epoch accessors — NCCL manages epoch internally; no-op
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getIntraEpoch() const { return 0; }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setIntraEpoch(uint64_t) {}
-  FLAGCX_DEVICE_INLINE_DECORATOR uint64_t getInterEpoch() const { return 0; }
-  FLAGCX_DEVICE_INLINE_DECORATOR void setInterEpoch(uint64_t) {}
 };
 #endif // NCCL_CHECK_CUDACC
 

flagcx/kernels/device_api_host_helpers.cu

@@ -28,6 +28,8 @@ static __global__ void flagcxDevNetConstructKernel(flagcxDevNet *nets,
 // synchronizing the stream afterwards.
 extern "C" void flagcxDevNetLaunchConstruct(void *devNets, void *devComm,
                                             int count, void *stream) {
+  if (count <= 0 || devNets == nullptr || devComm == nullptr)
+    return;
   int threads = (count < 256) ? count : 256;
   int blocks = (count + threads - 1) / threads;
   flagcxDevNetConstructKernel<<<blocks, threads, 0, (cudaStream_t)stream>>>(