Skip to content

Add opt-in reduced precision output for cuDNN MXFP8 norm - #3288

Open
sraman-rgb wants to merge 1 commit into
NVIDIA:mainfrom
sraman-rgb:fix/cudnn-mxfp8-norm-output-dtype
Open

Add opt-in reduced precision output for cuDNN MXFP8 norm#3288
sraman-rgb wants to merge 1 commit into
NVIDIA:mainfrom
sraman-rgb:fix/cudnn-mxfp8-norm-output-dtype

Conversation

@sraman-rgb

Copy link
Copy Markdown
Contributor

Description

Please include a brief summary of the changes, relevant motivation and context.

Fixes # (issue)

Type of change

  • Documentation change (change only to the documentation, either a fix or a new content)
  • Bug fix (non-breaking change which fixes an issue)
  • New feature (non-breaking change which adds functionality)
  • Breaking change (fix or feature that would cause existing functionality to not work as expected)
  • Infra/Build change
  • Code refactoring

Changes

Please list the changes introduced in this PR:

  • Change A
  • Change B

Checklist:

  • I have read and followed the contributing guidelines
  • The functionality is complete
  • I have commented my code, particularly in hard-to-understand areas
  • I have made corresponding changes to the documentation
  • My changes generate no new warnings
  • I have added tests that prove my fix is effective or that my feature works
  • New and existing unit tests pass locally with my changes

@github-actions github-actions Bot added the community-contribution PRs from external contributor outside the core maintainers, representing community-driven work. label Jul 30, 2026
Comment on lines +307 to 310
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));

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P2 Opt-in path lacks coverage

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-apps

greptile-apps Bot commented Jul 30, 2026

Copy link
Copy Markdown
Contributor

Greptile Summary

Adds an opt-in cuDNN MXFP8 normalization output mode.

  • Uses the FP16 or BF16 input datatype for the virtual normalization output when running MXFP8 block scaling with cuDNN 9.25 or later.
  • Adds the environment-variable helper declaration and documents its version constraints and default behavior.

Confidence Score: 5/5

The PR appears safe to merge.

No blocking failure remains.

Important Files Changed

Filename Overview
transformer_engine/common/normalization/common.cpp Gates the MXFP8 virtual normalization output datatype on the cuDNN version, scaling mode, and new environment variable.
transformer_engine/common/normalization/common.h Declares the internal environment-variable helper used by normalization graph construction.
docs/envvars.rst Documents the new opt-in variable, supported cuDNN version, datatype requirements, and fallback behavior.

Reviews (2): Last reviewed commit: "Add opt-in reduced precision output for ..." | Re-trigger Greptile

Signed-off-by: sraman-rgb <sraman@nvidia.com>
@sraman-rgb
sraman-rgb force-pushed the fix/cudnn-mxfp8-norm-output-dtype branch from 087adbe to c860887 Compare July 30, 2026 18:03
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

community-contribution PRs from external contributor outside the core maintainers, representing community-driven work.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant