From 5225ec0ad0ab9859f8fb84c527dd3e81892b1e3d Mon Sep 17 00:00:00 2001 From: sunlei1024 Date: Fri, 5 Dec 2025 11:23:48 +0000 Subject: [PATCH 01/13] feat(fmq): add ZMQ-based FMQ implementation and benchmark tools --- benchmarks/benchmark_fmq.py | 234 +++++++++++++ fastdeploy/inter_communicator/fmq.py | 336 +++++++++++++++++++ fastdeploy/inter_communicator/fmq_factory.py | 83 +++++ fastdeploy/utils.py | 1 + tests/inter_communicator/test_fmq.py | 92 +++++ tests/inter_communicator/test_fmq_factory.py | 91 +++++ 6 files changed, 837 insertions(+) create mode 100644 benchmarks/benchmark_fmq.py create mode 100644 fastdeploy/inter_communicator/fmq.py create mode 100644 fastdeploy/inter_communicator/fmq_factory.py create mode 100644 tests/inter_communicator/test_fmq.py create mode 100644 tests/inter_communicator/test_fmq_factory.py diff --git a/benchmarks/benchmark_fmq.py b/benchmarks/benchmark_fmq.py new file mode 100644 index 00000000000..f7a1b74031e --- /dev/null +++ b/benchmarks/benchmark_fmq.py @@ -0,0 +1,234 @@ +""" +# Copyright (c) 2025 PaddlePaddle Authors. All Rights Reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License" +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +""" + +import asyncio +import multiprocessing as mp +import os +import statistics +import time + +from tqdm import tqdm + +from fastdeploy.inter_communicator.fmq import FMQ + + +# ============================================================ +# Producer Task +# ============================================================ +async def producer_task(proc_id, msg_count, payload_size, shm_threshold, result_q): + fmq = FMQ() + q = fmq.queue("mp_bench_latency", role="producer") + payload = b"x" * payload_size + + # tqdm 进度条 + pbar = tqdm(total=msg_count, desc=f"Producer-{proc_id}", position=proc_id, leave=True, disable=False) + + t0 = time.perf_counter() + for i in range(msg_count): + send_ts = time.perf_counter() + await q.put(data={"pid": proc_id, "i": i, "send_ts": send_ts, "payload": payload}, shm_threshold=shm_threshold) + pbar.update(1) + # pbar.write(f"send {i}") + t1 = time.perf_counter() + result_q.put({"producer_id": proc_id, "count": msg_count, "time": t1 - t0}) + + pbar.close() + + # wait for 2 seconds before closing + await asyncio.sleep(5) + + +def producer_process(proc_id, msg_count, payload_size, shm_threshold, result_q): + async def run(): + await producer_task(proc_id, msg_count, payload_size, shm_threshold, result_q) + + asyncio.run(run()) + + +# ============================================================ +# Consumer Task +# ============================================================ +async def consumer_task(consumer_id, total_msgs, result_q, consumer_event): + fmq = FMQ() + q = fmq.queue("mp_bench_latency", role="consumer") + consumer_event.set() + + latencies = [] + recv = 0 + + # tqdm 显示进度 + pbar = tqdm(total=total_msgs, desc=f"Consumer-{consumer_id}", position=consumer_id + 1, leave=True, disable=False) + + first_recv = None + last_recv = None + + while recv < total_msgs: + msg = await q.get() + recv_ts = time.perf_counter() + if msg is None: + pbar.write("recv None") + continue + if first_recv is None: + first_recv = recv_ts + last_recv = recv_ts + send_ts = msg.payload["send_ts"] + latencies.append((recv_ts - send_ts) * 1000) # ms + pbar.update(1) + # pbar.write(f"recv {msg.payload['i']} = {msg.msg_id}") + recv += 1 + + pbar.close() + + result_q.put( + {"consumer_id": consumer_id, "latencies": latencies, "first_recv": first_recv, "last_recv": last_recv} + ) + + +def consumer_process(consumer_id, total_msgs, result_q, consumer_event): + async def run(): + await consumer_task(consumer_id, total_msgs, result_q, consumer_event) + + asyncio.run(run()) + + +# ============================================================ +# MAIN benchmark +# ============================================================ +def run_benchmark( + NUM_PRODUCERS=1, + NUM_CONSUMERS=1, + NUM_MESSAGES_PER_PRODUCER=1000, + PAYLOAD_SIZE=1 * 1024 * 1024, + SHM_THRESHOLD=1 * 1024 * 1024, +): + total_messages = NUM_PRODUCERS * NUM_MESSAGES_PER_PRODUCER + total_bytes = total_messages * PAYLOAD_SIZE + + print(f"\nFastDeploy Message Queue Benchmark, pid:{os.getpid()}") + print(f"Producers: {NUM_PRODUCERS}") + print(f"Consumers: {NUM_CONSUMERS}") + print(f"Messages per producer: {NUM_MESSAGES_PER_PRODUCER}") + print(f"Total bytes: {total_bytes / 1024 / 1024 / 1024:.2f} GB") + print(f"Total messages: {total_messages:,}") + print(f"Payload per message: {PAYLOAD_SIZE / 1024 / 1024:.2f} MB") + + mp.set_start_method("fork") + manager = mp.Manager() + result_q = manager.Queue() + + # 两个信号事件 + consumer_event = manager.Event() + + procs = [] + + # Start Consumers + msgs_per_consumer = total_messages // NUM_CONSUMERS + for i in range(NUM_CONSUMERS): + p = mp.Process(target=consumer_process, args=(i, msgs_per_consumer, result_q, consumer_event)) + procs.append(p) + p.start() + + consumer_event.wait() + + # Start Producers + for i in range(NUM_PRODUCERS): + p = mp.Process( + target=producer_process, args=(i, NUM_MESSAGES_PER_PRODUCER, PAYLOAD_SIZE, SHM_THRESHOLD, result_q) + ) + procs.append(p) + p.start() + + # Join + for p in procs: + p.join() + + # Collect results + producer_stats = [] + consumer_stats = {} + + while not result_q.empty(): + item = result_q.get() + if "producer_id" in item: + producer_stats.append(item) + if "consumer_id" in item: + consumer_stats[item["consumer_id"]] = item + + # Producer stats + print("\nProducer Stats:") + for p in producer_stats: + throughput = p["count"] / p["time"] + bandwidth = (p["count"] * PAYLOAD_SIZE) / (1024**2 * p["time"]) + print( + f"[Producer-{p['producer_id']}] Sent {p['count']:,} msgs " + f"in {p['time']:.3f} s | Throughput: {throughput:,.0f} msg/s | Bandwidth: {bandwidth:.2f} MB/s" + ) + + # Consumer latency stats + print("\nConsumer Latency Stats:") + all_latencies = [] + first_recv_times = [] + last_recv_times = [] + + for cid, data in consumer_stats.items(): + lats = data["latencies"] + if len(lats) == 0: + continue + all_latencies.extend(lats) + first_recv_times.append(data["first_recv"]) + last_recv_times.append(data["last_recv"]) + + avg = statistics.mean(lats) + p50 = statistics.median(lats) + p95 = statistics.quantiles(lats, n=20)[18] + p99 = statistics.quantiles(lats, n=100)[98] + + print( + f"[Consumer-{cid}] msgs={len(lats):5d} | avg={avg:.3f} ms | " + f"P50={p50:.3f} ms | P95={p95:.3f} ms | P99={p99:.3f} ms" + ) + + # Global summary + if first_recv_times and last_recv_times: + total_time = max(last_recv_times) - min(first_recv_times) + global_throughput = total_messages / total_time + global_bandwidth = total_bytes / (1024**2 * total_time) + + if all_latencies: + avg_latency = statistics.mean(all_latencies) + min_latency = min(all_latencies) + max_latency = max(all_latencies) + p50_latency = statistics.median(all_latencies) + p95_latency = statistics.quantiles(all_latencies, n=20)[18] + p99_latency = statistics.quantiles(all_latencies, n=100)[98] + else: + avg_latency = min_latency = max_latency = p50_latency = p95_latency = p99_latency = 0.0 + + print("\nGlobal Summary:") + print(f"Total messages : {total_messages:,}") + print(f"Total data : {total_bytes / 1024**2:.2f} MB") + print(f"Total time : {total_time:.3f} s") + print(f"Global throughput: {global_throughput:,.0f} msg/s") + print(f"Global bandwidth : {global_bandwidth:.2f} MB/s") + print( + f"Latency (ms) : avg={avg_latency:.3f} " + f"| min={min_latency:.3f} | max={max_latency:.3f} " + f"| P50={p50_latency:.3f} | P95={p95_latency:.3f} | P99={p99_latency:.3f}\n" + ) + + +# Entry +if __name__ == "__main__": + run_benchmark() diff --git a/fastdeploy/inter_communicator/fmq.py b/fastdeploy/inter_communicator/fmq.py new file mode 100644 index 00000000000..0b8f51d168b --- /dev/null +++ b/fastdeploy/inter_communicator/fmq.py @@ -0,0 +1,336 @@ +""" +# Copyright (c) 2025 PaddlePaddle Authors. All Rights Reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License" +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +""" + +import asyncio +import json +import os +import time +import uuid +from dataclasses import dataclass, field +from enum import Enum +from multiprocessing import shared_memory +from multiprocessing.reduction import ForkingPickler +from typing import Any, Callable, Dict, Optional + +import zmq +import zmq.asyncio + +from fastdeploy.utils import fmq_logger + +# ========================== +# Config & Enum Definitions +# ========================== + + +class EndpointType(Enum): + QUEUE = "queue" + TOPIC = "topic" + + +class Role(Enum): + PRODUCER = "producer" + CONSUMER = "consumer" + + +@dataclass +class SocketOptions: + sndhwm: int = 0 + rcvhwm: int = 0 + linger: int = -1 + sndbuf: int = 32 * 1024 * 1024 + rcvbuf: int = 32 * 1024 * 1024 + immediate: int = 1 + + def apply(self, socket: zmq.Socket, is_producer: bool): + # Apply socket-level configurations + socket.setsockopt(zmq.LINGER, self.linger) + socket.setsockopt(zmq.IMMEDIATE, self.immediate) + + if is_producer: + socket.setsockopt(zmq.SNDHWM, self.sndhwm) + socket.setsockopt(zmq.SNDBUF, self.sndbuf) + else: + socket.setsockopt(zmq.RCVHWM, self.rcvhwm) + socket.setsockopt(zmq.RCVBUF, self.rcvbuf) + + +@dataclass +class Endpoint: + # Represents a single endpoint with protocol, address, io_threads, and copy behavior + protocol: str + address: str + io_threads: int = 1 + copy: bool = False + + +@dataclass +class Config: + ipc_root: str = "/dev/shm" + io_threads: int = 1 + copy: bool = False + endpoints: Dict[str, Endpoint] = field(default_factory=dict) + socket_config: SocketOptions = SocketOptions() + + +# ========================== +# Endpoint Manager +# ========================== + + +class EndpointManager: + config: Config = Config() + + @classmethod + def load_config(cls, _ignored_file_path: str = None): + cfg_str = os.getenv("FMQ_CONFIG_JSON", None) + if cfg_str: + try: + custom_cfg = json.loads(cfg_str) + for key, value in vars(custom_cfg).items(): + if value is not None: + setattr(cls.config, key, value) + except Exception as e: + fmq_logger.error(f"Failed to load FMQ config: {e}") + fmq_logger.info(f"Loaded FMQ config: {cls.config}") + + @classmethod + def get_endpoint(cls, name: str) -> Endpoint: + # Retrieve endpoint object + if name in cls.config.endpoints: + return cls.config.endpoints[name] + + # Fallback: auto-generate endpoint + address = f"{cls.config.ipc_root}/fmq_{name}.ipc" + return Endpoint(protocol="ipc", address=address) + + +# ========================== +# Shared Memory Descriptor +# ========================== + + +@dataclass +class Descriptor: + shm_name: str + size: int + + @staticmethod + def create(data_bytes: bytes) -> "Descriptor": + # Create shared memory buffer and store payload + name = f"fmq_shm_{uuid.uuid4().hex}" + shm = shared_memory.SharedMemory(create=True, size=len(data_bytes), name=name) + shm.buf[: len(data_bytes)] = data_bytes + shm.close() + return Descriptor(shm_name=name, size=len(data_bytes)) + + def read_and_unlink(self) -> bytes: + # Read and cleanup shared memory + try: + shm = shared_memory.SharedMemory(name=self.shm_name) + data = bytes(shm.buf[: self.size]) + shm.close() + shm.unlink() + return data + except FileNotFoundError: + return b"" + + +# ========================== +# Message Wrapper +# ========================== + + +@dataclass +class Message: + payload: Any + msg_id: int = None + timestamp: float = field(default_factory=time.time) + descriptor: Optional[Descriptor] = None + + def serialize(self) -> bytes: + # Serialize message + return ForkingPickler.dumps(self) + + @staticmethod + def deserialize(data: bytes) -> "Message": + # Deserialize message + return ForkingPickler.loads(data) + + +# ========================== +# Base Component +# ========================== + + +class BaseComponent: + def __init__(self, context: zmq.asyncio.Context, endpoint: Endpoint): + self.context = context + self.endpoint = endpoint + self.socket = None + self.lock = asyncio.Lock() + + async def close(self): + # Close socket + if self.socket: + self.socket.close() + + +# ========================== +# FIFO Queue +# ========================== + + +class Queue(BaseComponent): + def __init__(self, context, name: str, role: str = "producer"): + endpoint = EndpointManager.get_endpoint(name) + super().__init__(context, endpoint) + + self.name = name + self.role = Role(role) + self.copy = endpoint.copy + self.socket_conf = EndpointManager.config.socket_config + self._msg_id = 0 + + full_ep = f"{endpoint.protocol}://{endpoint.address}" + + self.socket = self.context.socket(zmq.PUSH if self.role == Role.PRODUCER else zmq.PULL) + self.socket_conf.apply(self.socket, self.role == Role.PRODUCER) + + if self.role == Role.PRODUCER: + self.socket.connect(full_ep) + else: + self.socket.bind(full_ep) + + fmq_logger.info(f"Queue {name} initialized on {full_ep}") + + async def put(self, data: Any, shm_threshold: int = 1024 * 1024): + # Send data to queue + if self.role != Role.PRODUCER: + raise PermissionError("Only producers can send messages.") + + desc = None + payload = data + + if isinstance(data, bytes) and len(data) >= shm_threshold: + desc = Descriptor.create(data) + payload = None + + msg = Message(msg_id=self._msg_id, payload=payload, descriptor=desc) + raw = msg.serialize() + + async with self.lock: + await self.socket.send(raw, copy=self.copy) + self._msg_id += 1 + + async def get(self, timeout: int = None) -> Optional[Message]: + # Receive data from queue + if self.role != Role.CONSUMER: + raise PermissionError("Only consumers can get messages.") + + try: + if timeout: + raw = await asyncio.wait_for(self.socket.recv(), timeout / 1000) + else: + raw = await self.socket.recv(copy=self.copy) + except asyncio.TimeoutError: + fmq_logger.error(f"Timeout receiving message on {self.name}") + return None + + msg = Message.deserialize(raw) + if msg.descriptor: + msg.payload = msg.descriptor.read_and_unlink() + + self._msg_id += 1 + return msg + + +# ========================== +# Pub/Sub Topic +# ========================== + + +class Topic(BaseComponent): + def __init__(self, context, name: str): + endpoint = EndpointManager.get_endpoint(name) + super().__init__(context, endpoint) + self.name = name + self._pub_socket = None + self._sub_socket = None + self._task = None + + async def pub(self, data: Any): + # Publish a message + if not self._pub_socket: + ep = f"{self.endpoint.protocol}://{self.endpoint.address}" + self._pub_socket = self.context.socket(zmq.PUB) + self._pub_socket.bind(ep) + await asyncio.sleep(0.05) + + msg = Message(payload=data) + async with self.lock: + await self._pub_socket.send(msg.serialize()) + + async def sub(self, callback: Callable[[Message], Any]): + # Subscribe and handle messages + if not self._sub_socket: + ep = f"{self.endpoint.protocol}://{self.endpoint.address}" + self._sub_socket = self.context.socket(zmq.SUB) + self._sub_socket.connect(ep) + self._sub_socket.setsockopt_string(zmq.SUBSCRIBE, "") + + async def loop(): + while True: + raw = await self._sub_socket.recv() + msg = Message.deserialize(raw) + result = callback(msg) + if asyncio.iscoroutine(result): + await result + + self._task = asyncio.create_task(loop()) + + +# ========================== +# FMQ Main Interface +# ========================== + + +class FMQ: + _instance = None + _context = None + + def __new__(cls, config_path="fmq_config.json"): + if cls._instance is None: + cls._instance = super().__new__(cls) + EndpointManager.load_config() + + # Determine IO threads based on global defaults + io_threads = 1 + if EndpointManager.config.endpoints: + # Use max io_threads among all endpoints + io_threads = max(ep.io_threads for ep in EndpointManager.config.endpoints.values()) + + cls._context = zmq.asyncio.Context(io_threads=io_threads) + return cls._instance + + def queue(self, name: str, role="producer") -> Queue: + return Queue(self._context, name, role) + + def topic(self, name: str) -> Topic: + return Topic(self._context, name) + + async def destroy(self): + # Destroy ZeroMQ context + self._context.term() diff --git a/fastdeploy/inter_communicator/fmq_factory.py b/fastdeploy/inter_communicator/fmq_factory.py new file mode 100644 index 00000000000..d1c8e4dd244 --- /dev/null +++ b/fastdeploy/inter_communicator/fmq_factory.py @@ -0,0 +1,83 @@ +""" +# Copyright (c) 2025 PaddlePaddle Authors. All Rights Reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License" +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +""" + +from fastdeploy.inter_communicator.fmq import FMQ + + +class FMQFactory: + """ + Static factory for creating the four standard FMQ queues: + 1. q_a2e: api server --> engine + 2. q_e2w: engine --> worker + 3. q_w2e: worker --> engine + 4. q_e2a: engine --> api server + API Server: q_a2e producer / q_e2a consumer + Engine: q_a2e consumer / q_e2w producer / q_w2e consumer / q_e2a producer + Worker: q_e2w consumer / q_w2e producer + """ + + _fmq = FMQ() + + # ------------------------------ + # API → Engine + # ------------------------------ + @classmethod + def q_a2e_producer(cls): + return cls._fmq.queue("q_a2e", role="producer") + + @classmethod + def q_a2e_consumer(cls): + return cls._fmq.queue("q_a2e", role="consumer") + + # ------------------------------ + # Engine → Worker + # ------------------------------ + @classmethod + def q_e2w_producer(cls): + return cls._fmq.queue("q_e2w", role="producer") + + @classmethod + def q_e2w_consumer(cls): + return cls._fmq.queue("q_e2w", role="consumer") + + # ------------------------------ + # Worker → Engine + # ------------------------------ + @classmethod + def q_w2e_producer(cls): + return cls._fmq.queue("q_w2e", role="producer") + + @classmethod + def q_w2e_consumer(cls): + return cls._fmq.queue("q_w2e", role="consumer") + + # ------------------------------ + # Engine → API + # ------------------------------ + @classmethod + def q_e2a_producer(cls): + return cls._fmq.queue("q_e2a", role="producer") + + @classmethod + def q_e2a_consumer(cls): + return cls._fmq.queue("q_e2a", role="consumer") + + # ------------------------------ + # Destroy context + # ------------------------------ + @classmethod + async def destroy(cls): + await cls._fmq.destroy() diff --git a/fastdeploy/utils.py b/fastdeploy/utils.py index a0878fa7c73..97a975f4e84 100644 --- a/fastdeploy/utils.py +++ b/fastdeploy/utils.py @@ -1051,6 +1051,7 @@ def _bos_download(bos_client, link): zmq_client_logger = get_logger("zmq_client", "zmq_client.log") trace_logger = FastDeployLogger().get_trace_logger("trace_logger", "trace_logger.log") router_logger = get_logger("router", "router.log") +fmq_logger = get_logger("fmq", "fmq.log") def parse_type(return_type: Callable[[str], T]) -> Callable[[str], T]: diff --git a/tests/inter_communicator/test_fmq.py b/tests/inter_communicator/test_fmq.py new file mode 100644 index 00000000000..75011c2d7e4 --- /dev/null +++ b/tests/inter_communicator/test_fmq.py @@ -0,0 +1,92 @@ +""" +# Copyright (c) 2025 PaddlePaddle Authors. All Rights Reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License" +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +""" + +import asyncio +import json +import os +import unittest + +from fastdeploy.inter_communicator.fmq import FMQ, Message + +# Prepare environment config for testing +cfg = { + "ipc_root": "/dev/shm", + "io_threads": 1, + "copy": True, + "endpoints": { + "test_queue": {"protocol": "ipc", "address": "/dev/shm/fmq_test_queue.ipc", "io_threads": 1, "copy": True}, + "test_topic": {"protocol": "ipc", "address": "/dev/shm/fmq_test_topic.ipc", "io_threads": 1, "copy": True}, + }, +} +os.environ["FMQ_CONFIG_JSON"] = json.dumps(cfg) + + +class TestFMQ(unittest.TestCase): + + def setUp(self): + self.fmq = FMQ() + + def test_queue_send_receive(self): + async def run_test(): + producer = self.fmq.queue("test_queue", role="producer") + consumer = self.fmq.queue("test_queue", role="consumer") + + test_data = b"hello world" + await producer.put(test_data) + msg = await consumer.get(timeout=1000) + + self.assertIsNotNone(msg) + self.assertEqual(msg.payload, test_data) + + asyncio.run(run_test()) + + def test_queue_large_shm_transfer(self): + async def run_test(): + producer = self.fmq.queue("test_queue", role="producer") + consumer = self.fmq.queue("test_queue", role="consumer") + + large_data = b"x" * (2 * 1024 * 1024) # > 1MB + await producer.put(large_data) + msg = await consumer.get(timeout=1000) + + self.assertIsNotNone(msg) + self.assertEqual(msg.payload, large_data) + self.assertIsNotNone(msg.descriptor) + + asyncio.run(run_test()) + + def test_topic_pub_sub(self): + received = [] + + async def run_test(): + topic = self.fmq.topic("test_topic") + + async def callback(msg: Message): + received.append(msg.payload) + + await topic.sub(callback) + await asyncio.sleep(0.1) # allow SUB to connect + + await topic.pub("hello") + await asyncio.sleep(0.2) + + self.assertIn("hello", received) + + asyncio.run(run_test()) + + +if __name__ == "__main__": + unittest.main() diff --git a/tests/inter_communicator/test_fmq_factory.py b/tests/inter_communicator/test_fmq_factory.py new file mode 100644 index 00000000000..37a1a4193cd --- /dev/null +++ b/tests/inter_communicator/test_fmq_factory.py @@ -0,0 +1,91 @@ +""" +# Copyright (c) 2025 PaddlePaddle Authors. All Rights Reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License" +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +""" + +import unittest + +from fastdeploy.inter_communicator.fmq import Message +from fastdeploy.inter_communicator.fmq_factory import FMQFactory as factory + + +class TestFMQFactory(unittest.IsolatedAsyncioTestCase): + + async def test_create_queues(self): + """测试能否创建所有 producer / consumer 队列""" + q1 = factory.q_a2e_producer() + q2 = factory.q_a2e_consumer() + q3 = factory.q_e2w_producer() + q4 = factory.q_e2w_consumer() + q5 = factory.q_w2e_producer() + q6 = factory.q_w2e_consumer() + q7 = factory.q_e2a_producer() + q8 = factory.q_e2a_consumer() + + self.assertEqual(q1.name, "q_a2e") + self.assertEqual(q2.name, "q_a2e") + self.assertEqual(q3.name, "q_e2w") + self.assertEqual(q4.name, "q_e2w") + self.assertEqual(q5.name, "q_w2e") + self.assertEqual(q6.name, "q_w2e") + self.assertEqual(q7.name, "q_e2a") + self.assertEqual(q8.name, "q_e2a") + + # 同一进程内 context 应相同 + self.assertIs(q1.context, q2.context) + self.assertIs(q1.context, q3.context) + + async def test_message_roundtrip(self): + """测试 producer → consumer 消息流转""" + producer = factory.q_a2e_producer() + consumer = factory.q_a2e_consumer() + + payload = {"k": "v"} + + await producer.put(payload) + msg = await consumer.get(timeout=1500) + + self.assertIsInstance(msg, Message) + self.assertEqual(msg.payload, payload) + + async def test_multi_queue_independence(self): + """测试多个队列互不干扰""" + + prod_a2e = factory.q_a2e_producer() + cons_a2e = factory.q_a2e_consumer() + + prod_e2w = factory.q_e2w_producer() + cons_e2w = factory.q_e2w_consumer() + + await prod_a2e.put("msg_api") + await prod_e2w.put("msg_worker") + + msg1 = await cons_a2e.get(timeout=1500) + msg2 = await cons_e2w.get(timeout=1500) + + self.assertEqual(msg1.payload, "msg_api") + self.assertEqual(msg2.payload, "msg_worker") + + async def test_shared_context(self): + """验证 FMQFactory 始终返回同一个 context (单进程)""" + q1 = factory.q_a2e_producer() + q2 = factory.q_e2w_consumer() + q3 = factory.q_e2a_producer() + + self.assertIs(q1.context, q2.context) + self.assertIs(q1.context, q3.context) + + +if __name__ == "__main__": + unittest.main() From 9ffc4a980e4ad9d8ad9779b8a376bf6de0d657dd Mon Sep 17 00:00:00 2001 From: sunlei1024 Date: Fri, 5 Dec 2025 11:48:27 +0000 Subject: [PATCH 02/13] move FMQ_CONFIG_JSON to envs --- fastdeploy/envs.py | 1 + fastdeploy/inter_communicator/fmq.py | 4 ++-- tests/inter_communicator/test_fmq.py | 6 +++--- 3 files changed, 6 insertions(+), 5 deletions(-) diff --git a/fastdeploy/envs.py b/fastdeploy/envs.py index 93f135d09da..dc734af5eea 100644 --- a/fastdeploy/envs.py +++ b/fastdeploy/envs.py @@ -151,6 +151,7 @@ # "Number of tokens in the group for Mixture of Experts (MoE) computation processing on HPU" "FD_HPU_CHUNK_SIZE": lambda: int(os.getenv("FD_HPU_CHUNK_SIZE", "64")), "FD_PREFILL_WAIT_DECODE_RESOURCE_SECONDS": lambda: int(os.getenv("FD_PREFILL_WAIT_DECODE_RESOURCE_SECONDS", "30")), + "FMQ_CONFIG_JSON": lambda: os.getenv("FMQ_CONFIG_JSON", None), } diff --git a/fastdeploy/inter_communicator/fmq.py b/fastdeploy/inter_communicator/fmq.py index 0b8f51d168b..0915ce3e187 100644 --- a/fastdeploy/inter_communicator/fmq.py +++ b/fastdeploy/inter_communicator/fmq.py @@ -16,7 +16,6 @@ import asyncio import json -import os import time import uuid from dataclasses import dataclass, field @@ -28,6 +27,7 @@ import zmq import zmq.asyncio +from fastdeploy import envs from fastdeploy.utils import fmq_logger # ========================== @@ -95,7 +95,7 @@ class EndpointManager: @classmethod def load_config(cls, _ignored_file_path: str = None): - cfg_str = os.getenv("FMQ_CONFIG_JSON", None) + cfg_str = envs.FMQ_CONFIG_JSON if cfg_str: try: custom_cfg = json.loads(cfg_str) diff --git a/tests/inter_communicator/test_fmq.py b/tests/inter_communicator/test_fmq.py index 75011c2d7e4..a7d6a8153ff 100644 --- a/tests/inter_communicator/test_fmq.py +++ b/tests/inter_communicator/test_fmq.py @@ -25,10 +25,10 @@ cfg = { "ipc_root": "/dev/shm", "io_threads": 1, - "copy": True, + "copy": False, "endpoints": { - "test_queue": {"protocol": "ipc", "address": "/dev/shm/fmq_test_queue.ipc", "io_threads": 1, "copy": True}, - "test_topic": {"protocol": "ipc", "address": "/dev/shm/fmq_test_topic.ipc", "io_threads": 1, "copy": True}, + "test_queue": {"protocol": "ipc", "address": "/dev/shm/fmq_test_queue.ipc", "io_threads": 1, "copy": False}, + "test_topic": {"protocol": "ipc", "address": "/dev/shm/fmq_test_topic.ipc", "io_threads": 1, "copy": False}, }, } os.environ["FMQ_CONFIG_JSON"] = json.dumps(cfg) From 65abc55ab5815714667162dc19fc5bc964cfec7d Mon Sep 17 00:00:00 2001 From: GoldPancake <56388518+Deleter-D@users.noreply.github.com> Date: Fri, 5 Dec 2025 20:01:05 +0800 Subject: [PATCH 03/13] fix top_p_candidates (#5400) Co-authored-by: freeliuzc --- .../speculate_decoding/top_p_candidates.cu | 835 +++++++++--------- 1 file changed, 407 insertions(+), 428 deletions(-) diff --git a/custom_ops/gpu_ops/speculate_decoding/top_p_candidates.cu b/custom_ops/gpu_ops/speculate_decoding/top_p_candidates.cu index 0fced697d7e..bcafd45d08e 100644 --- a/custom_ops/gpu_ops/speculate_decoding/top_p_candidates.cu +++ b/custom_ops/gpu_ops/speculate_decoding/top_p_candidates.cu @@ -19,113 +19,113 @@ template __forceinline__ __device__ T CudaShuffleDownSync(unsigned mask, T val, int delta, int width = warpSize) { - return __shfl_down_sync(mask, val, static_cast(delta), width); + return __shfl_down_sync(mask, val, static_cast(delta), width); } template <> __forceinline__ __device__ phi::dtype::float16 CudaShuffleDownSync( unsigned mask, phi::dtype::float16 val, int delta, int width) { - return paddle::float16(__shfl_down_sync( - mask, val.to_half(), static_cast(delta), width)); + return paddle::float16(__shfl_down_sync( + mask, val.to_half(), static_cast(delta), width)); } template <> __forceinline__ __device__ phi::dtype::bfloat16 CudaShuffleDownSync( unsigned mask, phi::dtype::bfloat16 val, int delta, int width) { - return paddle::bfloat16(__shfl_down_sync( - mask, val.to_nv_bfloat16(), static_cast(delta), width)); + return paddle::bfloat16(__shfl_down_sync( + mask, val.to_nv_bfloat16(), static_cast(delta), width)); } struct BlockPrefixCallbackOp { - // Running prefix - float running_total; - // Constructor - __device__ BlockPrefixCallbackOp(float running_total) - : running_total(running_total) {} - // Callback operator to be entered by the first warp of threads in the - // block. Thread-0 is responsible for returning a value for seeding the - // block-wide scan. - __device__ float operator()(float block_aggregate) { - float old_prefix = running_total; - running_total += block_aggregate; - return old_prefix; - } + // Running prefix + float running_total; + // Constructor + __device__ BlockPrefixCallbackOp(float running_total) + : running_total(running_total) {} + // Callback operator to be entered by the first warp of threads in the + // block. Thread-0 is responsible for returning a value for seeding the + // block-wide scan. + __device__ float operator()(float block_aggregate) { + float old_prefix = running_total; + running_total += block_aggregate; + return old_prefix; + } }; #define FINAL_MASK 0xFFFFFFFF -#define FIXED_BLOCK_DIM_BASE(dim, ...) \ - case (dim): { \ - constexpr auto kBlockDim = (dim); \ - __VA_ARGS__; \ - } break - -#define FIXED_BLOCK_DIM(...) \ - FIXED_BLOCK_DIM_BASE(1024, ##__VA_ARGS__); \ - FIXED_BLOCK_DIM_BASE(512, ##__VA_ARGS__); \ - FIXED_BLOCK_DIM_BASE(256, ##__VA_ARGS__); \ - FIXED_BLOCK_DIM_BASE(128, ##__VA_ARGS__); \ - FIXED_BLOCK_DIM_BASE(64, ##__VA_ARGS__); \ - FIXED_BLOCK_DIM_BASE(32, ##__VA_ARGS__) - -#define FIXED_TOPK_BASE(topk, ...) \ - case (topk): { \ - constexpr auto kTopK = topk; \ - __VA_ARGS__; \ - } break - -#define FIXED_TOPK(...) \ - FIXED_TOPK_BASE(1, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(2, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(3, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(4, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(5, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(6, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(7, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(8, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(9, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(10, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(20, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(30, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(40, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(50, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(60, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(70, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(80, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(90, ##__VA_ARGS__); \ - FIXED_TOPK_BASE(100, ##__VA_ARGS__); +#define FIXED_BLOCK_DIM_BASE(dim, ...) \ + case (dim): { \ + constexpr auto kBlockDim = (dim); \ + __VA_ARGS__; \ + } break + +#define FIXED_BLOCK_DIM(...) \ + FIXED_BLOCK_DIM_BASE(1024, ##__VA_ARGS__); \ + FIXED_BLOCK_DIM_BASE(512, ##__VA_ARGS__); \ + FIXED_BLOCK_DIM_BASE(256, ##__VA_ARGS__); \ + FIXED_BLOCK_DIM_BASE(128, ##__VA_ARGS__); \ + FIXED_BLOCK_DIM_BASE(64, ##__VA_ARGS__); \ + FIXED_BLOCK_DIM_BASE(32, ##__VA_ARGS__) + +#define FIXED_TOPK_BASE(topk, ...) \ + case (topk): { \ + constexpr auto kTopK = topk; \ + __VA_ARGS__; \ + } break + +#define FIXED_TOPK(...) \ + FIXED_TOPK_BASE(1, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(2, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(3, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(4, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(5, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(6, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(7, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(8, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(9, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(10, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(20, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(30, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(40, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(50, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(60, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(70, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(80, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(90, ##__VA_ARGS__); \ + FIXED_TOPK_BASE(100, ##__VA_ARGS__); struct SegmentOffsetIter { - explicit SegmentOffsetIter(int num_cols) : num_cols_(num_cols) {} + explicit SegmentOffsetIter(int num_cols) : num_cols_(num_cols) {} - __host__ __device__ __forceinline__ int operator()(int idx) const { - return idx * num_cols_; - } + __host__ __device__ __forceinline__ int operator()(int idx) const { + return idx * num_cols_; + } - int num_cols_; + int num_cols_; }; inline int div_up(int a, int n) { return (a + n - 1) / n; } template __global__ void FillIndex(T* indices, T num_rows, T num_cols) { - int col_id = threadIdx.x; - int row_id = blockIdx.x; + int col_id = threadIdx.x; + int row_id = blockIdx.x; - for (T j = row_id; j < num_rows; j += gridDim.x) { - for (T i = col_id; i < num_cols; i += blockDim.x) { - indices[j * num_cols + i] = i; - } + for (T j = row_id; j < num_rows; j += gridDim.x) { + for (T i = col_id; i < num_cols; i += blockDim.x) { + indices[j * num_cols + i] = i; } + } } __global__ void SetCountIter(int* count_iter, int num) { - int tid = threadIdx.x; - int bid = blockIdx.x; - int idx = bid * blockDim.x + tid; - for (int i = idx; i < num; i += gridDim.x * blockDim.x) { - count_iter[i] = i; - } + int tid = threadIdx.x; + int bid = blockIdx.x; + int idx = bid * blockDim.x + tid; + for (int i = idx; i < num; i += gridDim.x * blockDim.x) { + count_iter[i] = i; + } } template @@ -137,148 +137,146 @@ __global__ void top_p_candidates_kernel(T* sorted_probs, const int vocab_size, const float topp, const int candidates_len) { - __shared__ int stop_shared; - __shared__ float rand_p; - const int tid = threadIdx.x; - const int bid = blockIdx.x; - constexpr int NUM_WARPS = BLOCK_SIZE / 32; - const int lane_id = tid % 32; - const int warp_id = tid / 32; - - typedef cub::BlockScan BlockScan; - typedef cub::BlockReduce BlockReduce; - __shared__ typename BlockScan::TempStorage temp_storage; - __shared__ typename BlockReduce::TempStorage temp_storage_reduce; - __shared__ uint32_t selected_shared[NUM_WARPS]; - - if (lane_id == 0) { - selected_shared[warp_id] = 0; + __shared__ int stop_shared; + __shared__ float rand_p; + const int tid = threadIdx.x; + const int bid = blockIdx.x; + constexpr int NUM_WARPS = BLOCK_SIZE / 32; + const int lane_id = tid % 32; + const int warp_id = tid / 32; + + typedef cub::BlockScan BlockScan; + typedef cub::BlockReduce BlockReduce; + __shared__ typename BlockScan::TempStorage temp_storage; + __shared__ typename BlockReduce::TempStorage temp_storage_reduce; + __shared__ uint32_t selected_shared[NUM_WARPS]; + + if (lane_id == 0) { + selected_shared[warp_id] = 0; + } + + // Initialize running total + BlockPrefixCallbackOp prefix_op(0); + + __syncthreads(); + + int offset = bid * vocab_size; + int end = ((vocab_size + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE; + int i_activate = 0; + float thread_offset = 0; + for (int i = tid; i < end; i += BLOCK_SIZE) { + float thread_count = + (i < vocab_size) ? static_cast(sorted_probs[offset + i]) : 0.f; + + BlockScan(temp_storage) + .InclusiveSum(thread_count, thread_offset, prefix_op); + + if (i < candidates_len) { + out_id[bid * candidates_len + i] = sorted_id[offset + i]; + out_val[bid * candidates_len + i] = sorted_probs[offset + i]; } - // Initialize running total - BlockPrefixCallbackOp prefix_op(0); - - __syncthreads(); - - int offset = bid * vocab_size; - int end = ((vocab_size + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE; - int i_activate = 0; - float thread_offset = 0; - for (int i = tid; i < end; i += BLOCK_SIZE) { - float thread_count = (i < vocab_size) - ? static_cast(sorted_probs[offset + i]) - : 0.f; - - BlockScan(temp_storage) - .InclusiveSum(thread_count, thread_offset, prefix_op); - - if (i < candidates_len) { - out_id[bid * candidates_len + i] = sorted_id[offset + i]; - out_val[bid * candidates_len + i] = sorted_probs[offset + i]; - } - - uint32_t activate_mask = - __ballot_sync(FINAL_MASK, topp <= thread_offset); - i_activate = i; - if (activate_mask != 0 || i >= candidates_len) { - if (lane_id == 0) { - atomicAdd(&stop_shared, 1); - selected_shared[warp_id] = activate_mask; - } - } - __syncthreads(); - if (stop_shared > 0) { - break; - } + uint32_t activate_mask = __ballot_sync(FINAL_MASK, topp <= thread_offset); + i_activate = i; + if (activate_mask != 0 || i >= candidates_len) { + if (lane_id == 0) { + atomicAdd(&stop_shared, 1); + selected_shared[warp_id] = activate_mask; + } } __syncthreads(); - bool skip = (selected_shared[warp_id] > 0) ? false : true; - for (int i = 0; i < warp_id; i++) { - if (selected_shared[i] != 0) { - // If the previous has stopped, skip the current warp - skip = true; - } + if (stop_shared > 0) { + break; } - if (!skip) { - int active_lane_id = - WARP_SIZE - __popc(selected_shared[warp_id]); // first not 0 - if (lane_id == active_lane_id) { - actual_candidates_lens[bid] = i_activate + 1; - } + } + __syncthreads(); + bool skip = (selected_shared[warp_id] > 0) ? false : true; + for (int i = 0; i < warp_id; i++) { + if (selected_shared[i] != 0) { + // If the previous has stopped, skip the current warp + skip = true; } - __syncthreads(); - if (tid == 0) { - // printf("actual_candidates_lens[%d] %d\n", bid, - // actual_candidates_lens[bid]); - if (actual_candidates_lens[bid] == 0) { - actual_candidates_lens[bid] = candidates_len; - } + } + if (!skip) { + int active_lane_id = + WARP_SIZE - __popc(selected_shared[warp_id]); // first not 0 + if (lane_id == active_lane_id) { + actual_candidates_lens[bid] = i_activate + 1; } + } + __syncthreads(); + if (tid == 0) { + // printf("actual_candidates_lens[%d] %d\n", bid, + // actual_candidates_lens[bid]); + if (actual_candidates_lens[bid] == 0) { + actual_candidates_lens[bid] = candidates_len; + } + } } template struct Pair { - __device__ __forceinline__ Pair() {} - __device__ __forceinline__ Pair(T value, int id) : v(value), id(id) {} - - __device__ __forceinline__ void set(T value, int id) { - this->v = value; - this->id = id; - } - - __device__ __forceinline__ void operator=(const Pair& in) { - v = in.v; - id = in.id; - } - - __device__ __forceinline__ bool operator<(const T value) const { - return (static_cast(v) < static_cast(value)); - } - - __device__ __forceinline__ bool operator>(const T value) const { - return (static_cast(v) > static_cast(value)); - } - __device__ __forceinline__ bool operator<(const Pair& in) const { - return (static_cast(v) < static_cast(in.v)) || - ((static_cast(v) == static_cast(in.v)) && - (id > in.id)); - } - - __device__ __forceinline__ bool operator>(const Pair& in) const { - return (static_cast(v) > static_cast(in.v)) || - ((static_cast(v) == static_cast(in.v)) && - (id < in.id)); - } - - T v; - int id; + __device__ __forceinline__ Pair() {} + __device__ __forceinline__ Pair(T value, int id) : v(value), id(id) {} + + __device__ __forceinline__ void set(T value, int id) { + this->v = value; + this->id = id; + } + + __device__ __forceinline__ void operator=(const Pair& in) { + v = in.v; + id = in.id; + } + + __device__ __forceinline__ bool operator<(const T value) const { + return (static_cast(v) < static_cast(value)); + } + + __device__ __forceinline__ bool operator>(const T value) const { + return (static_cast(v) > static_cast(value)); + } + __device__ __forceinline__ bool operator<(const Pair& in) const { + return (static_cast(v) < static_cast(in.v)) || + ((static_cast(v) == static_cast(in.v)) && + (id > in.id)); + } + + __device__ __forceinline__ bool operator>(const Pair& in) const { + return (static_cast(v) > static_cast(in.v)) || + ((static_cast(v) == static_cast(in.v)) && + (id < in.id)); + } + + T v; + int id; }; template __device__ __forceinline__ void AddTo(Pair topk[], const Pair& p, int beam_size) { - for (int k = beam_size - 2; k >= 0; k--) { - if (topk[k] < p) { - topk[k + 1] = topk[k]; - } else { - topk[k + 1] = p; - return; - } + for (int k = beam_size - 2; k >= 0; k--) { + if (topk[k] < p) { + topk[k + 1] = topk[k]; + } else { + topk[k + 1] = p; + return; } - topk[0] = p; + } + topk[0] = p; } template __device__ __forceinline__ void GetTopK( Pair topk[], const T* src, int idx, int dim, int beam_size) { - while (idx < dim) { - if (topk[beam_size - 1] < src[idx]) { - Pair tmp(src[idx], idx); - AddTo(topk, tmp, beam_size); - } - idx += BlockSize; + while (idx < dim) { + if (topk[beam_size - 1] < src[idx]) { + Pair tmp(src[idx], idx); + AddTo(topk, tmp, beam_size); } + idx += BlockSize; + } } template @@ -288,15 +286,15 @@ __device__ __forceinline__ void GetTopK(Pair topk[], int dim, const Pair& max, int beam_size) { - while (idx < dim) { - if (topk[beam_size - 1] < src[idx]) { - Pair tmp(src[idx], idx); - if (tmp < max) { - AddTo(topk, tmp, beam_size); - } - } - idx += BlockSize; + while (idx < dim) { + if (topk[beam_size - 1] < src[idx]) { + Pair tmp(src[idx], idx); + if (tmp < max) { + AddTo(topk, tmp, beam_size); + } } + idx += BlockSize; + } } template @@ -309,43 +307,43 @@ __device__ __forceinline__ void ThreadGetTopK(Pair topk[], Pair* max, int dim, const int tid) { - if (*beam > 0) { - int length = (*beam) < beam_size ? *beam : beam_size; - if (*firstStep) { - *firstStep = false; - GetTopK(topk, src, tid, dim, length); + if (*beam > 0) { + int length = (*beam) < beam_size ? *beam : beam_size; + if (*firstStep) { + *firstStep = false; + GetTopK(topk, src, tid, dim, length); + } else { + for (int k = 0; k < MaxLength; k++) { + if (k < MaxLength - (*beam)) { + topk[k] = topk[k + *beam]; } else { - for (int k = 0; k < MaxLength; k++) { - if (k < MaxLength - (*beam)) { - topk[k] = topk[k + *beam]; - } else { - topk[k].set(std::numeric_limits::min(), -1); - } - } - if (!(*is_empty)) { - GetTopK( - topk + MaxLength - *beam, src, tid, dim, *max, length); - } + topk[k].set(std::numeric_limits::min(), -1); } - - *max = topk[MaxLength - 1]; - if ((*max).id == -1) *is_empty = true; - *beam = 0; + } + if (!(*is_empty)) { + GetTopK( + topk + MaxLength - *beam, src, tid, dim, *max, length); + } } + + *max = topk[MaxLength - 1]; + if ((*max).id == -1) *is_empty = true; + *beam = 0; + } } template __forceinline__ __device__ Pair WarpReduce(Pair input) { #pragma unroll - for (int offset = 16; offset > 0; offset >>= 1) { - T tmp_val = CudaShuffleDownSync(FINAL_MASK, input.v, offset); - int tmp_id = CudaShuffleDownSync(FINAL_MASK, input.id, offset); - if (static_cast(input.v) < static_cast(tmp_val)) { - input.v = tmp_val; - input.id = tmp_id; - } + for (int offset = 16; offset > 0; offset >>= 1) { + T tmp_val = CudaShuffleDownSync(FINAL_MASK, input.v, offset); + int tmp_id = CudaShuffleDownSync(FINAL_MASK, input.id, offset); + if (static_cast(input.v) < static_cast(tmp_val)) { + input.v = tmp_val; + input.id = tmp_id; } - return input; + } + return input; } template @@ -358,52 +356,51 @@ __device__ __forceinline__ void BlockReduce(Pair shared_max[], const int tid, const int wid, const int lane) { - while (true) { - __syncthreads(); - Pair input_now = topk[0]; - input_now = WarpReduce(input_now); + while (true) { + __syncthreads(); + Pair input_now = topk[0]; + input_now = WarpReduce(input_now); - if (lane == 0) { - shared_max[wid] = input_now; - } - __syncthreads(); - input_now = (tid < BlockSize / 32) - ? shared_max[lane] - : Pair(std::numeric_limits::min(), -1); - if (wid == 0) { - input_now = WarpReduce(input_now); - if (lane == 0) shared_max[0] = input_now; - } - __syncthreads(); - if (tid == 0) { - beam_max[*count] = shared_max[0]; - (*count)++; - } - int tid_max = shared_max[0].id % BlockSize; - if (tid == tid_max) { - (*beam)++; - } - if (--(*k) == 0) break; - __syncthreads(); + if (lane == 0) { + shared_max[wid] = input_now; + } + __syncthreads(); + input_now = (tid < BlockSize / 32) + ? shared_max[lane] + : Pair(std::numeric_limits::min(), -1); + if (wid == 0) { + input_now = WarpReduce(input_now); + if (lane == 0) shared_max[0] = input_now; + } + __syncthreads(); + if (tid == 0) { + beam_max[*count] = shared_max[0]; + (*count)++; + } + int tid_max = shared_max[0].id % BlockSize; + if (tid == tid_max) { + (*beam)++; + } + if (--(*k) == 0) break; + __syncthreads(); - if (tid == tid_max) { - if (*beam < MaxLength) { - topk[0] = topk[*beam]; - } - } + if (tid == tid_max) { + if (*beam < MaxLength) { + topk[0] = topk[*beam]; + } + } - if (MaxLength < 5) { - if (*beam >= MaxLength) break; - } else { - unsigned mask = 0u; - mask = __ballot_sync(FINAL_MASK, true); - if (tid_max / 32 == wid) { - if (__shfl_down_sync(FINAL_MASK, *beam, tid_max % 32, 32) == - MaxLength) - break; - } - } + if (MaxLength < 5) { + if (*beam >= MaxLength) break; + } else { + unsigned mask = 0u; + mask = __ballot_sync(FINAL_MASK, true); + if (tid_max / 32 == wid) { + if (__shfl_down_sync(FINAL_MASK, *beam, tid_max % 32, 32) == MaxLength) + break; + } } + } } template @@ -417,70 +414,66 @@ __global__ void KeMatrixTopPBeamTopKFt( int vocab_size, const int max_cadidate_len, const int max_seq_len) { - const int tid = threadIdx.x; - const int wid = tid / 32; - const int lane = tid % 32; - const int token_id = blockIdx.x; - const int ori_token_id = token_id + output_padding_offset[token_id]; - const int bid = ori_token_id / max_seq_len; - - int top_num = TopPBeamTopK; - float top_p_value = static_cast(top_ps[bid]); - - __shared__ Pair shared_max[BlockSize / 32]; - __shared__ Pair beam_max[TopPBeamTopK]; - - Pair topk[MaxLength]; - int beam = MaxLength; - Pair max; - bool is_empty = false; - bool firststep = true; - __shared__ int count; - - if (tid == 0) { - count = 0; - } - - for (int j = 0; j < MaxLength; j++) { - topk[j].set(std::numeric_limits::min(), -1); - } - - while (top_num) { - ThreadGetTopK(topk, - &beam, - TopPBeamTopK, - src + token_id * vocab_size, - &firststep, - &is_empty, - &max, - vocab_size, - tid); - BlockReduce(shared_max, - topk, - beam_max, - &beam, - &top_num, - &count, - tid, - wid, - lane); - } - if (tid == 0) { - float sum_prob = 0.0f; - bool flag = false; - for (int i = 0; i < TopPBeamTopK; i++) { - out_id[token_id * max_cadidate_len + i] = - static_cast(beam_max[i].id); - out_val[token_id * max_cadidate_len + i] = beam_max[i].v; - float val = static_cast(beam_max[i].v); - sum_prob += val; - - if (sum_prob >= top_p_value) { - actual_candidates_lens[token_id] = i + 1; - break; - } - } + const int tid = threadIdx.x; + const int wid = tid / 32; + const int lane = tid % 32; + const int token_id = blockIdx.x; + const int ori_token_id = token_id + output_padding_offset[token_id]; + const int bid = ori_token_id / max_seq_len; + + int top_num = TopPBeamTopK; + float top_p_value = static_cast(top_ps[bid]); + + __shared__ Pair shared_max[BlockSize / 32]; + __shared__ Pair beam_max[TopPBeamTopK]; + + Pair topk[MaxLength]; + int beam = MaxLength; + Pair max; + bool is_empty = false; + bool firststep = true; + __shared__ int count; + + if (tid == 0) { + count = 0; + } + + for (int j = 0; j < MaxLength; j++) { + topk[j].set(std::numeric_limits::min(), -1); + } + + while (top_num) { + ThreadGetTopK(topk, + &beam, + TopPBeamTopK, + src + token_id * vocab_size, + &firststep, + &is_empty, + &max, + vocab_size, + tid); + BlockReduce( + shared_max, topk, beam_max, &beam, &top_num, &count, tid, wid, lane); + } + if (tid == 0) { + float sum_prob = 0.0f; + bool flag = false; + for (int i = 0; i < TopPBeamTopK; i++) { + out_id[token_id * max_cadidate_len + i] = + static_cast(beam_max[i].id); + out_val[token_id * max_cadidate_len + i] = beam_max[i].v; + float val = static_cast(beam_max[i].v); + sum_prob += val; + + if (sum_prob >= top_p_value) { + actual_candidates_lens[token_id] = i + 1; + break; + } } + } + if (top_p_value == 1.0 && actual_candidates_lens[token_id] == 0) { + actual_candidates_lens[token_id] = max_cadidate_len; + } } template @@ -495,30 +488,28 @@ void DispatchTopK(const T* src, const int cadidate_len, const int max_seq_len, const cudaStream_t& stream) { - int BlockSize = GetBlockSize(vocab_size); - switch (cadidate_len) { - FIXED_TOPK(switch (BlockSize) { - FIXED_BLOCK_DIM( - KeMatrixTopPBeamTopKFt - <<>>( - src, - top_ps, - output_padding_offset, - out_id, - out_val, - actual_candidates_lens_data, - vocab_size, - cadidate_len, - max_seq_len)); - default: - PD_THROW( - "Invalid max_candidate_len. Please set a value in [1,10] (step=1) " - "or [10,100] (step=10)." - ); - }); - default: - PD_THROW("the input topk is not implemented."); - } + int BlockSize = GetBlockSize(vocab_size); + switch (cadidate_len) { + FIXED_TOPK(switch (BlockSize) { + FIXED_BLOCK_DIM( + KeMatrixTopPBeamTopKFt + <<>>(src, + top_ps, + output_padding_offset, + out_id, + out_val, + actual_candidates_lens_data, + vocab_size, + cadidate_len, + max_seq_len)); + default: + PD_THROW( + "Invalid max_candidate_len. Please set a value in [1,10] (step=1) " + "or [10,100] (step=10)."); + }); + default: + PD_THROW("the input topk is not implemented."); + } } template @@ -528,38 +519,38 @@ std::vector LaunchTopPCandidates( const paddle::Tensor& output_padding_offset, const int candidates_len, const int max_seq_len) { - typedef PDTraits traits_; - typedef typename traits_::DataType DataType_; - typedef typename traits_::data_t data_t; - - std::vector input_shape = probs.shape(); - const int token_num = input_shape[0]; - const int vocab_size = input_shape[1]; - - auto verify_scores = - paddle::full({token_num, candidates_len}, 0, D, probs.place()); - auto verify_tokens = paddle::full( - {token_num, candidates_len}, 0, paddle::DataType::INT64, probs.place()); - auto actual_candidate_lens = - paddle::full({token_num}, 0, paddle::DataType::INT32, probs.place()); - - auto stream = probs.stream(); - - constexpr int TopKMaxLength = 2; - DispatchTopK( - reinterpret_cast(probs.data()), - reinterpret_cast(top_p.data()), - output_padding_offset.data(), - verify_tokens.data(), - reinterpret_cast(verify_scores.data()), - actual_candidate_lens.data(), - vocab_size, - token_num, - candidates_len, - max_seq_len, - stream); - - return {verify_scores, verify_tokens, actual_candidate_lens}; + typedef PDTraits traits_; + typedef typename traits_::DataType DataType_; + typedef typename traits_::data_t data_t; + + std::vector input_shape = probs.shape(); + const int token_num = input_shape[0]; + const int vocab_size = input_shape[1]; + + auto verify_scores = + paddle::full({token_num, candidates_len}, 0, D, probs.place()); + auto verify_tokens = paddle::full( + {token_num, candidates_len}, 0, paddle::DataType::INT64, probs.place()); + auto actual_candidate_lens = + paddle::full({token_num}, 0, paddle::DataType::INT32, probs.place()); + + auto stream = probs.stream(); + + constexpr int TopKMaxLength = 2; + DispatchTopK( + reinterpret_cast(probs.data()), + reinterpret_cast(top_p.data()), + output_padding_offset.data(), + verify_tokens.data(), + reinterpret_cast(verify_scores.data()), + actual_candidate_lens.data(), + vocab_size, + token_num, + candidates_len, + max_seq_len, + stream); + + return {verify_scores, verify_tokens, actual_candidate_lens}; } std::vector DispatchTopPCandidatesWithDtype( @@ -568,37 +559,25 @@ std::vector DispatchTopPCandidatesWithDtype( const paddle::Tensor& output_padding_offset, int candidates_len, int max_seq_len) { - switch (probs.type()) { - case paddle::DataType::BFLOAT16: - return LaunchTopPCandidates( - probs, - top_p, - output_padding_offset, - candidates_len, - max_seq_len); - break; - case paddle::DataType::FLOAT16: - return LaunchTopPCandidates( - probs, - top_p, - output_padding_offset, - candidates_len, - max_seq_len); - break; - case paddle::DataType::FLOAT32: - return LaunchTopPCandidates( - probs, - top_p, - output_padding_offset, - candidates_len, - max_seq_len); - break; - default: - PD_THROW( - "NOT supported data type. " - "Only bfloat16, float16 and float32 are supported. "); - break; - } + switch (probs.type()) { + case paddle::DataType::BFLOAT16: + return LaunchTopPCandidates( + probs, top_p, output_padding_offset, candidates_len, max_seq_len); + break; + case paddle::DataType::FLOAT16: + return LaunchTopPCandidates( + probs, top_p, output_padding_offset, candidates_len, max_seq_len); + break; + case paddle::DataType::FLOAT32: + return LaunchTopPCandidates( + probs, top_p, output_padding_offset, candidates_len, max_seq_len); + break; + default: + PD_THROW( + "NOT supported data type. " + "Only bfloat16, float16 and float32 are supported. "); + break; + } } std::vector TopPCandidates( @@ -607,8 +586,8 @@ std::vector TopPCandidates( const paddle::Tensor& output_padding_offset, int candidates_len, int max_seq_len) { - return DispatchTopPCandidatesWithDtype( - probs, top_p, output_padding_offset, candidates_len, max_seq_len); + return DispatchTopPCandidatesWithDtype( + probs, top_p, output_padding_offset, candidates_len, max_seq_len); } std::vector> TopPCandidatesInferShape( @@ -616,17 +595,17 @@ std::vector> TopPCandidatesInferShape( const std::vector& top_p_shape, const std::vector& output_padding_offset_shape, int max_candidates_len) { - int token_num = probs_shape[0]; - return {{token_num, max_candidates_len}, - {token_num, max_candidates_len}, - {token_num}}; + int token_num = probs_shape[0]; + return {{token_num, max_candidates_len}, + {token_num, max_candidates_len}, + {token_num}}; } std::vector TopPCandidatesInferDtype( const paddle::DataType& probs_dtype, const paddle::DataType& top_p_dtype, const paddle::DataType& output_padding_offset_dtype) { - return {probs_dtype, paddle::DataType::INT64, paddle::DataType::INT32}; + return {probs_dtype, paddle::DataType::INT64, paddle::DataType::INT32}; } PD_BUILD_STATIC_OP(top_p_candidates) From 55598eceafba00fb67d346bca8b20b9aff384f23 Mon Sep 17 00:00:00 2001 From: RAM Date: Fri, 5 Dec 2025 20:01:33 +0800 Subject: [PATCH 04/13] [RL] Support Rollout Routing Replay (#5321) * [RL] Support Rollout Routing Replay * add routing indices cache * fix config bug and moe forward bug * R3 Support GLM * support eb4.5 * fix merge bug * Apply suggestion from @Copilot Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Apply suggestion from @Copilot Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Apply suggestion from @Copilot Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Apply suggestion from @Copilot Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * add routing replay ci * support glm topk * support orther top_k * fix ci bug * pre-commit * only support chatcmpl --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> Co-authored-by: Yuanle Liu --- fastdeploy/config.py | 27 ++ fastdeploy/engine/args_utils.py | 22 ++ fastdeploy/engine/engine.py | 1 + fastdeploy/model_executor/forward_meta.py | 2 + .../backends/dcu/fused_moe_triton_backends.py | 5 + .../gcu/moe/fused_moe_method_gcu_backend.py | 6 + .../moe/fused_moe_triton_metax_backend.py | 6 + .../layers/moe/fused_moe_backend_base.py | 11 +- .../layers/moe/fused_moe_cutlass_backend.py | 17 + .../layers/moe/fused_moe_deepgemm_backend.py | 16 + .../layers/moe/fused_moe_marlin_backend.py | 6 + .../layers/moe/fused_moe_triton_backend.py | 18 + .../layers/moe/fused_moe_wint2_backend.py | 10 + fastdeploy/model_executor/layers/moe/moe.py | 60 ++- .../layers/moe/routing_indices_cache.py | 346 ++++++++++++++++++ fastdeploy/model_executor/models/glm4_moe.py | 7 +- fastdeploy/rl/rollout_config.py | 2 + fastdeploy/worker/gpu_model_runner.py | 32 ++ fastdeploy/worker/worker_process.py | 10 + tests/distributed/chunked_moe.py | 4 +- tests/e2e/test_EB_Lite_serving.py | 2 + tests/layers/test_fusedmoe.py | 2 + tests/layers/test_w4a8_moe.py | 2 + tests/layers/test_w4afp8_moe.py | 2 + 24 files changed, 592 insertions(+), 24 deletions(-) create mode 100644 fastdeploy/model_executor/layers/moe/routing_indices_cache.py diff --git a/fastdeploy/config.py b/fastdeploy/config.py index 58e7c4f3144..f1eb23852f8 100644 --- a/fastdeploy/config.py +++ b/fastdeploy/config.py @@ -1484,6 +1484,31 @@ def __str__(self) -> str: return json.dumps({key: value for key, value in self.__dict__.items()}) +class RoutingReplayConfig: + """Configuration for Routing Replay used in RL training""" + + def __init__(self, args) -> None: + self.enable_routing_replay: bool = False + self.routing_store_type: str = "local" + + # Local routing store + self.local_store_dir: str = "./routing_replay_output" + + # RDMA routing store + # TODO: Add RDMA routing store configuration attributes here when the feature is implemented. + + if args is not None: + for key, value in args.items(): + if hasattr(self, key) and value != "None": + setattr(self, key, value) + + def to_json_string(self): + """ + Convert routing replay config to json string. + """ + return json.dumps({key: value for key, value in self.__dict__.items()}) + + class FDConfig: """ The configuration class which contains all fastdeploy-related configuration. This @@ -1517,6 +1542,7 @@ def __init__( early_stop_config: Optional[Dict[str, Any]] = None, tool_parser: str = None, test_mode=False, + routing_replay_config: Optional[RoutingReplayConfig] = None, ): self.model_config: ModelConfig = model_config # type: ignore self.cache_config: CacheConfig = cache_config # type: ignore @@ -1533,6 +1559,7 @@ def __init__( self.plas_attention_config: Optional[PlasAttentionConfig] = plas_attention_config self.structured_outputs_config: StructuredOutputsConfig = structured_outputs_config self.router_config: RouterConfig = router_config + self.routing_replay_config = routing_replay_config # Initialize cuda graph capture list max_capture_shape = self.scheduler_config.max_num_seqs diff --git a/fastdeploy/engine/args_utils.py b/fastdeploy/engine/args_utils.py index 1eaf535498a..d2d7c6f908a 100644 --- a/fastdeploy/engine/args_utils.py +++ b/fastdeploy/engine/args_utils.py @@ -35,6 +35,7 @@ PlasAttentionConfig, PoolerConfig, RouterConfig, + RoutingReplayConfig, RunnerOption, SpeculativeConfig, StructuredOutputsConfig, @@ -491,6 +492,11 @@ class EngineArgs: Configuration for eplb. """ + routing_replay_config: Optional[Dict[str, Any]] = None + """ + Flag to rollout routing replay(r3) + """ + def __post_init__(self): """ Post-initialization processing to set default tokenizer if not provided. @@ -882,6 +888,12 @@ def add_cli_args(parser: FlexibleArgumentParser) -> FlexibleArgumentParser: default=EngineArgs.eplb_config, help="Config of eplb.", ) + parallel_group.add_argument( + "--routing-replay-config", + type=json.loads, + default=EngineArgs.routing_replay_config, + help="Flag of rollout routing replay(r3).", + ) parallel_group.add_argument( "--enable-chunked-moe", action="store_true", @@ -1235,6 +1247,14 @@ def create_eplb_config(self) -> EPLBConfig: eplb_args["enable_eplb"] = self.enable_eplb return EPLBConfig(eplb_args) + def create_routing_repaly_config(self) -> RoutingReplayConfig: + """ """ + routing_replay_args = asdict(self) + if self.routing_replay_config is not None: + for k, v in self.routing_replay_config.items(): + routing_replay_args[k] = v + return RoutingReplayConfig(routing_replay_args) + def create_engine_config(self, port_availability_check=True) -> FDConfig: """ Create and return a Config object based on the current settings. @@ -1278,6 +1298,7 @@ def create_engine_config(self, port_availability_check=True) -> FDConfig: graph_opt_cfg = self.create_graph_optimization_config() plas_attention_config = self.create_plas_attention_config() eplb_cfg = self.create_eplb_config() + routing_replay_config = self.create_routing_repaly_config() router_config = RouterConfig(all_dict) early_stop_cfg = self.create_early_stop_config() @@ -1310,4 +1331,5 @@ def create_engine_config(self, port_availability_check=True) -> FDConfig: graph_opt_config=graph_opt_cfg, plas_attention_config=plas_attention_config, early_stop_config=early_stop_cfg, + routing_replay_config=routing_replay_config, ) diff --git a/fastdeploy/engine/engine.py b/fastdeploy/engine/engine.py index 4a493843df7..fadf954679b 100644 --- a/fastdeploy/engine/engine.py +++ b/fastdeploy/engine/engine.py @@ -568,6 +568,7 @@ def _start_worker_service(self): f" --logprobs_mode {self.cfg.model_config.logprobs_mode}" f" --max_logprobs {self.cfg.model_config.max_logprobs}" f" --eplb_config '{self.cfg.eplb_config.to_json_string()}'" + f" --routing_replay_config '{self.cfg.routing_replay_config.to_json_string()}'" ) if self.cfg.structured_outputs_config.logits_processors is not None: arguments += f" --logits-processors {' '.join(self.cfg.structured_outputs_config.logits_processors)}" diff --git a/fastdeploy/model_executor/forward_meta.py b/fastdeploy/model_executor/forward_meta.py index 4e9df0d3ce3..787ec77c0eb 100644 --- a/fastdeploy/model_executor/forward_meta.py +++ b/fastdeploy/model_executor/forward_meta.py @@ -142,6 +142,8 @@ class ForwardMeta: caches: Optional[list[paddle.Tensor]] = None # Flag of profile run is_dummy_or_profile_run: bool = False + # Routing Replay table buffer + routing_replay_table: Optional[paddle.Tensor] = None # chunked MoE related moe_num_chunk: int = 1 diff --git a/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py b/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py index 918450c74f1..192c0b8833a 100644 --- a/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py +++ b/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -101,6 +103,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -117,6 +120,8 @@ def apply( scores += layer.gate_correction_bias topk_weights, topk_ids = paddle.topk(scores, k=top_k, axis=-1, sorted=False) topk_weights = topk_weights / topk_weights.sum(axis=-1, keepdim=True) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) intermediate_cache1 = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], diff --git a/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py b/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py index e67dd6dbdaf..2260d7caf7b 100644 --- a/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py +++ b/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py @@ -16,6 +16,7 @@ import multiprocessing import os +from typing import Callable import numpy as np import paddle @@ -182,6 +183,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle gcu compute Fused MoE. @@ -194,6 +196,7 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -205,6 +208,7 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -216,6 +220,7 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -381,6 +386,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle gcu compute Fused MoE. diff --git a/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py b/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py index 7b61d58b6f5..fbbfac277b8 100644 --- a/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py +++ b/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -245,6 +247,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -274,6 +277,9 @@ def apply( True, # apply_norm_weight False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py b/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py index b34291a96f4..a8bd70465ea 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py @@ -15,6 +15,7 @@ """ from abc import abstractmethod +from typing import Callable import paddle from paddle import nn @@ -163,6 +164,7 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -175,6 +177,7 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -187,6 +190,7 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -198,6 +202,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -207,13 +212,13 @@ def apply( if layer.fd_config.model_config.moe_phase.phase == "prefill": if layer.fd_config.scheduler_config.splitwise_role == "mixed" and is_moe_start_layer: self.ep_prefill_runner.clean_low_latency_buffer() - return self.apply_ep_prefill(layer, x, gate) + return self.apply_ep_prefill(layer, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) else: if layer.fd_config.scheduler_config.splitwise_role == "mixed" and is_moe_start_layer: self.ep_decoder_runner.clean_low_latency_buffer() - return self.apply_ep_decode(layer, x, gate) + return self.apply_ep_decode(layer, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) else: - return self.apply_tp(layer, x, gate) + return self.apply_tp(layer, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) class UnquantizedFusedMoEMethod(MoEMethodBase): diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py index e45ad63b19c..c3dbfc9ba5f 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn from paddle.nn.quant import weight_quantize @@ -132,6 +134,7 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -148,8 +151,13 @@ def apply_ep_prefill( handle, event, ) = self.ep_prefill_runner.dispatch(x, topk_idx, topk_weights) + + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + if self.ep_prefill_runner.ep_engine.async_finish: event.current_stream_wait() + token_all_num = sum(recv_num_tokens_per_expert_list) # 3. Compute ffn @@ -217,6 +225,7 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -225,6 +234,10 @@ def apply_ep_decode( estimate_total_token_nums = gate_out.shape[0] * layer.top_k # 1. Select topk experts and weights topk_idx, topk_weights = self.ep_decoder_runner.moe_select(layer, gate_out) + + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + expertwise_scale = None if hasattr(layer, "up_gate_proj_in_scale_all_experts"): # only use in w4a8 expertwise_scale = getattr(layer, "up_gate_proj_in_scale_all_experts", None) @@ -269,6 +282,7 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -369,6 +383,9 @@ def apply_tp( if hasattr(layer, "up_gate_proj_in_scale"): dequant_scale = None + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + if not layer.with_bias and self.moe_quant_type != "w4a8" and self.moe_quant_type != "w4afp8": # only w4a8 need expert_idx_per_token # Other need not this tensor, so we make it None. diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py index 1245cddcebc..881f9a22c4d 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn from paddle.distributed.communication import deep_ep @@ -139,6 +141,7 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -147,6 +150,10 @@ def apply_ep_prefill( # 1. Select topk experts and weights topk_idx, topk_weights = self.ep_prefill_runner.moe_select(layer, gate_out) + + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + # 2. Dynamic compute blockwise quantization scales x, x_scale_tensor = fastdeploy.model_executor.ops.gpu.per_token_quant( x, self.quant_config.weight_block_size[0] @@ -264,6 +271,7 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -271,6 +279,10 @@ def apply_ep_decode( gate_out = gate(x.cast("float32")) # 1. Select topk experts and weights topk_idx, topk_weights = self.ep_decoder_runner.moe_select(layer, gate_out) + + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + # 2. EP Dispatch permute_input, token_nums_per_expert, handle = self.ep_decoder_runner.dispatch( x, topk_idx, topk_weights, use_fp8=True @@ -335,6 +347,7 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Use DeepGemm compute Fused MoE. @@ -363,6 +376,9 @@ def apply_tp( False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + tmp = count_tokens_per_expert_func(topk_ids, layer.num_experts) recv_x, recv_x_scale = fastdeploy.model_executor.ops.gpu.per_token_quant(x, 128) diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py index 094d3df8f1a..cd836dbaf09 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -239,6 +241,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Marlin compute Fused MoE. @@ -273,6 +276,9 @@ def apply( False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + block_size_m = 64 for m in [8, 16, 32, 48, 64]: diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py index 3c14859375c..2861d96e8d3 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -282,6 +284,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -314,6 +317,10 @@ def apply( True, # apply_norm_weight, False, ) + + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, @@ -664,6 +671,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -724,6 +732,9 @@ def apply( * ceil_div(moe_intermediate_size * 2, config["BLOCK_SIZE_N"]), ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, @@ -953,6 +964,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -974,6 +986,9 @@ def apply( False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, @@ -1466,6 +1481,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -1488,6 +1504,8 @@ def apply( True, # apply_norm_weight False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) config = { "BLOCK_SIZE_M": 64, diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py index f75e36bcbdd..3c548ba57c8 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -261,6 +263,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Use Wint2 Triton Fusedmoe compute Fused MoE. @@ -288,6 +291,9 @@ def apply( topk_only_mode=False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + ffn_out = fastdeploy.model_executor.ops.gpu.moe_expert_ffn_wint2( permute_input, token_nums_per_expert, @@ -328,6 +334,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Use Wint2 Triton Fusedmoe compute Fused MoE. @@ -343,6 +350,9 @@ def apply( False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + num_tokens, K = x.shape E, _, N = layer.up_gate_proj_weight.shape M = num_tokens diff --git a/fastdeploy/model_executor/layers/moe/moe.py b/fastdeploy/model_executor/layers/moe/moe.py index 743e05031f6..5b1be52d183 100644 --- a/fastdeploy/model_executor/layers/moe/moe.py +++ b/fastdeploy/model_executor/layers/moe/moe.py @@ -14,7 +14,8 @@ # limitations under the License. """ -from typing import Optional +from functools import partial +from typing import Callable, Optional import paddle from paddle import nn @@ -26,6 +27,9 @@ tensor_model_parallel_all_reduce_custom, ) from fastdeploy.model_executor.forward_meta import ForwardMeta +from fastdeploy.model_executor.layers.moe.routing_indices_cache import ( + save_routing_to_buffer, +) from fastdeploy.model_executor.layers.utils import get_tensor from fastdeploy.model_executor.utils import h2d_copy, slice_fn from fastdeploy.platforms import current_platform @@ -226,7 +230,7 @@ def __init__( self.is_rearrange = False if self.ep_size > 1: self.quant_method.init_ep(self) - + self.enable_routing_replay = fd_config.routing_replay_config.enable_routing_replay # Merge normal and RL build model if gate_correction_bias is not None: self.gate_correction_bias = gate_correction_bias @@ -600,7 +604,7 @@ def load_state_dict(self, state_dict, is_rearrange: bool = False): else: self.quant_method.process_loaded_weights(self, state_dict) - def forward_split_allgather(self, x: paddle.Tensor, gate: nn.Layer): + def forward_split_allgather(self, x: paddle.Tensor, gate: nn.Layer, topk_ids_hookfunc: Callable = None): """ Forward split allgather function. """ @@ -615,14 +619,14 @@ def forward_split_allgather(self, x: paddle.Tensor, gate: nn.Layer): if end_offset > token_num: end_offset = token_num part_x[: (end_offset - start_offset), :] = x[start_offset:end_offset, :] - out = self.quant_method.apply(self, part_x, gate) + out = self.quant_method.apply(self, part_x, gate, topk_ids_hookfunc=topk_ids_hookfunc) multi_outs = paddle.zeros([token_num_per_rank * self.attn_tp_size, x.shape[1]], dtype=x.dtype) paddle.distributed.all_gather(multi_outs, out, self.tp_group) out = multi_outs[:token_num, :] return out - def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): + def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta = None): """ Defines the forward computation of the moe layer. @@ -633,6 +637,21 @@ def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): Tensor: Output tensor.s """ + topk_ids_hookfunc = None + if self.enable_routing_replay: + if forward_meta is not None: # forward_meta is None when execute empty_input_forward + topk_ids_hookfunc = partial( + save_routing_to_buffer, + routing_replay_table=forward_meta.routing_replay_table, + batch_id_per_token=forward_meta.batch_id_per_token, + seq_lens_decoder=forward_meta.seq_lens_decoder, + cu_seqlens_q=forward_meta.cu_seqlens_q, + layer_idx=self.layer_idx, + tp_size=self.fd_config.parallel_config.tensor_parallel_size, + ep_size=self.fd_config.parallel_config.expert_parallel_size, + tp_group=self.fd_config.parallel_config.tp_group, + ) + token_num = x.shape[0] if ( self.ep_size > 1 @@ -640,11 +659,16 @@ def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): and (not self.fd_config.parallel_config.use_sequence_parallel_moe) and token_num >= self.attn_tp_size ): - out = self.forward_split_allgather(x, gate) + out = self.forward_split_allgather(x, gate, topk_ids_hookfunc=topk_ids_hookfunc) elif self.fd_config.parallel_config.use_ep and self.fd_config.parallel_config.enable_chunked_moe: - out = self.forward_chunked_moe(x, gate, forward_meta) + out = self.forward_chunked_moe( + x, + gate, + forward_meta, + topk_ids_hookfunc=topk_ids_hookfunc, + ) else: - out = self.forward_normal(x, gate) + out = self.forward_normal(x, gate, forward_meta, topk_ids_hookfunc=topk_ids_hookfunc) if self.reduce_results and self.tp_size > 1: if current_platform.is_intel_hpu(): @@ -653,7 +677,9 @@ def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): out = tensor_model_parallel_all_reduce(out, self.tp_group) return out - def forward_chunked_moe(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): + def forward_chunked_moe( + self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta, topk_ids_hookfunc: Callable = None + ): """ Split input to multi chunk to reduce the memory usage of moe. @@ -677,21 +703,25 @@ def forward_chunked_moe(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: Fo for i in range(forward_meta.max_moe_num_chunk): if i < forward_meta.moe_num_chunk: - out_split_list[i] = self.quant_method.apply(self, x_split_list[i], gate) + out_split_list[i] = self.quant_method.apply( + self, x_split_list[i], gate, topk_ids_hookfunc=topk_ids_hookfunc + ) else: # just need to use real data to infer max_moe_num_chunk times. - self.quant_method.apply(self, fake_x, gate) + self.quant_method.apply(self, fake_x, gate, topk_ids_hookfunc=topk_ids_hookfunc) out = paddle.concat(out_split_list, axis=0) else: # when only one chunk, just need to use real data to infer once. - out = self.quant_method.apply(self, x, gate) + out = self.quant_method.apply(self, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) for i in range(forward_meta.max_moe_num_chunk - 1): - self.quant_method.apply(self, fake_x, gate) + self.quant_method.apply(self, fake_x, gate, topk_ids_hookfunc=topk_ids_hookfunc) return out - def forward_normal(self, x: paddle.Tensor, gate: nn.Layer): + def forward_normal( + self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta, topk_ids_hookfunc: Callable = None + ): """ Normal mode of forward. @@ -702,5 +732,5 @@ def forward_normal(self, x: paddle.Tensor, gate: nn.Layer): Tensor: Output tensor.s """ - out = self.quant_method.apply(self, x, gate) + out = self.quant_method.apply(self, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) return out diff --git a/fastdeploy/model_executor/layers/moe/routing_indices_cache.py b/fastdeploy/model_executor/layers/moe/routing_indices_cache.py new file mode 100644 index 00000000000..e95a3d8569f --- /dev/null +++ b/fastdeploy/model_executor/layers/moe/routing_indices_cache.py @@ -0,0 +1,346 @@ +""" +# Copyright (c) 2025 PaddlePaddle Authors. All Rights Reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +""" + +import copy +import os +import shutil +from abc import ABC, abstractmethod +from typing import Dict, List, Optional + +import paddle +import paddle.distributed as dist +import triton +import triton.language as tl + +from fastdeploy.config import FDConfig + + +@triton.jit +def _save_routing_kernel( + ROUTING_REPLAY_TABLE_PTR, + TOPK_IDS_PTR, + BATCH_ID_PER_TOKEN_PTR, + CU_SEQLENS_Q_PTR, + SEQ_LENS_DECODER_PTR, + LAYER_IDX, + TOKEN_NUM, + TOP_K, + NUM_HIDDEN_LAYERS, + MAX_MODEL_LEN, + BLOCK_SIZE_M: tl.constexpr, + BLOCK_SIZE_K: tl.constexpr, +): + pid_m = tl.program_id(axis=0) + + token_offsets = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) + token_mask = token_offsets < TOKEN_NUM + + k_offsets = tl.arange(0, BLOCK_SIZE_K) + + k_mask = k_offsets < TOP_K + + topk_ids_ptrs = TOPK_IDS_PTR + token_offsets[:, None] * TOP_K + k_offsets[None, :] + # [BLOCK_SIZE_M, BLOCK_SIZE_K] + + load_mask = token_mask[:, None] & k_mask[None, :] + topk_vals = tl.load(topk_ids_ptrs, mask=load_mask) + + batch_ids = tl.load(BATCH_ID_PER_TOKEN_PTR + token_offsets, mask=token_mask) + pad_mask = token_mask & (batch_ids != -1) + # [0, 3, 4, 10, 12][0, 0, 0, 0, 2, 2, 2, 2, 2, 2, 3, 3] + # -> [0, 0, 0, 0, 4, 4, 4, 4, 4, 4, 10, 10] + # [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11] - [0, 0, 0, 0, 4, 4, 4, 4, 4, 4, 10, 10] + # -> [0, 1, 2, 3, 0, 1, 2, 3, 4, 5, 0, 1] + start_offsets = tl.load(CU_SEQLENS_Q_PTR + batch_ids, mask=pad_mask) + token_relative_index = token_offsets - start_offsets + + # [BLOCK_SIZE_M] + len_decoder = tl.load(SEQ_LENS_DECODER_PTR + batch_ids, mask=pad_mask) + token_seq_pos = len_decoder + token_relative_index + + STRIDE_BUF_SEQ = NUM_HIDDEN_LAYERS * MAX_MODEL_LEN * TOP_K + STRIDE_BUF_LAYER = MAX_MODEL_LEN * TOP_K + STRIDE_BUF_TOKEN = TOP_K + + # [BLOCK_SIZE_M, BLOCK_SIZE_K] + output_ptrs = ( + ROUTING_REPLAY_TABLE_PTR + + batch_ids[:, None] * STRIDE_BUF_SEQ + + LAYER_IDX * STRIDE_BUF_LAYER + + token_seq_pos[:, None] * STRIDE_BUF_TOKEN + + k_offsets[None, :] + ) + + pos_mask = token_seq_pos < MAX_MODEL_LEN + pos_mask = pos_mask & pad_mask + + # [BLOCK_SIZE_M, BLOCK_SIZE_K] + pos_mask = pos_mask[:, None] & k_mask[None, :] + + final_mask = load_mask & pos_mask + + tl.store(output_ptrs, topk_vals, mask=final_mask) + + +def save_routing_to_buffer( + routing_replay_table: paddle.Tensor, # [max_num_seqs, num_layers, max_len, top_k] + topk_ids: paddle.Tensor, # [token_num, top_k] + batch_id_per_token: paddle.Tensor, # [token_num, 1] + seq_lens_decoder: paddle.Tensor, # [max_num_seqs, 1] + cu_seqlens_q: paddle.Tensor, # [max_num_seqs + 1, 1] + layer_idx: int, + tp_size: int, + ep_size: int, + tp_group: dist.communication.group.Group, +): + if tp_size > 1 and ep_size > 1: + token_num_per_rank = topk_ids.shape[0] + topk_ids_all = paddle.zeros([token_num_per_rank * tp_size, topk_ids.shape[1]], dtype=topk_ids.dtype) + paddle.distributed.all_gather(topk_ids_all, topk_ids, tp_group) + topk_ids = topk_ids_all[: batch_id_per_token.shape[0], :] + + token_num, top_k = topk_ids.shape + max_num_seqs, num_hidden_layers, max_model_len, _ = routing_replay_table.shape + assert token_num > 0 + + assert topk_ids.shape[1] == routing_replay_table.shape[3], (topk_ids.shape[1], routing_replay_table.shape[3]) + assert batch_id_per_token.shape[0] == token_num, (batch_id_per_token.shape[0], token_num) + assert seq_lens_decoder.shape[0] == max_num_seqs, (seq_lens_decoder.shape[0], max_num_seqs) + + BLOCK_SIZE_M = 128 + BLOCK_SIZE_K = triton.next_power_of_2(top_k) # top_k + + grid = (triton.cdiv(token_num, BLOCK_SIZE_M),) + _save_routing_kernel[grid]( + routing_replay_table, + topk_ids, + batch_id_per_token, + cu_seqlens_q, + seq_lens_decoder, + LAYER_IDX=layer_idx, + TOKEN_NUM=token_num, + TOP_K=top_k, + NUM_HIDDEN_LAYERS=num_hidden_layers, + MAX_MODEL_LEN=max_model_len, + BLOCK_SIZE_M=BLOCK_SIZE_M, + BLOCK_SIZE_K=BLOCK_SIZE_K, + ) + + +class RoutingReplayManager: + """Request level routing replay table manager""" + + def __init__( + self, + fd_config: FDConfig, + ): + self.max_num_seqs = fd_config.scheduler_config.max_num_seqs + self.max_model_len = fd_config.model_config.max_model_len + self.num_moe_layers = fd_config.model_config.num_hidden_layers - fd_config.model_config.moe_layer_start_index + + if fd_config.model_config.architectures[0] == "Glm4MoeForCausalLM": + self.moe_top_k = fd_config.model_config.num_experts_per_tok + else: + self.moe_top_k = fd_config.model_config.moe_k + self.tp_rank = fd_config.parallel_config.tensor_parallel_rank + + self.routing_store = get_routing_store(fd_config=fd_config) + self.routing_batch_to_request: Dict[int, str] = {} + self.routing_replay_table = paddle.full( + shape=[self.max_num_seqs, self.num_moe_layers, self.max_model_len, self.moe_top_k], + fill_value=-1, + dtype="int32", + ) + + def register_request(self, batch_id: int, request_id: str): + """ + Register a new request to routing replay table + Args: + batch_id: The batch ID of this request + request_id: The global ID of the request is usually executed by the training process in RL + """ + # Save requests that have been finished for the current slot + if batch_id in self.routing_batch_to_request: + pre_request_id = self._deregister_request(batch_id) + self._put_request_to_store(batch_id, pre_request_id) + # Register the new request + self.routing_batch_to_request[batch_id] = request_id + + def _deregister_request(self, batch_id: int) -> str: + """ + Deregister a request from routing replay table + """ + assert batch_id in self.routing_batch_to_request + return self.routing_batch_to_request.pop(batch_id) + + def _put_request_to_store( + self, + batch_id: int, + request_id: str, + ): + if self.tp_rank == 0: + batch_buffer = self.routing_replay_table[batch_id] + for layer_id in range(self.num_moe_layers): + layer_buffer = batch_buffer[layer_id] + rollout_id = self.split_request_id(request_id) + self.routing_store.put(routing_indices=layer_buffer, rollout_id=rollout_id, layer_idx=layer_id) + + self._clear_table_slot(batch_id) + + def put_table_to_store(self): + """Put the routing table""" + batch_ids = copy.deepcopy(list(self.routing_batch_to_request.keys())) + for batch_id in batch_ids: + request_id = self._deregister_request(batch_id) + self._put_request_to_store(batch_id, request_id) + + def _clear_table_slot(self, batch_id: int): + assert 0 <= batch_id < self.max_num_seqs + self.routing_replay_table[batch_id].fill_(-1) + + def clear_routing_table(self): + """Clear all slots of the routing replay table""" + self.routing_replay_table.fill_(-1) + + def _clear_store(self): + """Clear routing store""" + self.routing_store.clear_store() + + def _clear_request_of_store(self, request_id): + """Clear one request of routing store""" + rollout_id = self.split_request_id(request_id) + for layer_idx in range(self.num_moe_layers): + self.routing_store.clear(rollout_id=rollout_id, layer_idx=layer_idx) + + def get_request_from_store(self, request_id: str) -> List[paddle.Tensor]: + """Get the routing indices of the request from store""" + routing_list = [] + rollout_id = self.split_request_id(request_id) + for layer_idx in range(self.num_moe_layers): + one_layer_routing = self.routing_store.get(rollout_id, layer_idx) + routing_list.append(one_layer_routing) + + return routing_list + + def get_routing_table(self) -> paddle.Tensor: + return self.routing_replay_table + + def split_request_id(self, request_id: str): + """Split the request id to get rollout id""" + chat_type, tmp_str = request_id.split("-", 1) + # NOTE(gongshaotian): only support chatcmpl now + # assert chat_type == "chatcmpl" + reversed_tmp_str = tmp_str[::-1].split("-", 5) + rollout_id = reversed_tmp_str[-1][::-1] + return rollout_id + + +class RoutingStoreBase(ABC): + """Base class for routing store""" + + def __init__(self, fd_config: FDConfig) -> None: + self.fd_config = fd_config + + @abstractmethod + def put(self, routing_indices: paddle.Tensor, rollout_id: str, layer_idx: Optional[int] = None) -> None: + """Put the routing indices into store""" + raise NotImplementedError + + @abstractmethod + def get(self, rollout_id: str, layer_idx: Optional[int] = None) -> paddle.Tensor: + """Get the routing indices from store""" + raise NotImplementedError + + @abstractmethod + def clear(self, rollout_id: str, layer_idx: Optional[int] = None) -> None: + """Clear the routing indices of the request""" + raise NotImplementedError + + @abstractmethod + def clear_store( + self, + ): + """Clear the routing indices store""" + raise NotImplementedError + + +class RoutingStoreLocal(RoutingStoreBase): + """Routing Store using local memory""" + + def __init__(self, fd_config) -> None: + super().__init__(fd_config=fd_config) + self.local_store_dir = fd_config.routing_replay_config.local_store_dir + + def put(self, routing_indices: paddle.Tensor, rollout_id: str, layer_idx: int) -> None: + """Put the routing indices into store""" + dir_path = os.path.join(self.local_store_dir, f"{rollout_id}") + os.makedirs(dir_path, exist_ok=True) + file_path = os.path.join(dir_path, f"layer_{layer_idx}.pdtensor") + paddle.save(routing_indices, file_path) + + def get( + self, + rollout_id: str, + layer_idx: int = None, + ) -> paddle.Tensor: + """Get the routing indices from store""" + dir_path = os.path.join(self.local_store_dir, f"{rollout_id}") + file_path = os.path.join(dir_path, f"layer_{layer_idx}.pdtensor") + assert os.path.exists(file_path), f"File not found: {file_path}" + layer_routing_indices = paddle.load(file_path) + + return layer_routing_indices + + def clear( + self, + rollout_id: str, + layer_idx: int = None, + ) -> None: + """Clear the routing indices of the request""" + dir_path = os.path.join(self.local_store_dir, f"{rollout_id}") + file_path = os.path.join(dir_path, f"layer_{layer_idx}.pdtensor") + assert os.path.exists(file_path), f"File not found: {file_path}" + os.remove(file_path) + + # Delete empty directory + if len(os.listdir(dir_path)) == 0: + os.rmdir(dir_path) + + def clear_store(self): + """Clear the routing indices store""" + if os.path.isdir(self.local_store_dir): + for file_name in os.listdir(self.local_store_dir): + file_path = os.path.join(self.local_store_dir, file_name) + shutil.rmtree(file_path) + + +class RoutingStoreRDMA(RoutingStoreBase): + """Routing Store using RDMA""" + + def __init__(self) -> None: + super().__init__() + + +def get_routing_store(fd_config: FDConfig) -> RoutingStoreBase: + if fd_config.routing_replay_config.routing_store_type == "local": + return RoutingStoreLocal(fd_config=fd_config) + elif fd_config.routing_replay_config.routing_store_type == "rdma": + return RoutingStoreRDMA(fd_config=fd_config) + else: + raise ValueError( + f"Invalid routing store type: '{fd_config.routing_replay_config.routing_store_type}'. " + "Valid types are: 'local', 'rdma'" + ) diff --git a/fastdeploy/model_executor/models/glm4_moe.py b/fastdeploy/model_executor/models/glm4_moe.py index d5ad6e3916b..0cc7c4dae45 100644 --- a/fastdeploy/model_executor/models/glm4_moe.py +++ b/fastdeploy/model_executor/models/glm4_moe.py @@ -161,7 +161,7 @@ def __init__( reduce_results=False, ) - def forward(self, x, forward_meta): + def forward(self, x, forward_meta: ForwardMeta = None): shared_experts_out = self.shared_experts(x) out = self.experts(x, self.gate, forward_meta) out = out + shared_experts_out @@ -306,10 +306,7 @@ def forward( # Fully Connected hidden_states, residual = self.post_attention_layernorm(hidden_states, residual) - hidden_states = self.mlp( - hidden_states, - forward_meta, - ) + hidden_states = self.mlp(hidden_states, forward_meta) return hidden_states, residual diff --git a/fastdeploy/rl/rollout_config.py b/fastdeploy/rl/rollout_config.py index 6bd3c3bcb35..f7ff748fed7 100644 --- a/fastdeploy/rl/rollout_config.py +++ b/fastdeploy/rl/rollout_config.py @@ -65,6 +65,7 @@ def __init__( data_parallel_size: int = 1, num_nextn_predict_layers: int = 0, eplb_config: str = {}, + routing_replay_config: str = None, ): # Required parameters self.model = model_name_or_path @@ -113,6 +114,7 @@ def __init__( self.plas_attention_config = plas_attention_config self.num_nextn_predict_layers = num_nextn_predict_layers self.eplb_config = eplb_config + self.routing_replay_config = routing_replay_config def __str__(self): return "\n".join(f"{k}: {v}" for k, v in self.__dict__.items()) diff --git a/fastdeploy/worker/gpu_model_runner.py b/fastdeploy/worker/gpu_model_runner.py index a91611524ac..94c7a0b3f21 100644 --- a/fastdeploy/worker/gpu_model_runner.py +++ b/fastdeploy/worker/gpu_model_runner.py @@ -45,6 +45,9 @@ from fastdeploy.model_executor.layers.attention.base_attention_backend import ( AttentionBackend, ) +from fastdeploy.model_executor.layers.moe.routing_indices_cache import ( + RoutingReplayManager, +) from fastdeploy.model_executor.layers.rotary_embedding import get_rope, get_rope_3d from fastdeploy.model_executor.layers.sample.meta_data import SamplingMetadata from fastdeploy.model_executor.layers.sample.sampler import Sampler, SpeculativeSampler @@ -202,6 +205,11 @@ def __init__( os.environ["INFERENCE_MSG_QUEUE_ID"] = str(self.parallel_config.engine_worker_queue_port) logger.info(f"queue id is {str(self.parallel_config.engine_worker_queue_port)}") + # Rollout routing replay config + self.routing_replay_manager = None + if self.fd_config.routing_replay_config.enable_routing_replay: + self.routing_replay_manager = RoutingReplayManager(fd_config=self.fd_config) + self.zmq_client = None self.async_output_queue = None if envs.FD_USE_GET_SAVE_OUTPUT_V1: @@ -648,6 +656,7 @@ def insert_tasks_v1(self, req_dicts: List[Request], num_running_requests: int = self.share_inputs["step_seq_lens_decoder"][idx : idx + 1] = 0 self.share_inputs["prompt_lens"][idx : idx + 1] = len(input_ids) self.share_inputs["is_block_step"][idx : idx + 1] = False + self.share_inputs["is_chunk_step"][idx : idx + 1] = prefill_end_index < len(input_ids) self.share_inputs["step_idx"][idx : idx + 1] = ( len(request.output_token_ids) if prefill_end_index >= len(input_ids) else 0 ) @@ -656,6 +665,12 @@ def insert_tasks_v1(self, req_dicts: List[Request], num_running_requests: int = if request.sampling_params is not None and request.sampling_params.prompt_logprobs is not None: self.prompt_logprobs_reqs[request.request_id] = request has_prefill_task = True + + # Routing Replay + if self.fd_config.routing_replay_config.enable_routing_replay: + if prefill_start_index == 0: + self.routing_replay_manager.register_request(batch_id=idx, request_id=request.request_id) + if ( self.fd_config.scheduler_config.splitwise_role == "decode" ): # In PD, we continue to decode after P generate first token @@ -1148,6 +1163,7 @@ def _init_share_inputs(self, max_num_seqs: int): self.share_inputs["bad_tokens_len"] = paddle.full([max_num_seqs], 1, dtype="int64") self.share_inputs["next_tokens"] = paddle.full([max_num_seqs, 1], -1, dtype="int64") self.share_inputs["is_block_step"] = paddle.full([max_num_seqs], False, dtype="bool") + self.share_inputs["is_chunk_step"] = paddle.full([max_num_seqs], False, dtype="bool").cpu() self.share_inputs["encoder_block_lens"] = paddle.full([max_num_seqs], 0, dtype="int32") self.share_inputs["step_block_list"] = paddle.full([max_num_seqs], -1, dtype="int32") self.share_inputs["step_lens"] = paddle.full([1], 0, dtype="int32") @@ -1418,6 +1434,9 @@ def initialize_forward_meta(self, is_dummy_or_profile_run=False): Initialize forward meta, attention meta data and update some config. """ # Initialize forward meta + routing_replay_table = None + if self.routing_replay_manager is not None: + routing_replay_table = self.routing_replay_manager.get_routing_table() self.forward_meta = ForwardMeta( ids_remove_padding=self.share_inputs["ids_remove_padding"], rotary_embs=self.share_inputs["rope_emb"], @@ -1444,6 +1463,7 @@ def initialize_forward_meta(self, is_dummy_or_profile_run=False): kv_batch_ids=self.share_inputs["kv_batch_ids"], kv_tile_ids_per_batch=self.share_inputs["kv_tile_ids_per_batch"], kv_num_blocks_x_cpu=self.share_inputs["kv_num_blocks_x_cpu"], + routing_replay_table=routing_replay_table, ) dist_status = self.collect_distributed_status() @@ -1932,6 +1952,9 @@ def _dummy_run( if int((self.share_inputs["seq_lens_this_time"] > 0).sum()) == 0: break + if self.fd_config.routing_replay_config.enable_routing_replay: + self.routing_replay_manager.clear_routing_table() + def _update_chunked_prefill(self, tasks): """ Update chunked prefill related parameters @@ -2429,6 +2452,15 @@ class at the server level, which is too granular for ModelRunner. self.speculative_config.num_speculative_tokens, ) + # Routing replay + if self.fd_config.routing_replay_config.enable_routing_replay: + if ( + not self.exist_prefill() + and not self.exist_decode() + and self.share_inputs["is_block_step"].sum() == 0 + and self.share_inputs["is_chunk_step"].sum() == 0 + ): + self.routing_replay_manager.put_table_to_store() return None def _pool(self, hidden_states: paddle.Tensor, num_running_requests: int) -> Optional[ModelRunnerOutput]: diff --git a/fastdeploy/worker/worker_process.py b/fastdeploy/worker/worker_process.py index 02d66f4bc53..0c29ce4d757 100644 --- a/fastdeploy/worker/worker_process.py +++ b/fastdeploy/worker/worker_process.py @@ -38,6 +38,7 @@ ModelConfig, ParallelConfig, PlasAttentionConfig, + RoutingReplayConfig, SpeculativeConfig, StructuredOutputsConfig, ) @@ -885,6 +886,13 @@ def parse_args(): help="EPLB Configuration.", ) + parser.add_argument( + "--routing_replay_config", + type=json.loads, + default=None, + help="Configation of Rollout Routing Replay.", + ) + args = parser.parse_args() return args @@ -944,6 +952,7 @@ def initialize_fd_config(args, ranks: int = 1, local_rank: int = 0) -> FDConfig: eplb_config = EPLBConfig(args.eplb_config) structured_outputs_config: StructuredOutputsConfig = StructuredOutputsConfig(args=vars(args)) + routing_replay_config = RoutingReplayConfig(args.routing_replay_config) # Note(tangbinhan): used for load_checkpoint model_config.pretrained_config.tensor_parallel_rank = parallel_config.tensor_parallel_rank @@ -1003,6 +1012,7 @@ def initialize_fd_config(args, ranks: int = 1, local_rank: int = 0) -> FDConfig: plas_attention_config=plas_attention_config, structured_outputs_config=structured_outputs_config, eplb_config=eplb_config, + routing_replay_config=routing_replay_config, ) update_fd_config_for_mm(fd_config) if fd_config.load_config.load_choices == "default_v1" and not v1_loader_support(fd_config): diff --git a/tests/distributed/chunked_moe.py b/tests/distributed/chunked_moe.py index 0be645d38e2..ef41a610dcb 100644 --- a/tests/distributed/chunked_moe.py +++ b/tests/distributed/chunked_moe.py @@ -90,7 +90,7 @@ def init_attention_metadata(self, forward_meta): class MockQuantMethod: - def apply(self, layer, x, gate): + def apply(self, layer, x, gate, topk_ids_hookfunc=None): return x @@ -129,6 +129,7 @@ def setup_model_runner(self): model_runner.speculative_decoding = False model_runner._init_share_inputs(mock_fd_config.scheduler_config.max_num_seqs) model_runner.share_inputs["caches"] = None + model_runner.routing_replay_manager = None if dist.get_rank() == 0: model_runner.share_inputs["ids_remove_padding"] = paddle.ones([10]) @@ -148,6 +149,7 @@ def setup_fused_moe(self): fused_moe.fd_config = mock_fd_config fused_moe.quant_method = MockQuantMethod() + fused_moe.enable_routing_replay = None return fused_moe def run_model_runner(self): diff --git a/tests/e2e/test_EB_Lite_serving.py b/tests/e2e/test_EB_Lite_serving.py index bc27daab993..c71b7667260 100644 --- a/tests/e2e/test_EB_Lite_serving.py +++ b/tests/e2e/test_EB_Lite_serving.py @@ -78,6 +78,8 @@ def setup_and_run_server(): "wint4", "--graph-optimization-config", '{"cudagraph_capture_sizes": [1], "use_cudagraph":true}', + "--routing-replay-config", + '{"enable_routing_replay":true, "routing_store_type":"local", "local_store_dir":"./routing_replay_output"}', ] # Start subprocess in new process group diff --git a/tests/layers/test_fusedmoe.py b/tests/layers/test_fusedmoe.py index ed4fe5b28b6..346afc98fd3 100644 --- a/tests/layers/test_fusedmoe.py +++ b/tests/layers/test_fusedmoe.py @@ -31,6 +31,7 @@ LoadConfig, ModelConfig, ParallelConfig, + RoutingReplayConfig, ) from fastdeploy.model_executor.layers.moe.moe import FusedMoE from fastdeploy.model_executor.layers.quantization.block_wise_fp8 import ( @@ -476,6 +477,7 @@ def __init__( graph_opt_config=GraphOptimizationConfig({}), load_config=LoadConfig({}), ips=",".join(["0"] * nnodes), + routing_replay_config=RoutingReplayConfig({}), ) self.fd_config.parallel_config.tp_group = None self.fd_config.parallel_config.tensor_parallel_rank = tp_rank diff --git a/tests/layers/test_w4a8_moe.py b/tests/layers/test_w4a8_moe.py index dc6dab15427..f20c27b06bf 100644 --- a/tests/layers/test_w4a8_moe.py +++ b/tests/layers/test_w4a8_moe.py @@ -13,6 +13,7 @@ LoadConfig, ModelConfig, ParallelConfig, + RoutingReplayConfig, ) from fastdeploy.model_executor.layers.moe.moe import FusedMoE from fastdeploy.model_executor.layers.quantization.w4a8 import W4A8Config @@ -59,6 +60,7 @@ def __init__( graph_opt_config=GraphOptimizationConfig({}), load_config=LoadConfig({}), ips=",".join(["0"] * nnodes), + routing_replay_config=RoutingReplayConfig({}), ) self.fd_config.parallel_config.tp_group = None self.fd_config.parallel_config.tensor_parallel_rank = tp_rank diff --git a/tests/layers/test_w4afp8_moe.py b/tests/layers/test_w4afp8_moe.py index 65b7733172c..8f1ae79cd67 100644 --- a/tests/layers/test_w4afp8_moe.py +++ b/tests/layers/test_w4afp8_moe.py @@ -13,6 +13,7 @@ LoadConfig, ModelConfig, ParallelConfig, + RoutingReplayConfig, ) from fastdeploy.model_executor.layers.moe.moe import FusedMoE from fastdeploy.model_executor.layers.quantization.w4afp8 import W4AFP8Config @@ -65,6 +66,7 @@ def __init__( graph_opt_config=GraphOptimizationConfig({}), load_config=LoadConfig({}), ips=",".join(["0"] * nnodes), + routing_replay_config=RoutingReplayConfig({}), ) self.fd_config.parallel_config.tp_group = None self.fd_config.parallel_config.tensor_parallel_rank = tp_rank From 3ab2d53d590281a3e069cc2f2578f7ed6aeafdfd Mon Sep 17 00:00:00 2001 From: lizexu123 <39205361+lizexu123@users.noreply.github.com> Date: Fri, 5 Dec 2025 20:18:17 +0800 Subject: [PATCH 05/13] [Bug fix] Fix the multi-input accuracy issue in the pooling model. (#5374) * fix multi-inputs * fix threshold * fix threshold * fix --- fastdeploy/entrypoints/openai/protocol.py | 2 +- fastdeploy/input/text_processor.py | 16 +++- tests/pooling/test_Qwen3-Embedding_serving.py | 95 ++++++++++++++++++- 3 files changed, 107 insertions(+), 6 deletions(-) diff --git a/fastdeploy/entrypoints/openai/protocol.py b/fastdeploy/entrypoints/openai/protocol.py index 6a57cf47242..000861470dd 100644 --- a/fastdeploy/entrypoints/openai/protocol.py +++ b/fastdeploy/entrypoints/openai/protocol.py @@ -933,7 +933,7 @@ class EmbeddingChatRequest(BaseModel): ) add_special_tokens: bool = Field( - default=False, + default=True, description=( "If true, special tokens (e.g. BOS) will be added to the prompt " "on top of what is added by the chat template. " diff --git a/fastdeploy/input/text_processor.py b/fastdeploy/input/text_processor.py index a27d125ebc8..366244e5244 100644 --- a/fastdeploy/input/text_processor.py +++ b/fastdeploy/input/text_processor.py @@ -230,13 +230,16 @@ def process_request(self, request, max_model_len=None, **kwargs): if request.prompt_token_ids is None or len(request.prompt_token_ids) == 0: if request.prompt is not None: prompt = request.prompt + add_special_tokens = request.get("add_special_tokens", False) assert isinstance(prompt, str) or ( isinstance(prompt, list) and all([isinstance(t, int) for t in prompt]) ), f"prompt must be a string or a list of integers, but got {type(prompt)}" if isinstance(prompt, list): # if prompt is a token id list request.prompt_token_ids = prompt else: - request.prompt_token_ids = self.text2ids(request.prompt, max_model_len) + request.prompt_token_ids = self.text2ids( + request.prompt, max_model_len, add_special_tokens=add_special_tokens + ) elif request.messages is not None: if self.tokenizer.chat_template is None: raise ValueError("This model does not support chat_template.") @@ -305,13 +308,16 @@ def process_request_dict(self, request, max_model_len=None, **kwargs): if not request.get("prompt_token_ids"): if request.get("prompt"): prompt = request.get("prompt") + add_special_tokens = request.get("add_special_tokens", False) assert isinstance(prompt, str) or ( isinstance(prompt, list) and all([isinstance(t, int) for t in prompt]) ), f"prompt must be a string or a list of integers, but got {type(prompt)}" if isinstance(prompt, list): # if prompt is a token id list request["prompt_token_ids"] = prompt else: - request["prompt_token_ids"] = self.text2ids(request["prompt"], max_model_len).tolist() + request["prompt_token_ids"] = self.text2ids( + request["prompt"], max_model_len, add_special_tokens=add_special_tokens + ).tolist() elif request.get("messages"): if self.tokenizer.chat_template is None: raise ValueError("This model does not support chat_template.") @@ -503,7 +509,7 @@ def process_response_dict(self, response_dict, **kwargs): **kwargs, ) - def text2ids(self, text, max_model_len): + def text2ids(self, text, max_model_len, **kwargs): """ text to token ids @@ -513,6 +519,8 @@ def text2ids(self, text, max_model_len): Returns: List[int]: token ids list """ + + add_special_tokens = kwargs.get("add_special_tokens") if envs.FD_USE_HF_TOKENIZER: tokens = self.tokenizer( text, @@ -529,7 +537,7 @@ def text2ids(self, text, max_model_len): padding=True, truncation=True, max_length=max_model_len, - add_special_tokens=False, + add_special_tokens=add_special_tokens, ) return tokens["input_ids"][0] diff --git a/tests/pooling/test_Qwen3-Embedding_serving.py b/tests/pooling/test_Qwen3-Embedding_serving.py index 80a4410d996..74365bd8e54 100644 --- a/tests/pooling/test_Qwen3-Embedding_serving.py +++ b/tests/pooling/test_Qwen3-Embedding_serving.py @@ -237,4 +237,97 @@ def test_single_text_embedding(embedding_api_url, headers): save_embedding_baseline(embedding, baseline_file) else: print(f"Comparing with baseline: {baseline_file}") - check_embedding_against_baseline(embedding, baseline_file, threshold=0.01) + check_embedding_against_baseline(embedding, baseline_file, threshold=0.02) + + +def test_multi_text_embedding(embedding_api_url, headers): + """Test embedding generation for batch text inputs.""" + payload = { + "model": "default", + "input": ["北京天安门在哪里?", "上海东方明珠有多高?", "杭州西湖的面积是多少?"], + } + + resp = requests.post(embedding_api_url, headers=headers, json=payload) + assert resp.status_code == 200, f"Unexpected status code: {resp.status_code}, response: {resp.text}" + + result = resp.json() + assert "data" in result, "Response missing 'data' field" + assert len(result["data"]) == 3, f"Expected 3 embedding results, got {len(result['data'])}" + + # Validate each embedding in the batch + for idx, item in enumerate(result["data"]): + assert "embedding" in item, f"Item {idx} missing 'embedding' field" + assert "index" in item, f"Item {idx} missing 'index' field" + assert item["index"] == idx, f"Item index mismatch: expected {idx}, got {item['index']}" + + embedding = item["embedding"] + assert isinstance(embedding, list), f"Embedding {idx} should be a list" + assert len(embedding) > 0, f"Embedding {idx} vector should not be empty" + assert all(isinstance(x, (int, float)) for x in embedding), f"Embedding {idx} values should be numeric" + + print(f"Text {idx} embedding dimension: {len(embedding)}") + + # Verify all embeddings have the same dimension + dimensions = [len(item["embedding"]) for item in result["data"]] + assert len(set(dimensions)) == 1, f"All embeddings should have same dimension, got: {dimensions}" + + # Compare embeddings with baseline + base_path = os.getenv("MODEL_PATH", "") + baseline_filename = "test-Qwen3-Embedding-0.6B-multi-input-baseline.json" + + if base_path: + baseline_file = os.path.join(base_path, "torch", baseline_filename) + else: + baseline_file = baseline_filename + + # Save all embeddings to baseline + batch_embeddings = [item["embedding"] for item in result["data"]] + + if not os.path.exists(baseline_file): + print("Batch baseline file not found. Saving current embeddings as baseline...") + baseline_data = { + "embeddings": batch_embeddings, + "dimension": len(batch_embeddings[0]), + "count": len(batch_embeddings), + "inputs": payload["input"], + } + with open(baseline_file, "w", encoding="utf-8") as f: + json.dump(baseline_data, f, indent=2) + print(f"Batch baseline saved to: {baseline_file}") + else: + print(f"Comparing batch with baseline: {baseline_file}") + with open(baseline_file, "r", encoding="utf-8") as f: + baseline_data = json.load(f) + baseline_embeddings = baseline_data["embeddings"] + + assert len(batch_embeddings) == len( + baseline_embeddings + ), f"Embedding count mismatch: current={len(batch_embeddings)}, baseline={len(baseline_embeddings)}" + + # Compare each embedding + for idx, (current_emb, baseline_emb) in enumerate(zip(batch_embeddings, baseline_embeddings)): + print(f"\n--- Comparing embedding {idx}: '{payload['input'][idx]}' ---") + mean_abs_diff = compare_embeddings(current_emb, baseline_emb, threshold=0.05) + + if mean_abs_diff >= 0.05: + # Save current batch for debugging + temp_file = f"{baseline_file}.current" + print("temp_file", temp_file) + with open(temp_file, "w", encoding="utf-8") as f: + json.dump( + { + "embeddings": batch_embeddings, + "dimension": len(batch_embeddings[0]), + "count": len(batch_embeddings), + "inputs": payload["input"], + }, + f, + indent=2, + ) + + raise AssertionError( + f"Embedding {idx} differs from baseline by too much " + f"(mean_abs_diff={mean_abs_diff:.6f} >= 0.01):\n" + f"Current batch saved to: {temp_file}\n" + f"Please check the differences." + ) From f77e222c383ee1ce6897455bec5787d3adf22544 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E5=91=A8=E5=91=A8=E5=91=A8?= <39978853+zhoutianzi666@users.noreply.github.com> Date: Fri, 5 Dec 2025 20:19:01 +0800 Subject: [PATCH 06/13] [BugFix]remove _execute_empty_input (#5396) --- fastdeploy/config.py | 3 +++ fastdeploy/worker/gpu_model_runner.py | 27 ++++++++++++++++----------- 2 files changed, 19 insertions(+), 11 deletions(-) diff --git a/fastdeploy/config.py b/fastdeploy/config.py index f1eb23852f8..a820e8d94e6 100644 --- a/fastdeploy/config.py +++ b/fastdeploy/config.py @@ -1577,6 +1577,9 @@ def __init__( self.graph_opt_config._set_cudagraph_sizes(max_capture_size=max_capture_shape) self.graph_opt_config.init_with_cudagrpah_size(max_capture_size=max_capture_shape) + if self.parallel_config.use_ep: + self.graph_opt_config.cudagraph_capture_sizes = [0] + self.graph_opt_config.cudagraph_capture_sizes + self.tokenizer = tokenizer self.ips = ips self.tool_parser = tool_parser diff --git a/fastdeploy/worker/gpu_model_runner.py b/fastdeploy/worker/gpu_model_runner.py index 94c7a0b3f21..9b550f10438 100644 --- a/fastdeploy/worker/gpu_model_runner.py +++ b/fastdeploy/worker/gpu_model_runner.py @@ -1020,10 +1020,14 @@ def get_input_length_list( """ # NOTE(gongshaotian): The maximum decoding length is equal to the expected decoded tokens plus the eos token max_dec_len = expected_decode_len + 1 - input_length = min( - num_tokens // (1 if capture_prefill else batch_size), - self.model_config.max_model_len - max_dec_len, - ) + if batch_size == 0: + # Note(ZKK): divided by 0 is invalid, here we give a input_length = 1 + input_length = 1 + else: + input_length = min( + num_tokens // (1 if capture_prefill else batch_size), + self.model_config.max_model_len - max_dec_len, + ) # NOTE(wanglongzhi): When the full length is too large, DeepEP's buffer size will not be enough to cause the result to appear nan. # TODO(wanglongzhi): Figure out the accurate buffer size of DeepEP. @@ -2223,13 +2227,6 @@ class at the server level, which is too granular for ModelRunner. for proc in self.sampling_metadata.logits_processors: proc.update_state(self.share_inputs) - # NOTE(wufeisheng): If `not_need_stop`` is False, it means the current worker is in an idle state. - # This logic is not used in TP (Tensor Parallelism) mode. However, in EP (Expert Parallelism) mode, - # when there is data on other runner, the current runner is required to execute part of the model. - if not self.not_need_stop(): - self._execute_empty_input(self.forward_meta) - return None - # 2. Padding inputs for cuda graph self.padding_cudagraph_inputs() @@ -2245,6 +2242,14 @@ class at the server level, which is too granular for ModelRunner. self.forward_meta.ids_remove_padding, self.forward_meta, ) + + # NOTE(wufeisheng): If `not_need_stop`` is False, it means the current worker is in an idle state. + # This logic is not used in TP (Tensor Parallelism) mode. However, in EP (Expert Parallelism) mode, + # Then there is data on other runner, the current runner is required to execute part of the model. + # But not need to run the below code. + if not self.not_need_stop(): + return None + if self.use_cudagraph: model_output = model_output[: self.real_token_num] From e70351c8bdd19cbefb50e3c6e9a153c9372015c9 Mon Sep 17 00:00:00 2001 From: Jiang-Jia-Jun <163579578+Jiang-Jia-Jun@users.noreply.github.com> Date: Fri, 5 Dec 2025 20:19:39 +0800 Subject: [PATCH 07/13] Revert "[RL] Support Rollout Routing Replay (#5321)" (#5402) This reverts commit 96d2d4877b8a1f374d1e0bcb4f02b707b3d51d06. --- fastdeploy/config.py | 27 -- fastdeploy/engine/args_utils.py | 22 -- fastdeploy/engine/engine.py | 1 - fastdeploy/model_executor/forward_meta.py | 2 - .../backends/dcu/fused_moe_triton_backends.py | 5 - .../gcu/moe/fused_moe_method_gcu_backend.py | 6 - .../moe/fused_moe_triton_metax_backend.py | 6 - .../layers/moe/fused_moe_backend_base.py | 11 +- .../layers/moe/fused_moe_cutlass_backend.py | 17 - .../layers/moe/fused_moe_deepgemm_backend.py | 16 - .../layers/moe/fused_moe_marlin_backend.py | 6 - .../layers/moe/fused_moe_triton_backend.py | 18 - .../layers/moe/fused_moe_wint2_backend.py | 10 - fastdeploy/model_executor/layers/moe/moe.py | 60 +-- .../layers/moe/routing_indices_cache.py | 346 ------------------ fastdeploy/model_executor/models/glm4_moe.py | 7 +- fastdeploy/rl/rollout_config.py | 2 - fastdeploy/worker/gpu_model_runner.py | 32 -- fastdeploy/worker/worker_process.py | 10 - tests/distributed/chunked_moe.py | 4 +- tests/e2e/test_EB_Lite_serving.py | 2 - tests/layers/test_fusedmoe.py | 2 - tests/layers/test_w4a8_moe.py | 2 - tests/layers/test_w4afp8_moe.py | 2 - 24 files changed, 24 insertions(+), 592 deletions(-) delete mode 100644 fastdeploy/model_executor/layers/moe/routing_indices_cache.py diff --git a/fastdeploy/config.py b/fastdeploy/config.py index a820e8d94e6..63ac382d108 100644 --- a/fastdeploy/config.py +++ b/fastdeploy/config.py @@ -1484,31 +1484,6 @@ def __str__(self) -> str: return json.dumps({key: value for key, value in self.__dict__.items()}) -class RoutingReplayConfig: - """Configuration for Routing Replay used in RL training""" - - def __init__(self, args) -> None: - self.enable_routing_replay: bool = False - self.routing_store_type: str = "local" - - # Local routing store - self.local_store_dir: str = "./routing_replay_output" - - # RDMA routing store - # TODO: Add RDMA routing store configuration attributes here when the feature is implemented. - - if args is not None: - for key, value in args.items(): - if hasattr(self, key) and value != "None": - setattr(self, key, value) - - def to_json_string(self): - """ - Convert routing replay config to json string. - """ - return json.dumps({key: value for key, value in self.__dict__.items()}) - - class FDConfig: """ The configuration class which contains all fastdeploy-related configuration. This @@ -1542,7 +1517,6 @@ def __init__( early_stop_config: Optional[Dict[str, Any]] = None, tool_parser: str = None, test_mode=False, - routing_replay_config: Optional[RoutingReplayConfig] = None, ): self.model_config: ModelConfig = model_config # type: ignore self.cache_config: CacheConfig = cache_config # type: ignore @@ -1559,7 +1533,6 @@ def __init__( self.plas_attention_config: Optional[PlasAttentionConfig] = plas_attention_config self.structured_outputs_config: StructuredOutputsConfig = structured_outputs_config self.router_config: RouterConfig = router_config - self.routing_replay_config = routing_replay_config # Initialize cuda graph capture list max_capture_shape = self.scheduler_config.max_num_seqs diff --git a/fastdeploy/engine/args_utils.py b/fastdeploy/engine/args_utils.py index d2d7c6f908a..1eaf535498a 100644 --- a/fastdeploy/engine/args_utils.py +++ b/fastdeploy/engine/args_utils.py @@ -35,7 +35,6 @@ PlasAttentionConfig, PoolerConfig, RouterConfig, - RoutingReplayConfig, RunnerOption, SpeculativeConfig, StructuredOutputsConfig, @@ -492,11 +491,6 @@ class EngineArgs: Configuration for eplb. """ - routing_replay_config: Optional[Dict[str, Any]] = None - """ - Flag to rollout routing replay(r3) - """ - def __post_init__(self): """ Post-initialization processing to set default tokenizer if not provided. @@ -888,12 +882,6 @@ def add_cli_args(parser: FlexibleArgumentParser) -> FlexibleArgumentParser: default=EngineArgs.eplb_config, help="Config of eplb.", ) - parallel_group.add_argument( - "--routing-replay-config", - type=json.loads, - default=EngineArgs.routing_replay_config, - help="Flag of rollout routing replay(r3).", - ) parallel_group.add_argument( "--enable-chunked-moe", action="store_true", @@ -1247,14 +1235,6 @@ def create_eplb_config(self) -> EPLBConfig: eplb_args["enable_eplb"] = self.enable_eplb return EPLBConfig(eplb_args) - def create_routing_repaly_config(self) -> RoutingReplayConfig: - """ """ - routing_replay_args = asdict(self) - if self.routing_replay_config is not None: - for k, v in self.routing_replay_config.items(): - routing_replay_args[k] = v - return RoutingReplayConfig(routing_replay_args) - def create_engine_config(self, port_availability_check=True) -> FDConfig: """ Create and return a Config object based on the current settings. @@ -1298,7 +1278,6 @@ def create_engine_config(self, port_availability_check=True) -> FDConfig: graph_opt_cfg = self.create_graph_optimization_config() plas_attention_config = self.create_plas_attention_config() eplb_cfg = self.create_eplb_config() - routing_replay_config = self.create_routing_repaly_config() router_config = RouterConfig(all_dict) early_stop_cfg = self.create_early_stop_config() @@ -1331,5 +1310,4 @@ def create_engine_config(self, port_availability_check=True) -> FDConfig: graph_opt_config=graph_opt_cfg, plas_attention_config=plas_attention_config, early_stop_config=early_stop_cfg, - routing_replay_config=routing_replay_config, ) diff --git a/fastdeploy/engine/engine.py b/fastdeploy/engine/engine.py index fadf954679b..4a493843df7 100644 --- a/fastdeploy/engine/engine.py +++ b/fastdeploy/engine/engine.py @@ -568,7 +568,6 @@ def _start_worker_service(self): f" --logprobs_mode {self.cfg.model_config.logprobs_mode}" f" --max_logprobs {self.cfg.model_config.max_logprobs}" f" --eplb_config '{self.cfg.eplb_config.to_json_string()}'" - f" --routing_replay_config '{self.cfg.routing_replay_config.to_json_string()}'" ) if self.cfg.structured_outputs_config.logits_processors is not None: arguments += f" --logits-processors {' '.join(self.cfg.structured_outputs_config.logits_processors)}" diff --git a/fastdeploy/model_executor/forward_meta.py b/fastdeploy/model_executor/forward_meta.py index 787ec77c0eb..4e9df0d3ce3 100644 --- a/fastdeploy/model_executor/forward_meta.py +++ b/fastdeploy/model_executor/forward_meta.py @@ -142,8 +142,6 @@ class ForwardMeta: caches: Optional[list[paddle.Tensor]] = None # Flag of profile run is_dummy_or_profile_run: bool = False - # Routing Replay table buffer - routing_replay_table: Optional[paddle.Tensor] = None # chunked MoE related moe_num_chunk: int = 1 diff --git a/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py b/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py index 192c0b8833a..918450c74f1 100644 --- a/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py +++ b/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py @@ -14,8 +14,6 @@ # limitations under the License. """ -from typing import Callable - import paddle from paddle import nn @@ -103,7 +101,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -120,8 +117,6 @@ def apply( scores += layer.gate_correction_bias topk_weights, topk_ids = paddle.topk(scores, k=top_k, axis=-1, sorted=False) topk_weights = topk_weights / topk_weights.sum(axis=-1, keepdim=True) - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_ids) intermediate_cache1 = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], diff --git a/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py b/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py index 2260d7caf7b..e67dd6dbdaf 100644 --- a/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py +++ b/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py @@ -16,7 +16,6 @@ import multiprocessing import os -from typing import Callable import numpy as np import paddle @@ -183,7 +182,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle gcu compute Fused MoE. @@ -196,7 +194,6 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -208,7 +205,6 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -220,7 +216,6 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -386,7 +381,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle gcu compute Fused MoE. diff --git a/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py b/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py index fbbfac277b8..7b61d58b6f5 100644 --- a/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py +++ b/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py @@ -14,8 +14,6 @@ # limitations under the License. """ -from typing import Callable - import paddle from paddle import nn @@ -247,7 +245,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -277,9 +274,6 @@ def apply( True, # apply_norm_weight False, ) - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_ids) - up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py b/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py index a8bd70465ea..b34291a96f4 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py @@ -15,7 +15,6 @@ """ from abc import abstractmethod -from typing import Callable import paddle from paddle import nn @@ -164,7 +163,6 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -177,7 +175,6 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -190,7 +187,6 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -202,7 +198,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -212,13 +207,13 @@ def apply( if layer.fd_config.model_config.moe_phase.phase == "prefill": if layer.fd_config.scheduler_config.splitwise_role == "mixed" and is_moe_start_layer: self.ep_prefill_runner.clean_low_latency_buffer() - return self.apply_ep_prefill(layer, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) + return self.apply_ep_prefill(layer, x, gate) else: if layer.fd_config.scheduler_config.splitwise_role == "mixed" and is_moe_start_layer: self.ep_decoder_runner.clean_low_latency_buffer() - return self.apply_ep_decode(layer, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) + return self.apply_ep_decode(layer, x, gate) else: - return self.apply_tp(layer, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) + return self.apply_tp(layer, x, gate) class UnquantizedFusedMoEMethod(MoEMethodBase): diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py index c3dbfc9ba5f..e45ad63b19c 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py @@ -14,8 +14,6 @@ # limitations under the License. """ -from typing import Callable - import paddle from paddle import nn from paddle.nn.quant import weight_quantize @@ -134,7 +132,6 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -151,13 +148,8 @@ def apply_ep_prefill( handle, event, ) = self.ep_prefill_runner.dispatch(x, topk_idx, topk_weights) - - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_idx) - if self.ep_prefill_runner.ep_engine.async_finish: event.current_stream_wait() - token_all_num = sum(recv_num_tokens_per_expert_list) # 3. Compute ffn @@ -225,7 +217,6 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -234,10 +225,6 @@ def apply_ep_decode( estimate_total_token_nums = gate_out.shape[0] * layer.top_k # 1. Select topk experts and weights topk_idx, topk_weights = self.ep_decoder_runner.moe_select(layer, gate_out) - - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_idx) - expertwise_scale = None if hasattr(layer, "up_gate_proj_in_scale_all_experts"): # only use in w4a8 expertwise_scale = getattr(layer, "up_gate_proj_in_scale_all_experts", None) @@ -282,7 +269,6 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -383,9 +369,6 @@ def apply_tp( if hasattr(layer, "up_gate_proj_in_scale"): dequant_scale = None - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_idx) - if not layer.with_bias and self.moe_quant_type != "w4a8" and self.moe_quant_type != "w4afp8": # only w4a8 need expert_idx_per_token # Other need not this tensor, so we make it None. diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py index 881f9a22c4d..1245cddcebc 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py @@ -14,8 +14,6 @@ # limitations under the License. """ -from typing import Callable - import paddle from paddle import nn from paddle.distributed.communication import deep_ep @@ -141,7 +139,6 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -150,10 +147,6 @@ def apply_ep_prefill( # 1. Select topk experts and weights topk_idx, topk_weights = self.ep_prefill_runner.moe_select(layer, gate_out) - - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_idx) - # 2. Dynamic compute blockwise quantization scales x, x_scale_tensor = fastdeploy.model_executor.ops.gpu.per_token_quant( x, self.quant_config.weight_block_size[0] @@ -271,7 +264,6 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -279,10 +271,6 @@ def apply_ep_decode( gate_out = gate(x.cast("float32")) # 1. Select topk experts and weights topk_idx, topk_weights = self.ep_decoder_runner.moe_select(layer, gate_out) - - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_idx) - # 2. EP Dispatch permute_input, token_nums_per_expert, handle = self.ep_decoder_runner.dispatch( x, topk_idx, topk_weights, use_fp8=True @@ -347,7 +335,6 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Use DeepGemm compute Fused MoE. @@ -376,9 +363,6 @@ def apply_tp( False, ) - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_ids) - tmp = count_tokens_per_expert_func(topk_ids, layer.num_experts) recv_x, recv_x_scale = fastdeploy.model_executor.ops.gpu.per_token_quant(x, 128) diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py index cd836dbaf09..094d3df8f1a 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py @@ -14,8 +14,6 @@ # limitations under the License. """ -from typing import Callable - import paddle from paddle import nn @@ -241,7 +239,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Marlin compute Fused MoE. @@ -276,9 +273,6 @@ def apply( False, ) - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_ids) - block_size_m = 64 for m in [8, 16, 32, 48, 64]: diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py index 2861d96e8d3..3c14859375c 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py @@ -14,8 +14,6 @@ # limitations under the License. """ -from typing import Callable - import paddle from paddle import nn @@ -284,7 +282,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -317,10 +314,6 @@ def apply( True, # apply_norm_weight, False, ) - - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_ids) - up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, @@ -671,7 +664,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -732,9 +724,6 @@ def apply( * ceil_div(moe_intermediate_size * 2, config["BLOCK_SIZE_N"]), ) - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_ids) - up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, @@ -964,7 +953,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -986,9 +974,6 @@ def apply( False, ) - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_ids) - up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, @@ -1481,7 +1466,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -1504,8 +1488,6 @@ def apply( True, # apply_norm_weight False, ) - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_ids) config = { "BLOCK_SIZE_M": 64, diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py index 3c548ba57c8..f75e36bcbdd 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py @@ -14,8 +14,6 @@ # limitations under the License. """ -from typing import Callable - import paddle from paddle import nn @@ -263,7 +261,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Use Wint2 Triton Fusedmoe compute Fused MoE. @@ -291,9 +288,6 @@ def apply( topk_only_mode=False, ) - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_idx) - ffn_out = fastdeploy.model_executor.ops.gpu.moe_expert_ffn_wint2( permute_input, token_nums_per_expert, @@ -334,7 +328,6 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, - topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Use Wint2 Triton Fusedmoe compute Fused MoE. @@ -350,9 +343,6 @@ def apply( False, ) - if topk_ids_hookfunc is not None: - topk_ids_hookfunc(topk_ids=topk_ids) - num_tokens, K = x.shape E, _, N = layer.up_gate_proj_weight.shape M = num_tokens diff --git a/fastdeploy/model_executor/layers/moe/moe.py b/fastdeploy/model_executor/layers/moe/moe.py index 5b1be52d183..743e05031f6 100644 --- a/fastdeploy/model_executor/layers/moe/moe.py +++ b/fastdeploy/model_executor/layers/moe/moe.py @@ -14,8 +14,7 @@ # limitations under the License. """ -from functools import partial -from typing import Callable, Optional +from typing import Optional import paddle from paddle import nn @@ -27,9 +26,6 @@ tensor_model_parallel_all_reduce_custom, ) from fastdeploy.model_executor.forward_meta import ForwardMeta -from fastdeploy.model_executor.layers.moe.routing_indices_cache import ( - save_routing_to_buffer, -) from fastdeploy.model_executor.layers.utils import get_tensor from fastdeploy.model_executor.utils import h2d_copy, slice_fn from fastdeploy.platforms import current_platform @@ -230,7 +226,7 @@ def __init__( self.is_rearrange = False if self.ep_size > 1: self.quant_method.init_ep(self) - self.enable_routing_replay = fd_config.routing_replay_config.enable_routing_replay + # Merge normal and RL build model if gate_correction_bias is not None: self.gate_correction_bias = gate_correction_bias @@ -604,7 +600,7 @@ def load_state_dict(self, state_dict, is_rearrange: bool = False): else: self.quant_method.process_loaded_weights(self, state_dict) - def forward_split_allgather(self, x: paddle.Tensor, gate: nn.Layer, topk_ids_hookfunc: Callable = None): + def forward_split_allgather(self, x: paddle.Tensor, gate: nn.Layer): """ Forward split allgather function. """ @@ -619,14 +615,14 @@ def forward_split_allgather(self, x: paddle.Tensor, gate: nn.Layer, topk_ids_hoo if end_offset > token_num: end_offset = token_num part_x[: (end_offset - start_offset), :] = x[start_offset:end_offset, :] - out = self.quant_method.apply(self, part_x, gate, topk_ids_hookfunc=topk_ids_hookfunc) + out = self.quant_method.apply(self, part_x, gate) multi_outs = paddle.zeros([token_num_per_rank * self.attn_tp_size, x.shape[1]], dtype=x.dtype) paddle.distributed.all_gather(multi_outs, out, self.tp_group) out = multi_outs[:token_num, :] return out - def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta = None): + def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): """ Defines the forward computation of the moe layer. @@ -637,21 +633,6 @@ def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta = Tensor: Output tensor.s """ - topk_ids_hookfunc = None - if self.enable_routing_replay: - if forward_meta is not None: # forward_meta is None when execute empty_input_forward - topk_ids_hookfunc = partial( - save_routing_to_buffer, - routing_replay_table=forward_meta.routing_replay_table, - batch_id_per_token=forward_meta.batch_id_per_token, - seq_lens_decoder=forward_meta.seq_lens_decoder, - cu_seqlens_q=forward_meta.cu_seqlens_q, - layer_idx=self.layer_idx, - tp_size=self.fd_config.parallel_config.tensor_parallel_size, - ep_size=self.fd_config.parallel_config.expert_parallel_size, - tp_group=self.fd_config.parallel_config.tp_group, - ) - token_num = x.shape[0] if ( self.ep_size > 1 @@ -659,16 +640,11 @@ def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta = and (not self.fd_config.parallel_config.use_sequence_parallel_moe) and token_num >= self.attn_tp_size ): - out = self.forward_split_allgather(x, gate, topk_ids_hookfunc=topk_ids_hookfunc) + out = self.forward_split_allgather(x, gate) elif self.fd_config.parallel_config.use_ep and self.fd_config.parallel_config.enable_chunked_moe: - out = self.forward_chunked_moe( - x, - gate, - forward_meta, - topk_ids_hookfunc=topk_ids_hookfunc, - ) + out = self.forward_chunked_moe(x, gate, forward_meta) else: - out = self.forward_normal(x, gate, forward_meta, topk_ids_hookfunc=topk_ids_hookfunc) + out = self.forward_normal(x, gate) if self.reduce_results and self.tp_size > 1: if current_platform.is_intel_hpu(): @@ -677,9 +653,7 @@ def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta = out = tensor_model_parallel_all_reduce(out, self.tp_group) return out - def forward_chunked_moe( - self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta, topk_ids_hookfunc: Callable = None - ): + def forward_chunked_moe(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): """ Split input to multi chunk to reduce the memory usage of moe. @@ -703,25 +677,21 @@ def forward_chunked_moe( for i in range(forward_meta.max_moe_num_chunk): if i < forward_meta.moe_num_chunk: - out_split_list[i] = self.quant_method.apply( - self, x_split_list[i], gate, topk_ids_hookfunc=topk_ids_hookfunc - ) + out_split_list[i] = self.quant_method.apply(self, x_split_list[i], gate) else: # just need to use real data to infer max_moe_num_chunk times. - self.quant_method.apply(self, fake_x, gate, topk_ids_hookfunc=topk_ids_hookfunc) + self.quant_method.apply(self, fake_x, gate) out = paddle.concat(out_split_list, axis=0) else: # when only one chunk, just need to use real data to infer once. - out = self.quant_method.apply(self, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) + out = self.quant_method.apply(self, x, gate) for i in range(forward_meta.max_moe_num_chunk - 1): - self.quant_method.apply(self, fake_x, gate, topk_ids_hookfunc=topk_ids_hookfunc) + self.quant_method.apply(self, fake_x, gate) return out - def forward_normal( - self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta, topk_ids_hookfunc: Callable = None - ): + def forward_normal(self, x: paddle.Tensor, gate: nn.Layer): """ Normal mode of forward. @@ -732,5 +702,5 @@ def forward_normal( Tensor: Output tensor.s """ - out = self.quant_method.apply(self, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) + out = self.quant_method.apply(self, x, gate) return out diff --git a/fastdeploy/model_executor/layers/moe/routing_indices_cache.py b/fastdeploy/model_executor/layers/moe/routing_indices_cache.py deleted file mode 100644 index e95a3d8569f..00000000000 --- a/fastdeploy/model_executor/layers/moe/routing_indices_cache.py +++ /dev/null @@ -1,346 +0,0 @@ -""" -# Copyright (c) 2025 PaddlePaddle Authors. All Rights Reserved. -# -# Licensed under the Apache License, Version 2.0 (the "License"); -# you may not use this file except in compliance with the License. -# You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, software -# distributed under the License is distributed on an "AS IS" BASIS, -# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# See the License for the specific language governing permissions and -# limitations under the License. -""" - -import copy -import os -import shutil -from abc import ABC, abstractmethod -from typing import Dict, List, Optional - -import paddle -import paddle.distributed as dist -import triton -import triton.language as tl - -from fastdeploy.config import FDConfig - - -@triton.jit -def _save_routing_kernel( - ROUTING_REPLAY_TABLE_PTR, - TOPK_IDS_PTR, - BATCH_ID_PER_TOKEN_PTR, - CU_SEQLENS_Q_PTR, - SEQ_LENS_DECODER_PTR, - LAYER_IDX, - TOKEN_NUM, - TOP_K, - NUM_HIDDEN_LAYERS, - MAX_MODEL_LEN, - BLOCK_SIZE_M: tl.constexpr, - BLOCK_SIZE_K: tl.constexpr, -): - pid_m = tl.program_id(axis=0) - - token_offsets = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) - token_mask = token_offsets < TOKEN_NUM - - k_offsets = tl.arange(0, BLOCK_SIZE_K) - - k_mask = k_offsets < TOP_K - - topk_ids_ptrs = TOPK_IDS_PTR + token_offsets[:, None] * TOP_K + k_offsets[None, :] - # [BLOCK_SIZE_M, BLOCK_SIZE_K] - - load_mask = token_mask[:, None] & k_mask[None, :] - topk_vals = tl.load(topk_ids_ptrs, mask=load_mask) - - batch_ids = tl.load(BATCH_ID_PER_TOKEN_PTR + token_offsets, mask=token_mask) - pad_mask = token_mask & (batch_ids != -1) - # [0, 3, 4, 10, 12][0, 0, 0, 0, 2, 2, 2, 2, 2, 2, 3, 3] - # -> [0, 0, 0, 0, 4, 4, 4, 4, 4, 4, 10, 10] - # [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11] - [0, 0, 0, 0, 4, 4, 4, 4, 4, 4, 10, 10] - # -> [0, 1, 2, 3, 0, 1, 2, 3, 4, 5, 0, 1] - start_offsets = tl.load(CU_SEQLENS_Q_PTR + batch_ids, mask=pad_mask) - token_relative_index = token_offsets - start_offsets - - # [BLOCK_SIZE_M] - len_decoder = tl.load(SEQ_LENS_DECODER_PTR + batch_ids, mask=pad_mask) - token_seq_pos = len_decoder + token_relative_index - - STRIDE_BUF_SEQ = NUM_HIDDEN_LAYERS * MAX_MODEL_LEN * TOP_K - STRIDE_BUF_LAYER = MAX_MODEL_LEN * TOP_K - STRIDE_BUF_TOKEN = TOP_K - - # [BLOCK_SIZE_M, BLOCK_SIZE_K] - output_ptrs = ( - ROUTING_REPLAY_TABLE_PTR - + batch_ids[:, None] * STRIDE_BUF_SEQ - + LAYER_IDX * STRIDE_BUF_LAYER - + token_seq_pos[:, None] * STRIDE_BUF_TOKEN - + k_offsets[None, :] - ) - - pos_mask = token_seq_pos < MAX_MODEL_LEN - pos_mask = pos_mask & pad_mask - - # [BLOCK_SIZE_M, BLOCK_SIZE_K] - pos_mask = pos_mask[:, None] & k_mask[None, :] - - final_mask = load_mask & pos_mask - - tl.store(output_ptrs, topk_vals, mask=final_mask) - - -def save_routing_to_buffer( - routing_replay_table: paddle.Tensor, # [max_num_seqs, num_layers, max_len, top_k] - topk_ids: paddle.Tensor, # [token_num, top_k] - batch_id_per_token: paddle.Tensor, # [token_num, 1] - seq_lens_decoder: paddle.Tensor, # [max_num_seqs, 1] - cu_seqlens_q: paddle.Tensor, # [max_num_seqs + 1, 1] - layer_idx: int, - tp_size: int, - ep_size: int, - tp_group: dist.communication.group.Group, -): - if tp_size > 1 and ep_size > 1: - token_num_per_rank = topk_ids.shape[0] - topk_ids_all = paddle.zeros([token_num_per_rank * tp_size, topk_ids.shape[1]], dtype=topk_ids.dtype) - paddle.distributed.all_gather(topk_ids_all, topk_ids, tp_group) - topk_ids = topk_ids_all[: batch_id_per_token.shape[0], :] - - token_num, top_k = topk_ids.shape - max_num_seqs, num_hidden_layers, max_model_len, _ = routing_replay_table.shape - assert token_num > 0 - - assert topk_ids.shape[1] == routing_replay_table.shape[3], (topk_ids.shape[1], routing_replay_table.shape[3]) - assert batch_id_per_token.shape[0] == token_num, (batch_id_per_token.shape[0], token_num) - assert seq_lens_decoder.shape[0] == max_num_seqs, (seq_lens_decoder.shape[0], max_num_seqs) - - BLOCK_SIZE_M = 128 - BLOCK_SIZE_K = triton.next_power_of_2(top_k) # top_k - - grid = (triton.cdiv(token_num, BLOCK_SIZE_M),) - _save_routing_kernel[grid]( - routing_replay_table, - topk_ids, - batch_id_per_token, - cu_seqlens_q, - seq_lens_decoder, - LAYER_IDX=layer_idx, - TOKEN_NUM=token_num, - TOP_K=top_k, - NUM_HIDDEN_LAYERS=num_hidden_layers, - MAX_MODEL_LEN=max_model_len, - BLOCK_SIZE_M=BLOCK_SIZE_M, - BLOCK_SIZE_K=BLOCK_SIZE_K, - ) - - -class RoutingReplayManager: - """Request level routing replay table manager""" - - def __init__( - self, - fd_config: FDConfig, - ): - self.max_num_seqs = fd_config.scheduler_config.max_num_seqs - self.max_model_len = fd_config.model_config.max_model_len - self.num_moe_layers = fd_config.model_config.num_hidden_layers - fd_config.model_config.moe_layer_start_index - - if fd_config.model_config.architectures[0] == "Glm4MoeForCausalLM": - self.moe_top_k = fd_config.model_config.num_experts_per_tok - else: - self.moe_top_k = fd_config.model_config.moe_k - self.tp_rank = fd_config.parallel_config.tensor_parallel_rank - - self.routing_store = get_routing_store(fd_config=fd_config) - self.routing_batch_to_request: Dict[int, str] = {} - self.routing_replay_table = paddle.full( - shape=[self.max_num_seqs, self.num_moe_layers, self.max_model_len, self.moe_top_k], - fill_value=-1, - dtype="int32", - ) - - def register_request(self, batch_id: int, request_id: str): - """ - Register a new request to routing replay table - Args: - batch_id: The batch ID of this request - request_id: The global ID of the request is usually executed by the training process in RL - """ - # Save requests that have been finished for the current slot - if batch_id in self.routing_batch_to_request: - pre_request_id = self._deregister_request(batch_id) - self._put_request_to_store(batch_id, pre_request_id) - # Register the new request - self.routing_batch_to_request[batch_id] = request_id - - def _deregister_request(self, batch_id: int) -> str: - """ - Deregister a request from routing replay table - """ - assert batch_id in self.routing_batch_to_request - return self.routing_batch_to_request.pop(batch_id) - - def _put_request_to_store( - self, - batch_id: int, - request_id: str, - ): - if self.tp_rank == 0: - batch_buffer = self.routing_replay_table[batch_id] - for layer_id in range(self.num_moe_layers): - layer_buffer = batch_buffer[layer_id] - rollout_id = self.split_request_id(request_id) - self.routing_store.put(routing_indices=layer_buffer, rollout_id=rollout_id, layer_idx=layer_id) - - self._clear_table_slot(batch_id) - - def put_table_to_store(self): - """Put the routing table""" - batch_ids = copy.deepcopy(list(self.routing_batch_to_request.keys())) - for batch_id in batch_ids: - request_id = self._deregister_request(batch_id) - self._put_request_to_store(batch_id, request_id) - - def _clear_table_slot(self, batch_id: int): - assert 0 <= batch_id < self.max_num_seqs - self.routing_replay_table[batch_id].fill_(-1) - - def clear_routing_table(self): - """Clear all slots of the routing replay table""" - self.routing_replay_table.fill_(-1) - - def _clear_store(self): - """Clear routing store""" - self.routing_store.clear_store() - - def _clear_request_of_store(self, request_id): - """Clear one request of routing store""" - rollout_id = self.split_request_id(request_id) - for layer_idx in range(self.num_moe_layers): - self.routing_store.clear(rollout_id=rollout_id, layer_idx=layer_idx) - - def get_request_from_store(self, request_id: str) -> List[paddle.Tensor]: - """Get the routing indices of the request from store""" - routing_list = [] - rollout_id = self.split_request_id(request_id) - for layer_idx in range(self.num_moe_layers): - one_layer_routing = self.routing_store.get(rollout_id, layer_idx) - routing_list.append(one_layer_routing) - - return routing_list - - def get_routing_table(self) -> paddle.Tensor: - return self.routing_replay_table - - def split_request_id(self, request_id: str): - """Split the request id to get rollout id""" - chat_type, tmp_str = request_id.split("-", 1) - # NOTE(gongshaotian): only support chatcmpl now - # assert chat_type == "chatcmpl" - reversed_tmp_str = tmp_str[::-1].split("-", 5) - rollout_id = reversed_tmp_str[-1][::-1] - return rollout_id - - -class RoutingStoreBase(ABC): - """Base class for routing store""" - - def __init__(self, fd_config: FDConfig) -> None: - self.fd_config = fd_config - - @abstractmethod - def put(self, routing_indices: paddle.Tensor, rollout_id: str, layer_idx: Optional[int] = None) -> None: - """Put the routing indices into store""" - raise NotImplementedError - - @abstractmethod - def get(self, rollout_id: str, layer_idx: Optional[int] = None) -> paddle.Tensor: - """Get the routing indices from store""" - raise NotImplementedError - - @abstractmethod - def clear(self, rollout_id: str, layer_idx: Optional[int] = None) -> None: - """Clear the routing indices of the request""" - raise NotImplementedError - - @abstractmethod - def clear_store( - self, - ): - """Clear the routing indices store""" - raise NotImplementedError - - -class RoutingStoreLocal(RoutingStoreBase): - """Routing Store using local memory""" - - def __init__(self, fd_config) -> None: - super().__init__(fd_config=fd_config) - self.local_store_dir = fd_config.routing_replay_config.local_store_dir - - def put(self, routing_indices: paddle.Tensor, rollout_id: str, layer_idx: int) -> None: - """Put the routing indices into store""" - dir_path = os.path.join(self.local_store_dir, f"{rollout_id}") - os.makedirs(dir_path, exist_ok=True) - file_path = os.path.join(dir_path, f"layer_{layer_idx}.pdtensor") - paddle.save(routing_indices, file_path) - - def get( - self, - rollout_id: str, - layer_idx: int = None, - ) -> paddle.Tensor: - """Get the routing indices from store""" - dir_path = os.path.join(self.local_store_dir, f"{rollout_id}") - file_path = os.path.join(dir_path, f"layer_{layer_idx}.pdtensor") - assert os.path.exists(file_path), f"File not found: {file_path}" - layer_routing_indices = paddle.load(file_path) - - return layer_routing_indices - - def clear( - self, - rollout_id: str, - layer_idx: int = None, - ) -> None: - """Clear the routing indices of the request""" - dir_path = os.path.join(self.local_store_dir, f"{rollout_id}") - file_path = os.path.join(dir_path, f"layer_{layer_idx}.pdtensor") - assert os.path.exists(file_path), f"File not found: {file_path}" - os.remove(file_path) - - # Delete empty directory - if len(os.listdir(dir_path)) == 0: - os.rmdir(dir_path) - - def clear_store(self): - """Clear the routing indices store""" - if os.path.isdir(self.local_store_dir): - for file_name in os.listdir(self.local_store_dir): - file_path = os.path.join(self.local_store_dir, file_name) - shutil.rmtree(file_path) - - -class RoutingStoreRDMA(RoutingStoreBase): - """Routing Store using RDMA""" - - def __init__(self) -> None: - super().__init__() - - -def get_routing_store(fd_config: FDConfig) -> RoutingStoreBase: - if fd_config.routing_replay_config.routing_store_type == "local": - return RoutingStoreLocal(fd_config=fd_config) - elif fd_config.routing_replay_config.routing_store_type == "rdma": - return RoutingStoreRDMA(fd_config=fd_config) - else: - raise ValueError( - f"Invalid routing store type: '{fd_config.routing_replay_config.routing_store_type}'. " - "Valid types are: 'local', 'rdma'" - ) diff --git a/fastdeploy/model_executor/models/glm4_moe.py b/fastdeploy/model_executor/models/glm4_moe.py index 0cc7c4dae45..d5ad6e3916b 100644 --- a/fastdeploy/model_executor/models/glm4_moe.py +++ b/fastdeploy/model_executor/models/glm4_moe.py @@ -161,7 +161,7 @@ def __init__( reduce_results=False, ) - def forward(self, x, forward_meta: ForwardMeta = None): + def forward(self, x, forward_meta): shared_experts_out = self.shared_experts(x) out = self.experts(x, self.gate, forward_meta) out = out + shared_experts_out @@ -306,7 +306,10 @@ def forward( # Fully Connected hidden_states, residual = self.post_attention_layernorm(hidden_states, residual) - hidden_states = self.mlp(hidden_states, forward_meta) + hidden_states = self.mlp( + hidden_states, + forward_meta, + ) return hidden_states, residual diff --git a/fastdeploy/rl/rollout_config.py b/fastdeploy/rl/rollout_config.py index f7ff748fed7..6bd3c3bcb35 100644 --- a/fastdeploy/rl/rollout_config.py +++ b/fastdeploy/rl/rollout_config.py @@ -65,7 +65,6 @@ def __init__( data_parallel_size: int = 1, num_nextn_predict_layers: int = 0, eplb_config: str = {}, - routing_replay_config: str = None, ): # Required parameters self.model = model_name_or_path @@ -114,7 +113,6 @@ def __init__( self.plas_attention_config = plas_attention_config self.num_nextn_predict_layers = num_nextn_predict_layers self.eplb_config = eplb_config - self.routing_replay_config = routing_replay_config def __str__(self): return "\n".join(f"{k}: {v}" for k, v in self.__dict__.items()) diff --git a/fastdeploy/worker/gpu_model_runner.py b/fastdeploy/worker/gpu_model_runner.py index 9b550f10438..2d5f47def46 100644 --- a/fastdeploy/worker/gpu_model_runner.py +++ b/fastdeploy/worker/gpu_model_runner.py @@ -45,9 +45,6 @@ from fastdeploy.model_executor.layers.attention.base_attention_backend import ( AttentionBackend, ) -from fastdeploy.model_executor.layers.moe.routing_indices_cache import ( - RoutingReplayManager, -) from fastdeploy.model_executor.layers.rotary_embedding import get_rope, get_rope_3d from fastdeploy.model_executor.layers.sample.meta_data import SamplingMetadata from fastdeploy.model_executor.layers.sample.sampler import Sampler, SpeculativeSampler @@ -205,11 +202,6 @@ def __init__( os.environ["INFERENCE_MSG_QUEUE_ID"] = str(self.parallel_config.engine_worker_queue_port) logger.info(f"queue id is {str(self.parallel_config.engine_worker_queue_port)}") - # Rollout routing replay config - self.routing_replay_manager = None - if self.fd_config.routing_replay_config.enable_routing_replay: - self.routing_replay_manager = RoutingReplayManager(fd_config=self.fd_config) - self.zmq_client = None self.async_output_queue = None if envs.FD_USE_GET_SAVE_OUTPUT_V1: @@ -656,7 +648,6 @@ def insert_tasks_v1(self, req_dicts: List[Request], num_running_requests: int = self.share_inputs["step_seq_lens_decoder"][idx : idx + 1] = 0 self.share_inputs["prompt_lens"][idx : idx + 1] = len(input_ids) self.share_inputs["is_block_step"][idx : idx + 1] = False - self.share_inputs["is_chunk_step"][idx : idx + 1] = prefill_end_index < len(input_ids) self.share_inputs["step_idx"][idx : idx + 1] = ( len(request.output_token_ids) if prefill_end_index >= len(input_ids) else 0 ) @@ -665,12 +656,6 @@ def insert_tasks_v1(self, req_dicts: List[Request], num_running_requests: int = if request.sampling_params is not None and request.sampling_params.prompt_logprobs is not None: self.prompt_logprobs_reqs[request.request_id] = request has_prefill_task = True - - # Routing Replay - if self.fd_config.routing_replay_config.enable_routing_replay: - if prefill_start_index == 0: - self.routing_replay_manager.register_request(batch_id=idx, request_id=request.request_id) - if ( self.fd_config.scheduler_config.splitwise_role == "decode" ): # In PD, we continue to decode after P generate first token @@ -1167,7 +1152,6 @@ def _init_share_inputs(self, max_num_seqs: int): self.share_inputs["bad_tokens_len"] = paddle.full([max_num_seqs], 1, dtype="int64") self.share_inputs["next_tokens"] = paddle.full([max_num_seqs, 1], -1, dtype="int64") self.share_inputs["is_block_step"] = paddle.full([max_num_seqs], False, dtype="bool") - self.share_inputs["is_chunk_step"] = paddle.full([max_num_seqs], False, dtype="bool").cpu() self.share_inputs["encoder_block_lens"] = paddle.full([max_num_seqs], 0, dtype="int32") self.share_inputs["step_block_list"] = paddle.full([max_num_seqs], -1, dtype="int32") self.share_inputs["step_lens"] = paddle.full([1], 0, dtype="int32") @@ -1438,9 +1422,6 @@ def initialize_forward_meta(self, is_dummy_or_profile_run=False): Initialize forward meta, attention meta data and update some config. """ # Initialize forward meta - routing_replay_table = None - if self.routing_replay_manager is not None: - routing_replay_table = self.routing_replay_manager.get_routing_table() self.forward_meta = ForwardMeta( ids_remove_padding=self.share_inputs["ids_remove_padding"], rotary_embs=self.share_inputs["rope_emb"], @@ -1467,7 +1448,6 @@ def initialize_forward_meta(self, is_dummy_or_profile_run=False): kv_batch_ids=self.share_inputs["kv_batch_ids"], kv_tile_ids_per_batch=self.share_inputs["kv_tile_ids_per_batch"], kv_num_blocks_x_cpu=self.share_inputs["kv_num_blocks_x_cpu"], - routing_replay_table=routing_replay_table, ) dist_status = self.collect_distributed_status() @@ -1956,9 +1936,6 @@ def _dummy_run( if int((self.share_inputs["seq_lens_this_time"] > 0).sum()) == 0: break - if self.fd_config.routing_replay_config.enable_routing_replay: - self.routing_replay_manager.clear_routing_table() - def _update_chunked_prefill(self, tasks): """ Update chunked prefill related parameters @@ -2457,15 +2434,6 @@ class at the server level, which is too granular for ModelRunner. self.speculative_config.num_speculative_tokens, ) - # Routing replay - if self.fd_config.routing_replay_config.enable_routing_replay: - if ( - not self.exist_prefill() - and not self.exist_decode() - and self.share_inputs["is_block_step"].sum() == 0 - and self.share_inputs["is_chunk_step"].sum() == 0 - ): - self.routing_replay_manager.put_table_to_store() return None def _pool(self, hidden_states: paddle.Tensor, num_running_requests: int) -> Optional[ModelRunnerOutput]: diff --git a/fastdeploy/worker/worker_process.py b/fastdeploy/worker/worker_process.py index 0c29ce4d757..02d66f4bc53 100644 --- a/fastdeploy/worker/worker_process.py +++ b/fastdeploy/worker/worker_process.py @@ -38,7 +38,6 @@ ModelConfig, ParallelConfig, PlasAttentionConfig, - RoutingReplayConfig, SpeculativeConfig, StructuredOutputsConfig, ) @@ -886,13 +885,6 @@ def parse_args(): help="EPLB Configuration.", ) - parser.add_argument( - "--routing_replay_config", - type=json.loads, - default=None, - help="Configation of Rollout Routing Replay.", - ) - args = parser.parse_args() return args @@ -952,7 +944,6 @@ def initialize_fd_config(args, ranks: int = 1, local_rank: int = 0) -> FDConfig: eplb_config = EPLBConfig(args.eplb_config) structured_outputs_config: StructuredOutputsConfig = StructuredOutputsConfig(args=vars(args)) - routing_replay_config = RoutingReplayConfig(args.routing_replay_config) # Note(tangbinhan): used for load_checkpoint model_config.pretrained_config.tensor_parallel_rank = parallel_config.tensor_parallel_rank @@ -1012,7 +1003,6 @@ def initialize_fd_config(args, ranks: int = 1, local_rank: int = 0) -> FDConfig: plas_attention_config=plas_attention_config, structured_outputs_config=structured_outputs_config, eplb_config=eplb_config, - routing_replay_config=routing_replay_config, ) update_fd_config_for_mm(fd_config) if fd_config.load_config.load_choices == "default_v1" and not v1_loader_support(fd_config): diff --git a/tests/distributed/chunked_moe.py b/tests/distributed/chunked_moe.py index ef41a610dcb..0be645d38e2 100644 --- a/tests/distributed/chunked_moe.py +++ b/tests/distributed/chunked_moe.py @@ -90,7 +90,7 @@ def init_attention_metadata(self, forward_meta): class MockQuantMethod: - def apply(self, layer, x, gate, topk_ids_hookfunc=None): + def apply(self, layer, x, gate): return x @@ -129,7 +129,6 @@ def setup_model_runner(self): model_runner.speculative_decoding = False model_runner._init_share_inputs(mock_fd_config.scheduler_config.max_num_seqs) model_runner.share_inputs["caches"] = None - model_runner.routing_replay_manager = None if dist.get_rank() == 0: model_runner.share_inputs["ids_remove_padding"] = paddle.ones([10]) @@ -149,7 +148,6 @@ def setup_fused_moe(self): fused_moe.fd_config = mock_fd_config fused_moe.quant_method = MockQuantMethod() - fused_moe.enable_routing_replay = None return fused_moe def run_model_runner(self): diff --git a/tests/e2e/test_EB_Lite_serving.py b/tests/e2e/test_EB_Lite_serving.py index c71b7667260..bc27daab993 100644 --- a/tests/e2e/test_EB_Lite_serving.py +++ b/tests/e2e/test_EB_Lite_serving.py @@ -78,8 +78,6 @@ def setup_and_run_server(): "wint4", "--graph-optimization-config", '{"cudagraph_capture_sizes": [1], "use_cudagraph":true}', - "--routing-replay-config", - '{"enable_routing_replay":true, "routing_store_type":"local", "local_store_dir":"./routing_replay_output"}', ] # Start subprocess in new process group diff --git a/tests/layers/test_fusedmoe.py b/tests/layers/test_fusedmoe.py index 346afc98fd3..ed4fe5b28b6 100644 --- a/tests/layers/test_fusedmoe.py +++ b/tests/layers/test_fusedmoe.py @@ -31,7 +31,6 @@ LoadConfig, ModelConfig, ParallelConfig, - RoutingReplayConfig, ) from fastdeploy.model_executor.layers.moe.moe import FusedMoE from fastdeploy.model_executor.layers.quantization.block_wise_fp8 import ( @@ -477,7 +476,6 @@ def __init__( graph_opt_config=GraphOptimizationConfig({}), load_config=LoadConfig({}), ips=",".join(["0"] * nnodes), - routing_replay_config=RoutingReplayConfig({}), ) self.fd_config.parallel_config.tp_group = None self.fd_config.parallel_config.tensor_parallel_rank = tp_rank diff --git a/tests/layers/test_w4a8_moe.py b/tests/layers/test_w4a8_moe.py index f20c27b06bf..dc6dab15427 100644 --- a/tests/layers/test_w4a8_moe.py +++ b/tests/layers/test_w4a8_moe.py @@ -13,7 +13,6 @@ LoadConfig, ModelConfig, ParallelConfig, - RoutingReplayConfig, ) from fastdeploy.model_executor.layers.moe.moe import FusedMoE from fastdeploy.model_executor.layers.quantization.w4a8 import W4A8Config @@ -60,7 +59,6 @@ def __init__( graph_opt_config=GraphOptimizationConfig({}), load_config=LoadConfig({}), ips=",".join(["0"] * nnodes), - routing_replay_config=RoutingReplayConfig({}), ) self.fd_config.parallel_config.tp_group = None self.fd_config.parallel_config.tensor_parallel_rank = tp_rank diff --git a/tests/layers/test_w4afp8_moe.py b/tests/layers/test_w4afp8_moe.py index 8f1ae79cd67..65b7733172c 100644 --- a/tests/layers/test_w4afp8_moe.py +++ b/tests/layers/test_w4afp8_moe.py @@ -13,7 +13,6 @@ LoadConfig, ModelConfig, ParallelConfig, - RoutingReplayConfig, ) from fastdeploy.model_executor.layers.moe.moe import FusedMoE from fastdeploy.model_executor.layers.quantization.w4afp8 import W4AFP8Config @@ -66,7 +65,6 @@ def __init__( graph_opt_config=GraphOptimizationConfig({}), load_config=LoadConfig({}), ips=",".join(["0"] * nnodes), - routing_replay_config=RoutingReplayConfig({}), ) self.fd_config.parallel_config.tp_group = None self.fd_config.parallel_config.tensor_parallel_rank = tp_rank From daf3384da8e0ef133e7ad769026179eae917aea1 Mon Sep 17 00:00:00 2001 From: RAM Date: Fri, 5 Dec 2025 22:06:26 +0800 Subject: [PATCH 08/13] [New][RL] Support Rollout Routing Replay (#5405) * [RL] Support Rollout Routing Replay * add routing indices cache * fix config bug and moe forward bug * R3 Support GLM * support eb4.5 * fix merge bug * Apply suggestion from @Copilot Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Apply suggestion from @Copilot Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Apply suggestion from @Copilot Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Apply suggestion from @Copilot Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * add routing replay ci * support glm topk * support orther top_k * fix ci bug * pre-commit * only support chatcmpl * Revert "Revert "[RL] Support Rollout Routing Replay (#5321)" (#5402)" This reverts commit c45e064f3df5a84c98f5427d5fd8a9f6e7d26177. * Fix XPU and NPU bug --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> Co-authored-by: Yuanle Liu --- fastdeploy/config.py | 27 ++ fastdeploy/engine/args_utils.py | 22 ++ fastdeploy/engine/engine.py | 1 + fastdeploy/model_executor/forward_meta.py | 2 + .../backends/dcu/fused_moe_triton_backends.py | 5 + .../gcu/moe/fused_moe_method_gcu_backend.py | 6 + .../intel_hpu/moe/fused_moe_hpu_backend.py | 8 + .../moe/fused_moe_triton_metax_backend.py | 6 + .../layers/backends/xpu/moe/fused_moe.py | 8 + .../layers/moe/fused_moe_backend_base.py | 11 +- .../layers/moe/fused_moe_cutlass_backend.py | 17 + .../layers/moe/fused_moe_deepgemm_backend.py | 16 + .../layers/moe/fused_moe_marlin_backend.py | 6 + .../layers/moe/fused_moe_triton_backend.py | 18 + .../layers/moe/fused_moe_wint2_backend.py | 10 + fastdeploy/model_executor/layers/moe/moe.py | 60 ++- .../layers/moe/routing_indices_cache.py | 346 ++++++++++++++++++ fastdeploy/model_executor/models/glm4_moe.py | 7 +- fastdeploy/rl/rollout_config.py | 2 + fastdeploy/worker/gpu_model_runner.py | 32 ++ fastdeploy/worker/worker_process.py | 10 + tests/distributed/chunked_moe.py | 4 +- tests/e2e/test_EB_Lite_serving.py | 2 + tests/layers/test_fusedmoe.py | 2 + tests/layers/test_w4a8_moe.py | 2 + tests/layers/test_w4afp8_moe.py | 2 + 26 files changed, 608 insertions(+), 24 deletions(-) create mode 100644 fastdeploy/model_executor/layers/moe/routing_indices_cache.py diff --git a/fastdeploy/config.py b/fastdeploy/config.py index 63ac382d108..a820e8d94e6 100644 --- a/fastdeploy/config.py +++ b/fastdeploy/config.py @@ -1484,6 +1484,31 @@ def __str__(self) -> str: return json.dumps({key: value for key, value in self.__dict__.items()}) +class RoutingReplayConfig: + """Configuration for Routing Replay used in RL training""" + + def __init__(self, args) -> None: + self.enable_routing_replay: bool = False + self.routing_store_type: str = "local" + + # Local routing store + self.local_store_dir: str = "./routing_replay_output" + + # RDMA routing store + # TODO: Add RDMA routing store configuration attributes here when the feature is implemented. + + if args is not None: + for key, value in args.items(): + if hasattr(self, key) and value != "None": + setattr(self, key, value) + + def to_json_string(self): + """ + Convert routing replay config to json string. + """ + return json.dumps({key: value for key, value in self.__dict__.items()}) + + class FDConfig: """ The configuration class which contains all fastdeploy-related configuration. This @@ -1517,6 +1542,7 @@ def __init__( early_stop_config: Optional[Dict[str, Any]] = None, tool_parser: str = None, test_mode=False, + routing_replay_config: Optional[RoutingReplayConfig] = None, ): self.model_config: ModelConfig = model_config # type: ignore self.cache_config: CacheConfig = cache_config # type: ignore @@ -1533,6 +1559,7 @@ def __init__( self.plas_attention_config: Optional[PlasAttentionConfig] = plas_attention_config self.structured_outputs_config: StructuredOutputsConfig = structured_outputs_config self.router_config: RouterConfig = router_config + self.routing_replay_config = routing_replay_config # Initialize cuda graph capture list max_capture_shape = self.scheduler_config.max_num_seqs diff --git a/fastdeploy/engine/args_utils.py b/fastdeploy/engine/args_utils.py index 1eaf535498a..d2d7c6f908a 100644 --- a/fastdeploy/engine/args_utils.py +++ b/fastdeploy/engine/args_utils.py @@ -35,6 +35,7 @@ PlasAttentionConfig, PoolerConfig, RouterConfig, + RoutingReplayConfig, RunnerOption, SpeculativeConfig, StructuredOutputsConfig, @@ -491,6 +492,11 @@ class EngineArgs: Configuration for eplb. """ + routing_replay_config: Optional[Dict[str, Any]] = None + """ + Flag to rollout routing replay(r3) + """ + def __post_init__(self): """ Post-initialization processing to set default tokenizer if not provided. @@ -882,6 +888,12 @@ def add_cli_args(parser: FlexibleArgumentParser) -> FlexibleArgumentParser: default=EngineArgs.eplb_config, help="Config of eplb.", ) + parallel_group.add_argument( + "--routing-replay-config", + type=json.loads, + default=EngineArgs.routing_replay_config, + help="Flag of rollout routing replay(r3).", + ) parallel_group.add_argument( "--enable-chunked-moe", action="store_true", @@ -1235,6 +1247,14 @@ def create_eplb_config(self) -> EPLBConfig: eplb_args["enable_eplb"] = self.enable_eplb return EPLBConfig(eplb_args) + def create_routing_repaly_config(self) -> RoutingReplayConfig: + """ """ + routing_replay_args = asdict(self) + if self.routing_replay_config is not None: + for k, v in self.routing_replay_config.items(): + routing_replay_args[k] = v + return RoutingReplayConfig(routing_replay_args) + def create_engine_config(self, port_availability_check=True) -> FDConfig: """ Create and return a Config object based on the current settings. @@ -1278,6 +1298,7 @@ def create_engine_config(self, port_availability_check=True) -> FDConfig: graph_opt_cfg = self.create_graph_optimization_config() plas_attention_config = self.create_plas_attention_config() eplb_cfg = self.create_eplb_config() + routing_replay_config = self.create_routing_repaly_config() router_config = RouterConfig(all_dict) early_stop_cfg = self.create_early_stop_config() @@ -1310,4 +1331,5 @@ def create_engine_config(self, port_availability_check=True) -> FDConfig: graph_opt_config=graph_opt_cfg, plas_attention_config=plas_attention_config, early_stop_config=early_stop_cfg, + routing_replay_config=routing_replay_config, ) diff --git a/fastdeploy/engine/engine.py b/fastdeploy/engine/engine.py index 4a493843df7..fadf954679b 100644 --- a/fastdeploy/engine/engine.py +++ b/fastdeploy/engine/engine.py @@ -568,6 +568,7 @@ def _start_worker_service(self): f" --logprobs_mode {self.cfg.model_config.logprobs_mode}" f" --max_logprobs {self.cfg.model_config.max_logprobs}" f" --eplb_config '{self.cfg.eplb_config.to_json_string()}'" + f" --routing_replay_config '{self.cfg.routing_replay_config.to_json_string()}'" ) if self.cfg.structured_outputs_config.logits_processors is not None: arguments += f" --logits-processors {' '.join(self.cfg.structured_outputs_config.logits_processors)}" diff --git a/fastdeploy/model_executor/forward_meta.py b/fastdeploy/model_executor/forward_meta.py index 4e9df0d3ce3..787ec77c0eb 100644 --- a/fastdeploy/model_executor/forward_meta.py +++ b/fastdeploy/model_executor/forward_meta.py @@ -142,6 +142,8 @@ class ForwardMeta: caches: Optional[list[paddle.Tensor]] = None # Flag of profile run is_dummy_or_profile_run: bool = False + # Routing Replay table buffer + routing_replay_table: Optional[paddle.Tensor] = None # chunked MoE related moe_num_chunk: int = 1 diff --git a/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py b/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py index 918450c74f1..192c0b8833a 100644 --- a/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py +++ b/fastdeploy/model_executor/layers/backends/dcu/fused_moe_triton_backends.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -101,6 +103,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -117,6 +120,8 @@ def apply( scores += layer.gate_correction_bias topk_weights, topk_ids = paddle.topk(scores, k=top_k, axis=-1, sorted=False) topk_weights = topk_weights / topk_weights.sum(axis=-1, keepdim=True) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) intermediate_cache1 = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], diff --git a/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py b/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py index e67dd6dbdaf..2260d7caf7b 100644 --- a/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py +++ b/fastdeploy/model_executor/layers/backends/gcu/moe/fused_moe_method_gcu_backend.py @@ -16,6 +16,7 @@ import multiprocessing import os +from typing import Callable import numpy as np import paddle @@ -182,6 +183,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle gcu compute Fused MoE. @@ -194,6 +196,7 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -205,6 +208,7 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -216,6 +220,7 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -381,6 +386,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle gcu compute Fused MoE. diff --git a/fastdeploy/model_executor/layers/backends/intel_hpu/moe/fused_moe_hpu_backend.py b/fastdeploy/model_executor/layers/backends/intel_hpu/moe/fused_moe_hpu_backend.py index d47bfc86b93..8e4d7b1cc5e 100644 --- a/fastdeploy/model_executor/layers/backends/intel_hpu/moe/fused_moe_hpu_backend.py +++ b/fastdeploy/model_executor/layers/backends/intel_hpu/moe/fused_moe_hpu_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -48,6 +50,7 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate_out: paddle.Tensor, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -59,6 +62,7 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate_out: paddle.Tensor, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -70,6 +74,7 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle hpu Fused MoE. @@ -142,6 +147,7 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate_out: paddle.Tensor, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -153,6 +159,7 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate_out: paddle.Tensor, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -164,6 +171,7 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle hpu Fused MoE. diff --git a/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py b/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py index 7b61d58b6f5..fbbfac277b8 100644 --- a/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py +++ b/fastdeploy/model_executor/layers/backends/metax/moe/fused_moe_triton_metax_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -245,6 +247,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -274,6 +277,9 @@ def apply( True, # apply_norm_weight False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, diff --git a/fastdeploy/model_executor/layers/backends/xpu/moe/fused_moe.py b/fastdeploy/model_executor/layers/backends/xpu/moe/fused_moe.py index 3a14e28e305..4356f8cc442 100644 --- a/fastdeploy/model_executor/layers/backends/xpu/moe/fused_moe.py +++ b/fastdeploy/model_executor/layers/backends/xpu/moe/fused_moe.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -235,6 +237,7 @@ def apply_tp_fused_op( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply TP Fused Op. @@ -262,6 +265,7 @@ def apply_tp_scatter_op( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply TP Scatter Op. @@ -318,6 +322,7 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ apply tp @@ -368,6 +373,7 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -442,6 +448,7 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -488,6 +495,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ compute Fused MoE. diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py b/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py index b34291a96f4..a8bd70465ea 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_backend_base.py @@ -15,6 +15,7 @@ """ from abc import abstractmethod +from typing import Callable import paddle from paddle import nn @@ -163,6 +164,7 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -175,6 +177,7 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -187,6 +190,7 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -198,6 +202,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -207,13 +212,13 @@ def apply( if layer.fd_config.model_config.moe_phase.phase == "prefill": if layer.fd_config.scheduler_config.splitwise_role == "mixed" and is_moe_start_layer: self.ep_prefill_runner.clean_low_latency_buffer() - return self.apply_ep_prefill(layer, x, gate) + return self.apply_ep_prefill(layer, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) else: if layer.fd_config.scheduler_config.splitwise_role == "mixed" and is_moe_start_layer: self.ep_decoder_runner.clean_low_latency_buffer() - return self.apply_ep_decode(layer, x, gate) + return self.apply_ep_decode(layer, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) else: - return self.apply_tp(layer, x, gate) + return self.apply_tp(layer, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) class UnquantizedFusedMoEMethod(MoEMethodBase): diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py index e45ad63b19c..c3dbfc9ba5f 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_cutlass_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn from paddle.nn.quant import weight_quantize @@ -132,6 +134,7 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -148,8 +151,13 @@ def apply_ep_prefill( handle, event, ) = self.ep_prefill_runner.dispatch(x, topk_idx, topk_weights) + + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + if self.ep_prefill_runner.ep_engine.async_finish: event.current_stream_wait() + token_all_num = sum(recv_num_tokens_per_expert_list) # 3. Compute ffn @@ -217,6 +225,7 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -225,6 +234,10 @@ def apply_ep_decode( estimate_total_token_nums = gate_out.shape[0] * layer.top_k # 1. Select topk experts and weights topk_idx, topk_weights = self.ep_decoder_runner.moe_select(layer, gate_out) + + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + expertwise_scale = None if hasattr(layer, "up_gate_proj_in_scale_all_experts"): # only use in w4a8 expertwise_scale = getattr(layer, "up_gate_proj_in_scale_all_experts", None) @@ -269,6 +282,7 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Cutlass compute Fused MoE. @@ -369,6 +383,9 @@ def apply_tp( if hasattr(layer, "up_gate_proj_in_scale"): dequant_scale = None + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + if not layer.with_bias and self.moe_quant_type != "w4a8" and self.moe_quant_type != "w4afp8": # only w4a8 need expert_idx_per_token # Other need not this tensor, so we make it None. diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py index 1245cddcebc..881f9a22c4d 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn from paddle.distributed.communication import deep_ep @@ -139,6 +141,7 @@ def apply_ep_prefill( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP prefill method. @@ -147,6 +150,10 @@ def apply_ep_prefill( # 1. Select topk experts and weights topk_idx, topk_weights = self.ep_prefill_runner.moe_select(layer, gate_out) + + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + # 2. Dynamic compute blockwise quantization scales x, x_scale_tensor = fastdeploy.model_executor.ops.gpu.per_token_quant( x, self.quant_config.weight_block_size[0] @@ -264,6 +271,7 @@ def apply_ep_decode( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Apply the EP decoder method. @@ -271,6 +279,10 @@ def apply_ep_decode( gate_out = gate(x.cast("float32")) # 1. Select topk experts and weights topk_idx, topk_weights = self.ep_decoder_runner.moe_select(layer, gate_out) + + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + # 2. EP Dispatch permute_input, token_nums_per_expert, handle = self.ep_decoder_runner.dispatch( x, topk_idx, topk_weights, use_fp8=True @@ -335,6 +347,7 @@ def apply_tp( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Paddle Use DeepGemm compute Fused MoE. @@ -363,6 +376,9 @@ def apply_tp( False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + tmp = count_tokens_per_expert_func(topk_ids, layer.num_experts) recv_x, recv_x_scale = fastdeploy.model_executor.ops.gpu.per_token_quant(x, 128) diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py index 094d3df8f1a..cd836dbaf09 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_marlin_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -239,6 +241,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Marlin compute Fused MoE. @@ -273,6 +276,9 @@ def apply( False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + block_size_m = 64 for m in [8, 16, 32, 48, 64]: diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py index 3c14859375c..2861d96e8d3 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_triton_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -282,6 +284,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -314,6 +317,10 @@ def apply( True, # apply_norm_weight, False, ) + + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, @@ -664,6 +671,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -724,6 +732,9 @@ def apply( * ceil_div(moe_intermediate_size * 2, config["BLOCK_SIZE_N"]), ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, @@ -953,6 +964,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -974,6 +986,9 @@ def apply( False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + up_gate_proj_out = paddle.empty( [token_num * top_k, moe_intermediate_size * 2], dtype=x.dtype, @@ -1466,6 +1481,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Triton compute Fused MoE. @@ -1488,6 +1504,8 @@ def apply( True, # apply_norm_weight False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) config = { "BLOCK_SIZE_M": 64, diff --git a/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py b/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py index f75e36bcbdd..3c548ba57c8 100644 --- a/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py +++ b/fastdeploy/model_executor/layers/moe/fused_moe_wint2_backend.py @@ -14,6 +14,8 @@ # limitations under the License. """ +from typing import Callable + import paddle from paddle import nn @@ -261,6 +263,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Use Wint2 Triton Fusedmoe compute Fused MoE. @@ -288,6 +291,9 @@ def apply( topk_only_mode=False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_idx) + ffn_out = fastdeploy.model_executor.ops.gpu.moe_expert_ffn_wint2( permute_input, token_nums_per_expert, @@ -328,6 +334,7 @@ def apply( layer: nn.Layer, x: paddle.Tensor, gate: nn.Layer, + topk_ids_hookfunc: Callable = None, ) -> paddle.Tensor: """ Use Wint2 Triton Fusedmoe compute Fused MoE. @@ -343,6 +350,9 @@ def apply( False, ) + if topk_ids_hookfunc is not None: + topk_ids_hookfunc(topk_ids=topk_ids) + num_tokens, K = x.shape E, _, N = layer.up_gate_proj_weight.shape M = num_tokens diff --git a/fastdeploy/model_executor/layers/moe/moe.py b/fastdeploy/model_executor/layers/moe/moe.py index 743e05031f6..5b1be52d183 100644 --- a/fastdeploy/model_executor/layers/moe/moe.py +++ b/fastdeploy/model_executor/layers/moe/moe.py @@ -14,7 +14,8 @@ # limitations under the License. """ -from typing import Optional +from functools import partial +from typing import Callable, Optional import paddle from paddle import nn @@ -26,6 +27,9 @@ tensor_model_parallel_all_reduce_custom, ) from fastdeploy.model_executor.forward_meta import ForwardMeta +from fastdeploy.model_executor.layers.moe.routing_indices_cache import ( + save_routing_to_buffer, +) from fastdeploy.model_executor.layers.utils import get_tensor from fastdeploy.model_executor.utils import h2d_copy, slice_fn from fastdeploy.platforms import current_platform @@ -226,7 +230,7 @@ def __init__( self.is_rearrange = False if self.ep_size > 1: self.quant_method.init_ep(self) - + self.enable_routing_replay = fd_config.routing_replay_config.enable_routing_replay # Merge normal and RL build model if gate_correction_bias is not None: self.gate_correction_bias = gate_correction_bias @@ -600,7 +604,7 @@ def load_state_dict(self, state_dict, is_rearrange: bool = False): else: self.quant_method.process_loaded_weights(self, state_dict) - def forward_split_allgather(self, x: paddle.Tensor, gate: nn.Layer): + def forward_split_allgather(self, x: paddle.Tensor, gate: nn.Layer, topk_ids_hookfunc: Callable = None): """ Forward split allgather function. """ @@ -615,14 +619,14 @@ def forward_split_allgather(self, x: paddle.Tensor, gate: nn.Layer): if end_offset > token_num: end_offset = token_num part_x[: (end_offset - start_offset), :] = x[start_offset:end_offset, :] - out = self.quant_method.apply(self, part_x, gate) + out = self.quant_method.apply(self, part_x, gate, topk_ids_hookfunc=topk_ids_hookfunc) multi_outs = paddle.zeros([token_num_per_rank * self.attn_tp_size, x.shape[1]], dtype=x.dtype) paddle.distributed.all_gather(multi_outs, out, self.tp_group) out = multi_outs[:token_num, :] return out - def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): + def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta = None): """ Defines the forward computation of the moe layer. @@ -633,6 +637,21 @@ def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): Tensor: Output tensor.s """ + topk_ids_hookfunc = None + if self.enable_routing_replay: + if forward_meta is not None: # forward_meta is None when execute empty_input_forward + topk_ids_hookfunc = partial( + save_routing_to_buffer, + routing_replay_table=forward_meta.routing_replay_table, + batch_id_per_token=forward_meta.batch_id_per_token, + seq_lens_decoder=forward_meta.seq_lens_decoder, + cu_seqlens_q=forward_meta.cu_seqlens_q, + layer_idx=self.layer_idx, + tp_size=self.fd_config.parallel_config.tensor_parallel_size, + ep_size=self.fd_config.parallel_config.expert_parallel_size, + tp_group=self.fd_config.parallel_config.tp_group, + ) + token_num = x.shape[0] if ( self.ep_size > 1 @@ -640,11 +659,16 @@ def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): and (not self.fd_config.parallel_config.use_sequence_parallel_moe) and token_num >= self.attn_tp_size ): - out = self.forward_split_allgather(x, gate) + out = self.forward_split_allgather(x, gate, topk_ids_hookfunc=topk_ids_hookfunc) elif self.fd_config.parallel_config.use_ep and self.fd_config.parallel_config.enable_chunked_moe: - out = self.forward_chunked_moe(x, gate, forward_meta) + out = self.forward_chunked_moe( + x, + gate, + forward_meta, + topk_ids_hookfunc=topk_ids_hookfunc, + ) else: - out = self.forward_normal(x, gate) + out = self.forward_normal(x, gate, forward_meta, topk_ids_hookfunc=topk_ids_hookfunc) if self.reduce_results and self.tp_size > 1: if current_platform.is_intel_hpu(): @@ -653,7 +677,9 @@ def forward(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): out = tensor_model_parallel_all_reduce(out, self.tp_group) return out - def forward_chunked_moe(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta): + def forward_chunked_moe( + self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta, topk_ids_hookfunc: Callable = None + ): """ Split input to multi chunk to reduce the memory usage of moe. @@ -677,21 +703,25 @@ def forward_chunked_moe(self, x: paddle.Tensor, gate: nn.Layer, forward_meta: Fo for i in range(forward_meta.max_moe_num_chunk): if i < forward_meta.moe_num_chunk: - out_split_list[i] = self.quant_method.apply(self, x_split_list[i], gate) + out_split_list[i] = self.quant_method.apply( + self, x_split_list[i], gate, topk_ids_hookfunc=topk_ids_hookfunc + ) else: # just need to use real data to infer max_moe_num_chunk times. - self.quant_method.apply(self, fake_x, gate) + self.quant_method.apply(self, fake_x, gate, topk_ids_hookfunc=topk_ids_hookfunc) out = paddle.concat(out_split_list, axis=0) else: # when only one chunk, just need to use real data to infer once. - out = self.quant_method.apply(self, x, gate) + out = self.quant_method.apply(self, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) for i in range(forward_meta.max_moe_num_chunk - 1): - self.quant_method.apply(self, fake_x, gate) + self.quant_method.apply(self, fake_x, gate, topk_ids_hookfunc=topk_ids_hookfunc) return out - def forward_normal(self, x: paddle.Tensor, gate: nn.Layer): + def forward_normal( + self, x: paddle.Tensor, gate: nn.Layer, forward_meta: ForwardMeta, topk_ids_hookfunc: Callable = None + ): """ Normal mode of forward. @@ -702,5 +732,5 @@ def forward_normal(self, x: paddle.Tensor, gate: nn.Layer): Tensor: Output tensor.s """ - out = self.quant_method.apply(self, x, gate) + out = self.quant_method.apply(self, x, gate, topk_ids_hookfunc=topk_ids_hookfunc) return out diff --git a/fastdeploy/model_executor/layers/moe/routing_indices_cache.py b/fastdeploy/model_executor/layers/moe/routing_indices_cache.py new file mode 100644 index 00000000000..e95a3d8569f --- /dev/null +++ b/fastdeploy/model_executor/layers/moe/routing_indices_cache.py @@ -0,0 +1,346 @@ +""" +# Copyright (c) 2025 PaddlePaddle Authors. All Rights Reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +""" + +import copy +import os +import shutil +from abc import ABC, abstractmethod +from typing import Dict, List, Optional + +import paddle +import paddle.distributed as dist +import triton +import triton.language as tl + +from fastdeploy.config import FDConfig + + +@triton.jit +def _save_routing_kernel( + ROUTING_REPLAY_TABLE_PTR, + TOPK_IDS_PTR, + BATCH_ID_PER_TOKEN_PTR, + CU_SEQLENS_Q_PTR, + SEQ_LENS_DECODER_PTR, + LAYER_IDX, + TOKEN_NUM, + TOP_K, + NUM_HIDDEN_LAYERS, + MAX_MODEL_LEN, + BLOCK_SIZE_M: tl.constexpr, + BLOCK_SIZE_K: tl.constexpr, +): + pid_m = tl.program_id(axis=0) + + token_offsets = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) + token_mask = token_offsets < TOKEN_NUM + + k_offsets = tl.arange(0, BLOCK_SIZE_K) + + k_mask = k_offsets < TOP_K + + topk_ids_ptrs = TOPK_IDS_PTR + token_offsets[:, None] * TOP_K + k_offsets[None, :] + # [BLOCK_SIZE_M, BLOCK_SIZE_K] + + load_mask = token_mask[:, None] & k_mask[None, :] + topk_vals = tl.load(topk_ids_ptrs, mask=load_mask) + + batch_ids = tl.load(BATCH_ID_PER_TOKEN_PTR + token_offsets, mask=token_mask) + pad_mask = token_mask & (batch_ids != -1) + # [0, 3, 4, 10, 12][0, 0, 0, 0, 2, 2, 2, 2, 2, 2, 3, 3] + # -> [0, 0, 0, 0, 4, 4, 4, 4, 4, 4, 10, 10] + # [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11] - [0, 0, 0, 0, 4, 4, 4, 4, 4, 4, 10, 10] + # -> [0, 1, 2, 3, 0, 1, 2, 3, 4, 5, 0, 1] + start_offsets = tl.load(CU_SEQLENS_Q_PTR + batch_ids, mask=pad_mask) + token_relative_index = token_offsets - start_offsets + + # [BLOCK_SIZE_M] + len_decoder = tl.load(SEQ_LENS_DECODER_PTR + batch_ids, mask=pad_mask) + token_seq_pos = len_decoder + token_relative_index + + STRIDE_BUF_SEQ = NUM_HIDDEN_LAYERS * MAX_MODEL_LEN * TOP_K + STRIDE_BUF_LAYER = MAX_MODEL_LEN * TOP_K + STRIDE_BUF_TOKEN = TOP_K + + # [BLOCK_SIZE_M, BLOCK_SIZE_K] + output_ptrs = ( + ROUTING_REPLAY_TABLE_PTR + + batch_ids[:, None] * STRIDE_BUF_SEQ + + LAYER_IDX * STRIDE_BUF_LAYER + + token_seq_pos[:, None] * STRIDE_BUF_TOKEN + + k_offsets[None, :] + ) + + pos_mask = token_seq_pos < MAX_MODEL_LEN + pos_mask = pos_mask & pad_mask + + # [BLOCK_SIZE_M, BLOCK_SIZE_K] + pos_mask = pos_mask[:, None] & k_mask[None, :] + + final_mask = load_mask & pos_mask + + tl.store(output_ptrs, topk_vals, mask=final_mask) + + +def save_routing_to_buffer( + routing_replay_table: paddle.Tensor, # [max_num_seqs, num_layers, max_len, top_k] + topk_ids: paddle.Tensor, # [token_num, top_k] + batch_id_per_token: paddle.Tensor, # [token_num, 1] + seq_lens_decoder: paddle.Tensor, # [max_num_seqs, 1] + cu_seqlens_q: paddle.Tensor, # [max_num_seqs + 1, 1] + layer_idx: int, + tp_size: int, + ep_size: int, + tp_group: dist.communication.group.Group, +): + if tp_size > 1 and ep_size > 1: + token_num_per_rank = topk_ids.shape[0] + topk_ids_all = paddle.zeros([token_num_per_rank * tp_size, topk_ids.shape[1]], dtype=topk_ids.dtype) + paddle.distributed.all_gather(topk_ids_all, topk_ids, tp_group) + topk_ids = topk_ids_all[: batch_id_per_token.shape[0], :] + + token_num, top_k = topk_ids.shape + max_num_seqs, num_hidden_layers, max_model_len, _ = routing_replay_table.shape + assert token_num > 0 + + assert topk_ids.shape[1] == routing_replay_table.shape[3], (topk_ids.shape[1], routing_replay_table.shape[3]) + assert batch_id_per_token.shape[0] == token_num, (batch_id_per_token.shape[0], token_num) + assert seq_lens_decoder.shape[0] == max_num_seqs, (seq_lens_decoder.shape[0], max_num_seqs) + + BLOCK_SIZE_M = 128 + BLOCK_SIZE_K = triton.next_power_of_2(top_k) # top_k + + grid = (triton.cdiv(token_num, BLOCK_SIZE_M),) + _save_routing_kernel[grid]( + routing_replay_table, + topk_ids, + batch_id_per_token, + cu_seqlens_q, + seq_lens_decoder, + LAYER_IDX=layer_idx, + TOKEN_NUM=token_num, + TOP_K=top_k, + NUM_HIDDEN_LAYERS=num_hidden_layers, + MAX_MODEL_LEN=max_model_len, + BLOCK_SIZE_M=BLOCK_SIZE_M, + BLOCK_SIZE_K=BLOCK_SIZE_K, + ) + + +class RoutingReplayManager: + """Request level routing replay table manager""" + + def __init__( + self, + fd_config: FDConfig, + ): + self.max_num_seqs = fd_config.scheduler_config.max_num_seqs + self.max_model_len = fd_config.model_config.max_model_len + self.num_moe_layers = fd_config.model_config.num_hidden_layers - fd_config.model_config.moe_layer_start_index + + if fd_config.model_config.architectures[0] == "Glm4MoeForCausalLM": + self.moe_top_k = fd_config.model_config.num_experts_per_tok + else: + self.moe_top_k = fd_config.model_config.moe_k + self.tp_rank = fd_config.parallel_config.tensor_parallel_rank + + self.routing_store = get_routing_store(fd_config=fd_config) + self.routing_batch_to_request: Dict[int, str] = {} + self.routing_replay_table = paddle.full( + shape=[self.max_num_seqs, self.num_moe_layers, self.max_model_len, self.moe_top_k], + fill_value=-1, + dtype="int32", + ) + + def register_request(self, batch_id: int, request_id: str): + """ + Register a new request to routing replay table + Args: + batch_id: The batch ID of this request + request_id: The global ID of the request is usually executed by the training process in RL + """ + # Save requests that have been finished for the current slot + if batch_id in self.routing_batch_to_request: + pre_request_id = self._deregister_request(batch_id) + self._put_request_to_store(batch_id, pre_request_id) + # Register the new request + self.routing_batch_to_request[batch_id] = request_id + + def _deregister_request(self, batch_id: int) -> str: + """ + Deregister a request from routing replay table + """ + assert batch_id in self.routing_batch_to_request + return self.routing_batch_to_request.pop(batch_id) + + def _put_request_to_store( + self, + batch_id: int, + request_id: str, + ): + if self.tp_rank == 0: + batch_buffer = self.routing_replay_table[batch_id] + for layer_id in range(self.num_moe_layers): + layer_buffer = batch_buffer[layer_id] + rollout_id = self.split_request_id(request_id) + self.routing_store.put(routing_indices=layer_buffer, rollout_id=rollout_id, layer_idx=layer_id) + + self._clear_table_slot(batch_id) + + def put_table_to_store(self): + """Put the routing table""" + batch_ids = copy.deepcopy(list(self.routing_batch_to_request.keys())) + for batch_id in batch_ids: + request_id = self._deregister_request(batch_id) + self._put_request_to_store(batch_id, request_id) + + def _clear_table_slot(self, batch_id: int): + assert 0 <= batch_id < self.max_num_seqs + self.routing_replay_table[batch_id].fill_(-1) + + def clear_routing_table(self): + """Clear all slots of the routing replay table""" + self.routing_replay_table.fill_(-1) + + def _clear_store(self): + """Clear routing store""" + self.routing_store.clear_store() + + def _clear_request_of_store(self, request_id): + """Clear one request of routing store""" + rollout_id = self.split_request_id(request_id) + for layer_idx in range(self.num_moe_layers): + self.routing_store.clear(rollout_id=rollout_id, layer_idx=layer_idx) + + def get_request_from_store(self, request_id: str) -> List[paddle.Tensor]: + """Get the routing indices of the request from store""" + routing_list = [] + rollout_id = self.split_request_id(request_id) + for layer_idx in range(self.num_moe_layers): + one_layer_routing = self.routing_store.get(rollout_id, layer_idx) + routing_list.append(one_layer_routing) + + return routing_list + + def get_routing_table(self) -> paddle.Tensor: + return self.routing_replay_table + + def split_request_id(self, request_id: str): + """Split the request id to get rollout id""" + chat_type, tmp_str = request_id.split("-", 1) + # NOTE(gongshaotian): only support chatcmpl now + # assert chat_type == "chatcmpl" + reversed_tmp_str = tmp_str[::-1].split("-", 5) + rollout_id = reversed_tmp_str[-1][::-1] + return rollout_id + + +class RoutingStoreBase(ABC): + """Base class for routing store""" + + def __init__(self, fd_config: FDConfig) -> None: + self.fd_config = fd_config + + @abstractmethod + def put(self, routing_indices: paddle.Tensor, rollout_id: str, layer_idx: Optional[int] = None) -> None: + """Put the routing indices into store""" + raise NotImplementedError + + @abstractmethod + def get(self, rollout_id: str, layer_idx: Optional[int] = None) -> paddle.Tensor: + """Get the routing indices from store""" + raise NotImplementedError + + @abstractmethod + def clear(self, rollout_id: str, layer_idx: Optional[int] = None) -> None: + """Clear the routing indices of the request""" + raise NotImplementedError + + @abstractmethod + def clear_store( + self, + ): + """Clear the routing indices store""" + raise NotImplementedError + + +class RoutingStoreLocal(RoutingStoreBase): + """Routing Store using local memory""" + + def __init__(self, fd_config) -> None: + super().__init__(fd_config=fd_config) + self.local_store_dir = fd_config.routing_replay_config.local_store_dir + + def put(self, routing_indices: paddle.Tensor, rollout_id: str, layer_idx: int) -> None: + """Put the routing indices into store""" + dir_path = os.path.join(self.local_store_dir, f"{rollout_id}") + os.makedirs(dir_path, exist_ok=True) + file_path = os.path.join(dir_path, f"layer_{layer_idx}.pdtensor") + paddle.save(routing_indices, file_path) + + def get( + self, + rollout_id: str, + layer_idx: int = None, + ) -> paddle.Tensor: + """Get the routing indices from store""" + dir_path = os.path.join(self.local_store_dir, f"{rollout_id}") + file_path = os.path.join(dir_path, f"layer_{layer_idx}.pdtensor") + assert os.path.exists(file_path), f"File not found: {file_path}" + layer_routing_indices = paddle.load(file_path) + + return layer_routing_indices + + def clear( + self, + rollout_id: str, + layer_idx: int = None, + ) -> None: + """Clear the routing indices of the request""" + dir_path = os.path.join(self.local_store_dir, f"{rollout_id}") + file_path = os.path.join(dir_path, f"layer_{layer_idx}.pdtensor") + assert os.path.exists(file_path), f"File not found: {file_path}" + os.remove(file_path) + + # Delete empty directory + if len(os.listdir(dir_path)) == 0: + os.rmdir(dir_path) + + def clear_store(self): + """Clear the routing indices store""" + if os.path.isdir(self.local_store_dir): + for file_name in os.listdir(self.local_store_dir): + file_path = os.path.join(self.local_store_dir, file_name) + shutil.rmtree(file_path) + + +class RoutingStoreRDMA(RoutingStoreBase): + """Routing Store using RDMA""" + + def __init__(self) -> None: + super().__init__() + + +def get_routing_store(fd_config: FDConfig) -> RoutingStoreBase: + if fd_config.routing_replay_config.routing_store_type == "local": + return RoutingStoreLocal(fd_config=fd_config) + elif fd_config.routing_replay_config.routing_store_type == "rdma": + return RoutingStoreRDMA(fd_config=fd_config) + else: + raise ValueError( + f"Invalid routing store type: '{fd_config.routing_replay_config.routing_store_type}'. " + "Valid types are: 'local', 'rdma'" + ) diff --git a/fastdeploy/model_executor/models/glm4_moe.py b/fastdeploy/model_executor/models/glm4_moe.py index d5ad6e3916b..0cc7c4dae45 100644 --- a/fastdeploy/model_executor/models/glm4_moe.py +++ b/fastdeploy/model_executor/models/glm4_moe.py @@ -161,7 +161,7 @@ def __init__( reduce_results=False, ) - def forward(self, x, forward_meta): + def forward(self, x, forward_meta: ForwardMeta = None): shared_experts_out = self.shared_experts(x) out = self.experts(x, self.gate, forward_meta) out = out + shared_experts_out @@ -306,10 +306,7 @@ def forward( # Fully Connected hidden_states, residual = self.post_attention_layernorm(hidden_states, residual) - hidden_states = self.mlp( - hidden_states, - forward_meta, - ) + hidden_states = self.mlp(hidden_states, forward_meta) return hidden_states, residual diff --git a/fastdeploy/rl/rollout_config.py b/fastdeploy/rl/rollout_config.py index 6bd3c3bcb35..f7ff748fed7 100644 --- a/fastdeploy/rl/rollout_config.py +++ b/fastdeploy/rl/rollout_config.py @@ -65,6 +65,7 @@ def __init__( data_parallel_size: int = 1, num_nextn_predict_layers: int = 0, eplb_config: str = {}, + routing_replay_config: str = None, ): # Required parameters self.model = model_name_or_path @@ -113,6 +114,7 @@ def __init__( self.plas_attention_config = plas_attention_config self.num_nextn_predict_layers = num_nextn_predict_layers self.eplb_config = eplb_config + self.routing_replay_config = routing_replay_config def __str__(self): return "\n".join(f"{k}: {v}" for k, v in self.__dict__.items()) diff --git a/fastdeploy/worker/gpu_model_runner.py b/fastdeploy/worker/gpu_model_runner.py index 2d5f47def46..9b550f10438 100644 --- a/fastdeploy/worker/gpu_model_runner.py +++ b/fastdeploy/worker/gpu_model_runner.py @@ -45,6 +45,9 @@ from fastdeploy.model_executor.layers.attention.base_attention_backend import ( AttentionBackend, ) +from fastdeploy.model_executor.layers.moe.routing_indices_cache import ( + RoutingReplayManager, +) from fastdeploy.model_executor.layers.rotary_embedding import get_rope, get_rope_3d from fastdeploy.model_executor.layers.sample.meta_data import SamplingMetadata from fastdeploy.model_executor.layers.sample.sampler import Sampler, SpeculativeSampler @@ -202,6 +205,11 @@ def __init__( os.environ["INFERENCE_MSG_QUEUE_ID"] = str(self.parallel_config.engine_worker_queue_port) logger.info(f"queue id is {str(self.parallel_config.engine_worker_queue_port)}") + # Rollout routing replay config + self.routing_replay_manager = None + if self.fd_config.routing_replay_config.enable_routing_replay: + self.routing_replay_manager = RoutingReplayManager(fd_config=self.fd_config) + self.zmq_client = None self.async_output_queue = None if envs.FD_USE_GET_SAVE_OUTPUT_V1: @@ -648,6 +656,7 @@ def insert_tasks_v1(self, req_dicts: List[Request], num_running_requests: int = self.share_inputs["step_seq_lens_decoder"][idx : idx + 1] = 0 self.share_inputs["prompt_lens"][idx : idx + 1] = len(input_ids) self.share_inputs["is_block_step"][idx : idx + 1] = False + self.share_inputs["is_chunk_step"][idx : idx + 1] = prefill_end_index < len(input_ids) self.share_inputs["step_idx"][idx : idx + 1] = ( len(request.output_token_ids) if prefill_end_index >= len(input_ids) else 0 ) @@ -656,6 +665,12 @@ def insert_tasks_v1(self, req_dicts: List[Request], num_running_requests: int = if request.sampling_params is not None and request.sampling_params.prompt_logprobs is not None: self.prompt_logprobs_reqs[request.request_id] = request has_prefill_task = True + + # Routing Replay + if self.fd_config.routing_replay_config.enable_routing_replay: + if prefill_start_index == 0: + self.routing_replay_manager.register_request(batch_id=idx, request_id=request.request_id) + if ( self.fd_config.scheduler_config.splitwise_role == "decode" ): # In PD, we continue to decode after P generate first token @@ -1152,6 +1167,7 @@ def _init_share_inputs(self, max_num_seqs: int): self.share_inputs["bad_tokens_len"] = paddle.full([max_num_seqs], 1, dtype="int64") self.share_inputs["next_tokens"] = paddle.full([max_num_seqs, 1], -1, dtype="int64") self.share_inputs["is_block_step"] = paddle.full([max_num_seqs], False, dtype="bool") + self.share_inputs["is_chunk_step"] = paddle.full([max_num_seqs], False, dtype="bool").cpu() self.share_inputs["encoder_block_lens"] = paddle.full([max_num_seqs], 0, dtype="int32") self.share_inputs["step_block_list"] = paddle.full([max_num_seqs], -1, dtype="int32") self.share_inputs["step_lens"] = paddle.full([1], 0, dtype="int32") @@ -1422,6 +1438,9 @@ def initialize_forward_meta(self, is_dummy_or_profile_run=False): Initialize forward meta, attention meta data and update some config. """ # Initialize forward meta + routing_replay_table = None + if self.routing_replay_manager is not None: + routing_replay_table = self.routing_replay_manager.get_routing_table() self.forward_meta = ForwardMeta( ids_remove_padding=self.share_inputs["ids_remove_padding"], rotary_embs=self.share_inputs["rope_emb"], @@ -1448,6 +1467,7 @@ def initialize_forward_meta(self, is_dummy_or_profile_run=False): kv_batch_ids=self.share_inputs["kv_batch_ids"], kv_tile_ids_per_batch=self.share_inputs["kv_tile_ids_per_batch"], kv_num_blocks_x_cpu=self.share_inputs["kv_num_blocks_x_cpu"], + routing_replay_table=routing_replay_table, ) dist_status = self.collect_distributed_status() @@ -1936,6 +1956,9 @@ def _dummy_run( if int((self.share_inputs["seq_lens_this_time"] > 0).sum()) == 0: break + if self.fd_config.routing_replay_config.enable_routing_replay: + self.routing_replay_manager.clear_routing_table() + def _update_chunked_prefill(self, tasks): """ Update chunked prefill related parameters @@ -2434,6 +2457,15 @@ class at the server level, which is too granular for ModelRunner. self.speculative_config.num_speculative_tokens, ) + # Routing replay + if self.fd_config.routing_replay_config.enable_routing_replay: + if ( + not self.exist_prefill() + and not self.exist_decode() + and self.share_inputs["is_block_step"].sum() == 0 + and self.share_inputs["is_chunk_step"].sum() == 0 + ): + self.routing_replay_manager.put_table_to_store() return None def _pool(self, hidden_states: paddle.Tensor, num_running_requests: int) -> Optional[ModelRunnerOutput]: diff --git a/fastdeploy/worker/worker_process.py b/fastdeploy/worker/worker_process.py index 02d66f4bc53..0c29ce4d757 100644 --- a/fastdeploy/worker/worker_process.py +++ b/fastdeploy/worker/worker_process.py @@ -38,6 +38,7 @@ ModelConfig, ParallelConfig, PlasAttentionConfig, + RoutingReplayConfig, SpeculativeConfig, StructuredOutputsConfig, ) @@ -885,6 +886,13 @@ def parse_args(): help="EPLB Configuration.", ) + parser.add_argument( + "--routing_replay_config", + type=json.loads, + default=None, + help="Configation of Rollout Routing Replay.", + ) + args = parser.parse_args() return args @@ -944,6 +952,7 @@ def initialize_fd_config(args, ranks: int = 1, local_rank: int = 0) -> FDConfig: eplb_config = EPLBConfig(args.eplb_config) structured_outputs_config: StructuredOutputsConfig = StructuredOutputsConfig(args=vars(args)) + routing_replay_config = RoutingReplayConfig(args.routing_replay_config) # Note(tangbinhan): used for load_checkpoint model_config.pretrained_config.tensor_parallel_rank = parallel_config.tensor_parallel_rank @@ -1003,6 +1012,7 @@ def initialize_fd_config(args, ranks: int = 1, local_rank: int = 0) -> FDConfig: plas_attention_config=plas_attention_config, structured_outputs_config=structured_outputs_config, eplb_config=eplb_config, + routing_replay_config=routing_replay_config, ) update_fd_config_for_mm(fd_config) if fd_config.load_config.load_choices == "default_v1" and not v1_loader_support(fd_config): diff --git a/tests/distributed/chunked_moe.py b/tests/distributed/chunked_moe.py index 0be645d38e2..ef41a610dcb 100644 --- a/tests/distributed/chunked_moe.py +++ b/tests/distributed/chunked_moe.py @@ -90,7 +90,7 @@ def init_attention_metadata(self, forward_meta): class MockQuantMethod: - def apply(self, layer, x, gate): + def apply(self, layer, x, gate, topk_ids_hookfunc=None): return x @@ -129,6 +129,7 @@ def setup_model_runner(self): model_runner.speculative_decoding = False model_runner._init_share_inputs(mock_fd_config.scheduler_config.max_num_seqs) model_runner.share_inputs["caches"] = None + model_runner.routing_replay_manager = None if dist.get_rank() == 0: model_runner.share_inputs["ids_remove_padding"] = paddle.ones([10]) @@ -148,6 +149,7 @@ def setup_fused_moe(self): fused_moe.fd_config = mock_fd_config fused_moe.quant_method = MockQuantMethod() + fused_moe.enable_routing_replay = None return fused_moe def run_model_runner(self): diff --git a/tests/e2e/test_EB_Lite_serving.py b/tests/e2e/test_EB_Lite_serving.py index bc27daab993..c71b7667260 100644 --- a/tests/e2e/test_EB_Lite_serving.py +++ b/tests/e2e/test_EB_Lite_serving.py @@ -78,6 +78,8 @@ def setup_and_run_server(): "wint4", "--graph-optimization-config", '{"cudagraph_capture_sizes": [1], "use_cudagraph":true}', + "--routing-replay-config", + '{"enable_routing_replay":true, "routing_store_type":"local", "local_store_dir":"./routing_replay_output"}', ] # Start subprocess in new process group diff --git a/tests/layers/test_fusedmoe.py b/tests/layers/test_fusedmoe.py index ed4fe5b28b6..346afc98fd3 100644 --- a/tests/layers/test_fusedmoe.py +++ b/tests/layers/test_fusedmoe.py @@ -31,6 +31,7 @@ LoadConfig, ModelConfig, ParallelConfig, + RoutingReplayConfig, ) from fastdeploy.model_executor.layers.moe.moe import FusedMoE from fastdeploy.model_executor.layers.quantization.block_wise_fp8 import ( @@ -476,6 +477,7 @@ def __init__( graph_opt_config=GraphOptimizationConfig({}), load_config=LoadConfig({}), ips=",".join(["0"] * nnodes), + routing_replay_config=RoutingReplayConfig({}), ) self.fd_config.parallel_config.tp_group = None self.fd_config.parallel_config.tensor_parallel_rank = tp_rank diff --git a/tests/layers/test_w4a8_moe.py b/tests/layers/test_w4a8_moe.py index dc6dab15427..f20c27b06bf 100644 --- a/tests/layers/test_w4a8_moe.py +++ b/tests/layers/test_w4a8_moe.py @@ -13,6 +13,7 @@ LoadConfig, ModelConfig, ParallelConfig, + RoutingReplayConfig, ) from fastdeploy.model_executor.layers.moe.moe import FusedMoE from fastdeploy.model_executor.layers.quantization.w4a8 import W4A8Config @@ -59,6 +60,7 @@ def __init__( graph_opt_config=GraphOptimizationConfig({}), load_config=LoadConfig({}), ips=",".join(["0"] * nnodes), + routing_replay_config=RoutingReplayConfig({}), ) self.fd_config.parallel_config.tp_group = None self.fd_config.parallel_config.tensor_parallel_rank = tp_rank diff --git a/tests/layers/test_w4afp8_moe.py b/tests/layers/test_w4afp8_moe.py index 65b7733172c..8f1ae79cd67 100644 --- a/tests/layers/test_w4afp8_moe.py +++ b/tests/layers/test_w4afp8_moe.py @@ -13,6 +13,7 @@ LoadConfig, ModelConfig, ParallelConfig, + RoutingReplayConfig, ) from fastdeploy.model_executor.layers.moe.moe import FusedMoE from fastdeploy.model_executor.layers.quantization.w4afp8 import W4AFP8Config @@ -65,6 +66,7 @@ def __init__( graph_opt_config=GraphOptimizationConfig({}), load_config=LoadConfig({}), ips=",".join(["0"] * nnodes), + routing_replay_config=RoutingReplayConfig({}), ) self.fd_config.parallel_config.tp_group = None self.fd_config.parallel_config.tensor_parallel_rank = tp_rank From 7ab6b98ca430247ee0dd08d3a85b299b2c98c74c Mon Sep 17 00:00:00 2001 From: bukejiyu <52310069+bukejiyu@users.noreply.github.com> Date: Fri, 5 Dec 2025 22:23:30 +0800 Subject: [PATCH 09/13] bf16 deepseek (#5379) --- fastdeploy/model_executor/layers/linear.py | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/fastdeploy/model_executor/layers/linear.py b/fastdeploy/model_executor/layers/linear.py index e126aed2ba1..687db766acf 100644 --- a/fastdeploy/model_executor/layers/linear.py +++ b/fastdeploy/model_executor/layers/linear.py @@ -367,11 +367,14 @@ def weight_loader(self, param, loaded_weight, loaded_shard_id: Optional[str] = N # loaded_shard_id == "kv_a" param_shard_offset = self.output_sizes[0] param_shard_size = self.output_sizes[1] - param_output_dim = True if hasattr(param, "tensor_track"): - param_output_dim = param.tensor_track.output_dim param.tensor_track.mark(start=param_shard_offset, end=param_shard_offset + param_shard_size) - param = slice_fn(param, param_output_dim, start=param_shard_offset, end=param_shard_offset + param_shard_size) + param = slice_fn( + param, + (self.fd_config.model_config.model_format == "torch") ^ True, + start=param_shard_offset, + end=param_shard_offset + param_shard_size, + ) assert param.shape == loaded_weight.shape, ( f" Attempted to load weight ({loaded_weight.shape}) " f"into parameter ({param.shape})" ) From 31159f7d4914404170b3959f1d50a356f876bfad Mon Sep 17 00:00:00 2001 From: bukejiyu <52310069+bukejiyu@users.noreply.github.com> Date: Sat, 6 Dec 2025 00:45:48 +0800 Subject: [PATCH 10/13] fix deepseek (#5410) --- fastdeploy/model_executor/layers/linear.py | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/fastdeploy/model_executor/layers/linear.py b/fastdeploy/model_executor/layers/linear.py index 687db766acf..14d1e0dcc0c 100644 --- a/fastdeploy/model_executor/layers/linear.py +++ b/fastdeploy/model_executor/layers/linear.py @@ -977,7 +977,12 @@ def __init__( def process_weights_after_loading(self): if self.fd_config.load_config.dynamic_load_weight: return - w = self.kv_b_proj.weight.reshape( + w = ( + self.kv_b_proj.weight.transpose([1, 0]) + if self.fd_config.model_config.model_format == "torch" + else self.kv_b_proj.weight + ) + w = w.reshape( [ self.kv_lora_rank, self.num_heads_per_partition, From bef4ea87eab81aa74bd352fbfdd126e5369d4da7 Mon Sep 17 00:00:00 2001 From: SunLei Date: Mon, 8 Dec 2025 13:26:27 +0800 Subject: [PATCH 11/13] Update tests/inter_communicator/test_fmq_factory.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> --- tests/inter_communicator/test_fmq_factory.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/inter_communicator/test_fmq_factory.py b/tests/inter_communicator/test_fmq_factory.py index 37a1a4193cd..50da82f4fae 100644 --- a/tests/inter_communicator/test_fmq_factory.py +++ b/tests/inter_communicator/test_fmq_factory.py @@ -23,7 +23,7 @@ class TestFMQFactory(unittest.IsolatedAsyncioTestCase): async def test_create_queues(self): - """测试能否创建所有 producer / consumer 队列""" + """Test whether all producer/consumer queues can be created.""" q1 = factory.q_a2e_producer() q2 = factory.q_a2e_consumer() q3 = factory.q_e2w_producer() From 717bb0c16fd11c307c239584ef919d18bd472951 Mon Sep 17 00:00:00 2001 From: SunLei Date: Mon, 8 Dec 2025 13:26:54 +0800 Subject: [PATCH 12/13] Update benchmarks/benchmark_fmq.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> --- benchmarks/benchmark_fmq.py | 1 - 1 file changed, 1 deletion(-) diff --git a/benchmarks/benchmark_fmq.py b/benchmarks/benchmark_fmq.py index f7a1b74031e..3878f790cc6 100644 --- a/benchmarks/benchmark_fmq.py +++ b/benchmarks/benchmark_fmq.py @@ -87,7 +87,6 @@ async def consumer_task(consumer_id, total_msgs, result_q, consumer_event): send_ts = msg.payload["send_ts"] latencies.append((recv_ts - send_ts) * 1000) # ms pbar.update(1) - # pbar.write(f"recv {msg.payload['i']} = {msg.msg_id}") recv += 1 pbar.close() From 8dbb0871c753fa0a68b3dbf8e81eb42e90a1d56f Mon Sep 17 00:00:00 2001 From: SunLei Date: Mon, 8 Dec 2025 13:27:20 +0800 Subject: [PATCH 13/13] Update fastdeploy/inter_communicator/fmq.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> --- fastdeploy/inter_communicator/fmq.py | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) diff --git a/fastdeploy/inter_communicator/fmq.py b/fastdeploy/inter_communicator/fmq.py index 0915ce3e187..f2c98196c99 100644 --- a/fastdeploy/inter_communicator/fmq.py +++ b/fastdeploy/inter_communicator/fmq.py @@ -217,7 +217,18 @@ def __init__(self, context, name: str, role: str = "producer"): fmq_logger.info(f"Queue {name} initialized on {full_ep}") async def put(self, data: Any, shm_threshold: int = 1024 * 1024): - # Send data to queue + """ + Send data to the queue. + + Args: + data: The data to send. Can be any serializable object or bytes. + shm_threshold: Size threshold in bytes. If the data is of type bytes and its size is + greater than or equal to this threshold, shared memory will be used to send the message. + Default is 1MB (1024 * 1024 bytes). + + Raises: + PermissionError: If called by a non-producer role. + """ if self.role != Role.PRODUCER: raise PermissionError("Only producers can send messages.")