@@ -984,10 +984,8 @@ def route_tokens_to_experts(router_logits):
984984 .sum (dim = - 1 )
985985 )
986986 group_idx = torch .topk (group_scores , k = topk_group , dim = - 1 , sorted = False )[1 ]
987- # one_hot + sum instead of scatter_: torch scatter_ lowers to a
988- # stablehlo.scatter whose dim-0 row index is a token iota; Shardy
989- # shards that iota WITHOUT the per-shard offset, so the group mask
990- # marks only mesh-row 0's tokens and zeros routing for rows 1-3.
987+ # one_hot + sum instead of scatter_: scatter_'s token-iota row index
988+ # isn't offset per shard under Shardy, zeroing routing for mesh-rows 1-3.
991989 group_mask = (
992990 (
993991 group_idx .unsqueeze (- 1 )
@@ -1007,11 +1005,8 @@ def route_tokens_to_experts(router_logits):
10071005 topk_indices = torch .topk (scores_for_choice , k = top_k , dim = - 1 , sorted = False )[
10081006 1
10091007 ]
1010- # one_hot + einsum instead of gather: torch.gather lowers to
1011- # ttir.embedding over the all-gathered [tokens*E] score table with a
1012- # flat index token*E+expert computed at fp16-class precision; for
1013- # mesh-rows 1-3 the large index rounds off the expert bits -> wrong
1014- # weights. einsum keeps the small expert index implicit and exact.
1008+ # one_hot + einsum instead of gather: gather's flat token*E+expert
1009+ # index rounds off at fp16 for mesh-rows 1-3, picking wrong weights.
10151010 _ohw = (
10161011 topk_indices .unsqueeze (- 1 )
10171012 == torch .arange (n_routed_experts , device = topk_indices .device )
0 commit comments