diff --git a/src/ext/ep/CMakeLists.txt b/src/ext/ep/CMakeLists.txt
index 36acf022..ae775d05 100644
--- a/src/ext/ep/CMakeLists.txt
+++ b/src/ext/ep/CMakeLists.txt
@@ -6,7 +6,7 @@
 #
 # Two backends share one module, both with a torch-free, raw-pointer (uintptr_t)
 # nanobind API so the module never links libtorch:
-#  - Low-latency (LL): moe_runtime.cc + kernels/low_latency.cu.
+#  - Low-latency (LL): moe_runtime.cc + low_latency/{dispatch,combine}.cu.
 #  - High-throughput (HT): ht_runtime.cc + ht/kernels/*.cu, a DeepEP-style
 #    runtime de-torched to the same pointer boundary as the LL runtime (dynamic
 #    recv sizing via the two-phase notify -> allocate -> dispatch API).
diff --git a/src/ext/ep/README.md b/src/ext/ep/README.md
index 47d902dc..5e0b5c8e 100644
--- a/src/ext/ep/README.md
+++ b/src/ext/ep/README.md
@@ -4,7 +4,7 @@ A torch-free nanobind extension for MoE `dispatch` / `combine` primitives in
 MSCCL++. The module builds two active backends:
 
 - **Low-latency (LL)**: `MoERuntime` / `MoECommunicator(mode=LOW_LATENCY)`,
-  backed by `kernels/low_latency.cu`.
+  backed by `low_latency/dispatch.cu` and `low_latency/combine.cu`.
 - **High-throughput (HT)**: `ExpertParallelRuntime` /
   `MoECommunicator(mode=HIGH_THROUGHPUT)`, backed by `ht_runtime.cc` and
   `ht/kernels/*`.
@@ -488,11 +488,19 @@ Env knobs:
 
 | Backend | Python API | C++ runtime | Kernel sources | Layout |
 |---------|------------|-------------|----------------|--------|
-| LL | `MoECommunicator(mode=MoEMode.LOW_LATENCY)` / `MoERuntime` | `moe_runtime.cc` | `kernels/low_latency.cu` | `EXPERT_MAJOR` |
+| LL | `MoECommunicator(mode=MoEMode.LOW_LATENCY)` / `MoERuntime` | `moe_runtime.cc` | `low_latency/{dispatch,combine}.cu` | `EXPERT_MAJOR` |
 | HT | `MoECommunicator(mode=MoEMode.HIGH_THROUGHPUT)` / `ExpertParallelRuntime` | `ht_runtime.cc` | `ht/kernels/*` | `FLAT` |
 
-The `ht/` directory is active source code for the HT backend in the current
-build.
+Shared internal headers live in `include/`. The previous LL implementation is
+kept in `legacy/low_latency.cu` for reference and is not compiled.
+
+The LL runtime uses one `low_latency_num_blocks` setting. Its default is 130:
+dispatch launches 128 workers plus scheduler/notify blocks, while combine
+launches 128 workers. `RANK_LOCAL_REDUCE` is the default combine mode;
+`DIRECT_SEND` preserves bit-exact top-k reduction order.
+
+The LL payload layout reserves optional scale storage for future quantization,
+but the active LL kernels currently accept BF16 inputs and expert outputs only.
 
 ### MSCCL++ EP LL vs NCCL EP LL vs DeepEP elastic dispatch payloads
 
diff --git a/src/ext/ep/config.hpp b/src/ext/ep/config.hpp
index b92b6563..3419c4fb 100644
--- a/src/ext/ep/config.hpp
+++ b/src/ext/ep/config.hpp
@@ -143,7 +143,6 @@ struct PayloadView {
 struct Buffer {
   void* dispatchData_;
   void* combineData_;
-  mscclpp::LL8Packet* combineReadyPackets_;
 };
 
 struct Layout {
@@ -154,15 +153,12 @@ struct Layout {
     const PayloadView<__bfloat16> bf16Payload(hidden, numTopk);
     const PayloadView<__nv_fp8_storage_t, float> fp8Payload(hidden, numTopk, 128);
     const size_t dispatchMetadataBytes =
-        configAlign<size_t>(static_cast<size_t>(2 * numRanks + numExperts) * sizeof(uint64_t), 128);
+        configAlign<size_t>(static_cast<size_t>(numRanks + numExperts) * sizeof(uint64_t), 128);
     const size_t dispatchPayloadStride =
         configAlign<size_t>(std::max(bf16Payload.numBytes_, fp8Payload.numBytes_), 128);
     const size_t dispatchBufferBytes =
         dispatchMetadataBytes + static_cast<size_t>(numRanks) * maxTokensPerRank * dispatchPayloadStride;
-    const size_t combineControlBytes =
-        configAlign<size_t>(static_cast<size_t>(numRanks) * sizeof(mscclpp::LL8Packet), 128);
-    const size_t combineBufferBytes =
-        combineControlBytes + static_cast<size_t>(numExperts) * maxTokensPerRank * hidden * sizeof(__bfloat16);
+    const size_t combineBufferBytes = static_cast<size_t>(numExperts) * maxTokensPerRank * hidden * sizeof(__bfloat16);
     const size_t bufferBytes = configAlign<size_t>(std::max(dispatchBufferBytes, combineBufferBytes), 128);
     totalBytes_ = 2 * bufferBytes;
 
@@ -172,8 +168,7 @@ struct Layout {
         auto* bufferBase = base + static_cast<size_t>(bufferIdx) * bufferBytes;
         buffers_[bufferIdx] = {
             .dispatchData_ = bufferBase,
-            .combineData_ = bufferBase + combineControlBytes,
-            .combineReadyPackets_ = reinterpret_cast<mscclpp::LL8Packet*>(bufferBase),
+            .combineData_ = bufferBase,
         };
       }
     }
diff --git a/src/ext/ep/ht_runtime.cc b/src/ext/ep/ht_runtime.cc
index ffbabfc5..9a56d441 100644
--- a/src/ext/ep/ht_runtime.cc
+++ b/src/ext/ep/ht_runtime.cc
@@ -1272,7 +1272,7 @@ void MoEHighThroughputRuntime::intranodeCombine(void* combinedX, float* combined
 // -----------------------------------------------------------------------------
 // Internode (NVLink + RDMA) high-throughput path. Ported from DeepEP
 // `csrc/deep_ep.cpp`; the kernels it drives are in
-// `src/ext/ep/kernels/internode.cu`. Validated end-to-end on 2 x H100 x 8
+// `src/ext/ep/ht/kernels/internode.cu`. Validated end-to-end on 2 x H100 x 8
 // via `test/python/ep/test_internode_multirank.py`. De-torched the same
 // way as the intranode path: tensor params became raw pointers + size ints,
 // output tensors became caller pointers, the EventHandle / async / record_stream
diff --git a/src/ext/ep/include/api.cuh b/src/ext/ep/include/api.cuh
index 3d0d6634..d6792962 100644
--- a/src/ext/ep/include/api.cuh
+++ b/src/ext/ep/include/api.cuh
@@ -222,6 +222,8 @@ struct Workload {
 struct CommContext {
   /// Base address of the locally-registered RDMA buffer.
   void* rdmaBufferBase_;
+  /// Base memory channel handles used only for signal/wait synchronization.
+  mscclpp::BaseMemoryChannelDeviceHandle* baseMemoryChannels_;
   /// Peer-mapped base addresses.
   void* const* peerBases_;
   /// Maximum shared memory available to one block after opt-in.
@@ -236,6 +238,12 @@ struct CommContext {
   int numRanks_;
 };
 
+/// Return the optimized low-latency workspace size.
+/// @param[in] numRanks Total number of ranks.
+/// @param[in] numExperts Total number of experts.
+/// @return Required workspace bytes.
+size_t workspaceSize(int numRanks, int numExperts);
+
 /// Low-latency dispatch that distributes tokens to experts across ranks.
 /// @param[out] output Expert-major packed output
 /// [num_local_experts, num_ranks * max_tokens_per_rank, hidden].
@@ -264,7 +272,6 @@ void dispatch(void* output, int* outputSrcInfo, int64_t* outputLayout, int* outp
 /// @param[in] layoutRange Per-[local expert, source rank] packed count and offset.
 /// @param[in] workload Per-call workload dimensions.
 /// @param[in,out] recvBuffer Current symmetric ping-pong buffer receiving partials or expert rows.
-/// @param[in,out] readyPackets One readiness packet per sender rank in the current symmetric buffer.
 /// @param[in] dispatchRecvBuffer Previous dispatch buffer containing rewritten routing metadata.
 /// @param[in] comm Persistent communication context.
 /// @param[in,out] workspace Persistent dispatch metadata plus the combine device barrier.
@@ -272,9 +279,8 @@ void dispatch(void* output, int* outputSrcInfo, int64_t* outputLayout, int* outp
 /// @param[in] mode Combine algorithm.
 /// @param[in] stream CUDA stream.
 void combine(void* output, const void* input, const int64_t* topkIdx, const float* topkWeights, const int* srcInfo,
-             const int64_t* layoutRange, const Workload& workload, void* recvBuffer, mscclpp::LL8Packet* readyPackets,
-             void* dispatchRecvBuffer, const CommContext& comm, void* workspace, int numBlocks, CombineMode mode,
-             cudaStream_t stream);
+             const int64_t* layoutRange, const Workload& workload, void* recvBuffer, void* dispatchRecvBuffer,
+             const CommContext& comm, void* workspace, int numBlocks, CombineMode mode, cudaStream_t stream);
 
 }  // namespace low_latency
 
diff --git a/src/ext/ep/include/device_helpers.cuh b/src/ext/ep/include/device_helpers.cuh
index c026024e..4784a3ff 100644
--- a/src/ext/ep/include/device_helpers.cuh
+++ b/src/ext/ep/include/device_helpers.cuh
@@ -66,17 +66,6 @@ __device__ __forceinline__ void memory_fence_gpu() { asm volatile("fence.acq_rel
 
 __device__ __forceinline__ void memory_fence_cta() { asm volatile("fence.acq_rel.cta;" ::: "memory"); }
 
-MSCCLPP_DEVICE_INLINE void publishLl8Packet(mscclpp::LL8Packet *packet, uint32_t value, uint32_t flag) {
-  memory_fence();
-  packet->write(value, flag);
-}
-
-MSCCLPP_DEVICE_INLINE uint32_t waitLl8Packet(const mscclpp::LL8Packet *packet, uint32_t flag) {
-  const uint32_t value = packet->read(flag, -1);
-  memory_fence();
-  return value;
-}
-
 __device__ __forceinline__ void *peerBufferPtr(void *localBuffer, void *localBufferBase, void *peerBufferBase) {
   if (localBufferBase == nullptr) return peerBufferBase;
   const auto offset = reinterpret_cast<uint8_t *>(localBuffer) - reinterpret_cast<uint8_t *>(localBufferBase);
diff --git a/src/ext/ep/low_latency/combine.cu b/src/ext/ep/low_latency/combine.cu
index 383f805c..dc186f27 100644
--- a/src/ext/ep/low_latency/combine.cu
+++ b/src/ext/ep/low_latency/combine.cu
@@ -10,7 +10,8 @@
 
 namespace mscclpp {
 namespace ep {
-namespace low_latency_opt {
+namespace low_latency {
+namespace detail {
 
 constexpr int CombineNWarps = 32;
 constexpr int CombineNThreads = CombineNWarps * WARP_SIZE;
@@ -42,8 +43,8 @@ MSCCLPP_DEVICE_INLINE int4 reduceWeightedBf16x8(const void* expertOutput, int ro
     const int sourceRowOffset = warpBroadcast(rowOffset, topkLane);
     if (sourceRowOffset < 0) continue;
     const float sourceWeight = warpBroadcast(weight, topkLane);
-    const int4 packed = ld_nc_global(reinterpret_cast<const int4*>(expertOutput) +
-                                     static_cast<size_t>(sourceRowOffset) * HiddenInt4 + hiddenIdx);
+    const int4 packed =
+        reinterpret_cast<const int4*>(expertOutput)[static_cast<size_t>(sourceRowOffset) * HiddenInt4 + hiddenIdx];
     const auto* values = reinterpret_cast<const mscclpp::bf16x2*>(&packed);
 #pragma unroll
     for (int pairIdx = 0; pairIdx < Bf16PairsPerInt4; ++pairIdx) {
@@ -70,9 +71,8 @@ MSCCLPP_DEVICE_INLINE int4 reduceRankPartialsBf16x8(const void* combineRecvBuffe
   for (int topkLane = 0; topkLane < nTopk; ++topkLane) {
     const int partialRank = warpBroadcast(partialRankCandidate, topkLane);
     if (partialRank < 0) continue;
-    const int4 packed =
-        ld_nc_global(reinterpret_cast<const int4*>(combineRecvBuffer) +
-                     (static_cast<size_t>(partialRank) * maxTokensPerRank + tokenIdx) * HiddenInt4 + hiddenIdx);
+    const int4 packed = reinterpret_cast<const int4*>(
+        combineRecvBuffer)[(static_cast<size_t>(partialRank) * maxTokensPerRank + tokenIdx) * HiddenInt4 + hiddenIdx];
     const auto* values = reinterpret_cast<const mscclpp::bf16x2*>(&packed);
 #pragma unroll
     for (int pairIdx = 0; pairIdx < Bf16PairsPerInt4; ++pairIdx) {
@@ -128,8 +128,8 @@ MSCCLPP_DEVICE_INLINE void sendRankReducedPartials(const void* expertOutput, int
       const auto* sourcePayload =
           reinterpret_cast<const uint8_t*>(dispatchRecvBuffer) + dispatchMetadataSize +
           (static_cast<size_t>(sourceRank) * maxTokensPerRank + sourceTokenSlot) * payloadStride;
-      const int rowOffset = laneId < nTopk ? ld_nc_global(payloadView.topKIndices(sourcePayload) + laneId) : -1;
-      const float weight = laneId < nTopk ? ld_nc_global(payloadView.topKValues(sourcePayload) + laneId) : 0.0f;
+      const int rowOffset = laneId < nTopk ? payloadView.topKIndices(sourcePayload)[laneId] : -1;
+      const float weight = laneId < nTopk ? payloadView.topKValues(sourcePayload)[laneId] : 0.0f;
       if (rowOffset >= 0) EP_DEVICE_ASSERT(rowOffset < nExpertOutputRows);
 
       int4 reduced[ChunksPerThread] = {};
@@ -154,8 +154,7 @@ MSCCLPP_DEVICE_INLINE void sendRankReducedPartials(const void* expertOutput, int
 
       if (threadId == 0) {
         fenceProxyAsyncSharedCta();
-        const int sourceTokenIdx =
-            ld_nc_global(payloadView.srcTokenGlobalIdx(sourcePayload)) - sourceRank * maxTokensPerRank;
+        const int sourceTokenIdx = *payloadView.srcTokenGlobalIdx(sourcePayload) - sourceRank * maxTokensPerRank;
         EP_DEVICE_ASSERT(sourceTokenIdx >= 0 && sourceTokenIdx < maxTokensPerRank);
         void* destinationBuffer = sourceRank == rank
                                       ? combineRecvBuffer
@@ -222,7 +221,7 @@ MSCCLPP_DEVICE_INLINE void sendExpertRowsDirect(const void* expertOutput, const
       }
       EP_DEVICE_ASSERT(sourceRank < nRanks);
       const int inputRowOffset = localExpertIdx * nOutputSlotsPerExpert + expertTokenIdx;
-      const int sourceTokenIdx = ld_nc_global(srcInfo + inputRowOffset);
+      const int sourceTokenIdx = srcInfo[inputRowOffset];
       EP_DEVICE_ASSERT(sourceTokenIdx >= 0 && sourceTokenIdx < maxTokensPerRank);
       const auto* inputRow =
           reinterpret_cast<const uint8_t*>(expertOutput) + static_cast<size_t>(inputRowOffset) * HiddenBytes;
@@ -243,21 +242,24 @@ MSCCLPP_DEVICE_INLINE void sendExpertRowsDirect(const void* expertOutput, const
   }
 }
 
-MSCCLPP_DEVICE_INLINE void combineSynchronize(mscclpp::LL8Packet* readyPackets, void* rdmaBufferBase,
-                                              void* const* peerRecvBuffers, int rank, int nRanks, uint32_t readyFlag) {
+MSCCLPP_DEVICE_INLINE void combineSynchronize(mscclpp::BaseMemoryChannelDeviceHandle* baseMemoryChannels,
+                                              mscclpp::DeviceSemaphore* localReady, int rank, int nRanks) {
   const int threadId = static_cast<int>(threadIdx.x);
   if (blockIdx.x == 0 && threadId < nRanks) {
     const int peerRank = threadId;
-    void* destination =
-        peerRank == rank ? readyPackets : peerBufferPtr(readyPackets, rdmaBufferBase, peerRecvBuffers[peerRank]);
-    publishLl8Packet(reinterpret_cast<mscclpp::LL8Packet*>(destination) + rank, 1, readyFlag);
-    waitLl8Packet(readyPackets + peerRank, readyFlag);
+    if (peerRank == rank) {
+      localReady->release();
+      localReady->acquire();
+    } else {
+      baseMemoryChannels[peerRank].signal();
+      baseMemoryChannels[peerRank].wait(-1);
+    }
   }
 }
 
 template <int Hidden>
-MSCCLPP_DEVICE_INLINE void recvRankLocalPartials(void* output, const int64_t* topkIndices, int nTokens, int nTopk,
-                                                 int nExperts, int nRanks, int maxTokensPerRank,
+MSCCLPP_DEVICE_INLINE void recvRankLocalPartials(void* output, const int64_t* __restrict__ topkIndices, int nTokens,
+                                                 int nTopk, int nExperts, int nRanks, int maxTokensPerRank,
                                                  const void* combineRecvBuffer, uint8_t* sharedMemory) {
   const int threadId = static_cast<int>(threadIdx.x);
   const int laneId = get_lane_id();
@@ -273,7 +275,7 @@ MSCCLPP_DEVICE_INLINE void recvRankLocalPartials(void* output, const int64_t* to
        tokenIdx += static_cast<int>(gridDim.x), ++tokenIteration) {
     const int stage = tokenIteration % CombineNStages;
     auto* outputTile = reinterpret_cast<int4*>(outputTiles + static_cast<size_t>(stage) * HiddenBytes);
-    const int globalExpertIdx = laneId < nTopk ? static_cast<int>(__ldg(topkIndices + tokenIdx * nTopk + laneId)) : -1;
+    const int globalExpertIdx = laneId < nTopk ? static_cast<int>(topkIndices[tokenIdx * nTopk + laneId]) : -1;
     const int destinationRank = globalExpertIdx >= 0 ? globalExpertIdx / nLocalExperts : -1;
     const bool firstLaneForRank = isFirstLaneForRank(destinationRank, laneId);
     const int partialRank = destinationRank >= 0 && firstLaneForRank ? destinationRank : -1;
@@ -308,9 +310,9 @@ MSCCLPP_DEVICE_INLINE void recvRankLocalPartials(void* output, const int64_t* to
 }
 
 template <int Hidden>
-MSCCLPP_DEVICE_INLINE void recvExpertRowsDirect(void* output, const int64_t* topkIndices, const float* topkWeights,
-                                                int nTokens, int nTopk, int maxTokensPerRank,
-                                                const void* combineRecvBuffer) {
+MSCCLPP_DEVICE_INLINE void recvExpertRowsDirect(void* output, const int64_t* __restrict__ topkIndices,
+                                                const float* __restrict__ topkWeights, int nTokens, int nTopk,
+                                                int maxTokensPerRank, const void* combineRecvBuffer) {
   constexpr int Bf16PerInt4 = sizeof(int4) / sizeof(__bfloat16);
   constexpr int HiddenInt4 = Hidden / Bf16PerInt4;
   const int threadId = static_cast<int>(threadIdx.x);
@@ -319,8 +321,8 @@ MSCCLPP_DEVICE_INLINE void recvExpertRowsDirect(void* output, const int64_t* top
     int regTopkIndices[CombineMaxNTopk];
     float regTopkWeights[CombineMaxNTopk];
     for (int topkIdx = 0; topkIdx < nTopk; ++topkIdx) {
-      regTopkIndices[topkIdx] = static_cast<int>(__ldg(topkIndices + tokenIdx * nTopk + topkIdx));
-      regTopkWeights[topkIdx] = topkWeights == nullptr ? 1.0f : __ldg(topkWeights + tokenIdx * nTopk + topkIdx);
+      regTopkIndices[topkIdx] = static_cast<int>(topkIndices[tokenIdx * nTopk + topkIdx]);
+      regTopkWeights[topkIdx] = topkWeights == nullptr ? 1.0f : topkWeights[tokenIdx * nTopk + topkIdx];
     }
 
 #pragma unroll
@@ -331,7 +333,7 @@ MSCCLPP_DEVICE_INLINE void recvExpertRowsDirect(void* output, const int64_t* top
         if (expertIdx < 0) continue;
         const auto* expertRow = reinterpret_cast<const int4*>(combineRecvBuffer) +
                                 (static_cast<size_t>(expertIdx) * maxTokensPerRank + tokenIdx) * HiddenInt4;
-        const int4 packed = ld_nc_global(expertRow + hiddenIdx);
+        const int4 packed = expertRow[hiddenIdx];
         const auto* values = reinterpret_cast<const __bfloat16*>(&packed);
 #pragma unroll
         for (int elemIdx = 0; elemIdx < Bf16PerInt4; ++elemIdx) {
@@ -353,13 +355,13 @@ MSCCLPP_DEVICE_INLINE void recvExpertRowsDirect(void* output, const int64_t* top
 
 template <low_latency::CombineMode Mode, int Hidden>
 __global__ __launch_bounds__(CombineNThreads, 1) void combineKernel(
-    void* output, const void* expertOutput, const int64_t* topkIndices, const float* topkWeights, const int* srcInfo,
-    const int64_t* layoutRange, int nTokens, int nExperts, int rank, int nRanks, int nTopk, int maxTokensPerRank,
-    void* combineRecvBuffer, const void* dispatchRecvBuffer, void* rdmaBufferBase, void* const* peerRecvBuffers,
-    mscclpp::LL8Packet* readyPackets, void* workspace) {
+    void* output, const void* expertOutput, const int64_t* __restrict__ topkIndices,
+    const float* __restrict__ topkWeights, const int* srcInfo, const int64_t* layoutRange, int nTokens, int nExperts,
+    int rank, int nRanks, int nTopk, int maxTokensPerRank, void* combineRecvBuffer, const void* dispatchRecvBuffer,
+    void* rdmaBufferBase, void* const* peerRecvBuffers, mscclpp::BaseMemoryChannelDeviceHandle* baseMemoryChannels,
+    void* workspace) {
   extern __shared__ __align__(128) uint8_t sharedMemory[];
   WorkspaceView workspaceView(workspace, nRanks, nExperts);
-  const uint32_t readyFlag = *workspaceView.metadataEpoch_;
 
   if constexpr (Mode == low_latency::CombineMode::RANK_LOCAL_REDUCE) {
     sendRankReducedPartials<Hidden>(expertOutput, nExperts, rank, nRanks, nTopk, maxTokensPerRank, combineRecvBuffer,
@@ -370,7 +372,7 @@ __global__ __launch_bounds__(CombineNThreads, 1) void combineKernel(
   }
 
   workspaceView.combineSyncer_->sync(gridDim.x);
-  combineSynchronize(readyPackets, rdmaBufferBase, peerRecvBuffers, rank, nRanks, readyFlag);
+  combineSynchronize(baseMemoryChannels, workspaceView.localPayloadReady_, rank, nRanks);
   workspaceView.combineSyncer_->sync(gridDim.x);
 
   if constexpr (Mode == low_latency::CombineMode::RANK_LOCAL_REDUCE) {
@@ -384,9 +386,9 @@ __global__ __launch_bounds__(CombineNThreads, 1) void combineKernel(
 template <low_latency::CombineMode Mode, int Hidden>
 inline void combineHiddenMode(void* output, const void* expertOutput, const int64_t* topkIndices,
                               const float* topkWeights, const int* srcInfo, const int64_t* layoutRange,
-                              const low_latency::Workload& workload, void* recvBuffer, mscclpp::LL8Packet* readyPackets,
-                              void* dispatchRecvBuffer, const low_latency::CommContext& comm, void* workspace,
-                              int numBlocks, cudaStream_t stream) {
+                              const low_latency::Workload& workload, void* recvBuffer, void* dispatchRecvBuffer,
+                              const low_latency::CommContext& comm, void* workspace, int numBlocks,
+                              cudaStream_t stream) {
   static_assert(Hidden == 4096 || Hidden == 7168 || Hidden == 8192 || Hidden == 9216);
   if constexpr (Mode == low_latency::CombineMode::DIRECT_SEND) {
     static_assert(tmaWorkerCount<Hidden, DirectSendMaxNWorkers>() > 0);
@@ -407,31 +409,30 @@ inline void combineHiddenMode(void* output, const void* expertOutput, const int6
 
   combineKernel<Mode, Hidden><<<dim3(numBlocks), dim3(CombineNThreads), sharedBytes, stream>>>(
       output, expertOutput, topkIndices, topkWeights, srcInfo, layoutRange, nTokens, nExperts, rank, nRanks, nTopk,
-      maxTokensPerRank, recvBuffer, dispatchRecvBuffer, comm.rdmaBufferBase_, comm.peerBases_, readyPackets, workspace);
+      maxTokensPerRank, recvBuffer, dispatchRecvBuffer, comm.rdmaBufferBase_, comm.peerBases_, comm.baseMemoryChannels_,
+      workspace);
   CUDA_CHECK(cudaGetLastError());
 }
 
 template <int Hidden>
 inline void combineHidden(void* output, const void* expertOutput, const int64_t* topkIndices, const float* topkWeights,
                           const int* srcInfo, const int64_t* layoutRange, const low_latency::Workload& workload,
-                          void* recvBuffer, mscclpp::LL8Packet* readyPackets, void* dispatchRecvBuffer,
-                          const low_latency::CommContext& comm, void* workspace, int numBlocks,
-                          low_latency::CombineMode mode, cudaStream_t stream) {
+                          void* recvBuffer, void* dispatchRecvBuffer, const low_latency::CommContext& comm,
+                          void* workspace, int numBlocks, low_latency::CombineMode mode, cudaStream_t stream) {
   if (mode == low_latency::CombineMode::RANK_LOCAL_REDUCE) {
     return combineHiddenMode<low_latency::CombineMode::RANK_LOCAL_REDUCE, Hidden>(
-        output, expertOutput, topkIndices, topkWeights, srcInfo, layoutRange, workload, recvBuffer, readyPackets,
-        dispatchRecvBuffer, comm, workspace, numBlocks, stream);
+        output, expertOutput, topkIndices, topkWeights, srcInfo, layoutRange, workload, recvBuffer, dispatchRecvBuffer,
+        comm, workspace, numBlocks, stream);
   }
   return combineHiddenMode<low_latency::CombineMode::DIRECT_SEND, Hidden>(
-      output, expertOutput, topkIndices, topkWeights, srcInfo, layoutRange, workload, recvBuffer, readyPackets,
-      dispatchRecvBuffer, comm, workspace, numBlocks, stream);
+      output, expertOutput, topkIndices, topkWeights, srcInfo, layoutRange, workload, recvBuffer, dispatchRecvBuffer,
+      comm, workspace, numBlocks, stream);
 }
 
 inline void combine(void* output, const void* expertOutput, const int64_t* topkIndices, const float* topkWeights,
                     const int* srcInfo, const int64_t* layoutRange, const low_latency::Workload& workload,
-                    void* recvBuffer, mscclpp::LL8Packet* readyPackets, void* dispatchRecvBuffer,
-                    const low_latency::CommContext& comm, void* workspace, int numBlocks, low_latency::CombineMode mode,
-                    cudaStream_t stream) {
+                    void* recvBuffer, void* dispatchRecvBuffer, const low_latency::CommContext& comm, void* workspace,
+                    int numBlocks, low_latency::CombineMode mode, cudaStream_t stream) {
   const int nExperts = workload.numExperts_;
   const int rank = comm.rank_;
   const int nRanks = comm.numRanks_;
@@ -440,10 +441,10 @@ inline void combine(void* output, const void* expertOutput, const int64_t* topkI
   EP_HOST_ASSERT(expertOutput != nullptr);
   EP_HOST_ASSERT(topkIndices != nullptr);
   EP_HOST_ASSERT(recvBuffer != nullptr);
-  EP_HOST_ASSERT(readyPackets != nullptr);
   EP_HOST_ASSERT(dispatchRecvBuffer != nullptr);
   EP_HOST_ASSERT(comm.rdmaBufferBase_ != nullptr);
   EP_HOST_ASSERT(comm.peerBases_ != nullptr);
+  EP_HOST_ASSERT(comm.baseMemoryChannels_ != nullptr);
   EP_HOST_ASSERT(workspace != nullptr);
   EP_HOST_ASSERT(nRanks > 0 && nRanks <= 2 * WARP_SIZE);
   EP_HOST_ASSERT(nExperts > 0 && nExperts % nRanks == 0);
@@ -460,35 +461,28 @@ inline void combine(void* output, const void* expertOutput, const int64_t* topkI
   switch (workload.hidden_) {
     case 4096:
       return combineHidden<4096>(output, expertOutput, topkIndices, topkWeights, srcInfo, layoutRange, workload,
-                                 recvBuffer, readyPackets, dispatchRecvBuffer, comm, workspace, numBlocks, mode,
-                                 stream);
+                                 recvBuffer, dispatchRecvBuffer, comm, workspace, numBlocks, mode, stream);
     case 7168:
       return combineHidden<7168>(output, expertOutput, topkIndices, topkWeights, srcInfo, layoutRange, workload,
-                                 recvBuffer, readyPackets, dispatchRecvBuffer, comm, workspace, numBlocks, mode,
-                                 stream);
+                                 recvBuffer, dispatchRecvBuffer, comm, workspace, numBlocks, mode, stream);
     case 8192:
       return combineHidden<8192>(output, expertOutput, topkIndices, topkWeights, srcInfo, layoutRange, workload,
-                                 recvBuffer, readyPackets, dispatchRecvBuffer, comm, workspace, numBlocks, mode,
-                                 stream);
+                                 recvBuffer, dispatchRecvBuffer, comm, workspace, numBlocks, mode, stream);
     case 9216:
       return combineHidden<9216>(output, expertOutput, topkIndices, topkWeights, srcInfo, layoutRange, workload,
-                                 recvBuffer, readyPackets, dispatchRecvBuffer, comm, workspace, numBlocks, mode,
-                                 stream);
+                                 recvBuffer, dispatchRecvBuffer, comm, workspace, numBlocks, mode, stream);
     default:
       EP_HOST_ASSERT(false && "unsupported optimized low-latency hidden size");
   }
 }
 
-}  // namespace low_latency_opt
-
-namespace low_latency {
+}  // namespace detail
 
 void combine(void* output, const void* input, const int64_t* topkIdx, const float* topkWeights, const int* srcInfo,
-             const int64_t* layoutRange, const Workload& workload, void* recvBuffer, mscclpp::LL8Packet* readyPackets,
-             void* dispatchRecvBuffer, const CommContext& comm, void* workspace, int numBlocks, CombineMode mode,
-             cudaStream_t stream) {
-  low_latency_opt::combine(output, input, topkIdx, topkWeights, srcInfo, layoutRange, workload, recvBuffer,
-                           readyPackets, dispatchRecvBuffer, comm, workspace, numBlocks, mode, stream);
+             const int64_t* layoutRange, const Workload& workload, void* recvBuffer, void* dispatchRecvBuffer,
+             const CommContext& comm, void* workspace, int numBlocks, CombineMode mode, cudaStream_t stream) {
+  detail::combine(output, input, topkIdx, topkWeights, srcInfo, layoutRange, workload, recvBuffer, dispatchRecvBuffer,
+                  comm, workspace, numBlocks, mode, stream);
 }
 
 }  // namespace low_latency
diff --git a/src/ext/ep/low_latency/config.cuh b/src/ext/ep/low_latency/config.cuh
index 5386e8b0..f6f71234 100644
--- a/src/ext/ep/low_latency/config.cuh
+++ b/src/ext/ep/low_latency/config.cuh
@@ -12,7 +12,8 @@
 
 namespace mscclpp {
 namespace ep {
-namespace low_latency_opt {
+namespace low_latency {
+namespace detail {
 
 constexpr int DispatchNWarps = 16;
 constexpr int DispatchMinNWarpsPerGroup = 8;
@@ -22,15 +23,13 @@ constexpr int DispatchMaxNRecvTmaWorkers = DispatchNWarps;
 constexpr size_t OptimizedDynamicSharedMemoryBytes = 226 * 1024;
 constexpr size_t TmaWorkerControlBytes = DispatchMaxNWarpGroups * WARP_SIZE * sizeof(int);
 static_assert(DispatchNWarps % DispatchMinNWarpsPerGroup == 0);
+static_assert(sizeof(mscclpp::DeviceSemaphore) == sizeof(int));
+static_assert(alignof(mscclpp::DeviceSemaphore) <= alignof(int));
 static_assert(sizeof(mscclpp::DeviceSyncer) % sizeof(int) == 0);
 static_assert(alignof(mscclpp::DeviceSyncer) <= alignof(int));
 
 MSCCLPP_HOST_DEVICE_INLINE size_t dispatchMetadataBytes(int nRanks, int nExperts) {
-  return configAlign<size_t>(static_cast<size_t>(2 * nRanks + nExperts) * sizeof(mscclpp::LL8Packet), 128);
-}
-
-MSCCLPP_DEVICE_INLINE mscclpp::LL8Packet* dispatchReadyPackets(void* recvBuffer, int nRanks, int nExperts) {
-  return reinterpret_cast<mscclpp::LL8Packet*>(recvBuffer) + nRanks + nExperts;
+  return configAlign<size_t>(static_cast<size_t>(nRanks + nExperts) * sizeof(mscclpp::LL8Packet), 128);
 }
 
 MSCCLPP_HOST_DEVICE_INLINE size_t dispatchPayloadStride(int hidden, int nTopk) {
@@ -47,11 +46,14 @@ struct RecvTask {
   int tokenBegin_;
   int tokenEnd_;
 };
+static_assert(sizeof(RecvTask) % sizeof(int) == 0);
+static_assert(alignof(RecvTask) <= alignof(int));
 
 struct WorkspaceView {
   uint32_t* metadataEpoch_;
   int* rankPayloadSlots_;
   int* rankPayloadCompletions_;
+  mscclpp::DeviceSemaphore* localPayloadReady_;
   int* recvExpertCopiedCounts_;
   uint32_t* rankReadyEpochs_;
   RecvTask* recvTasks_;
@@ -66,20 +68,28 @@ struct WorkspaceView {
     cursor += nRanks;
     rankPayloadCompletions_ = cursor;
     cursor += nRanks;
+    localPayloadReady_ = reinterpret_cast<mscclpp::DeviceSemaphore*>(cursor++);
     recvExpertCopiedCounts_ = cursor;
     cursor += nExperts;
     rankReadyEpochs_ = reinterpret_cast<uint32_t*>(cursor);
     cursor += nRanks;
     recvTasks_ = reinterpret_cast<RecvTask*>(cursor);
-    cursor += 3 * low_latency::MaxWorkerBlocks;
+    cursor += static_cast<size_t>(MaxWorkerBlocks) * sizeof(RecvTask) / sizeof(int);
     tasksAssignedEpoch_ = reinterpret_cast<uint32_t*>(cursor++);
     nRecvTasks_ = cursor++;
     combineSyncer_ = reinterpret_cast<mscclpp::DeviceSyncer*>(cursor);
   }
 
   MSCCLPP_HOST_DEVICE_INLINE static size_t numBytes(int nRanks, int nExperts) {
-    return static_cast<size_t>(3 * nRanks + nExperts + 3 * low_latency::MaxWorkerBlocks + 3) * sizeof(int) +
-           sizeof(mscclpp::DeviceSyncer);
+    return sizeof(uint32_t) +                                                            // metadataEpoch_
+           static_cast<size_t>(nRanks) * sizeof(int) +                                   // rankPayloadSlots_
+           static_cast<size_t>(nRanks) * sizeof(int) +                                   // rankPayloadCompletions_
+           sizeof(mscclpp::DeviceSemaphore) +                                            // localPayloadReady_
+           static_cast<size_t>(nExperts) * sizeof(int) +                                 // recvExpertCopiedCounts_
+           static_cast<size_t>(nRanks) * sizeof(uint32_t) +                              // rankReadyEpochs_
+           static_cast<size_t>(MaxWorkerBlocks) * sizeof(RecvTask) + sizeof(uint32_t) +  // tasksAssignedEpoch_
+           sizeof(int) +                                                                 // nRecvTasks_
+           sizeof(mscclpp::DeviceSyncer);                                                // combineSyncer_
   }
 };
 
@@ -90,7 +100,7 @@ struct KernelConfigCache {
 };
 
 template <typename Kernel>
-inline int configureKernel(Kernel kernel, int nThreads, size_t dynamicSharedBytes, const low_latency::CommContext& comm,
+inline int configureKernel(Kernel kernel, int nThreads, size_t dynamicSharedBytes, const CommContext& comm,
                            KernelConfigCache& cache) {
   if (cache.deviceId_ != comm.deviceId_ || cache.dynamicSharedBytes_ < dynamicSharedBytes) {
     cudaFuncAttributes attributes;
@@ -148,6 +158,7 @@ MSCCLPP_HOST_DEVICE_INLINE size_t dispatchSharedBytes(int nRanks, int nExperts,
   return tmaBytes > metadataBytes ? tmaBytes : metadataBytes;
 }
 
-}  // namespace low_latency_opt
+}  // namespace detail
+}  // namespace low_latency
 }  // namespace ep
 }  // namespace mscclpp
diff --git a/src/ext/ep/low_latency/dispatch.cu b/src/ext/ep/low_latency/dispatch.cu
index d744cccf..51401010 100644
--- a/src/ext/ep/low_latency/dispatch.cu
+++ b/src/ext/ep/low_latency/dispatch.cu
@@ -9,17 +9,20 @@
 
 namespace mscclpp {
 namespace ep {
-namespace low_latency_opt {
+namespace low_latency {
+namespace detail {
 
 template <int Hidden>
-MSCCLPP_DEVICE_INLINE void dispatchSend(const void* inputTokens, int nExperts, int rank, int nRanks,
-                                        const int64_t* topkIndices, const float* topkWeights, int nTokens, int nTopk,
+MSCCLPP_DEVICE_INLINE void dispatchSend(const void* inputTokens,
+                                        mscclpp::BaseMemoryChannelDeviceHandle* baseMemoryChannels, int nExperts,
+                                        int rank, int nRanks, const int64_t* __restrict__ topkIndices,
+                                        const float* __restrict__ topkWeights, int nTokens, int nTopk,
                                         int maxTokensPerRank, void* recvBuffer, void* const* peerRecvBuffers,
                                         void* rdmaBufferBase, void* workspace, uint32_t metadataFlag, int* sharedMem) {
   const int threadId = static_cast<int>(threadIdx.x);
   const int warpId = threadId / WARP_SIZE;
   const int laneId = get_lane_id();
-  const int nWorkerBlocks = static_cast<int>(gridDim.x) - low_latency::DispatchControlBlocks;
+  const int nWorkerBlocks = static_cast<int>(gridDim.x) - DispatchControlBlocks;
   const int notifyBlockIdx = nWorkerBlocks + 1;
   const int nLocalExperts = nExperts / nRanks;
   const size_t metadataBytes = dispatchMetadataBytes(nRanks, nExperts);
@@ -55,8 +58,7 @@ MSCCLPP_DEVICE_INLINE void dispatchSend(const void* inputTokens, int nExperts, i
         if (laneId == 0) {
           issueTmaLoad(inputData, stagedPayload, tmaBarrier, static_cast<uint32_t>(HiddenBytes));
         }
-        const int routedExpertIdx =
-            laneId < nTopk ? static_cast<int>(__ldg(topkIndices + tokenIdx * nTopk + laneId)) : -1;
+        const int routedExpertIdx = laneId < nTopk ? static_cast<int>(topkIndices[tokenIdx * nTopk + laneId]) : -1;
         const int dstRank = routedExpertIdx >= 0 ? routedExpertIdx / nLocalExperts : -1;
         const bool firstLaneForRank = isFirstLaneForRank(dstRank, laneId);
         const bool shouldSend = dstRank >= 0 && firstLaneForRank;
@@ -69,7 +71,7 @@ MSCCLPP_DEVICE_INLINE void dispatchSend(const void* inputTokens, int nExperts, i
           destinationSlots[laneId] = destinationSlot;
           payloadView.topKIndices(stagedPayload)[laneId] = routedExpertIdx;
           payloadView.topKValues(stagedPayload)[laneId] =
-              topkWeights == nullptr ? 1.0f : __ldg(topkWeights + tokenIdx * nTopk + laneId);
+              topkWeights == nullptr ? 1.0f : topkWeights[tokenIdx * nTopk + laneId];
         }
         if (laneId == 0) {
           *payloadView.srcTokenGlobalIdx(stagedPayload) = rank * maxTokensPerRank + tokenIdx;
@@ -108,8 +110,7 @@ MSCCLPP_DEVICE_INLINE void dispatchSend(const void* inputTokens, int nExperts, i
     }
     __syncthreads();
     for (int tokenIdx = warpId; tokenIdx < nTokens; tokenIdx += DispatchNWarps) {
-      const int routedExpertIdx =
-          laneId < nTopk ? static_cast<int>(__ldg(topkIndices + tokenIdx * nTopk + laneId)) : -1;
+      const int routedExpertIdx = laneId < nTopk ? static_cast<int>(topkIndices[tokenIdx * nTopk + laneId]) : -1;
       const int dstRank = routedExpertIdx >= 0 ? routedExpertIdx / nLocalExperts : -1;
       if (routedExpertIdx >= 0) {
         atomicAdd_block(sharedExpertTokenCounts + routedExpertIdx, 1);
@@ -143,10 +144,12 @@ MSCCLPP_DEVICE_INLINE void dispatchSend(const void* inputTokens, int nExperts, i
       }
       workspaceView.rankPayloadSlots_[dstRank] = 0;
       workspaceView.rankPayloadCompletions_[dstRank] = 0;
-      void* destinationBuffer =
-          dstRank == rank ? recvBuffer : peerBufferPtr(recvBuffer, rdmaBufferBase, peerRecvBuffers[dstRank]);
-      auto* readyPackets = dispatchReadyPackets(destinationBuffer, nRanks, nExperts);
-      publishLl8Packet(readyPackets + rank, 1, metadataFlag);
+      if (expectedPayloadCount == 0) continue;
+      if (dstRank == rank) {
+        workspaceView.localPayloadReady_->release();
+      } else {
+        baseMemoryChannels[dstRank].signal();
+      }
     }
     __syncthreads();
   }
@@ -156,9 +159,10 @@ MSCCLPP_DEVICE_INLINE int proportionalTaskBoundary(int nTokens, int nTasks, int
   return nTotalTokens == 0 ? 0 : static_cast<int>(static_cast<int64_t>(nTokens) * nTasks / nTotalTokens);
 }
 
-MSCCLPP_DEVICE_INLINE void dispatchRecvScheduler(int64_t* outputLayout, int* outputCount, int nExperts, int rank,
-                                                 int nRanks, void* recvBuffer, void* workspace, uint32_t metadataFlag,
-                                                 int* sharedMem) {
+MSCCLPP_DEVICE_INLINE void dispatchRecvScheduler(int64_t* outputLayout, int* outputCount,
+                                                 mscclpp::BaseMemoryChannelDeviceHandle* baseMemoryChannels,
+                                                 int nExperts, int rank, int nRanks, void* recvBuffer, void* workspace,
+                                                 uint32_t metadataFlag, int* sharedMem) {
   const int threadId = static_cast<int>(threadIdx.x);
   const int warpId = threadId / WARP_SIZE;
   const int laneId = get_lane_id();
@@ -238,7 +242,11 @@ MSCCLPP_DEVICE_INLINE void dispatchRecvScheduler(int64_t* outputLayout, int* out
     }
 
     if (sourceRank < nRanks && nRankTokens > 0) {
-      waitLl8Packet(dispatchReadyPackets(recvBuffer, nRanks, nExperts) + sourceRank, metadataFlag);
+      if (sourceRank == rank) {
+        workspaceView.localPayloadReady_->acquire();
+      } else {
+        baseMemoryChannels[sourceRank].wait(-1);
+      }
       mscclpp::atomicStore<uint32_t, mscclpp::scopeDevice>(workspaceView.rankReadyEpochs_ + sourceRank, metadataFlag,
                                                            mscclpp::memoryOrderRelease);
     }
@@ -331,13 +339,12 @@ MSCCLPP_DEVICE_INLINE void dispatchRecvWorker(void* output, int* outputSrcInfo,
     }
     __syncwarp();
 
-    const int routedExpertIdx = laneId < nTopk ? ld_nc_global(payloadView.topKIndices(sourcePayload) + laneId) : -1;
+    const int routedExpertIdx = laneId < nTopk ? payloadView.topKIndices(sourcePayload)[laneId] : -1;
     const int localExpertIdx = routedExpertIdx >= globalExpertBase && routedExpertIdx < globalExpertEnd
                                    ? routedExpertIdx - globalExpertBase
                                    : -1;
     const int sourceTokenIdx = warpBroadcast(
-        laneId == 0 ? ld_nc_global(payloadView.srcTokenGlobalIdx(sourcePayload)) - sourceRank * maxTokensPerRank : 0,
-        0);
+        laneId == 0 ? *payloadView.srcTokenGlobalIdx(sourcePayload) - sourceRank * maxTokensPerRank : 0, 0);
     int outputTokenIdx = -1;
     int combineInputOffset = -1;
     if (localExpertIdx >= 0) {
@@ -376,21 +383,24 @@ MSCCLPP_DEVICE_INLINE void dispatchRecvWorker(void* output, int* outputSrcInfo,
 
 template <int Hidden>
 __global__ __launch_bounds__(DispatchNThreads, 1) void dispatchKernel(
-    void* output, int* outputSrcInfo, int64_t* outputLayout, int* outputCount, int nExperts, int rank, int nRanks,
-    const int64_t* topkIndices, const float* topkWeights, const void* inputTokens, int nTokens, int nTopk,
-    int maxTokensPerRank, void* recvBuffer, void* rdmaBufferBase, void* const* peerRecvBuffers, void* workspace) {
+    void* output, int* outputSrcInfo, int64_t* outputLayout, int* outputCount,
+    mscclpp::BaseMemoryChannelDeviceHandle* baseMemoryChannels, int nExperts, int rank, int nRanks,
+    const int64_t* __restrict__ topkIndices, const float* __restrict__ topkWeights, const void* inputTokens,
+    int nTokens, int nTopk, int maxTokensPerRank, void* recvBuffer, void* rdmaBufferBase, void* const* peerRecvBuffers,
+    void* workspace) {
   extern __shared__ __align__(128) uint8_t sharedMemory[];
   auto* sharedMem = reinterpret_cast<int*>(sharedMemory);
-  const int nWorkerBlocks = static_cast<int>(gridDim.x) - low_latency::DispatchControlBlocks;
+  const int nWorkerBlocks = static_cast<int>(gridDim.x) - DispatchControlBlocks;
   WorkspaceView workspaceView(workspace, nRanks, nExperts);
   const uint32_t metadataFlag = *workspaceView.metadataEpoch_ + 1;
 
-  dispatchSend<Hidden>(inputTokens, nExperts, rank, nRanks, topkIndices, topkWeights, nTokens, nTopk, maxTokensPerRank,
-                       recvBuffer, peerRecvBuffers, rdmaBufferBase, workspace, metadataFlag, sharedMem);
+  dispatchSend<Hidden>(inputTokens, baseMemoryChannels, nExperts, rank, nRanks, topkIndices, topkWeights, nTokens,
+                       nTopk, maxTokensPerRank, recvBuffer, peerRecvBuffers, rdmaBufferBase, workspace, metadataFlag,
+                       sharedMem);
 
   if (static_cast<int>(blockIdx.x) == 0) {
-    dispatchRecvScheduler(outputLayout, outputCount, nExperts, rank, nRanks, recvBuffer, workspace, metadataFlag,
-                          sharedMem);
+    dispatchRecvScheduler(outputLayout, outputCount, baseMemoryChannels, nExperts, rank, nRanks, recvBuffer, workspace,
+                          metadataFlag, sharedMem);
   } else if (static_cast<int>(blockIdx.x) <= nWorkerBlocks) {
     dispatchRecvWorker<Hidden>(output, outputSrcInfo, outputLayout, nExperts, rank, nRanks, nTopk, maxTokensPerRank,
                                recvBuffer, workspace, metadataFlag, sharedMem);
@@ -420,8 +430,9 @@ inline void dispatchHidden(void* output, int* outputSrcInfo, int64_t* outputLayo
       configureKernel(dispatchKernel<Hidden>, DispatchNThreads, dynamicSharedBytes, comm, kernelConfig);
   EP_HOST_ASSERT(residentBlocks >= numBlocks);
   dispatchKernel<Hidden><<<dim3(numBlocks), dim3(DispatchNThreads), dynamicSharedBytes, stream>>>(
-      output, outputSrcInfo, outputLayout, outputCount, nExperts, rank, nRanks, topkIdx, topkWeights, input, nTokens,
-      nTopk, workload.maxTokensPerRank_, recvBuffer, comm.rdmaBufferBase_, comm.peerBases_, workspace);
+      output, outputSrcInfo, outputLayout, outputCount, comm.baseMemoryChannels_, nExperts, rank, nRanks, topkIdx,
+      topkWeights, input, nTokens, nTopk, workload.maxTokensPerRank_, recvBuffer, comm.rdmaBufferBase_, comm.peerBases_,
+      workspace);
   CUDA_CHECK(cudaGetLastError());
 }
 
@@ -432,16 +443,17 @@ inline void dispatch(void* output, int* outputSrcInfo, int64_t* outputLayout, in
   const int nExperts = workload.numExperts_;
   const int rank = comm.rank_;
   const int nRanks = comm.numRanks_;
-  const int numWorkerBlocks = numBlocks - low_latency::DispatchControlBlocks;
+  const int numWorkerBlocks = numBlocks - DispatchControlBlocks;
 
   EP_HOST_ASSERT(nRanks > 0);
   EP_HOST_ASSERT(nExperts > 0);
   EP_HOST_ASSERT(nExperts % nRanks == 0);
   EP_HOST_ASSERT(rank >= 0 && rank < nRanks);
+  EP_HOST_ASSERT(comm.baseMemoryChannels_ != nullptr);
   EP_HOST_ASSERT(workload.numTokens_ >= 0);
   EP_HOST_ASSERT(workload.numTopk_ > 0 && workload.numTopk_ <= WARP_SIZE);
   EP_HOST_ASSERT(nRanks <= 2 * WARP_SIZE);
-  EP_HOST_ASSERT(numWorkerBlocks >= nRanks && numWorkerBlocks <= low_latency::MaxWorkerBlocks);
+  EP_HOST_ASSERT(numWorkerBlocks >= nRanks && numWorkerBlocks <= MaxWorkerBlocks);
   EP_HOST_ASSERT(output != nullptr);
   EP_HOST_ASSERT(outputSrcInfo != nullptr);
   EP_HOST_ASSERT(outputLayout != nullptr);
@@ -471,15 +483,15 @@ inline void dispatch(void* output, int* outputSrcInfo, int64_t* outputLayout, in
   }
 }
 
-}  // namespace low_latency_opt
+}  // namespace detail
 
-namespace low_latency {
+size_t workspaceSize(int numRanks, int numExperts) { return detail::workspaceBytes(numRanks, numExperts); }
 
 void dispatch(void* output, int* outputSrcInfo, int64_t* outputLayout, int* outputCount, const void* input,
               const int64_t* topkIdx, const float* topkWeights, const Workload& workload, void* recvBuffer,
               const CommContext& comm, void* workspace, int numBlocks, cudaStream_t stream) {
-  low_latency_opt::dispatch(output, outputSrcInfo, outputLayout, outputCount, input, topkIdx, topkWeights, workload,
-                            recvBuffer, comm, workspace, numBlocks, stream);
+  detail::dispatch(output, outputSrcInfo, outputLayout, outputCount, input, topkIdx, topkWeights, workload, recvBuffer,
+                   comm, workspace, numBlocks, stream);
 }
 
 }  // namespace low_latency
diff --git a/src/ext/ep/moe_runtime.cc b/src/ext/ep/moe_runtime.cc
index c4cdb797..d5f62ad0 100644
--- a/src/ext/ep/moe_runtime.cc
+++ b/src/ext/ep/moe_runtime.cc
@@ -114,6 +114,7 @@ void MoERuntime::setup() {
   communicator_->bootstrap()->barrier();
   CUDA_CHECK(cudaDeviceSynchronize());
 
+  const mscclpp::EndpointConfig ipcConfig(ipcTransport);
   const int ipcDomainSize = useFabricIpcAlloc ? numRanks_ : numNvlRanks_;
   auto isIpcPeer = [&](int peer) {
     return peer != rank_ && ipcDomainSize > 1 && rank_ / ipcDomainSize == peer / ipcDomainSize;
@@ -123,28 +124,39 @@ void MoERuntime::setup() {
   peerRdmaMemories_.resize(numRanks_);
   peerRdmaMemories_[rank_] = communicator_->registerMemory(rdmaBufferPtr_, numRdmaBytes_, ipcTransport);
   std::vector<std::shared_future<mscclpp::RegisteredMemory>> remoteFutures(numRanks_);
+  std::vector<std::shared_future<mscclpp::Connection>> connectionFutures(numRanks_);
   for (int r = 0; r < numRanks_; ++r) {
     if (!isIpcPeer(r)) continue;
     communicator_->sendMemory(peerRdmaMemories_[rank_], r, IpcTag);
     remoteFutures[r] = communicator_->recvMemory(r, IpcTag);
+    connectionFutures[r] = communicator_->connect(ipcConfig, r, IpcTag);
   }
 
   peerRdmaBases_.assign(numRanks_, nullptr);
   peerRdmaBases_[rank_] = rdmaBufferPtr_;
+  std::vector<mscclpp::BaseMemoryChannelDeviceHandle> baseMemoryChannelHandles(numRanks_);
   for (int r = 0; r < numRanks_; ++r) {
     if (!isIpcPeer(r)) continue;
     peerRdmaMemories_[r] = remoteFutures[r].get();
     peerRdmaBases_[r] = peerRdmaMemories_[r].data();
+    auto semaphore =
+        std::make_shared<mscclpp::MemoryDevice2DeviceSemaphore>(*communicator_, connectionFutures[r].get());
+    baseMemoryChannels_.emplace_back(semaphore);
+    baseMemoryChannelHandles[r] = baseMemoryChannels_.back().deviceHandle();
   }
 
   CUDA_CHECK(cudaMalloc(&peerRdmaBasesGpu_, sizeof(void*) * numRanks_));
   CUDA_CHECK(cudaMemcpy(peerRdmaBasesGpu_, peerRdmaBases_.data(), sizeof(void*) * numRanks_, cudaMemcpyHostToDevice));
+  baseMemoryChannelHandles_ = mscclpp::detail::gpuCallocShared<mscclpp::BaseMemoryChannelDeviceHandle>(numRanks_);
+  mscclpp::gpuMemcpy<mscclpp::BaseMemoryChannelDeviceHandle>(
+      baseMemoryChannelHandles_.get(), baseMemoryChannelHandles.data(), numRanks_, cudaMemcpyHostToDevice);
 
   int maxSharedMemoryPerBlock;
   int numSms;
   CUDA_CHECK(cudaDeviceGetAttribute(&maxSharedMemoryPerBlock, cudaDevAttrMaxSharedMemoryPerBlockOptin, deviceId_));
   CUDA_CHECK(cudaDeviceGetAttribute(&numSms, cudaDevAttrMultiProcessorCount, deviceId_));
   commContext_ = {.rdmaBufferBase_ = rdmaBufferPtr_,
+                  .baseMemoryChannels_ = baseMemoryChannelHandles_.get(),
                   .peerBases_ = peerRdmaBasesGpu_,
                   .maxSharedMemoryPerBlock_ = maxSharedMemoryPerBlock,
                   .numSms_ = numSms,
@@ -174,8 +186,7 @@ void MoERuntime::dispatch(void* output, int* outputSrcInfo, int64_t* outputLayou
                                        .numTopk_ = numTopk,
                                        .numExperts_ = numExperts,
                                        .maxTokensPerRank_ = maxTokensPerRank};
-  const size_t workspaceBytes =
-      static_cast<size_t>(3 * numRanks_ + numExperts + 3 * 128 + 3) * sizeof(int) + sizeof(mscclpp::DeviceSyncer);
+  const size_t workspaceBytes = low_latency::workspaceSize(numRanks_, numExperts);
   EP_HOST_ASSERT(workspaceBytes <= NUM_WORKSPACE_BYTES);
   low_latency::dispatch(output, outputSrcInfo, outputLayout, outputCount, input, topkIdx, topkWeights, workload,
                         recvBuffer, commContext_, workspace_, numBlocks, stream);
@@ -192,9 +203,7 @@ void MoERuntime::combine(void* output, const void* input, const int64_t* topkIdx
 
   low_latency::Layout layout(rdmaBufferPtr_, maxTokensPerRank, hidden, numRanks_, numExperts, numTopk);
   EP_HOST_ASSERT(layout.totalBytes_ <= static_cast<size_t>(numRdmaBytes_));
-  auto& combineBuffer = layout.buffers_[lowLatencyBufferIdx_];
-  void* recvBuffer = combineBuffer.combineData_;
-  auto* readyPackets = combineBuffer.combineReadyPackets_;
+  void* recvBuffer = layout.buffers_[lowLatencyBufferIdx_].combineData_;
   lowLatencyBufferIdx_ ^= 1;
   void* dispatchRecvBuffer = layout.buffers_[lowLatencyBufferIdx_].dispatchData_;
 
@@ -203,7 +212,7 @@ void MoERuntime::combine(void* output, const void* input, const int64_t* topkIdx
                                        .numTopk_ = numTopk,
                                        .numExperts_ = numExperts,
                                        .maxTokensPerRank_ = maxTokensPerRank};
-  low_latency::combine(output, input, topkIdx, topkWeights, srcInfo, layoutRange, workload, recvBuffer, readyPackets,
+  low_latency::combine(output, input, topkIdx, topkWeights, srcInfo, layoutRange, workload, recvBuffer,
                        dispatchRecvBuffer, commContext_, workspace_, numBlocks, mode, stream);
 }
 
diff --git a/src/ext/ep/moe_runtime.hpp b/src/ext/ep/moe_runtime.hpp
index 68df3fb8..933e7983 100644
--- a/src/ext/ep/moe_runtime.hpp
+++ b/src/ext/ep/moe_runtime.hpp
@@ -8,6 +8,7 @@
 #include <memory>
 #include <mscclpp/core.hpp>
 #include <mscclpp/gpu_utils.hpp>
+#include <mscclpp/memory_channel.hpp>
 #include <string>
 #include <vector>
 
@@ -60,6 +61,8 @@ class MoERuntime {
   std::vector<void*> peerRdmaBases_;
   std::vector<mscclpp::RegisteredMemory> peerRdmaMemories_;
   void** peerRdmaBasesGpu_ = nullptr;
+  std::vector<mscclpp::BaseMemoryChannel> baseMemoryChannels_;
+  std::shared_ptr<mscclpp::BaseMemoryChannelDeviceHandle> baseMemoryChannelHandles_;
 
   void setup();
 };
diff --git a/test/mscclpp-test/CMakeLists.txt b/test/mscclpp-test/CMakeLists.txt
index 485b9a85..63d207d9 100644
--- a/test/mscclpp-test/CMakeLists.txt
+++ b/test/mscclpp-test/CMakeLists.txt
@@ -21,12 +21,6 @@ add_mscclpp_test_executable(sendrecv_test_perf sendrecv_test.cu)
 add_mscclpp_test_executable(allgather_test_perf allgather_test.cu)
 add_mscclpp_test_executable(allreduce_test_perf allreduce_test.cu)
 add_mscclpp_test_executable(alltoall_test_perf alltoall_test.cu)
-add_mscclpp_test_executable(ep_low_latency_opt_signal_perf ep_low_latency_opt_signal_perf.cu)
-target_sources(ep_low_latency_opt_signal_perf PRIVATE
-    ${PROJECT_SOURCE_DIR}/src/ext/ep/low_latency/dispatch.cu)
-target_include_directories(ep_low_latency_opt_signal_perf PRIVATE
-    ${PROJECT_SOURCE_DIR}/src/ext/ep
-    ${PROJECT_SOURCE_DIR}/src/ext/ep/include)
 if(MSCCLPP_USE_CUDA)
     set(_mscclpp_ep_test_gpu_archs "")
     foreach(_arch IN LISTS CMAKE_CUDA_ARCHITECTURES)
@@ -39,9 +33,6 @@ if(MSCCLPP_USE_CUDA)
             endif()
         endif()
     endforeach()
-    set_target_properties(ep_low_latency_opt_signal_perf PROPERTIES
-        CUDA_ARCHITECTURES "${_mscclpp_ep_test_gpu_archs}")
-
     add_executable(tma_pipeline_perf tma_pipeline_perf.cu)
     target_link_libraries(tma_pipeline_perf PRIVATE ${GPU_LIBRARIES})
     target_include_directories(tma_pipeline_perf ${TEST_INC_COMMON})