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,23 @@ struct CommTraits<NcclBackend> {
       return flagcxInternalError;
     }
 
-    // --- get stub (fallback-only, vendor has no RDMA READ) ---
+    // --- One-sided: get (unsupported by the NCCL backend) ---
     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 (coop.threadRank() == 0) {
+        printf("FLAGCX WARN: NCCL Device Get is unsupported; operation is a "
+               "no-op\n");
+      }
+      (void)team;
+      (void)peer;
+      (void)src;
+      (void)srcOff;
+      (void)dst;
+      (void)dstOff;
+      (void)bytes;
+    }
   };
 #endif // FLAGCX_CHECK_DEVICE_CC
 };

test/kernel/Makefile

@@ -16,8 +16,14 @@ all:
 		USE_SHMEM=$(USE_SHMEM) SHMEM_HOME=$(SHMEM_HOME) \
 		FORCE_DEFAULT_PATH=$(FORCE_DEFAULT_PATH)
 
+ifeq ($(strip $(KERNEL_PLATFORM)),)
+clean:
+	@$(MAKE) -C $(CURDIR)/nvidia clean
+	@$(MAKE) -C $(CURDIR)/du clean
+else
 clean:
 	@$(MAKE) -C $(KERNEL_PLATFORM_DIR) clean
+endif
 
 $(FORWARD_TARGETS):
 	@$(MAKE) -C $(KERNEL_PLATFORM_DIR) $@ \

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 =