perf(gdn): preserve Qwen3.5 H100 integration stack

[yaoyu-33]

Purpose

This draft preserves the exact public Megatron-Core integration used by the
Qwen3.5-35B-A3B 16xH100 model-card benchmark in
NVIDIA-NeMo/Megatron-Bridge#5098. It is a reproducibility ref while the source
upstream changes remain in flight.

Measured runtime commit:
606c01046bf411c54e730739844da9e613f84a76.

Current PR head:
323c5bdf72ac62e6249308fe5adfd22e9a1e243a.
Descendants of the measured commit update tests and repair the deterministic
torch fallback's FLA 0.4.2 l2norm call signature. The measured FlashQLA
performance path is unchanged.

Integrated upstream work

• refactor: extract and split common logic between GDN & GDN2 #5843: merged GDN/GDN2 refactor and integration base
• Bump FLA to 0.5.1 and integrate GDN2 kernels #5765: GDN2/FLA integration ancestry
• [dev] moe(perf): Add FlashQLA backend for Hopper/SM90 gated delta net (GDN). #5564: FlashQLA Hopper backend
• [main] moe(perf): Pre-GDR kernel fusion #5532: streamed pre-GDR fusion

The branch also carries adaptations of #5982 and #5396. They are inactive in
the accepted no-recompute, pre-GDR-fused path and are not credited for the
measured result.

Frozen runtime contract

• flash-linear-attention==0.4.2
• flash-qla==0.1.2
• transitive tilelang==0.1.9 and apache-tvm-ffi==0.1.9
• lazy FlashQLA loading fails closed unless the runtime is exactly 0.1.2
• fused gated RMSNorm is adapted to the refactored GDN path

The complete kernel stack is exposed through the opt-in
qwen35-h100-perf extra. It is explicitly incompatible with MCore's existing
ssm extra because that runtime currently requires TileLang 0.1.8.

Reproducibility scope

Bridge #5449 pins the current test-hardened PR head and supplies a hash-locked
container overlay;
stacked Bridge #5098 owns the recipe and model-card evidence. The benchmark
uses mock data and forced expert balancing, so it is performance stability
evidence rather than convergence evidence.

Validation

• exact dependency pins and public artifacts are represented in uv.lock
• focused FlashQLA version-contract tests passed on H100
• Black formatting issue from the first CI attempt was fixed at the frozen
  head
• deterministic-mode coverage explicitly selects its required torch backend
• deterministic GDN/GDN2 fallbacks match the pinned FLA 0.4.2 l2norm
  interface
• current-head GDN unit coverage, the broad SSM suite, and GB200 unit coverage
  pass
• all four affected current-head GDN functional jobs pass:
  gpt3_mcore_te_tp2_pp1_gdn, gdn_no_nvrx_async,
  gdn_no_nvrx_async_mcore, and gdn_no_nvrx_sync
• Black, isort, Ruff, DCO, installation, wheel, docs, and copyright checks
  pass; code-owner approval remains expected for merge

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[yaoyu-33]

/ok to test ca330d1

[yaoyu-33]

/ok to test 9f347b3

[yaoyu-33]

/ok to test 323c5bd

[yaoyu-33]

Freeze/validation note for current HEAD 323c5bdf72:

• All executable CI tests completed successfully, including current-head GDN unit coverage, the broad SSM suite, GB200 unit coverage, and all four affected GDN functional jobs.
• The remaining failures are merge-policy gates: code-owner approval, multi-approval, and Codecov patch coverage (25.74% vs 80%).
• The Codecov result reflects the large integrated upstream diff preserved by this reproducibility PR. It is not a failure of the Qwen3.5 model-card reproduction or measured FlashQLA path.
• Per the campaign freeze, no further performance exploration or unrelated coverage expansion is being added here.