Fix minor issues

Author: MC952-arch

flagcx/adaptor/ccl/nccl_adaptor.cc

@@ -420,6 +420,18 @@ flagcxResult_t ncclAdaptorDevCommCreate(flagcxInnerComm_t comm,
   ncclReqs.ginContextCount = reqs->interContextCount;
   ncclReqs.ginCounterCount = reqs->interCounterCount;
 
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2, 29, 7)
+  // GIN resources require an explicit connection type starting with 2.29.7.
+  const bool needsFullGin = reqs->interForceEnable ||
+                            reqs->interSignalCount > 0 ||
+                            reqs->interCounterCount > 0;
+  const bool needsRailGin =
+      reqs->barrierCount > 0 || reqs->interBarrierCount > 0;
+  ncclReqs.ginConnectionType = needsFullGin   ? NCCL_GIN_CONNECTION_FULL
+                               : needsRailGin ? NCCL_GIN_CONNECTION_RAIL
+                                              : NCCL_GIN_CONNECTION_NONE;
+#endif
+
   flagcxResult_t ret =
       ncclDevCommCreateHelper(comm->base, &ncclReqs, &inner->base);
   if (ret != flagcxSuccess) {

flagcx/adaptor/include/device_api/nccl_comm_traits.h

@@ -11,6 +11,8 @@
 #ifndef FLAGCX_NCCL_COMM_TRAITS_H_
 #define FLAGCX_NCCL_COMM_TRAITS_H_
 
+#include <stdio.h>
+
 #include "nccl.h"
 #ifndef __CUDACC__
 #include "flagcx.h"
@@ -461,10 +463,28 @@ struct CommTraits<NcclBackend> {
       return flagcxInternalError;
     }
 
-    // --- get stub (fallback-only, vendor has no RDMA READ) ---
+    // --- One-sided: get ---
     template <typename Coop>
-    FLAGCX_DEVICE_INLINE_DECORATOR void get(Team, int, Window, size_t, Window,
-                                            size_t, size_t, Coop) const {}
+    FLAGCX_DEVICE_INLINE_DECORATOR void
+    get(Team team, int peer, Window src, size_t srcOff, Window dst,
+        size_t dstOff, size_t bytes, Coop coop) const {
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2, 30, 7)
+      _gin.get((ncclTeam_t)team, peer, src._impl, srcOff, dst._impl, dstOff,
+               bytes, coop._impl);
+#else
+      if (coop.threadRank() == 0) {
+        printf("FLAGCX WARN: NCCL Device Get requires NCCL >= 2.30.7; "
+               "operation is a no-op\n");
+      }
+      (void)team;
+      (void)peer;
+      (void)src;
+      (void)srcOff;
+      (void)dst;
+      (void)dstOff;
+      (void)bytes;
+#endif
+    }
   };
 #endif // FLAGCX_CHECK_DEVICE_CC
 };

test/unittest/device_api/test_device_ir_unified_inter.cpp

@@ -88,6 +88,7 @@ int main(int argc, char *argv[]) {
   // Inter suite uses 6 combinations (INTER + WORLD) for most tests
   // S18 uses 8 signal slots (includes extra BLOCK single-leader patterns)
   flagcxDevCommRequirements reqs = FLAGCX_DEV_COMM_REQUIREMENTS_INITIALIZER;
+  reqs.barrierCount = FLAGCX_DEVICE_CTA_COUNT;
   reqs.intraBarrierCount = FLAGCX_DEVICE_CTA_COUNT;
   reqs.interBarrierCount = FLAGCX_DEVICE_CTA_COUNT;
   reqs.interSignalCount =

test/unittest/device_api/test_device_ir_unified_intra.cpp

@@ -84,6 +84,7 @@ int main(int argc, char *argv[]) {
   // Intra suite uses 6 combinations (INTRA + WORLD) for most tests
   // S18 uses 8 signal slots (includes extra BLOCK single-leader patterns)
   flagcxDevCommRequirements reqs = FLAGCX_DEV_COMM_REQUIREMENTS_INITIALIZER;
+  reqs.barrierCount = FLAGCX_DEVICE_CTA_COUNT;
   reqs.intraBarrierCount = FLAGCX_DEVICE_CTA_COUNT;
   reqs.interBarrierCount = FLAGCX_DEVICE_CTA_COUNT;
   reqs.interSignalCount =