Uh oh!
There was an error while loading. Please reload this page.
Add opt-in reduced precision output for cuDNN MXFP8 norm - #3288
Conversation
| const bool use_input_dtype = cudnnGetVersion() >= 92500 && _fp8_out && _ndim_scale_block == 1 && | ||
| use_cudnn_mxfp8_norm_output_in_input_dtype(); | ||
| const auto ZDtype = use_input_dtype ? itype : (_fp8_out ? ctype : otype); | ||
| _z->set_output(!_fp8_out).set_data_type(get_cudnn_fe_dtype(ZDtype)); |
There was a problem hiding this comment.
The existing MXFP8 normalization tests leave NVTE_CUDNN_MXFP8_NORM_OUTPUT_IN_INPUT_DTYPE disabled, so CI does not exercise graph construction or output correctness when this branch changes the virtual normalization output from FP32 to FP16 or BF16.
Knowledge Base Used:Common core (transformer_engine/common)
Note: If this suggestion doesn't match your team's coding style, reply to this and let me know. I'll remember it for next time!
Greptile SummaryAdds an opt-in reduced-precision intermediate output for cuDNN MXFP8 normalization.
Confidence Score: 5/5The PR appears safe to merge. No blocking failure remains. Important Files Changed
Reviews (4): Last reviewed commit: "Scope cuDNN MXFP8 dtype override to FP8 ..." | Re-trigger Greptile |
Signed-off-by: sraman-rgb <sraman@nvidia.com>
087adbe to
c860887CompareUh oh!
There was an error while loading. Please reload this page.
Signed-off-by: sraman-rgb <sraman@nvidia.com>
Signed-off-by: sraman-rgb <sraman@nvidia.com>
Uh oh!
There was an error while loading. Please reload this page.
Description
Please include a brief summary of the changes, relevant motivation and context.
Fixes # (issue)
Type of change
Changes
Please list the changes introduced in this PR:
Checklist: