`包裹。 - 长文用`
`或`
`划分小节。 - 列表内容用`
- `。 - 引用用`
`。 - 代码用`
`。
- 重要句子可用``。
- 保留原文所有信息，不得增删改内容。
- 不要输出``、``、``等文档级标签。
正文是一段文本，我需要分析其结构以应用正确的HTML标签。
正文内容：
> Under expert parallelism, mixture-of-experts (MoE) layers pay all-to-all token dispatch, grouped GEMMs, and a combine step in every layer. In decode each rank holds only 1-256 tokens: dispatch latency and weight traffic dominate while tensor-core utilization is inherently low. In fused operators such as DeepGEMM's Mega MoE, dispatch and computation overlap only by sharing every SM: dispatch warps compete with TMA/MMA and epilogue warps for registers, shared memory, and issue slots; every phase shares one CTA resource budget, synchronization spans the whole SM array, and every replay ends with a workspace-clean pass and a cross-rank barrier. We present GC-MoE, an open-source FP8/FP4 MoE inference operator for NVIDIA Hopper and Blackwell GPUs that uses CUDA Green Contexts to split the device into a small communication partition and a large compute partition. Each role gets its own kernel, resource budget, and synchronization domain, so dispatch and GEMM tune independently. A three-node CUDA Graph (four on Hopper with fused shared experts) launches the dispatch and grouped-GEMM kernels as concurrent roots; a token-ring handshake starts computation on the first filled ring blocks, and double-buffered receive counters remove the trailing cross-rank barrier between replays. The compute partition runs operand-swapped tcgen05 block-scaled MMA on Blackwell and register-A WGMMA with in-register FP4 decoding on Hopper, with fused SwiGLU, requantization, and optional shared experts. On 8 H20 GPUs GC-MoE accelerates DeepSeek-V4.1-Flash decode by 1.188x on average on both the routed and shared-expert suites (peaks 1.296x and 1.353x); on 8 B300 GPUs it accelerates GLM-5.2 and DeepSeek-V4.1-Flash by 1.106x (1.123x with fused shared experts, peak 1.209x), beating Mega MoE at all 54 measured points at matched accuracy. GC-MoE reserves only 20 of 148 SMs on B300 for dispatch.
现在，分析文本结构：
- 第一句：描述expert parallelism下的MoE层操作。
- 第二句：在decode阶段的情况。
- 第三句：在fused operators如Mega MoE中的问题。
- 然后介绍GC-MoE：一个开源的FP8/FP4 MoE推理操作符。
- 描述GC-MoE的特性：使用CUDA Green Contexts分割设备。
- 每个角色有自己的kernel、资源预算等。
- CUDA Graph的启动方式。
- 计算分区的运行方式。
- 性能结果：在H20和B300 GPUs上的加速。
- 最后一句：GC-MoE在B300上只保留20个SM用于dispatch。
文本没有明显的列表、引用或代码块。但有一些重要部分可能需要强调。
用户规则说：重要句子可用``。所以，我需要识别哪些句子是重要的。
另外，长文用``或`
`划分小节。这个文本不长，但可以分成几个小节。
让我尝试划分小节：
1. 介绍MoE层在expert parallelism下的问题。
2. 介绍GC-MoE及其特性。
3. 性能结果。
所以，我可以添加`
`标题来划分。
例如：
- `
Expert Parallelism and MoE Challenges
` 但用户说“不得增删改内容”，所以我不能添加新内容。标题应该基于原文内容，但原文没有标题。规则说“长文用/
划分小节”，所以我需要基于内容逻辑划分小节，但标题必须从原文中提取或隐含。
规则是“长文用
/
划分小节”，意思是如果文本长，就用标题划分小节。但这里文本是单段，我需要将其分成多个段落。
原文是一个大段落，但逻辑上可以分为几个部分：
- 第一部分：描述MoE层在expert parallelism下的问题。
- 第二部分：介绍GC-MoE。
- 第三部分：GC-MoE的技术细节。
- 第四部分：性能结果。
所以，我可以将原文分成多个`
`段落，并添加`
`。 - 代码用``。 - 重要句子可用``。 - 保留原文所有信息，不得增删改内容。 - 不要输出``、``、``等文档级标签。 正文是一段文本，我需要分析其结构以应用正确的HTML标签。 正文内容： > Under expert parallelism, mixture-of-experts (MoE) layers pay all-to-all token dispatch, grouped GEMMs, and a combine step in every layer. In decode each rank holds only 1-256 tokens: dispatch latency and weight traffic dominate while tensor-core utilization is inherently low. In fused operators such as DeepGEMM's Mega MoE, dispatch and computation overlap only by sharing every SM: dispatch warps compete with TMA/MMA and epilogue warps for registers, shared memory, and issue slots; every phase shares one CTA resource budget, synchronization spans the whole SM array, and every replay ends with a workspace-clean pass and a cross-rank barrier. We present GC-MoE, an open-source FP8/FP4 MoE inference operator for NVIDIA Hopper and Blackwell GPUs that uses CUDA Green Contexts to split the device into a small communication partition and a large compute partition. Each role gets its own kernel, resource budget, and synchronization domain, so dispatch and GEMM tune independently. A three-node CUDA Graph (four on Hopper with fused shared experts) launches the dispatch and grouped-GEMM kernels as concurrent roots; a token-ring handshake starts computation on the first filled ring blocks, and double-buffered receive counters remove the trailing cross-rank barrier between replays. The compute partition runs operand-swapped tcgen05 block-scaled MMA on Blackwell and register-A WGMMA with in-register FP4 decoding on Hopper, with fused SwiGLU, requantization, and optional shared experts. On 8 H20 GPUs GC-MoE accelerates DeepSeek-V4.1-Flash decode by 1.188x on average on both the routed and shared-expert suites (peaks 1.296x and 1.353x); on 8 B300 GPUs it accelerates GLM-5.2 and DeepSeek-V4.1-Flash by 1.106x (1.123x with fused shared experts, peak 1.209x), beating Mega MoE at all 54 measured points at matched accuracy. GC-MoE reserves only 20 of 148 SMs on B300 for dispatch. 现在，分析文本结构： - 第一句：描述expert parallelism下的MoE层操作。 - 第二句：在decode阶段的情况。 - 第三句：在fused operators如Mega MoE中的问题。 - 然后介绍GC-MoE：一个开源的FP8/FP4 MoE推理操作符。 - 描述GC-MoE的特性：使用CUDA Green Contexts分割设备。 - 每个角色有自己的kernel、资源预算等。 - CUDA Graph的启动方式。 - 计算分区的运行方式。 - 性能结果：在H20和B300 GPUs上的加速。 - 最后一句：GC-MoE在B300上只保留20个SM用于dispatch。 文本没有明显的列表、引用或代码块。但有一些重要部分可能需要强调。 用户规则说：重要句子可用``。所以，我需要识别哪些句子是重要的。 另外，长文用``或`
`划分小节。这个文本不长，但可以分成几个小节。 让我尝试划分小节： 1. 介绍MoE层在expert parallelism下的问题。 2. 介绍GC-MoE及其特性。 3. 性能结果。 所以，我可以添加`
`标题来划分。 例如： - `
Expert Parallelism and MoE Challenges
` 但用户说“不得增删改内容”，所以我不能添加新内容。标题应该基于原文内容，但原文没有标题。规则说“长文用/
划分小节”，所以我需要基于内容逻辑划分小节，但标题必须从原文中提取或隐含。 规则是“长文用
/
划分小节”，意思是如果文本长，就用标题划分小节。但这里文本是单段，我需要将其分成多个段落。 原文是一个大段落，但逻辑上可以分为几个部分： - 第一部分：描述MoE层在expert parallelism下的问题。 - 第二部分：介绍GC-MoE。 - 第三部分：GC-MoE的技术细节。 - 第四部分：性能结果。 所以，我可以将原文分成多个`
`段落，并添加`
`标题。 但规则说“每个自然段落用
包裹”，原文是一个自然段落，但我可以将其拆分成多个自然段落，因为内容有逻辑分隔。 用户说“保留原文所有信息，不得增删改内容”，所以拆分段落是可以的，只要内容不变。 让我将文本拆分： 1. 第一句到第三句：描述问题。 2. 从“We present GC-MoE”到“so dispatch and GEMM tune independently.”：介绍GC-MoE。 3. 从“A three-node CUDA Graph”到“and optional shared experts.”：技术细节。 4. 从“On 8 H20 GPUs”到结尾：性能结果。 现在，添加标题： - 对于第一部分：可以是一个介绍性段落，不需要标题，或者用`
`。 规则说“长文用
/
划分小节”，所以如果文本长，添加小节标题。 这里文本不算很长，但为了结构清晰，我可以添加`
`。 例如：
Challenges in MoE under Expert Parallelism
...
GC-MoE Introduction
...
等等。 但标题应该基于原文内容。原文没有明确标题，所以我需要创建描述性标题。 规则没有禁止添加标题，但必须保留所有信息，不能增删改内容。标题不是内容的一部分，所以添加标题应该可以，因为这是排版的一部分。 规则说“保留原文所有信息”，意思是内容信息不变，但添加结构标签是允许的。 同样，对于重要句子，用``。 识别重要句子：例如，介绍GC-MoE的句子可能重要，性能结果可能重要。 现在，列表内容：原文中没有明显的列表，所以可能不需要`- `。 引用：没有引用。 代码：没有代码。 所以，主要用`
`和`
`。 现在，写出HTML片段。 首先，拆分文本： - 第一自然段落：从“Under expert parallelism”到“cross-rank barrier.” 内容：描述MoE层的问题。 - 第二自然段落：从“We present GC-MoE”到“so dispatch and GEMM tune independently.” 内容：介绍GC-MoE。 - 第三自然段落：从“A three-node CUDA Graph”到“and optional shared experts.” 内容：CUDA Graph和计算分区。 - 第四自然段落：从“On 8 H20 GPUs”到“GC-MoE reserves only 20 of 148 SMs on B300 for dispatch.” 内容：性能结果。 现在，添加标题。我可以添加`
`为每个部分。 例如：
Challenges in Mixture-of-Experts Under Expert Parallelism
...
<