Enabling grouped MLP fusion breaks previously working 128-aligned inputs
メンテナーはふだん 2 日以内に返信
まだ誰も着手していません。
評価
- 難易度
- 4/5
- 見積もり時間
- 3〜5日
- 初心者へのやさしさ
- 54/100
- issue の種類
- バグ
- 明瞭さ
- 明確に書かれている
- 活発さ
- 活発
調査の方向性
Save the supplied reproducer as repro_alignment.py and run it with NVTE_CUTEDSL_FUSED_GROUPED_MLP=1; compare the failing 128-row case with the padded and fusion-disabled runs. Start by tracing how grouped MLP fusion is selected and where per-expert row counts are checked. Done means misaligned groups no longer launch the failing fused kernel, and regression coverage verifies the fallback or other chosen handling.
索引モデルが issue の本文から書いたものです。
説明
My agent says:
The same MXFP8 grouped MLP inputs with 128-aligned expert sizes work correctly without fusion, but produce incorrect results when fusion is enabled. The fused cuDNN backend requires each expert’s row count to be a multiple of 256, while TE checks only the total row count and still selects the fusion.
On B100, 128 rows/expert produce incorrect, nondeterministic outputs and gradients; 384 rows/expert can cause illegal memory access. Padding each expert to 256 or disabling fusion passes the reproducer.
Tests missed this because the main fused numerical tests already generate 256-aligned groups. TE should enforce the requirement, pad, or fall back before launching the kernel.
Environment: B100, TE 2.21.0.dev0, PyTorch 2.15.0a0, CUDA 13.4, cuDNN frontend 1.29.0, CUTLASS DSL 4.8.0.dev0.
Tested reproducer (expand)
Save as repro_alignment.py; run each command in a fresh process:
export CUDA_VISIBLE_DEVICES=0 NVTE_FLASH_ATTN_V4=0 OMP_NUM_THREADS=8
NVTE_CUTEDSL_FUSED_GROUPED_MLP=1 python repro_alignment.py # fails
NVTE_CUTEDSL_FUSED_GROUPED_MLP=1 python repro_alignment.py --pad-to 256 # passes
NVTE_CUTEDSL_FUSED_GROUPED_MLP=0 python repro_alignment.py # passes
import argparse
import importlib.metadata
import os
os.environ.setdefault("NVTE_CUTEDSL_FUSED_GROUPED_MLP", "1")
os.environ["NVTE_GROUPED_LINEAR_SINGLE_PARAM"] = "0"
os.environ["NVTE_ALLOW_NONDETERMINISTIC_ALGO"] = "1"
import torch
import transformer_engine
import transformer_engine.pytorch as te
from transformer_engine.common.recipe import MXFP8BlockScaling
parser = argparse.ArgumentParser()
parser.add_argument("--rows", type=int, default=128)
parser.add_argument("--pad-to", type=int, choices=(128, 256), default=128)
args = parser.parse_args()
assert torch.cuda.is_available()
print("torch:", torch.__version__, "CUDA:", torch.version.cuda,
"TE:", transformer_engine.__version__,
"cuDNN frontend:", importlib.metadata.version("nvidia-cudnn-frontend"),
"CUTLASS DSL:", importlib.metadata.version("nvidia-cutlass-dsl"),
"GPU:", torch.cuda.get_device_name(), flush=True)
groups, hidden, ffn = 8, 2048, 1024
physical_rows = (args.rows + args.pad_to - 1) // args.pad_to * args.pad_to
torch.manual_seed(0)
common = dict(bias=False, dtype=torch.bfloat16, device="cuda")
model = te.ops.Sequential(
te.ops.GroupedLinear(groups, hidden, 2 * ffn, **common),
te.ops.ScaledSwiGLU(glu_interleave_size=32),
te.ops.GroupedLinear(groups, ffn, hidden, **common),
)
torch.manual_seed(2)
data = torch.randn(groups, args.rows, hidden, dtype=torch.bfloat16, device="cuda")
probabilities = torch.rand(groups, args.rows, dtype=torch.bfloat16, device="cuda")
splits = torch.full((groups,), physical_rows, dtype=torch.int64, device="cuda")
baseline = None
for iteration in range(5):
x = data.clone().requires_grad_()
probs = probabilities.clone().requires_grad_()
padded_x = torch.nn.functional.pad(x, (0, 0, 0, physical_rows - args.rows))
padded_probs = torch.nn.functional.pad(probs, (0, physical_rows - args.rows))
model.zero_grad(set_to_none=True)
with te.autocast(enabled=True, recipe=MXFP8BlockScaling()):
output = model(padded_x.reshape(-1, hidden), splits, padded_probs.reshape(-1), splits)
output = output.reshape(groups, physical_rows, hidden)[:, :args.rows]
output.backward(torch.ones_like(output))
torch.cuda.synchronize()
if iteration == 0:
selected = [type(op[0]).__name__ for op in model._module_groups[0]._forward_ops]
print("Selected ops:", selected, "Rows per expert:", physical_rows, flush=True)
expected_fused = os.environ["NVTE_CUTEDSL_FUSED_GROUPED_MLP"] == "1"
assert ("GroupedMLP_CuTeGEMMGLU" in selected) == expected_fused
current = {"output": output.detach(), "dx": x.grad, "dprob": probs.grad}
current.update({name: parameter.grad for name, parameter in model.named_parameters()})
assert all(value is not None and torch.isfinite(value).all() for value in current.values())
if baseline is None:
baseline = {name: value.clone() for name, value in current.items()}
else:
differences = {
name: (value.float() - baseline[name].float()).abs().max().item()
for name, value in current.items() if not torch.equal(value, baseline[name])
}
print("Repeat", iteration, "max absolute differences:", differences, flush=True)
assert not differences, "Identical inputs/weights produced different outputs or gradients"
print("PASS: all five forward/backward passes are bitwise identical", flush=True)
- 主要言語
- Python
- スター
- 3.6k
- フォーク
- 851
- 平均マージ
- 4日 15時間
- マージ済み PR(30日)
- 54
環境構築
- Dockerfile・Docker Compose ファイルなし
- プルリクエストのテンプレートあり
- コントリビューションガイドを読む
はじめの一歩
- issue を最後まで読み、次にプロジェクトのコントリビューションガイドを読みます。
- 着手することを issue にコメントします — 二人が同じ作業をするのを防げます。
- リポジトリをフォークし、ブランチを切って変更します。
- issue 番号を参照したプルリクエストを送ります。
NVIDIA/TransformerEngine のほかの issue
-
[Bug] group_quantize_fp8_blockwise: mbarrier invalidated before other threads finish waiting on itオープン
難易度 2/5 1〜3時間 初心者へのやさしさ 82/100
NVIDIA/TransformerEngine#3647 ·
メンテナーはふだん 2 日以内に返信
-
[PyTorch] fp8_cs_quantize fake implementation returns a vector inverse scale instead of a scalar対応中かも @sanjana658 が 4 日前に担当しました。 オープン
難易度 2/5 1〜3時間 初心者へのやさしさ 82/100
NVIDIA/TransformerEngine#3636 · コメント 2 件 ·
メンテナーはふだん 2 日以内に返信
-
[Bug] Backend selection picks FA3 for training with head_dim_qk=192 / v_head_dim=128, but FA3 backward cannot run it再び着手できるかも @modelpath-dev が 38 日前に担当しましたが、オープン中のプルリクエストはありません。 オープンattention
難易度 2/5 1〜3時間 初心者へのやさしさ 85/100
NVIDIA/TransformerEngine#3481 · コメント 4 件 ·
メンテナーはふだん 2 日以内に返信
-
bug
難易度 2/5 1〜3時間 初心者へのやさしさ 68/100
NVIDIA/TransformerEngine#2189 · コメント 7 件 · リアクション 5 件 ·
メンテナーはふだん 2 日以内に返信
-
難易度 4/5 3〜5日 初心者へのやさしさ 48/100
NVIDIA/TransformerEngine#3668 ·
メンテナーはふだん 2 日以内に返信
NVIDIA/TransformerEngine の issue をすべて見る
似ている issue
-
難易度 2/5 1〜3時間 初心者へのやさしさ 76/100
spec-kitty/spec-kitty#6112 ·
メンテナーはふだん 1 日以内に返信
-
[Bug]: OpenRouter streaming chunks report the requested router name instead of the model OpenRouter served対応中かも @Mrun25 が今日担当しました。 オープンbug llm translation
難易度 2/5 1〜3時間 初心者へのやさしさ 75/100
BerriAI/litellm#46095 · コメント 1 件 ·
メンテナーはふだん 1 日以内に返信
-
bug
難易度 2/5 1〜3時間 初心者へのやさしさ 72/100
uberdeveloper/fastbt#29 ·
-
bug
難易度 2/5 1〜3時間 初心者へのやさしさ 78/100
andrew-blake/melcloudhome#368 ·
メンテナーはふだん 1 日以内に返信
-
lane:discover
難易度 2/5 1〜3時間 初心者へのやさしさ 72/100
SPARDiffData/differential-data-collection#33 ·
メンテナーはふだん 1 日以内に返信