Triton current scaling: avoid casting amax input

[matthiasdiener]

Description

Suggested by @ipanfilo

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

[Copilot]

Pull request overview

This PR adjusts Triton current-scaling amax reduction to avoid casting the loaded input tile to float32, instead casting only the reduced tile_amax before atomic/store. This targets current-scaling behavior in the Triton cast+transpose path.

Changes:

• Remove tl.float32 cast on tl.load(...) in amax-reduction kernels.
• Cast tile_amax to tl.float32 only at the atomic/store sites.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[Copilot]

PR description still lists placeholder items (“Change A”, “Change B”) and doesn’t describe the actual change (moving float32 cast from load to the reduced amax result). Please update the description to match what this PR does so reviewers/users can understand intent and impact.

[wenchenvincent]

Please run level 3 CI tests before merging.

[matthiasdiener]

Please run level 3 CI tests before merging.

I just started a level 3 test here: https://github.com/ROCm/TransformerEngine/actions/runs/22327178479