Refine build system and update documentation for v0.6 release (flagos-ai#291)

* Modify Makefile and update v0.6 release notes

* Fix a bug

Author: MC952-arch

Makefile

@@ -17,6 +17,7 @@ USE_DU ?= 0
 USE_MPI ?= 0
 USE_UCX ?= 0
 USE_IBUC ?= 0
+COMPILE_KERNEL ?= 0
 
 # set to empty if not provided
 DEVICE_HOME ?=
@@ -116,6 +117,7 @@ UCX_LIB =
 UCX_INCLUDE =
 UCX_LINK =
 NET_ADAPTOR_FLAG =
+COMPILE_KERNEL_FLAG =
 ifeq ($(USE_NVIDIA), 1)
 	DEVICE_LIB = $(DEVICE_HOME)/lib64
 	DEVICE_INCLUDE = $(DEVICE_HOME)/include
@@ -242,6 +244,10 @@ ifeq ($(USE_IBUC), 1)
 	NET_ADAPTOR_FLAG += -DUSE_IBUC
 endif
 
+ifeq ($(COMPILE_KERNEL), 1)
+	COMPILE_KERNEL_FLAG = -DCOMPILE_KERNEL
+endif
+
 LIBDIR := $(BUILDDIR)/lib
 OBJDIR := $(BUILDDIR)/obj
 
@@ -263,11 +269,12 @@ LIBSRCFILES:= \
 	$(wildcard flagcx/adaptor/tuner/*.cc) \
 	$(wildcard flagcx/service/*.cc)
 
+ifeq ($(COMPILE_KERNEL), 1)
 DEVSRCFILES:= \
 	$(wildcard flagcx/kernels/*.$(DEVICE_FILE_EXTENSION))
-
-LIBOBJ:= $(LIBSRCFILES:%.cc=$(OBJDIR)/%.o)
 DEVOBJ:= $(DEVSRCFILES:%.$(DEVICE_FILE_EXTENSION)=$(OBJDIR)/%.o)
+endif
+LIBOBJ:= $(LIBSRCFILES:%.cc=$(OBJDIR)/%.o)
 
 TARGET = libflagcx.so
 all: $(LIBDIR)/$(TARGET)
@@ -288,6 +295,7 @@ print_var:
 	@echo "USE_MUSA: $(USE_MUSA)"
 	@echo "USE_DU: $(USE_DU)"
 	@echo "USE_AMD: $(USE_AMD)"
+	@echo "COMPILE_KERNEL: $(COMPILE_KERNEL)"
 	@echo "DEVICE_LIB: $(DEVICE_LIB)"
 	@echo "DEVICE_INCLUDE: $(DEVICE_INCLUDE)"
 	@echo "CCL_LIB: $(CCL_LIB)"
@@ -303,26 +311,37 @@ print_var:
 	@echo "USE_IBUC: $(USE_IBUC)"
 	@echo "NET_ADAPTOR_FLAG: $(NET_ADAPTOR_FLAG)"
 
-$(LIBDIR)/$(TARGET): $(LIBOBJ) $(DEVOBJ) $(OBJDIR)/kernel_dlink.o
+ifeq ($(COMPILE_KERNEL), 1)
+DEVOBJS = $(DEVOBJ) $(OBJDIR)/kernel_dlink.o
+else
+DEVOBJS =
+endif
+
+$(LIBDIR)/$(TARGET): $(LIBOBJ) $(DEVOBJS)
 	@mkdir -p `dirname $@`
 	@echo "Linking   $@"
 	@g++ $^ -o $@ -L$(CCL_LIB) -L$(DEVICE_LIB) -L$(HOST_CCL_LIB) -L$(UCX_LIB) -shared -fvisibility=default -Wl,--no-as-needed -Wl,-rpath,$(LIBDIR) -Wl,-rpath,$(CCL_LIB) -Wl,-rpath,$(HOST_CCL_LIB) -Wl,-rpath,$(UCX_LIB) -lpthread -lrt -ldl $(CCL_LINK) $(DEVICE_LINK) $(HOST_CCL_LINK) $(UCX_LINK) -g
 
-
 $(OBJDIR)/%.o: %.cc
 	@mkdir -p `dirname $@`
 	@echo "Compiling $@"
 	@g++ $< -o $@ $(foreach dir,$(INCLUDEDIR),-I$(dir)) -I$(CCL_INCLUDE) -I$(DEVICE_INCLUDE) -I$(HOST_CCL_INCLUDE) -I$(UCX_INCLUDE) $(ADAPTOR_FLAG) $(HOST_CCL_ADAPTOR_FLAG) $(NET_ADAPTOR_FLAG) -c -fPIC -fvisibility=default -Wvla -Wno-unused-function -Wno-sign-compare -Wall -MMD -MP -g
 
+ifeq ($(COMPILE_KERNEL), 1)
 $(OBJDIR)/kernel_dlink.o: $(DEVOBJ)
 	@$(DEVICE_LINKER) $^ -o $@ $(DEVICE_LINK) $(DEVICE_LINK_FLAG)
 
 $(OBJDIR)/%.o: %.$(DEVICE_FILE_EXTENSION)
 	@mkdir -p `dirname $@`
 	@echo "Compiling $@ ($(DEVICE_RUNTIME))"
-	@$(DEVICE_COMPILER) $< -o $@ $(foreach dir,$(INCLUDEDIR),-I$(dir)) -I$(CCL_INCLUDE) -I$(DEVICE_INCLUDE) -I$(HOST_CCL_INCLUDE) -I$(UCX_INCLUDE) $(ADAPTOR_FLAG) $(HOST_CCL_ADAPTOR_FLAG) $(NET_ADAPTOR_FLAG) $(DEVICE_COMPILE_FLAG)
+	@$(DEVICE_COMPILER) $< -o $@ $(foreach dir,$(INCLUDEDIR),-I$(dir)) -I$(CCL_INCLUDE) -I$(DEVICE_INCLUDE) -I$(HOST_CCL_INCLUDE) -I$(UCX_INCLUDE) $(ADAPTOR_FLAG) $(HOST_CCL_ADAPTOR_FLAG) $(NET_ADAPTOR_FLAG) $(DEVICE_COMPILE_FLAG) $(COMPILE_KERNEL_FLAG) -g
+endif
 
+ifeq ($(COMPILE_KERNEL), 1)
 -include $(LIBOBJ:.o=.d) $(DEVOBJ:.o=.d)
+else
+-include $(LIBOBJ:.o=.d)
+endif
 
 clean:
 	@rm -rf $(LIBDIR)/$(TARGET) $(OBJDIR)

README.md

@@ -1,30 +1,34 @@
 [<img src="docs/images/flagopen.png">](https://flagopen.baai.ac.cn/)
 
 ## Latest News
+- **[2025/10]** Released [v0.6](https://github.com/FlagOpen/FlagCX/tree/release/v0.6):
+  - Supported device-buffer P2P communication to achieve intra-node SendRecv operations.
+  - Introduced Device-initiated, Host-launched device-side primitives, enabling kernel-based communication directly from the device.
+  - Enhanced automatic tuning functionality, achieving up to 50% performance improvement on Metax platforms for the AllReduce operation.
 - **[2025/09]** Released [v0.5](https://github.com/FlagOpen/FlagCX/tree/release/v0.5):
-  - Adds AMD support (hipAdaptor and rcclAdaptor).
-  - Introduces flagcxNetAdaptor to unify network backends, currently supporting SOCKET, IBRC, UCX and IBUC (experimently).
-  - Enables zero-copy device-buffer RDMA (user-buffer RDMA) to boost small-message performance.
-  - Supports automatic tuning in homogeneous scenarios via flagcxTuner.
-  - Integrates automated PyTorch API tests into CI/CD.
+  - Added AMD support (hipAdaptor and rcclAdaptor).
+  - Introduced flagcxNetAdaptor to unify network backends, currently supporting SOCKET, IBRC, UCX and IBUC (experimently).
+  - Enabled zero-copy device-buffer RDMA (user-buffer RDMA) to boost small-message performance.
+  - Supported automatic tuning in homogeneous scenarios via flagcxTuner.
+  - Integrated automated PyTorch API tests into CI/CD.
 - **[2025/08]** Released [v0.4](https://github.com/FlagOpen/FlagCX/tree/release/v0.4):
-  - Supports heterogeneous training of ERNIE4.5 on Nvidia and Iluvatar GPUs with Paddle + FlagCX.
-  - Enables more robust and flexible deployments with full support of heterogeneous communication across arbitrary NIC configurations (bug fixes). 
-  - Introduces an early experimental net plugin interface extending its support for both IBRC and SOCKET, along with the ability to register device buffers via DMA-BUF.
-  - Adds an InterOp-level DSL to allow users designing customized C2C algorithms.
-  - Provides usage documentation under docs/.
+  - Supported heterogeneous training of ERNIE4.5 on Nvidia and Iluvatar GPUs with Paddle + FlagCX.
+  - Enabled more robust and flexible deployments with full support of heterogeneous communication across arbitrary NIC configurations (bug fixes). 
+  - Introduced an early experimental net plugin interface extending its support for both IBRC and SOCKET, along with the ability to register device buffers via DMA-BUF.
+  - Added an InterOp-level DSL to allow users designing customized C2C algorithms.
+  - Provided usage documentation under docs/.
 - **[2025/07]** Released [v0.3](https://github.com/FlagOpen/FlagCX/tree/release/v0.3):
-  - Integrates three additional native communication libraries: HCCL, MUSACCL and MPI.
-  - Enhances heterogeneous collective communication operations with pipeline optimizations. 
-  - Introduces a device-side function mechanism to enable device-buffer RDMA, complementing the original host-side function mechanism.
-  - Delivers a full-stack open-source solution, FlagScale + FlagCX, for efficient heterogeneous prefilling-decoding disaggregation.
+  - Integrated three additional native communication libraries: HCCL, MUSACCL and MPI.
+  - Enhanced heterogeneous collective communication operations with pipeline optimizations. 
+  - Introduced a device-side function mechanism to enable device-buffer RDMA, complementing the original host-side function mechanism.
+  - Delivered a full-stack open-source solution, FlagScale + FlagCX, for efficient heterogeneous prefilling-decoding disaggregation.
 - **[2025/05]** Released [v0.2](https://github.com/FlagOpen/FlagCX/tree/release/v0.2):
-  - Integrates three additional native communications libraries, including MCCL, XCCL and DUCCL.
-  - Improves 11 heterogeneous collective communication operations with automatic topology detection, fully supporting both single-NIC and multi-NIC environments.
+  - Integrated three additional native communications libraries, including MCCL, XCCL and DUCCL.
+  - Improved 11 heterogeneous collective communication operations with automatic topology detection, fully supporting both single-NIC and multi-NIC environments.
 - **[2025/04]** Released [v0.1](https://github.com/FlagOpen/FlagCX/tree/release/v0.1):
-  - Integrates five native communications libraries including NCCL, IXCCL, CNCL, BOOTSTRAP and GLOO.
-  - Supports 11 heterogeneous collective communication operations using the originally proposed C2C (Cluster-to-Cluster) algorithm.
-  - Provides a full-stack open-source solution, FlagScale + FlagCX, for efficient heterogeneous training.
+  - Integrated five native communications libraries including NCCL, IXCCL, CNCL, BOOTSTRAP and GLOO.
+  - Supported 11 heterogeneous collective communication operations using the originally proposed C2C (Cluster-to-Cluster) algorithm.
+  - Provided a full-stack open-source solution, FlagScale + FlagCX, for efficient heterogeneous training.
   - Natively integrated into PaddlePaddle [v3.0.0](https://github.com/PaddlePaddle/Paddle/tree/v3.0.0), with support for both dynamic and static graphs.
 
 ## About
@@ -87,11 +91,11 @@ FlagCX also integrates with upper-layer applications such as PyTorch and PaddleP
 | all_to_all_single                | ✓                            | ✓            |
 | barrier                          | ✓                            | ✓            |
 
-The FlagCX PyTorch plugin has undergone comprehensive validation across multiple communication backends, as summarized below:
+In particular, PyTorch support is enabled via the FlagCX Torch Plugin, which provides native integration with the PyTorch distributed backend. This plugin has undergone comprehensive validation across diverse communication backends and hardware platforms, ensuring robust functionality, consistent performance, and compatibility in heterogeneous multi-chip environments, as summarized below:
 
-| FlagCX Backend         | NCCL | IXCCL | CNCL | MCCL | XCCL | DUCCL | HCCL | MUSACCL | RCCL |
-| ---------------------- | ---- | ----- | ---- | ---- | ---- | ----- | ---- | ------- | ---- |
-| PyTorch Plugin Support | ✓    | ✓     | ✓    | ✓    | ✓    | ✓     | ✓    | ☓       | ✓    |
+| FlagCX Backend   | NCCL | IXCCL | CNCL | MCCL | XCCL | DUCCL | HCCL | MUSACCL | RCCL |
+|:-----------------|:-----|:------|:-----|:-----|:-----|:------|:-----|:--------|:-----|
+| PyTorch Support  | ✓    | ✓     | ✓    | ✓    | ✓    | ✓     | ✓    | ☓       | ✓    |
 
 To enable heterogeneous cross-chip communication using the PyTorch DDP FlagCX backend, it is recommended to use identical PyTorch versions across all nodes. Mismatched versions may lead to initialization failures during process group setup. Further compatibility and performance tests will be conducted in future releases, and we warmly welcome community contributions to help expand and strengthen the validation matrix.
 

flagcx/core/flagcx_kernel_host.cc

@@ -0,0 +1,31 @@
+#include "comm.h"
+#include "flagcx.h"
+#include "flagcx_kernel.h"
+
+FLAGCX_HOST_DECORATOR flagcxResult_t dequeue(void *fifoBuffer,
+                                             flagcxDeviceTrigger_t trigger) {
+  int idx = -1;
+  unsigned long long int *buffer = (unsigned long long int *)fifoBuffer;
+  int capacity = buffer[0];
+  int distance = buffer[2] - buffer[1];
+  if (distance > 0) {
+    idx = buffer[1] % capacity;
+    buffer[1] = buffer[1] + 1;
+    __sync_synchronize();
+  }
+  if (idx > -1) {
+    memcpy((void *)trigger,
+           (void *)(buffer + 3 +
+                    sizeof(flagcxDeviceTrigger) /
+                        sizeof(unsigned long long int) * idx),
+           sizeof(flagcxDeviceTrigger));
+  } else {
+    memset((void *)trigger, 0, sizeof(flagcxDeviceTrigger));
+  }
+  return flagcxSuccess;
+}
+
+// __host__ flagcxResult_t flagcxFifo::enqueue(flagcxReduceTrigger trigger) {
+//   // to be implemented
+//   return flagcxNotSupported;
+// }

flagcx/include/flagcx_kernel.h

@@ -107,14 +107,15 @@ struct flagcxFifo {
 };
 typedef struct flagcxFifo *flagcxFifo_t;
 
+FLAGCX_HOST_DECORATOR flagcxResult_t dequeue(void *fifoBuffer,
+                                             flagcxDeviceTrigger_t trigger);
+#ifdef COMPILE_KERNEL
 // device-producer + host-consumer APIs
 FLAGCX_DEVICE_DECORATOR flagcxResult_t enqueue(void *fifoBuffer, uint64_t addr,
                                                uint64_t count,
                                                uint64_t peerRank,
                                                uint64_t datatype,
                                                uint64_t type);
-FLAGCX_HOST_DECORATOR flagcxResult_t dequeue(void *fifoBuffer,
-                                             flagcxDeviceTrigger_t trigger);
 // host-producer + device-consumer APIs
 // FLAGCX_HOST_DECORATOR flagcxResult_t enqueue(flagcxReduceTrigger trigger);
 // FLAGCX_HOST_DECORATOR  flagcxResult_t dequeue(flagcxReduceTrigger_t trigger);
@@ -134,4 +135,5 @@ FLAGCX_DEVICE_DECORATOR flagcxResult_t flagcxDeviceWait(void *fifoBuffer);
 void flagcxP2pDemo(const void *sendbuff, void *recvbuff, size_t count,
                    flagcxDataType_t datatype, int sendPeer, int recvPeer,
                    flagcxComm_t comm, flagcxStream_t stream);
+#endif // COMPILE_KERNEL
 #endif

flagcx/kernels/flagcx_kernel.cu flagcx/kernels/flagcx_kernel_device.cuflagcx/kernels/flagcx_kernel.cu renamed to flagcx/kernels/flagcx_kernel_device.cu

@@ -102,29 +102,6 @@ FLAGCX_DEVICE_DECORATOR flagcxResult_t enqueue(void *fifoBuffer, uint64_t addr,
   return flagcxSuccess;
 }
 
-FLAGCX_HOST_DECORATOR flagcxResult_t dequeue(void *fifoBuffer, flagcxDeviceTrigger_t trigger) {
-  int idx = -1;
-  unsigned long long int *buffer = (unsigned long long int *)fifoBuffer;
-  int capacity = buffer[0];
-  int distance = buffer[2] - buffer[1];
-  if (distance > 0) {
-    idx = buffer[1] % capacity;
-    buffer[1] = buffer[1] + 1;
-    __sync_synchronize();
-  }
-  if (idx > -1) {
-    memcpy((void *)trigger, (void *)(buffer + 3 + sizeof(flagcxDeviceTrigger) / sizeof(unsigned long long int) * idx), sizeof(flagcxDeviceTrigger));
-  } else {
-    memset((void *)trigger, 0, sizeof(flagcxDeviceTrigger));
-  }
-  return flagcxSuccess;
-}
-
-// __host__ flagcxResult_t flagcxFifo::enqueue(flagcxReduceTrigger trigger) {
-//   // to be implemented
-//   return flagcxNotSupported;
-// }
-
 // __device__ flagcxResult_t flagcxFifo::dequeue(flagcxReduceTrigger_t trigger) {
 //   // to be implemented
 //   return flagcxNotSupported;

test/kernel/Makefile

@@ -25,7 +25,7 @@ all: test-kernel-demo
 
 test-kernel-demo: test_kernel_demo.cpp
 	@echo "Compiling $@"
-	@$(COMPILER) $(EXTRA_COMPILER_FLAG) -o test_kernel_demo test_kernel_demo.cpp $(LIBSRCFILES) -I../../flagcx/include -I../../flagcx/adaptor/include -I../../flagcx/service/ -I../../flagcx/core/ -I$(INCLUDEDIR) -I$(MPI_INCLUDE) -I$(DEVICE_INCLUDE) -L../../build/lib -L$(MPI_LIB) -L$(DEVICE_LIB) -lflagcx $(MPI_LINK) $(DEVICE_LINK)
+	@$(COMPILER) $(EXTRA_COMPILER_FLAG) -o test_kernel_demo test_kernel_demo.cpp $(LIBSRCFILES) -I../../flagcx/include -I../../flagcx/adaptor/include -I../../flagcx/service/ -I../../flagcx/core/ -I$(INCLUDEDIR) -I$(MPI_INCLUDE) -I$(DEVICE_INCLUDE) -L../../build/lib -L$(MPI_LIB) -L$(DEVICE_LIB) -lflagcx $(MPI_LINK) $(DEVICE_LINK) -DCOMPILE_KERNEL
 
 clean:
 	@rm -f test_kernel_demo