diff --git a/tests/pytorch/test_grouped_mlp.py b/tests/pytorch/test_grouped_mlp.py index 8c3bf5b2c0..b1e13fc5e9 100644 --- a/tests/pytorch/test_grouped_mlp.py +++ b/tests/pytorch/test_grouped_mlp.py @@ -1804,7 +1804,7 @@ def test_grouped_mlp_glu_mxfp8_real_cudnn_fusion( "a_dtype": torch.float8_e4m3fn, "b_dtype": torch.float8_e4m3fn, "alpha_dtype": torch.bfloat16, - "prob_dtype": torch.bfloat16, + "prob_dtype": torch.float32, "c_dtype": torch.bfloat16, "d_dtype": torch.float8_e4m3fn, "act_func": act_func, diff --git a/transformer_engine/pytorch/ops/fused/grouped_mlp.py b/transformer_engine/pytorch/ops/fused/grouped_mlp.py index 84edc056fb..05ad6d178c 100644 --- a/transformer_engine/pytorch/ops/fused/grouped_mlp.py +++ b/transformer_engine/pytorch/ops/fused/grouped_mlp.py @@ -1589,11 +1589,11 @@ def fuser_forward( fc1_prob_tensor = None if not unit_activation_scale: - fc1_prob_tensor = ( - scales.detach() - .to(dtype=torch.float32 if use_nvfp4 else dtype) - .reshape((-1,) if canonical else (-1, 1, 1)) - ) + fc1_prob_tensor = scales.detach().reshape((-1,) if canonical else (-1, 1, 1)) + if use_nvfp4: + # cuDNN NVFP4 kernel only accepts fp32 scales. + fc1_prob_tensor = fc1_prob_tensor.to(torch.float32) + if use_nvfp4: # cuDNN receives NVFP4 block-scaled inputs without TE's per-group # global scale factors, so alpha supplies the product of the two @@ -2701,9 +2701,6 @@ def fuser_backward( else: fc2_bias_grads = [fc2_dbias_packed[idx] for idx in range(num_groups)] - if grad_scales is not None: - grad_scales = grad_scales.to(dtype=dtype) - fc1_bias_grads: Optional[list[Optional[torch.Tensor]]] = None fc1_bias_grad_packed: Optional[torch.Tensor] = None if fc1_op.has_bias: