You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
The reason will be displayed to describe this comment to others. Learn more.
Pull request overview
Refactors the MoE dispatch/combine implementation to support multi-backend autotuning, and introduces a new ROCm Mori EP backend alongside shared MoE utilities and updated distributed tests.
Changes:
Added MoEDispatchCombineAutoTuner with shape-keyed caching and handle-bound reuse across dispatch/combine.
Added MoriEPBackend and BackendType.MORI, plus ROCm-specific dispatch token-count handling.
Extracted shared helpers into moe_utils.py and updated dispatcher tests to use NCCL process groups.
Reviewed changes
Copilot reviewed 5 out of 5 changed files in this pull request and generated 10 comments.
Show a summary per file
File
Description
tests/pytorch/modules/test_token_dispatcher.py
Adds autotune sweep coverage and switches tests to self.pg + NCCL init behavior.
The reason will be displayed to describe this comment to others. Learn more.
MoEDispatchCombineAutoTuner keys _handle_cache by id(handle) only. Python can reuse object ids after GC, so a later, unrelated handle could collide and incorrectly reuse a stale tuned result. To make this robust, consider storing a strong reference to the handle in the cache entry and verifying stored_handle is handle on lookup (bounded by the existing LRU), or using a weakref-based approach where the handle type supports it.
The reason will be displayed to describe this comment to others. Learn more.
set_buffer_global_config(backend=None) only updates _buffer_config_per_backend keys that already exist (derived from _DEFAULT_BUFFER_CONFIG_PER_BACKEND). If a new backend is registered later via register_ep_backend(), it won’t receive the global config unless callers remember to pass backend=<name>. Consider iterating over _BACKEND_REGISTRY.keys() (or taking the union) when backend is None so newly registered backends are included.
The reason will be displayed to describe this comment to others. Learn more.
detect_group_topology() performs an all_gather_object every time it’s called. Since group topology is static, calling this from per-forward paths (e.g., Mori init/config resolution) adds an unnecessary synchronous collective overhead. Consider caching (node_idx, num_nodes) per ProcessGroup (e.g., keyed by id(group) with a weakref) to avoid repeated collectives.
The reason will be displayed to describe this comment to others. Learn more.
Stale comment
Security review (automation)
Reviewed the diff for injection, authn/authz, permission-boundary, secret-leakage, SSRF/XSS/request-forgery, path-traversal, unsafe-deserialization, and supply-chain risks.
Scope
primus_turbo/pytorch/core/backend.py (optional mori import + enum)
Env vars (PRIMUS_TURBO_MORI_NUM_QP_PER_PE, NODE_RANK/GROUP_RANK/SLURM_NODEID) feed only int() casts and rank-id comparisons — not shell, path, SQL, or deserialization sinks.
tempfile.NamedTemporaryFile(suffix=".json") uses secure defaults; the JSON is written by prof.export_chrome_trace in the same process (no attacker input).
suppress_stdout_stderr manipulates only its own file descriptors.
dist.all_gather_object in detect_group_topology pickles over the internal torch.distributed process group (trusted channel).
Triton compute_expert_token_info_kernel clamps OOB rows/cols with tl.where + load_mask/flat_mask and clamps invalid/out-of-range expert ids before tl.histogram / tl.atomic_add, preventing OOB reads/writes and illegal bin counts.
inplace_unique pads negatives to num_slots and drops the padding bin, keeping scatter_add_ in-range.
MoEDispatchCombineAutoTuner._handle_cache is LRU-bounded (_HANDLE_CACHE_MAX=1024); shape cache uses TuneCache(capacity=1024).
torch._C._distributed_c10d._register_process_group("mori", group) receives only internally constructed arguments.
The new mori dependency is an optional import (guarded by try/except ImportError); no code runs unless installed.
Result
No medium+ severity security issue introduced or exposed by this PR. No new inline findings.
Env vars (PRIMUS_TURBO_MORI_NUM_QP_PER_PE, NODE_RANK/GROUP_RANK/SLURM_NODEID) feed only int() casts and rank-id comparisons — no shell, path, SQL, or deserialization sinks.
tempfile.NamedTemporaryFile(suffix=".json") uses secure defaults; content is written by the in-process prof.export_chrome_trace (no attacker input).
suppress_stdout_stderr manipulates only its own file descriptors.
dist.all_gather_object in detect_group_topology pickles over an internal torch.distributed group (trusted training peers only).
Triton compute_expert_token_info_kernel clamps OOB rows/cols via tl.where + load_mask/flat_mask and clamps invalid/out-of-range expert ids before tl.histogram/tl.atomic_add — no OOB reads/writes or illegal bin counts.
inplace_unique pads negatives into num_slots and drops the padding bin, keeping scatter_add_ in-range.
MoEDispatchCombineAutoTuner._handle_cache is LRU-bounded (_HANDLE_CACHE_MAX=1024) and now identity-checks the stored handle (stored_handle is not handle) to guard against id() reuse; shape cache uses TuneCache(capacity=1024).
torch._C._distributed_c10d._register_process_group("mori", group) receives only internally constructed arguments.
The new mori dependency is an optional import (guarded by try/except ImportError); no code runs unless installed.
Result
No medium+ severity security issue introduced or exposed by this PR. No inline findings.
The reason will be displayed to describe this comment to others. Learn more.
test_autotune_sweep enables PRIMUS_TURBO_AUTO_TUNE=1 and then runs two full dispatch+combine passes at the default test shape (4096×4096 with topk=8). With the new autotuner this triggers a full config sweep using torch.profiler/Kineto inside the test, which is likely to be very slow and can make CI flaky/time out. Please consider reducing the problem size for this test (smaller NUM_TOKENS/HIDDEN_SIZE), lowering the sweep iteration counts for test mode, or marking/gating the test as “slow” so it doesn’t run in the default unit test shard.
The reason will be displayed to describe this comment to others. Learn more.
_get_backend_name() now only consults GlobalBackendManager.get_moe_dispatch_combine_backend(). That helper intentionally returns None when PRIMUS_TURBO_MOE_DISPATCH_COMBINE_BACKEND is set to a custom (non-BackendType) backend name, so this change forces a silent fallback to TURBO and breaks the documented “custom names like UCCL_EP” behavior. Please restore the previous behavior by reading ENV_MOE_DISPATCH_COMBINE_BACKEND directly when GlobalBackendManager returns None, and using the raw env value (uppercased/stripped) as the backend registry lookup key.
The reason will be displayed to describe this comment to others. Learn more.
_DeepEPLikeBackend.init_buffer() computes hidden_bytes as hidden_size * 2 for all non-FP8 inputs (element_size is hard-coded to 2). This can under-allocate buffers when x is float32 (4 bytes/elem) and conflicts with the previous logic that sized by max(inp.element_size(), 2). Either (a) size by the actual input element size (pass dtype/element_size into init_buffer), or (b) explicitly restrict dispatch/combine to BF16/FP16 and raise a clear error for FP32 to avoid silent buffer overruns.
The reason will be displayed to describe this comment to others. Learn more.
MoEDispatchCombineAutoTuner._handle_cache stores strong references to the full handle object (cache[hid] = (handle, result)). If a backend’s handle captures large tensors or GPU resources (e.g., MoriEPBackend’s handle currently includes recv_topk_idx_i32), this cache can retain significant memory until eviction, which is risky in long-running training. Consider storing only the tuned result keyed by id(handle) (accepting the small risk of id reuse), or redesigning handles to be lightweight/opaque IDs so the cache doesn’t keep heavyweight objects alive.
The reason will be displayed to describe this comment to others. Learn more.
Stale comment
Security review (automation)
Re-validated after the latest push (8bd040b). The incremental change since the last assessment is a mechanical rename (PRIMUS_TURBO_MOE_DISPATCH_COMBINE_BACKEND → PRIMUS_TURBO_EP_BACKEND, get_moe_dispatch_combine_backend → get_ep_backend) and the removal of eager deep_ep / mori imports in primus_turbo/pytorch/core/backend.py (both are now lazily loaded inside the EP backend registry). No trust boundary, permission check, or supply-chain surface changes.
Env vars (PRIMUS_TURBO_MORI_NUM_QP_PER_PE, NODE_RANK/GROUP_RANK/SLURM_NODEID, PRIMUS_TURBO_EP_BACKEND) feed only int() casts, bounded registry lookups (_BACKEND_REGISTRY), and rank/host token equality — no shell, path, SQL, or deserialization sinks.
tempfile.NamedTemporaryFile(suffix=".json") uses secure defaults; content produced by the in-process prof.export_chrome_trace (no attacker input).
suppress_stdout_stderr manipulates only its own file descriptors.
dist.all_gather_object in detect_group_topology runs over the trusted torch.distributed group.
Triton compute_expert_token_info_kernel clamps OOB rows/cols via tl.where + load_mask/flat_mask and clamps invalid/out-of-range expert ids before tl.histogram / tl.atomic_add — no OOB reads/writes or illegal bins.
inplace_unique pads negatives into num_slots and drops the padding bin, keeping scatter_add_ in-range.
MoEDispatchCombineAutoTuner._handle_cache is LRU-bounded (_HANDLE_CACHE_MAX=1024) and identity-checks the stored handle (stored_handle is not handle) to defend against id() reuse; shape cache uses TuneCache(capacity=1024).
torch._C._distributed_c10d._register_process_group("mori", group) receives only internally constructed arguments.
mori / deep_ep remain optional dependencies (guarded by try/except ImportError in each backend's is_available() / _get_module()).
No secret, token, payload, or model-input values are logged.
Result
No medium+ severity security issue introduced or exposed by this PR. No inline findings.
The reason will be displayed to describe this comment to others. Learn more.
Stale comment
Security review (automation)
Re-validated after the latest push (b96afca). The incremental changes since the last assessment are: (1) CK GEMM/Grouped GEMM directory restructure (file moves only); (2) numerical correctness fixes in csrc/kernels/quantization/quantization_mxfp8.cu (cast to int64_t to avoid int32 offset overflow), csrc/kernels/reduce/reduce_row.cuh (partial-tile indexing), and csrc/kernels/gemm/turbo/turbo_gemm_mxfp8_kernel.h (__builtin_amdgcn_s_barrier() WAR fence); (3) attention BHSD layout support in attention_aiter_impl.py / flash_attn_interface.py / flash_attn_usp_interface.py / attention_ring.py, with torch.library.custom_op wrappers and matching register_fake shape implementations; (4) a new deepep_use_comm_stream=False knob in DeepEPTokenDispatcher that calls clear_backend_instances() and sets the constant PRIMUS_TURBO_EP_FORCE_CURRENT_STREAM env var to "1"; (5) a _format_kwargs(...) helper that renders tensor shape/dtype (not data) and repr() of scalars/enums into backend mismatch error messages; (6) the nits commit dropping dispatch_config/combine_config from a debug log line.
_format_kwargs only emits Tensor(shape=..., dtype=...) for tensors and repr() for scalar/enum kwargs — no tensor contents, weights, gradients, or attacker payload values are logged. Error messages stay free of secrets/PII.
os.environ[ENV_EP_FORCE_CURRENT_STREAM] = "1" writes a hard-coded constant; no user-controlled string flows into os.environ or any shell. The C++ side reads the same env var via std::getenv and passes it through std::stoi, with the value bounded to a bool.
clear_backend_instances() only mutates the module-private _backend_instances dict — no FS, network, or trust-boundary surface.
torch.library.custom_op(..., mutates_args=("dq","dk","dv"), device_types="cuda") wrappers receive only Tensors / scalars from in-process callers; register_fake impls allocate empty tensors and don't dereference user data.
BHSD layout handling is restricted to qkv_format in ("bshd","sbhd","bhsd") with explicit assert/raise ValueError on unknown values; no path/format string is interpolated into a sink.
Sink-attention fast-path constraints (head_dim_qk == head_dim_v, power-of-2) moved into can_handle() — same-or-stronger validation as before.
mxfp8 int64_t cast and reduce_row partial-tile fix close potential signed-int32 offset overflows that could have produced OOB stores on extreme shapes; this is a security/safety improvement, not a regression.
MoEDispatchCombineAutoTuner._handle_cache remains LRU-bounded (_HANDLE_CACHE_MAX=1024) with id() reuse defended by stored_handle is not handle identity check; TuneCache(capacity=1024) shape cache unchanged.
mori / deep_ep remain optional, lazily imported inside each backend's is_available() / _get_module().
No new env vars feed into shell/path/SQL/deserialization sinks; PRIMUS_TURBO_EP_FORCE_CURRENT_STREAM and other env reads (int() casts, registry lookups, rank/host token equality) stay numeric/bounded.
Result
No medium+ severity security issue introduced or exposed by this PR. No inline findings.
primus_turbo/pytorch/kernels/moe/moe_dispatch_combine_impl.py: bumps EPBufferConfig.num_sms default from 32 to 64, adds a "UCCL_EP" entry to _DEFAULT_BUFFER_CONFIG_PER_BACKEND, and adds a new UCCLEPBackend(_DeepEPLikeBackend) whose is_available() / _get_module() lazily import uccl.ep. The new class is not wired into _BACKEND_REGISTRY, so it is only reachable if a caller invokes register_ep_backend("UCCL_EP", UCCLEPBackend) explicitly.
setup.py: adds a mandatory install_requires entry uccl @ git+https://github.com/uccl-project/uccl.git@9bae94c59229f82efbbd5b78a7f222ce19e74e86.
Key checks
Supply chain: the new uccl dependency is fetched from the upstream uccl-project/uccl repo (already referenced by the project's existing DeepEP fallback documentation) and is pinned to a fixed commit SHA, so installs cannot silently move to a different revision via tag/branch retargeting. No version-range or floating-ref injection risk introduced. The constant UCCL_COMMIT is a hard-coded literal — no env-var or user input flows into the install URL.
UCCLEPBackend._get_module() only does a lazy import uccl.ep; no attacker-controlled string is passed to import and the import target is fixed.
_make_buffer_kwargs() introspects the upstream Buffer constructor with inspect.signature; only the boolean is_intranode (derived from group.size()) is forwarded — no untrusted kwargs propagation.
EPBufferConfig default change (num_sms 32 → 64) and new UCCL_EP defaults are pure numeric tuning constants; no trust-boundary impact.
The reason will be displayed to describe this comment to others. Learn more.
Stale comment
Security review (automation)
Re-validated after the latest push (c4906b5). Incremental delta since the last assessment (aab8575):
primus_turbo/pytorch/kernels/moe/moe_dispatch_combine_impl.py: a single-line registry entry ("UCCL": UCCLEPBackend) added to _BACKEND_REGISTRY. The UCCLEPBackend class itself (added in the prior commit) follows the same optional-import + _DeepEPLikeBackend pattern already used by DeepEPBackend (lazy import uccl.ep, no eager dependency, identical buffer-kwargs introspection).
primus_turbo/pytorch/deep_ep/buffer.py and csrc/pytorch/deep_ep/deep_ep.hpp (force-current-stream toggle).
setup.py (new uccl install requirement pinned to commit 9bae94c5...).
csrc/kernels/... (CK GEMM/Grouped-GEMM directory restructure + numerical-correctness fixes; no new I/O sinks).
Attention/USP refactor and tests.
Threat checks performed
Injection: env values (PRIMUS_TURBO_EP_BACKEND, PRIMUS_TURBO_MORI_NUM_QP_PER_PE, ENV_EP_FORCE_CURRENT_STREAM) are parsed via BackendType[...]/int(...) with defensive fallbacks; no shell or SQL sinks. The Mori SHMEM process-group name ("mori") is a hard-coded constant.
Authn/authz: no auth surfaces touched.
Path traversal / unsafe deserialization: bench_kineto writes to a tempfile.NamedTemporaryFile (random path) and json.loads only the file it just wrote; no attacker-controlled path or payload.
SSRF / request forgery / XSS: no network or HTML sinks introduced.
Secret leakage / insecure logging: log lines emit shapes and tuning metadata only; no credentials or tokens.
Supply chain: the new uccl dependency in setup.py is pinned to a specific commit hash on github.com/uccl-project/uccl, which mitigates floating-tag tampering risk and matches existing patterns (e.g. pinned AITER_COMMIT). UCCLEPBackend only imports uccl.ep lazily and is selected via the existing EP backend registry.
Process-wide env mutation in DeepEPTokenDispatcher.__init__ (os.environ[ENV_EP_FORCE_CURRENT_STREAM]="1" + clear_backend_instances()): a usability concern previously raised by Copilot, but not a security boundary issue (caller already controls the process), so out of scope here.
Conclusion: no medium-or-higher security findings. No new finding comments posted.
No new medium or higher severity vulnerabilities are introduced. Notes on items examined and judged not to rise to that bar:
setup.py adds a Git-URL install_requires for uccl but pins to a specific commit SHA, so it is not an unpinned/floating dependency.
All env-var reads (ENV_EP_BACKEND, ENV_MORI_NUM_QP_PER_PE, NCCL_* → MORI_*/UCCL_*) are constrained: backend name is dereferenced through an in-process registry (_BACKEND_REGISTRY[name], raises on unknown), num_qp_per_pe is coerced via int(...), and the NCCL→MORI/UCCL fallback only copies env values that are already present in the same process's environment to a different env-var name. No shell, no eval, no subprocess.
_align_mori_hip_with_torch derives the .so path from os.path.dirname(torch.__file__); not attacker-controlled.
compute_expert_token_info_kernel (Triton) guards loads/stores with explicit mask= operands and clamps OOB rows/experts before address arithmetic.
detect_group_topology uses dist.all_gather_object (pickle under the hood) over a token sourced from NODE_RANK/GROUP_RANK/SLURM_NODEID/socket.gethostname(). Peers in an EP process group share the same trust boundary as the training job, so this is not a new attack surface relative to existing torch.distributed idioms in this codebase.
bench_kineto writes traces through tempfile.NamedTemporaryFile; no untrusted-path / TOCTOU exposure.
Prior threads from earlier automation runs
Prior security-review summaries from this automation all reached the same "no findings" verdict. They are being collapsed as outdated; this comment supersedes them.
The reason will be displayed to describe this comment to others. Learn more.
Stale comment
Security review (automation)
Re-validated for the latest push (d094e9f). No medium+ severity security findings.
Incremental delta since the last assessment (5ba14d2):
EPBufferConfig split into primus_turbo/pytorch/kernels/moe/_backend/_config.py (refactor, no behavior change).
_backend/__init__.py re-exports updated; no new optional-dep imports at module load time.
setup.py: uccl @ git+https://github.com/uccl-project/uccl.git@<pinned-commit> moved from install_requires to extras_require["uccl"]. This is a supply-chain hygiene improvement — the unsigned external GitHub source is now opt-in instead of pulled by every default install. The commit hash remains pinned (immutable reference).
Minor docstring/comment edits in backend.py, constants.py.
token_dispatcher.py, ops/moe/moe_dispatch_combine.py — graph-capture handling.
Checks performed and outcome:
Env-var ingestion (ENV_EP_BACKEND, ENV_MORI_NUM_QP_PER_PE, NCCL_* -> UCCL_* / MORI_* fallback in _apply_env_with_nccl_fallback): values are only set into the current process env or parsed as int; no shell expansion, no path/command injection.
ctypes.CDLL(torch_hip) in _align_mori_hip_with_torch: path is derived from torch.__file__ (trusted Python install), gated by os.path.isfile. Not attacker-controlled.
Triton kernel compute_expert_token_info_kernel uses bounds masks (load_mask, footprint_mask, bin_mask) and clamps invalid lanes before atomic add; no OOB write path from tensor shapes.
tempfile.NamedTemporaryFile(suffix=".json") in bench_kineto: secure temp creation, no path traversal.
dist.all_gather_object in detect_group_topology: pickle is used inside an already-trusted EP process group (standard PyTorch distributed pattern, not new attack surface introduced by this PR).
lru_cache on _build_mori_op and OrderedDict handle cache: keyed on integers / id() with is re-check; no DoS or aliasing exploit path.
Supply chain: UCCL_COMMIT and AITER_COMMIT are pinned to immutable commit hashes; UCCL is now optional.
Result: No injection, authn/authz, secret-leakage, SSRF/XSS, path-traversal, or unsafe-deserialization issues introduced by this PR. No new comments to post.
This test still passes dist.group.WORLD into _run_dispatch_combine while the rest of the test suite now uses self.pg (NCCL). Keeping a single test on a different group can make failures backend-/launcher-dependent (e.g., if WORLD isn't NCCL under the harness).
with patch.dict(os.environ, {"PRIMUS_TURBO_EP_BACKEND": backend}):
_run_dispatch_combine(
self.rank,
dist.group.WORLD,
deepep_use_cuda_num_tokens_per_expert=True,
- Pluggable EP backends under kernels/moe/_backend (turbo, deep_ep, mori, uccl)
selected via PRIMUS_TURBO_EP_BACKEND; shared paired dispatch->combine bench +
sweep_configs autotune harness with per-iter cross-rank barriers.
- deep_ep/TURBO: int64 Buffer/AsymBuffer/SymBuffer offsets (fix int32 overflow on
large NVL buffers); default EP kernels to the current stream.
- MORI backend (ROCm Mori SHMEM); inter-node mori is not benchmark-safe, so skip
tuning there and use the base config (live builds one fresh op).
- autotuner: restrict the sweep to a pinned backend; release other backends'
comm buffers before use (one live EP backend per process).
- uccl: is_available() checks the Buffer API surface, not just the import.
- tests/bench: token-dispatcher tests + _build_tp_groups; deep_ep intranode bench
supports relative/script imports.
setup.py is left identical to main (mori/uccl are not added as install deps).
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- Add primus_turbo/common/mori_utils.py: get_mori() lazily imports the optional
amd_mori dependency with a clear install hint and a once-only version-pin
warning (pinned to ROCm/mori v1.2.0; built with BUILD_UMBP=OFF).
- mori backend uses get_mori() in is_available()/init_buffer()/tune_configs so a
missing/mismatched mori surfaces a clear message at use time.
- README: document the amd_mori install (optional, MORI EP backend).
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…inter-node
Backend capabilities:
- _EPCapabilities mixin: can_release(will_reinit) + is_feasible(group); each
backend overrides only what differs (uccl can't release — destroy corrupts
HIP; mori can't reinit — inter-node free+rebuild hangs). Inherited by both
_DeepEPLikeBackend and the standalone MoriEPBackend.
- DeepEP(uccl)/UCCL auto-configure UCCL_* RDMA env from NCCL_* + NIC defaults.
Autotuner:
- Filter candidates by is_feasible so an unpinned sweep auto-excludes backends
that can't run the topology (TURBO inter-node SIGABRTs) — no env needed.
- Free non-winner tuning buffers post-sweep (skip uccl) to avoid wasted HBM;
keep the winner for live reuse.
mori:
- Enable inter-node autotune: one op, never rebuilt; sweep (block,warp,rdma)
via per-call overrides; live reuses the tuned op (rebuild hangs).
- Defaults aligned to mori's official gfx950 tuning (intra 256/16, inter 128/8/64).
- Opt-in MORI_EP_LAUNCH_CONFIG_MODE=AUTO delegates config to mori's tuning JSON.
- Split the triton compute_expert_token_info kernel to expert_token_count.py.
Add common/uccl_utils.py (lazy uccl import + version pin) and README install lines.
Validated: test_token_dispatcher.py 25/25; DeepSeek-V3 EP8 + EP16 (2 nodes)
autotune across mori / uccl(DEEP_EP) / turbo.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
test_pad_multiple passes dist.group.WORLD as the EP group, while other tests in this class pass self.pg. This inconsistency can break under MultiProcContinuousTest if self.pg is not the default WORLD group (and is also inconsistent with the stated switch to self.pg).
with patch.dict(os.environ, {"PRIMUS_TURBO_EP_BACKEND": backend}):
_run_dispatch_combine(
self.rank,
dist.group.WORLD,
deepep_use_cuda_num_tokens_per_expert=True,
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
![cursor[bot]](https://avatars.githubusercontent.com/in/1210556?s=60&v=4){kind=small}
Description
Refactor MoE dispatch/combine to support multi-backend autotuning and add a new Mori EP backend on ROCm.
Fixes # (issue)
Type of change
Changes
MoEDispatchCombineAutoTuner: shape-keyed tune + handle-bound reuse for paired dispatch/combine, enabled viaPRIMUS_TURBO_AUTO_TUNE=1.MoriEPBackend(ROCm) and registerBackendType.MORI; keep Turbo/DeepEP unchanged.moe_utils.py(bench_kineto,detect_group_topology,inplace_unique).self.pgwithnccl.Checklist: