AMD GPU性能暴涨7倍，优化算法首次开源！高效MoE支持任意专家数量

新智元报道

编辑：LRST

【新智元导读】通过完全启用并发多块执行，支持任意专家数量（MAX_EXPERT_NUMBER==256），并积极利用共享内存（5kB LDS）和寄存器（52 VGPRs，48 SGPRs），MoE Align & Sort逻辑被精心设计，实现了显著的性能提升：A100提升3倍，H200提升3倍，MI100提升10倍，MI300X/MI300A提升7倍...

MoE（Mixture of Experts）模型模仿了人脑的低功耗运作模式：功能被划分为多个独立的部分，在思考时通过自适应路由部分激活，从而提高计算效率。

牛津大学研究论文中的人脑皮层示意图，来源于互联网

首个可在CUDA真正可行的版本是Switch Transformer[1]，随后通过循环利用(Up Cycling)稠密模型Mistral[2]进一步优化了该设计。

SwitchTransformer-MoE

随后，DeepSeek V2/V3/R1[3][4][5]通过引入共享专家[3]和门控偏差（gating bias）[4][5]进一步改进了MoE，最终实现了无辅助损失（auxiliary loss free）的MoE模型 [4][5]。这一优化本质上归因于一个关键事实：当使用共享专家（DeepSeek团队选择的值为1）时，可以通过在较大的专家池（256个上施加偏差分数的惩罚，从而缓解专家路由的不均衡问题[11]。

MoE层本质上是由多个专家前馈网络（FFN）组成的层，其中包含门控函数（gating functions），用于根据Top-K门控分数（DeepSeek V3/R1中引入偏差）进行激活路由，并在所选的FFN层上通过Group GEMM计算logits。

该功能在很大程度上依赖于基数排序（radix sort）逻辑。借助MoE Align & Sort，机器学习研究人员和实践者可以按照专家ID对tokens进行排序。

在某些应用中，例如TransformerEngine[6][7]，该操作最初是通过已废弃的cub::DeviceRadixSort实现的，而增加的permute操作用于记录源（左）到目标（右）的映射，其梯度操作为unpermute。

MoE Permute示例

尽管cub::DeviceRadixSort大量使用共享内存，相比于基于__shfl_xor_sync（仅使用线程本地内存）的实现略慢，但它不支持对齐排序（alignment sorting）。

对齐排序对于Group GEMM的效率至关重要，因为它允许专家以块（block 为单位处理tokens。

SGLang 中的MoE Align & Sort算法采用了对齐排序，但在支持多达256个专家的大规模prefill操作时效率并不理想。该问题已在issue#2732中被确认。

目前的实现将MoE Align & Sort拆分为两个kernel启动（kernel launches）：

对齐（alignment）：在单个block内执行传统基数排序算法对齐后的偏移计算（alignment-based offsets computation）;

相关推荐：男人边吻奶边挵进去AV片无码 1982年美国孕妇 暗夜官网入口app下载安装