MoRI v1.2.4

MoRI v1.2.4

Release since v1.2.3.post1 (70 commits). Major themes: EPv2 gains a preview internode path over CCO/GDA and makes the HIP backend the default on gfx125x; CCO compiles its SDMA path by default and adds a fused GEMM+all-reduce over SDMA; MORI-IO gets RDMA retry tunables, a fault-injection seam and honest peer-failure reporting; UMBP adds NUMA-aware DRAM placement; mlx5 gains dma-buf support.

Highlights

• EPv2 internode (preview): dispatch/combine across nodes over CCO/GDA (#625), with DSV4 (hidden 7168) launch geometry (#676).
• EPv2 on gfx1250: the HIP kernel backend is now the default on gfx125x (#677). EP8 tuning (#723, #745), push-mode fp4 combine (#719, #720), fp4 pack with no-drain metadata (#706), and a DSR1 EP4 topk-9 schedule (#725, #726, #737).
• CCO: the SDMA path is compiled by default (#710), with SDMA micro-optimizations (#603). gemm_ar fuses DSV4-Pro's wo_b GEMM with its all-reduce over SDMA (#662), and adds MXFP8 and a GEMM usable without a collective (#690).
• MORI-IO: RDMA retry tunables (#670), a verbs seam with a compile-time fault injector (#731), and peer failures are reported instead of peers always being reported alive (#716).

Features

• feat(EPv2): [preview] internode dispatch/combine over CCO/GDA (#625)
• feat(EPv2): default to the HIP kernel backend on gfx125x (#677)
• feat(cco): fuse DSV4-Pro wo_b's GEMM with its all-reduce over cco SDMA (#662)
• feat(cco): gemm_ar MXFP8 support, and the GEMM usable without a collective (#690)
• feat(ccl): support HIP graph capture in AllGather SDMA (#652)
• feat(io): RDMA retry tunables (#670); verbs seam + compile-time fault injector for MORI-IO RDMA (#731)
• feat(umbp): NUMA-aware DRAM placement (#707); share storage configuration across deployment modes (#722)
• feat(mlx5): DMA buffer support (#654)
• refactor(ep/v2): make TokOffExt public for out-of-op callers (#708)

Performance

• perf(EPv2/gfx1250): EP8 tuning (#723, #745); push-mode fp4 combine (#719); fp4 pack and no-drain metadata (#706)
• perf(EPv2/gfx1250): EP4 topk-9 dispatch+combine schedule and retune (#725, #726, #737)
• perf(EPv2): hide the dispatch drain read behind the grid barrier on gfx1250 (#642); fix 512-token performance on gfx1250 A1B0 (#667)
• perf(ep): skip redundant nodeRecvTokenNum re-reads in DispatchInterNodeLLRecv, V1+V2 (#702); skip the rest of an exhausted node in DispatchInterNodeLLRecv (#705)
• perf(ep): MI355X IntraNodeLL combine tuning (#647); DSV4 (hidden 7168) small-token InterNodeV1LL tuning on MI308X (#689)
• perf(cco): CCO SDMA micro-optimizations (#603)
• perf(umbp): remove three bottlenecks from the distributed ranged read path (#656); resolve a key set once instead of per layer group (#644)

Bug Fixes

• fix(cco): skip window MRs on a single node and pad >2 GiB allocations to 2 MiB, which fixes ENOMEM from ionic MR registration on ROCm 7.2.x (#704)
• fix(cco): group lanes by CQ before PSD non-CCQE poll (#639); cco/shmem index wraparound (#653)
• fix(EPv2): last-arrival-resets for the combine wide-EP grid barrier (#709); FP4 push reads input (#720); drop a redundant total_recv.zero_() on gfx1250 intranode (#735); aggReserve state as a module __device__ array (#746)
• fix(ep): slice recv tokens in DispatchInterNodeRecv by the actual chunk index (#621); restore replay-mode payload dispatch for intranode training (#675)
• fix(sdma): avoid fused-off odd SDMA engines on MI308X (#693)
• fix(mlx5): stop mixing dma-buf and peermem on GPU control structures (#740)
• fix(io/rdma): export VMM memory by allocation handle so deregistered regions are freed (#729); report peer failures (#716)
• fix(allocator): link c10_hip in the JIT path for getCurrentHIPStream (#666)
• fix: migrate GPU topology from rocm_smi to amd_smi for ROCm 10.1+ (#663)
• fix: strided offsets in run_single_once when batch_contiguous is disabled (#612)
• revert: the dispatch completion-spin skip from v1.2.3.post1 (#645) is reverted (#661)

Build / CI / Tests / Docs

• build(cco): compile the SDMA path by default (#710)
• fix(build): lift the setuptools<84 cap and require Cython for the cco extension (#701); let torch pick the C++ standard for mori_torch_symm (#713)
• ci: ROCm 10 intranode gate (#651); fix ionic CCQE hangs and MI300X_MLX CI failures (#728); MI355X intranode CI migration (#669); re-run PR CI against current main (#738, #744); move jobs off the TW runner label (#684, #686)
• test(EPv2): JSON result table with per-point clock/power telemetry (#648); end-to-end dispatch+combine latency in the bench (#660); graph timing (#673); per-token scales on by default for fp8/fp4 in bench_ep (#697); warm the internode bench with the timed loop's tensor lifetime (#665)
• test(jit): check EpCfg's field count against VisitFields (#727)
• refactor(EPv2): name the dispatch transport fp8 or fp4x2 (#668)
• refactor(gemm_ar): drop the fused epilogue's unused fence knob (#691)
• io: default bench is the C++ bench (#573)
• tools: env-check Mellanox QoS without sudo (#658)
• docs: refresh installation, EP backend guidance and site checks (#680); fix the -W build and the EP fp4 target list (#688); cluster-network-topology skill (#615)

Known issues

• mori.ops, mori.io, mori.shmem and mori.umbp need the gRPC runtime on the host. libmori_pybinds.so links gRPC through the UMBP bindings, so without it these imports fail with ImportError: libgrpc++.so.1. On Ubuntu 22.04, install it with apt install libgrpc++1. mori.cco is not affected. This was already the case in v1.2.3.post1.

Wheels: cp310 (manylinux_2_35_x86_64) and cp312 (manylinux_2_39_x86_64), published to PyPI as amd-mori==1.2.4.

Full changelog: v1.2.3.post1...v1.2.4