Merge branch 'main' into epv1_fp8_noquant

Author: maning00

.github/workflows/ci.yml

@@ -6,6 +6,7 @@ jobs:
   build:
     name: mori build
     runs-on: label-1
+    timeout-minutes: 10
     steps:
       - name: clone source code
         run: |
@@ -27,14 +28,19 @@ jobs:
           #pip install --target /apps/mori-ci/packages -r requirements-build.txt
           pip install -r requirements-build.txt
           #git submodule update --init --recursive
+          #it indicates the use of a mellanox NIC by default, it has been verified on MI300X/MI355X.
           #pip3 install . --no-build-isolation
+          #The string "BNXT" indicates the use of a Thor 2 NIC, it has been verified on MI300X/MI355X.
+          #USE_BNXT=ON pip3 install . --no-build-isolation 
+          #The string "IONIC" indicates the use of a AINIC, it has been verified on MI300X/MI355X.
           USE_IONIC=ON pip3 install . --no-build-isolation
           # --target /apps/mori-ci/python-packages
           scp -P 2233 -r /root/actions-runner/_work/mori/mori/mori root@smci355-ccs-aus-n08-33:/root/actions-runner/_work/mori/mori/
   test:
     name: mori test
     needs: build
     runs-on: label-1
+    timeout-minutes: 10
     steps:
         #- uses: actions/checkout@v4
         #- name: Set up Python

CMakeLists.txt

@@ -31,6 +31,7 @@ option(BUILD_PYBINDS "Whether to build mori python bindings" ON)
 option(BUILD_TESTS "Whether to build mori CPP tests" ON)
 option(ENABLE_PROFILER "Enable kernel profiling" OFF)
 option(ENABLE_DEBUG_PRINTF "Enable debug printf in device kernels" OFF)
+option(ENABLE_STANDARD_MOE_ADAPT "Enable standard moe adapt" OFF)
 
 set(CMAKE_LIBRARY_PATH "/usr/local/lib;${CMAKE_LIBRARY_PATH}")
 find_library(BNXT_RE_LIB NAMES bnxt_re)
@@ -76,6 +77,11 @@ if(ENABLE_DEBUG_PRINTF)
   add_compile_definitions(MORI_ENABLE_DEBUG_PRINTF)
 endif()
 
+message(STATUS "ENABLE_STANDARD_MOE_ADAPT = ${ENABLE_STANDARD_MOE_ADAPT}")
+if(ENABLE_STANDARD_MOE_ADAPT)
+  add_compile_definitions(ENABLE_STANDARD_MOE_ADAPT)
+endif()
+
 if(NOT DEFINED WARP_ACCUM_UNROLL)
   set(WARP_ACCUM_UNROLL 1)
 endif()

README.md

@@ -4,11 +4,11 @@
 
 **MORI** (**Mo**dular **R**DMA **I**nterface) is a **bottom-up, modular, and composable framework** for building high-performance communication applications with a strong focus on **RDMA + GPU integration**. Inspired by the role of MLIR in compiler infrastructure, MORI provides reusable and extensible building blocks that make it **easier for developers to adopt advanced techniques** such as IBGDA (Infiniband GPUDirect Async) and GDS (GPUDirect Storage).
 
-To help developers get started quickly, MORI also includes a suite of optimized libraries—**MORI-EP** (MoE dispatch & combine kernels), **MORI-IO** (p2p communication for KVCache transfer), and **MORI-CCL** (collective communication)—that deliver out-of-the-box performance.
+To help developers get started quickly, MORI also includes a suite of optimized libraries—**MORI-EP** (MoE dispatch & combine kernels), **MORI-IO** (p2p communication for KVCache transfer), and **MORI-CCL** (collective communication)—that deliver out-of-the-box performance, with support for AMD `Pensando DSC`, Broadcom `Thor2`, and NVIDIA Mellanox `ConnectX-7` NICs.
 
 Feature summary:
 - Applications
-    - MORI-EP: intra and inter-node dispatch/combine kernels with SOTA performance
+    - MORI-EP: intra and inter-node dispatch/combine kernels with SOTA performance.
     - MORI-IO: point-to-point communication library with ultra-low overhead
     - MORI-CCL: lightweight and flexible collective communication library designed for highly customized use cases such as latency-sensitive or resource-constrained environment
 - Framework

examples/dist_rdma_ops/dist_write.cpp

@@ -431,7 +431,7 @@ int GetGpuidByNicName(std::string nic_name) {
       auto* gpuPci = pciSys->Node(gpu->busId);
       auto* nicPci = pciSys->Node(nic->busId);
 #if 0
-      if (!path) {     
+      if (!path) {
         printf("gpu %s nic %s no direct link\n", gpu->busId.String().c_str(),
                nic->busId.String().c_str());
       } else {

examples/ops/dispatch_combine/test_dispatch_combine.cpp

@@ -403,10 +403,10 @@ class EpDispatchCombineTestCase {
       InitializeHandle();
       SystemBarrier();
 
-      handle.LaunchDispatch(runConfig.kernelType, -1, -1, stream);
+      handle.LaunchDispatch(runConfig.kernelType, -1, -1, -1, stream);
       CopyDispatchOutAsCombineInp();
       SystemBarrier();
-      handle.LaunchCombine(runConfig.kernelType, -1, -1, stream);
+      handle.LaunchCombine(runConfig.kernelType, -1, -1, -1, -1, stream);
       if (handle.config.rank == 0) std::cout << "Warmup Done" << std::endl;
     }
 
@@ -420,7 +420,7 @@ class EpDispatchCombineTestCase {
       SystemBarrier();
 
       HIP_RUNTIME_CHECK(hipEventRecord(events[0]));
-      handle.LaunchDispatch(runConfig.kernelType, -1, -1, stream);
+      handle.LaunchDispatch(runConfig.kernelType, -1, -1, -1, stream);
       HIP_RUNTIME_CHECK(hipEventRecord(events[1]));
 
       HIP_RUNTIME_CHECK(hipMemcpy(&dispatchTotalRecvTokenNum, handle.totalRecvTokenNum, sizeof(int),
@@ -429,7 +429,7 @@ class EpDispatchCombineTestCase {
       SystemBarrier();
 
       HIP_RUNTIME_CHECK(hipEventRecord(events[2]));
-      handle.LaunchCombine(runConfig.kernelType, -1, -1, stream);
+      handle.LaunchCombine(runConfig.kernelType, -1, -1, -1, -1, stream);
       HIP_RUNTIME_CHECK(hipEventRecord(events[3]));
 
       float dispatch, combine;