agentsclimarketplace

Mega moe fusion with symmetric memory

Skill kjuhwa/skills-hub/skills/moe-optimization/mega-moe-fusion-with-symmetric-memory

Self-correcting knowledge corpus for Claude Code — 9 stable shape clusters, bias-correction pipeline baked into contribution flow. 47 papers, 45 techniques, 1.1k skills.

Install
npx -y skills add kjuhwa/skills-hub --skill mega-moe-fusion-with-symmetric-memory

Assembled from the repository path, not quoted from the project. Check it against their README if it does not work.

One thing to look at

  • 0 stars0 stars. Stars are a popularity signal and not a quality one, but at this level it is likely that nobody has read this closely except its author, and you would be relying on your own review.

What its author says it does

Copied from the file, not written here

Fuse EP dispatch, FP8×FP4 GEMM×2, SwiGLU, and EP combine into a single kernel that reads across ranks via symmetric-memory (same VA on every rank) to overlap NVLink collectives with compute.

SKILL.md

3.1 KB, 638 tokens by cl100k_base, as published. Nobody here has run it

Mega MoE Fusion With Symmetric Memory

What / Why

A classic EP (expert-parallel) MoE step is five discrete ops: allgather tokens, gather-by-expert, GEMM1, SwiGLU, GEMM2, scatter-combine. Each has its own launch + barrier + memory pass. Mega MoE fuses all of them into one kernel that reads from peer-rank buffers using symmetric memory (same virtual address on every rank, backed by CUDA IPC). Compute on local experts overlaps with NVLink traffic to/from peers — no explicit allgather + compute phase boundaries.

Procedure

  1. Require PyTorch ≥ 2.9 for the torch.cuda.symmetric_memory API (same VA across ranks via IPC).
  2. Allocate a symmetric buffer. get_symm_buffer_for_mega_moe(group, num_experts, num_max_tokens_per_rank, num_topk, hidden, intermediate_hidden) returns one MegaMoEBuffer per rank, all sharing virtual addresses.
  3. Layout the buffer with named regions: x (activations FP8), x_sf (UE8M0 scales), topk_idx, topk_weights, weights (FP4, read-only), output.
  4. Select kernel config via mega_moe.hpp heuristic: picks block_m (tokens per expert), block_n (hidden per expert), block_k; accounts for a 2× imbalance factor when estimating per-expert token count.
  5. Kernel pipeline stages.
    • Phase 1: warp-group 0 issues cross-rank TMA loads of peer-rank x into smem.
    • Phase 2: warp-groups 1-N run GEMM1 (FP8 × FP4) on already-arrived tiles.
    • Phase 3: fused SwiGLU (gate × silu(up)).
    • Phase 4: GEMM2.
    • Phase 5: cross-rank TMA stores into peer-rank output — the EP combine.
  6. Barrier layout. Dependency tracking via lightweight mbarriers; no __syncthreads() across the whole block.

Key design points

  • Symmetric memory is the critical enabler — without it, you need host-side gather, which serializes the pipeline.
  • The imbalance factor (2×) is empirical for typical top-k routing. Heavily skewed models should override mk_alignment_for_contiguous_layout().
  • Keep FP8 activations and FP4 weights — no mid-kernel dequantization to FP32. Accumulator stays BF16 / FP32 per tensor-core protocol.

References

  • csrc/jit_kernels/heuristics/mega_moe.hpp — config search with imbalance factor.
  • deep_gemm/mega/__init__.py — Python API + symmetric buffer helper.
  • tests/test_mega_moe.py — reference correctness check against DeepEP baseline.

Keep looking

Skills are one crate of 328,083. Ordering is by how many stacks a row turns up in, so the top of any crate is what has actually been picked rather than what has the most stars.