Skip to content

Ifu dev 20260706 v2.17#660

Draft
AllenFarcas wants to merge 190 commits into
devfrom
IFU-dev-20260706-v2.17
Draft

Ifu dev 20260706 v2.17#660
AllenFarcas wants to merge 190 commits into
devfrom
IFU-dev-20260706-v2.17

Conversation

@AllenFarcas

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

ptrendx and others added 30 commits April 20, 2026 11:45
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Reduced padding kernel compilation time

Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com>

* Completely removed unroll for better performance

Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com>

---------

Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com>
…(#2905)

Zero out padded region when swizzling via group quantize

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
…2879)

* fix broken links in README

Signed-off-by: Santosh Bhavani <santosh.bhavani@live.com>

* update README to modernize description and standardize terminology

Signed-off-by: Santosh Bhavani <santosh.bhavani@live.com>

* Update README.rst

Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com>
Signed-off-by: Santosh Bhavani <santosh.bhavani@live.com>

* Removed the duplicate line

Signed-off-by: Przemek Tredak <ptredak@nvidia.com>

---------

Signed-off-by: Santosh Bhavani <santosh.bhavani@live.com>
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com>
Co-authored-by: Przemek Tredak <ptredak@nvidia.com>
* initial implementation for mxfp8

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* semi-working FP8; broken F16

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* clean up last commit

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* comment out F16 pass

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* pull in grouped_quantize for MXFP8

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* grouped tensor - pytorch

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* quantize mxfp8

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix shapes/strides

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix unfused; clean up

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* split d to d_qk/d_v; attempt at bwd

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix last merge

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* update FE

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* attempt at SWA/MLA

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove prints

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove leftover prints

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* Revert "update FE"

This reverts commit d9ff5662aa4b4b6267c77baf614aada6602fa133.

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* update FE

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix MLA O strides; add bottom_right_diagonal

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* attempt at bwd

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix get_quantizers; attempt at bwd

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix fprop; add o_format

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* attempt at bwd with o_format/d_out_format/dqkv_layout

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix dtype/o_format/etc in bwd calls

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix generateMatrixStridesWithFormats and _v1; fix padding for mxfp8

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix upon last commit for paddedsizes

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add mxfp8 env var

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* disable FA for mxfp8

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add mha test

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* attempt at bwd; force determinism; fix shapes

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove prints

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* update FE

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* update FE from pre-merge branch to post-merge develop

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* allow MXFP8 linear + f16 attn

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* test cp a2a

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove prints temporarily

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* test cp p2p

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* minor fixes for mla

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* open up a2a for mla

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* test ag

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* tweaks for last commit

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* enable mla ag

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix merge

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix merge

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* revert to main grouped tensor impl

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* minor tweaks to return to main

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove prints

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix combine_and_quantize for f16

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* minor tweaks

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* tweak tests

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix ds descale_o

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* Revert "fix ds descale_o"

This reverts commit cd0bd82e239ff01210338b4e34cb8784109d22ec.

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* minor fixes for p2p and ag

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* tweak cp test skips

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* update FE

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix bwd KV tensors

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* tweak recipe control and backend selection

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* tweak quantizer logic

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* minor fixes after last two commits

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* improve generate strides

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* minor fixes for previous commit

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix bwd for current/delayed

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* tweak test configs

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix dO/dO_f16 strides

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix tests: SWA logic/test configs

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix ag

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add fp8 sink attn

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix a2a comm for F16

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove nan/inf print in test

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix fa a2a

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix fa a2a+p2p f16

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* update FE to include new fixes

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix thd for bwd

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* refactor a2a for fu/fa

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* update FE to fix d64

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* refactor ag

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* refactor p2p/a2a+p2p; mostly regarding shapes

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add shadow f16 fwd

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* update FE to fix SWA/BRCM

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* switch to GH FE temporarily

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* switch back to GL FE

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* update FE to latest commit

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* update group tensor usage after merge main

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* env vars for qdq(q,k), o_f16 tests

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* allow other recipes than mxfp8

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix grouped tensor for MLA

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* change cp test configs

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* add shadow f16 bwd

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix a2a+p2p for sbhd

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix last commit and causal flag for fa

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* enable fp8 sink and disable fp8_mha

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* minor cleanup for cp/non-cp

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* update FE for FP8 sink

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix TE for FP8 sink

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* temporary: random sink/print sink

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* Revert "temporary: random sink/print sink"

This reverts commit 706095f802e04cbdd5d88ee53849cc5ec938203f.

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* replace d_out_format with do_format

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix compare_and_assert for None cases

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* remove logic for b and simplify logic for dqkv types

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* minor fix for ndim_q/kv

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add explanation of fp8_output/grad in MHA

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* tidy up FP8 checks for bhsd/learnable

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove leading underscores in nvte_convert_qkv_format

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* simplify logic in generateMatrixStridesWithLayout

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* clean up strides/ifelse-recipe logic

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* tweak checks in utils.py

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* tweak UnfusedDPA

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* enable testing for ag+swa and disable fp8_mha

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* tweak FusedAttn, fp8/f16 tensor naming/docstring

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* replace d_out_format with do_format

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix lint

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* clean up a2a

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* clean up ag

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* clean up p2p/a2a+p2p

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* tweak test configs

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* qdq dO in bwd shadow f16 path

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* tweak qdq dO logic

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove prints in shadow paths

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* update FE to allow non-determinism

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fuse qkv transposes; first pass

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remap parallelism to grid(bh, splits, 3) block(s/splits x d); use nvec = 128 bits

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* allocate contiguous block for qkv

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix grouped tensor row/col scale_inv offsets

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* use fused permute kernels

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* quantize row/col as needed in fwd/bwd, non-cp/cp

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* Revert "quantize row/col as needed in fwd/bwd, non-cp/cp"

This reverts commit ca5376956e8b8f662c7fa88661695b3e9eda4f8f.

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* Reapply "quantize row/col as needed in fwd/bwd, non-cp/cp"

This reverts commit f19e852be3463210f2b3be5839ae8931e5ad92d0.

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix v_col format when row is quantized

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add back necessary bwd quants for shadow paths/cp a2a

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove ZInv for all layouts except T3HD

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix cp p2p with zinv

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* temporarily switch to GH FE main

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* switch back to GL FE

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix ag after merge main

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add condition for qdq(do) to not affect other tests

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix custom_mha_fp8 test

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix amax dqkv

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix fp8_recipe in DPA utils

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove use of amax for mxfp8

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* add o_format/do_format/dqkv_layout to cache indicators for fp8 and f16

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* enable sink attn + FP8 in CP

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* update FE to GH v1.22.0

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix for inconsistent kwarg name in permute to grouped tensor

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add TMA permute

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* Revert "add TMA permute"

This reverts commit 2532a50e829144bee290fc94acb8f3f154a62ea9.
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* TMA load for bhsd transposes

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix some lint

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* temp: quant+perm+swizzle, rope, perm_fused

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove mla_rope for now; clean up quant+permute+pad_swizzle; create multi_tensor_swizzle

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix last commit

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* implement narrow-m for col swizzle; reorder to pad+perm+swizzle

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fused pad into perm; remove at::zeros as zeros done in perm kernels

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove shadow code

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* minor fix for permute shapes

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* check smem size before entering narrow-k/m kernels

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* expand permute to multi_tensor_

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* refactor qkv/do quant; create a fast_path call

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix lint

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* cleanup grouped tensor fix

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove _with_amax for create_unquantized_tensor

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix last commit

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* reimplement inplace_multi_tensor_swizzle

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix last commit; set swizzled flag in python

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove permute_to_grouped_tensor_bwd; clean up fwd

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add doxygen for multi_tensor_swizzle

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* clean up nvte_convert_qkv_format

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fixes based on code review

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* group layouts/formats in APIs

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* rename nvte_convert_qkv_format to nvte_convert_qkv_shape

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove MXFP8 create_unquantized_tensor

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* rename permute_to_grouped_tensor to transpose_to_bhsd

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add multi_tensor_swizzle_xx_unchecked and split the calls/paths

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* straighten up indexing for multi_tensor_pad

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* batch up kernel calls per-16-tensors for pad and permute

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* remove nvec128; rename nvec64 back to nvec

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add Macros/arch specifics for compilation

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix lint

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* attempt 1: MLA RoPE

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* Revert "attempt 1: MLA RoPE"

This reverts commit 79229248718d26a0ae7029206adc26c687bb42a7.
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix kv_cache tests for Fused, is_page=True

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* attempt 2: MLA RoPE

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* use DIVUP/_TO_MULTIPLE for pad_s_d_for_mxfp8

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove CUDNN_VERSION 8900 macros

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add narrow-k/m swizzle tests

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* compile flash_attn.cu with special archs

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* Revert "attempt 2: MLA RoPE"

This reverts commit 3b854b29a3677de2005fecff821d801ccd9bf5d4.
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* make contiguous instead of check is_contiguous

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove unused s_q/s_kv

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* remove unused issue_tma_store_strided

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* add version gate for mxfp8 for CPP users

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* replace nvte_get_qkv_shape with AttentionShape

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* populate nvte_ changes to Jax

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* update FE to 1.22.1

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix minor merge issue

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* revert to FE 1.21 since it's what mxfp8 needs

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* udpate jax attention shapes

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Revert "revert to FE 1.21 since it's what mxfp8 needs"

This reverts commit f09961a03bd7f5a316474b5d77b8292c7a49c1a6.
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* pick FE 1.22 to support mxfp8 and avoid rng issue in 1.22.1

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* fix CP AG test on Hopper

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

---------

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
* starting grouped linear integration, not tested, grouped_bias_add optimized and uses scales now

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* all changes

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* optimize grouped bias add kernel to 4TB/s handling load imbalance

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* kernel optimized + review comments

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Remove uneecessary load from the kernel

Remove unnecessary pointer reinterpretation for bias input.

Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>

* No need for text for grouped bias add

Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>

* Remove lru_cache from _is_deterministic_mode

Removed lru_cache decorator from the _is_deterministic_mode function.

Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>

* address more review comments

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* for better perf

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* remove unecessary comments

Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* rename a bit

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

---------

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>
Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
* Add better ordering enforcment to split_overlap_rs gemms.

This adds a short delay kernel to the split_overlap_rs function, which
ensures that the gemms are properly ordered when run with cuda graphs.

Signed-off-by: Chase Block <cblock@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

---------

Signed-off-by: Chase Block <cblock@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Björn Buschkämper <bjoern.buschkaemper@gmail.com>
…caches (#2805)

* fixing mem leaks

Signed-off-by: Peter St. John <pstjohn@nvidia.com>

* update xfail message

Signed-off-by: Peter St. John <pstjohn@nvidia.com>

* addressing greptile comments

Signed-off-by: Peter St. John <pstjohn@nvidia.com>

* remove fsdp_safe

Signed-off-by: Peter St. John <pstjohn@nvidia.com>

* fix Float8BlockScaling backward override, unused imports, xfail MXFP8 HSDP

Fix three issues:

1. LayerNormLinear weight quantizer was requesting columnwise usage even
   when backward_override is set. The FSDP2 refactor on this branch
   changed the columnwise condition to `is_grad_enabled and not is_fsdp2`
   but omitted the `backward_override is None` guard. This caused the
   weight to carry unnecessary columnwise data, which the dequantized
   backward path rejects. Also harden Float8BlockQuantizer::create_tensor
   to pass py::none() based on usage flags rather than relying on
   at::Tensor::defined() which is unreliable for default-constructed
   tensors in some PyTorch builds.

2. Remove unused imports: Float8Tensor from linear.py, and the entire
   float8_tensor import line from layernorm_linear.py (Float8Quantizer,
   Float8CurrentScalingQuantizer, Float8Tensor were all unused).

3. xfail MXFP8BlockScaling + fp8_init + HSDP in FSDP2 model tests.
   Pre-existing bug (confirmed on main) where fsdp_post_all_gather
   receives fewer output tensors than fsdp_pre_all_gather sent when the
   HSDP shard dimension is trivial (size 1).

Signed-off-by: Peter St. John <pstjohn@nvidia.com>

* address review comments

Signed-off-by: Peter St. John <pstjohn@nvidia.com>

* address greptile review

Signed-off-by: Peter St. John <pstjohn@nvidia.com>

* no need

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* revert grouped linear changes

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* fix all tests

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* fix cudnn commit

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* fix lint

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* revert

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* fix lint

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

---------

Signed-off-by: Peter St. John <pstjohn@nvidia.com>
Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>
Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>
Co-authored-by: Varun Thumbe <vthumbe@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>
…rentScalingQuantizer (#2816)

* Make Float8CurrentScalingQuantizer stateless (no amax/scale members)

Remove amax, scale, and use_existing_amax from
Float8CurrentScalingQuantizer on both C++ and Python sides.
All amax/scale allocations are now ad-hoc at quantization time:
- quantize() allocates a combined 2-element tensor for amax+scale
- quantize_with_amax() accepts amax as a parameter
- create_unquantized_tensor_with_amax() returns amax in a tuple
- set_quantization_params() is now a no-op

Update all call sites in activation.cpp, bias.cpp, normalization.cpp,
attention.cpp, and cast.cpp to propagate the amax buffer.

For FusedAdam FP8 kernel: when scale_ptr is null (current scaling),
derive scale from scale_inv and skip writing amax/scale_inv metadata.
Python side passes empty(0) tensors for scale/amax to signal this.

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
Made-with: Cursor

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Guard quantizer.amax/scale access with recipe.delayed() checks

Add delayed() guards in context_parallel.py for amax writeback and
scale cloning, since Float8CurrentScalingQuantizer no longer has
these attributes. Allocate scratch scale buffer in
_cast_master_weights_to_fp8_current_scaling instead of reading
quantizer.scale.

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
Made-with: Cursor

* Fix pylint unused-argument warnings in Float8CurrentScalingQuantizer

The device, use_existing_amax, scale, and amax parameters are kept
for backward compatibility but not used internally.

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
Made-with: Cursor

* Clear dangling scale/amax pointers in Float8CurrentScalingQuantizer::quantize_impl

amax_buf and scale_buf are caller-owned local tensors whose storage is
released as soon as quantize() returns, leaving raw pointers stored in
`out` dangling. No current caller reads out.scale()/out.amax() after
quantize_impl returns, so this is currently safe, but it is a silent
invariant that could turn into a use-after-free if new callers are added.

Defensively clear both pointers at the end of quantize_impl (and in the
empty-input early return), mirroring the existing set_amax(nullptr, ...)
call already present before nvte_quantize_v2.

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
Made-with: Cursor

* Route current-scaling Float8Tensor through FP32 master path in FusedAdam

Float8Tensor with Float8CurrentScalingQuantizer now goes through the
FP32 master + requantize path (same as MXFP8/NVFP4/blockwise) instead
of the fused FP8 Adam kernel. The fused FP8 kernel stays for delayed-
scaling Float8Tensor only.

Also revert adam.cu to upstream — current scaling no longer needs the
scale_inv-derived path in the kernel.

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
Made-with: Cursor

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Restore intent/perf comments in FusedAdam quantized dispatch

- Note about a possible fused Adam+requantize kernel removing the
  FP32 round-trip for the QuantizedTensor path.
- Justification for casting BF16/FP16 grads to FP32 before the FP32
  Adam kernel.

Also drop a stray f-string prefix on a literal-only RuntimeError part.

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
Made-with: Cursor

* tests skip

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>

* tests skip

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>

* Warn on deprecated kwargs in Float8CurrentScalingQuantizer

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>

* Guard remaining quantizer.scale/amax accesses with recipe.delayed()

Followup to e06253d7. Three more spots accessed Float8CurrentScaling-
Quantizer.scale/.amax (which no longer exist after the stateless
refactor):

- AttnFuncWithCPAndKVAllGather.forward used `not mxfp8()` instead of
  `delayed()` when cloning .scale, causing AttributeError under
  Float8CurrentScaling.
- AttnFuncWithCPAndKVP2P fwd/bwd assigned .amax on per-step quantizer
  copies under `not mxfp8()`; harmless for current scaling (Python
  attaches a dynamic attribute that nobody reads) but inconsistent with
  the delayed()-guarded amax aggregation that follows.
- print_quantizers (debug-only) read .scale/.amax for both DS and CS;
  restrict to DS only.

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>

---------

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
fix fp8 and is_bwd_fp8 relationship

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>
… (#2728)

* fix: scope get_full_cu_seqlens cache key by device and inference mode

Signed-off-by: Dongmin Ra <dongmin.ra@navercorp.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

---------

Signed-off-by: Dongmin Ra <dongmin.ra@navercorp.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>
Fix FA4 selection when FA3 is unavailable.

Signed-off-by: Björn Buschkämper <bjoern.buschkaemper@gmail.com>
* Add AIR TopK support to TE JAX extension

Adds a custom AIR TopK implementation (header-only, vendored into
transformer_engine/common/util/) exposed as a JAX FFI custom call
via the TE JAX extension.

Key changes:
- transformer_engine/common/util/air_topk.cu: AIR TopK CUDA kernel
- transformer_engine/common/util/standalone_air_topk.cuh: vendored header
- transformer_engine/common/include/transformer_engine/air_topk.h: C API
- transformer_engine/jax/csrc/extensions/air_topk.cpp: JAX FFI binding
- transformer_engine/jax/cpp_extensions/air_topk.py: Python wrapper
- CMakeLists.txt: compile new kernel; use CCCL from CUDA toolkit
- CMakeLists.txt: fix SM100 arch handling when all arches are special-cased

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: dcampora <961215+dcampora@users.noreply.github.com>
Signed-off-by: Daniel Campora <961215+dcampora@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

Signed-off-by: Daniel Campora <961215+dcampora@users.noreply.github.com>

* Address PR review comments: fix namespace pollution, unused var, missing export, cache sm_cnt

- Move WARP_SIZE/WARP_BITS/FULL_WARP_MASK/VECTORIZED_READ_SIZE into namespace nv
- Remove unused keys_element_bytes variable in AirTopkFFI; collapse switch to dtype validation
- Add missing `from .air_topk import *` export in jax/cpp_extensions/__init__.py
- Cache sm_cnt per device with static vars to avoid repeated cudaGetDevice/cudaDeviceGetAttribute calls
- Add CMAKE_BUILD_WITH_INSTALL_RPATH=ON to build_ext.py

Signed-off-by: dcampora <961215+dcampora@users.noreply.github.com>
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Daniel Campora <961215+dcampora@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

Signed-off-by: Daniel Campora <961215+dcampora@users.noreply.github.com>

* Rename air_topk -> topk throughout JAX extension

Remove the `air_` prefix from all TopK-related identifiers: file names,
C API functions (nvte_air_topk -> nvte_topk), FFI handler/primitive names
(te_air_topk_ffi -> te_topk_ffi), Python symbols, and the internal
`air_topk` namespace in standalone_topk.cuh.  No functional changes.

Signed-off-by: Diego Campora <dcampora@nvidia.com>
Signed-off-by: dcampora <961215+dcampora@users.noreply.github.com>
Signed-off-by: Daniel Campora <961215+dcampora@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

Signed-off-by: Daniel Campora <961215+dcampora@users.noreply.github.com>

* Address ptrendx review comments and fix CI lint issues

- Follow TE workspace convention: remove nvte_get_topk_workspace_bytes()
  and implement empty-workspace size-query pattern in nvte_topk() instead
- Remove unnecessary nv_detail::float_to_T helper; replace usages with
  static_cast<T>() directly
- Remove unrelated CMAKE_CUDA_ARCHITECTURES OFF block from CMakeLists.txt
- Fix cpplint errors in standalone_topk.cuh: replace unsigned long long int
  with uint64_t, add NOLINT for constexpr-sized arrays and else-with-comment

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Daniel Campora <961215+dcampora@users.noreply.github.com>

* Add assertion for 2D input in topk

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Daniel Campora <961215+dcampora@users.noreply.github.com>

---------

Signed-off-by: dcampora <961215+dcampora@users.noreply.github.com>
Signed-off-by: Daniel Campora <961215+dcampora@users.noreply.github.com>
Signed-off-by: Diego Campora <dcampora@nvidia.com>
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
…ed quantization kernel (#2921)

Fix the race in the dbias computation

Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* remove ctype to eliminate memory usage from the cudnn kernel

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* Remove c_dtype from fusible ops test

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>

---------

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* [Common, PyTorch] Add triton mHC kernels & pytorch operators

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* fix

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* nit

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* make linter happy

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* nit

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* ah OK

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* new configs to improve perf

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* add APIs to docs

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* fix typos, check deterministic, refactor

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* fix

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* reset rng for all tests

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* add docstring

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* fix api doc

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* whoops

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* grad_x doesn't have to zero

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* nit

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* nit

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* force pytorch to not use bf16 for reduction

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* use TE's general_gemm instead

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Looks like this is how to make TE use fp32 acc

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

---------

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
…True (#2936)

* fix

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* add test

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* cleanup

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* zero_out should also be tested

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

---------

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: root <root@gb-nvl-059-compute03.nvidia.com>
… (#2924)

* Fix contiguous path for k=2880

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>

* format

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>

* Review suggestion from @Oleg-Goncharov

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>

* Add test for swizzle + padding fusion

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>

* Address review comments

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>

---------

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
…942)

* accumulate bias in fp32 instead of bf16 in ref impl dbias to avoid accumulated numerical error

Signed-off-by: tdophung <tdophung@nvidia.com>
* loosen up thresholds. Also only check min loss of last 10% of steps to avoid failing by noise near convergence

Signed-off-by: tdophung <tdophung@nvidia.com>

* add deterministic flag for mnist run 

Signed-off-by: tdophung <tdophung@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
* feat: add support for grouped GEMM swizzling with variable shapes and update C++ operator interface
Signed-off-by: Abhishek <abhi.dtu11@gmail.com>

* Added confirmation with uniformity in one of the dimensions
Signed-off-by: Abhishek <abhi.dtu11@gmail.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Using single kernel for variable m and k
Signed-off-by: Abhishek <abhi.dtu11@gmail.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Cached blocks per sm for device and removed redundant checks
Signed-off-by: Abhishek <abhi.dtu11@gmail.com>

* Updated the code with newer changes in main
Signed-off-by: Abhishek <abhi.dtu11@gmail.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

---------

Signed-off-by: Abhishek <abhi.dtu11@gmail.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: vthumbe1503 <vthumbe@nvidia.com>
…929)

* Avoid removing usages from quantized weight in linear op

Quantized weight tensor may be used across steps, so removing a usage is not safe.

Signed-off-by: Tim Moon <tmoon@nvidia.com>

* Tweak test to catch bug when alternating train and infer steps

Signed-off-by: Tim Moon <tmoon@nvidia.com>

* Avoid removing usages from quantized weights in grouped linear op

Signed-off-by: Tim Moon <tmoon@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Restore pre-forward quantizer config in ops

Turns out we still need this in case the quantizer is used before the forward, e.g. in previous ops or CPU offloading.

Signed-off-by: Tim Moon <tmoon@nvidia.com>

* Blindly preserve quantizer usages in quantized weight params.

Signed-off-by: Tim Moon <tmoon@nvidia.com>

---------

Signed-off-by: Tim Moon <tmoon@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
… (#2943)

* Patch for NVFP4 test suite

Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com>

* C++ tests fix

Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com>

* Cleanup

Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Removed dead code

Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com>

* Set the golden value for amax in tests

Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com>

* Fixed memory leakage

Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com>

---------

Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
… (#2948)

* Switch to cuDNN-FE min version 1.23.0 to enable fused grouped MLP

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>

* Fix tests

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>

---------

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
…n expert (#2947)

Add workaround for cuteDSL stride requirement for zero token expert

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Remove unused NVFP4 quantize kernel

Signed-off-by: Tim Moon <tmoon@nvidia.com>
Co-authored-by: Oleg Goncharov <64355998+Oleg-Goncharov@users.noreply.github.com>
@github-actions

Copy link
Copy Markdown

Claude Walkthrough

Intent. IFU (integrate-from-upstream) of NVIDIA TransformerEngine release_v2.17 (156 commits, v2.15 → v2.17) into the ROCm dev branch. The bulk of the diff (~453 files, +70k/-18k) is the upstream merge itself; the rest is a build-fix commit plus ~25 targeted ROCm-specific patches to make v2.17 compile, run, and pass CI on gfx942/gfx950, plus a CI infrastructure refactor.

Key changes.

  • Upstream v2.17 merge (fbcaca7a): 78 conflicts resolved. Notable reconciliations — MXFP4 enum values shifted to 10/11 so upstream NVFP4 can take 8/9; _Linear refactored onto upstream's LinearFwdArgs/LinearBwdArgs dataclasses with ROCm fields folded in; the test Tensor helper adopted upstream's Buffer/CudaPtr model; new upstream NCCL Expert-Parallelism subsystem force-disabled on ROCm.
  • Build-fix pass (0cf6b69c): guards for CUDA-only kernels newly introduced upstream — row-scaled NVFP4 amax (transformer_engine/common/cast/dispatch/quantize.cuh), variable-shape grouped swizzle (transformer_engine/common/swizzle/swizzle.cu), TMA multi-tensor transpose in transformer_engine/common/fused_attn/flash_attn.cu, <cuda_pipeline.h> in transformer_engine/common/fused_router/async_loader.h. Restores sum()/max() helpers dropped by upstream's warp-reduce rewrite and gates warp-shuffle widths for wave64. Reconciles nvte_fused_attn_fwd/bwd with new format-descriptor params. MXFP4Quantizer updated to the new Quantizer virtuals.
  • CI refactor (71eb59d6): the sGPU test runner previously inlined in .github/workflows/rocm-ci.yml is extracted into .github/scripts/run_parallel_sgpu.sh + .github/scripts/ci_sgpu_jobs.conf (see .github/workflows/rocm-ci.yml:188); adds progress logging, persistent JAX compilation cache, and per-suite .rc files instead of FAIL_* sentinels.
  • NVFP4 correctness fixes: portable nvfp4::compute_rowwise_amax extracted so ROCm can enable row-scaled NVFP4 quantization (transformer_engine/common/cast/nvfp4/rowwise_amax_nvfp4.cuh, 93f88883); arch-aware E4M3 max restored in the fnuz dequant scale (transformer_engine/common/cast/nvfp4/dequantize_nvfp4.cuh, 4e94fd7a); fnuz-aware reference global scale (transformer_engine/pytorch/custom_recipes/quantization_ref_nvfp4.py, be3feea9); inplace_swizzle_scale_for_gemm ROCm early-return to fix NoneType scale_inv crash (transformer_engine/pytorch/csrc/extensions/swizzle.cpp, bebd2475); CPU-offload dequant recursion fix (9c2bbdad); accept NVFP4TensorStorage in update_quantized (cdd4e9ca).
  • Fused-attention fixes: neutralize non-finite per-step LSE/grads under is_hip_extension() guard for gfx942 striped ring-attention with sliding window (transformer_engine/jax/cpp_extensions/attention.py, 19d3daec); bwd layout fix for CPU-offload reload (a6bbb7c9); don't force FlashAttention THD pad-between-seqs leg on ROCm (990bf8bf).
  • Other correctness: FSDP2 all-gather thread scale_inv through metadata for Float8 current-scaling (transformer_engine/pytorch/tensor/fsdp2_allgather_tensor.py, 2f2b5fc6); route grouped-tensor GEMM through te_general_grouped_gemm on ROCm (306a371c); add missing roles= param to MXFP4BlockScalingRecipeState (c99b4d74); Triton current-scaling quantize fix (transformer_engine/pytorch/triton_kernels/cast.py, fdc9510a); circular import fix in transformer_engine/pytorch/quantized_tensor.py (f88abf74).
  • Test guards / xfails for ROCm-unsupported v2.17 features: 4over6 NVFP4, grouped NVFP4, Blackwell-only RHT swizzled-SF, standalone JAX TopK, NVFP4 async DCP parity, primary-FP8 + FSDP2 forward-memory accumulation, large-M grouped-GEMM allocations, cpp operator tests missing scale accessors.

Walkthrough.

The overwhelming majority of the diff is the mechanical upstream merge; the interesting parts are what the ROCm fork had to change on top.

C++ common library. The v2.17 merge brings row-scaled NVFP4, a new grouped-swizzle kernel with variable shapes, a TMA-based multi-tensor transpose path, and a new fused-router API. The build-fix commit guards these behind USE_ROCM/__HIP_DEVICE_COMPILE__ where the underlying primitives (Blackwell CUtensorMap, wave32-only shuffles, <cuda_pipeline.h>) have no HIP equivalent. Row-scaled NVFP4 is then unlocked in a follow-up (93f88883) by extracting the amax kernel out of the Blackwell-only quantize-transpose header into rowwise_amax_nvfp4.cuh, since the ROCm blockwise quantize/dequantize kernels already handle row-scaled inputs — this was the piece answering a review comment about "did we not support rowwise NVFP4?" (see the tip commit 6369177f). Separately, the fnuz dequant fix (4e94fd7a) is subtle: upstream hardcoded E4M3_MAX=448 in the dequant factor, but fnuz archs have maxNorm=240, so the un-fixed code was scaling outputs by ~1.87×; the fix restores detail::TypeExtrema<fp8e4m3>::max on HIP device compilation.

Fused attention. transformer_engine/common/fused_attn_rocm/fused_attn.cpp was reconciled with upstream's new format-descriptor params on nvte_fused_attn_fwd/bwd (strict on layout, permissive on FP8 scale-inv). transformer_engine/common/include/transformer_engine/fused_attn.h retains the NVTE_F16_max512_seqlen enumerator with a restored explanatory comment (6369177f) — upstream v2.17 dropped the max512 impl/dispatch but the enumerator is kept so the ROCm fork's CUDA-path bindings still compile. The JAX ring-attention NaN fix (19d3daec) neutralizes non-finite per-step output/LSE (fwd) and grads (bwd), gated on is_hip_extension() and a no-op on finite values, fixing 5 test_context_parallel_ring_attn failures with window_size=(20,0) + THD_SEPARATE + PADDING_CAUSAL.

PyTorch bindings. MXFP4Quantizer::create_tensor / create_grouped_tensor were rewritten to match the new base Quantizer virtual signatures (device/pin_memory, precomputed tensor_offsets). Multiple entry points that only exist upstream (grouped-tensor GEMM, fused grouped-linear path, primary-FP8 + FSDP2 fwd memory accumulation) are guarded so ROCm falls through to the general grouped-gemm path.

FSDP2 all-gather. transformer_engine/pytorch/tensor/fsdp2_allgather_tensor.py previously reconstructed scale_inv via 1 / quantizer.scale, which crashed on Float8CurrentScalingQuantizer because current scaling computes its scale in-kernel and exposes no .scale attribute. Since fsdp_pre_all_gather already runs an amax reduction (every rank's per-tensor scale_inv is identical), the fix threads the scale_inv through the all-gather metadata and uses it in post_all_gather for the current-scaling path; delayed scaling is unchanged.

CI infrastructure. The old inline heredoc in rocm-ci.yml that fanned out Torch/JAX/examples/core on separate GPUs is replaced by a general driver run_parallel_sgpu.sh reading a job table from ci_sgpu_jobs.conf, plus run_examples.sh for the examples slot. Per-job .rc files replace the FAIL_* sentinel scheme. Timeouts adjusted (sGPU 360→270 min but inner step 180→240 min; mGPU 360→210 min). JAX picks up a persistent compilation cache.

Testing. No new test functionality is added; the delta is mostly guards — skips and xfails for ROCm-unsupported v2.17 features (grouped NVFP4, 4over6, Blackwell-only RHT, standalone JAX TopK, primary-FP8+FSDP2 fwd mem accumulation, NVFP4 async DCP parity under torch 2.8), plus stale -k filter fixes after upstream's recipe-id rename (bd8c4042). Row-scaled NVFP4 is stated verified by OperatorTestRowScaled/FusedCastTransposeNVFP4TestSuite (32/32). The upstream merge commit explicitly notes it was static-validated only (no GPU in the resolution env); the follow-up ~25 commits are the observed-failure fixes from actually running CI.

Notes for reviewers.

  • Base is dev, not main. This large IFU will land on dev; expect a subsequent dev → main promotion after downstream soak.
  • ROCm guard style is inconsistent by design: some guards are USE_ROCM (C++ preprocessor), some are IS_HIP_EXTENSION (build-time Python), some are is_hip_extension() (runtime). The build-fix commit normalizes guard comments to "Disabled on ROCm" but does not homogenize the guard mechanism — follow each file's existing convention per CLAUDE.md.
  • MXFP4 enum renumbering (8/9 → 10/11) is a wire-format concern: any external consumer serializing NVTEDType values will need to be aware. NVFP4 now occupies the old MXFP4 slots.
  • New CUDA-only subsystems (NCCL Expert Parallelism, upstream's grouped-tensor GEMM path, TMA transpose in flash_attn.cu) are stubbed/disabled — porting them to ROCm is deferred.
  • The upstream merge commit's own note ("full ROCm HIP build + GPU runtime tests not yet run") is now outdated; subsequent CI runs on this branch drove the ~25 fix commits listed above. Treat CI-green on this PR as the actual signal.

Generated by Claude. To request a code review, comment /claude review.

Comment thread transformer_engine/common/cast/nvfp4/rowwise_amax_nvfp4.cuh
Comment thread transformer_engine/common/fused_router/async_loader.h
Comment thread transformer_engine/jax/csrc/extensions/topk.cpp
Comment thread tests/cpp/operator/test_multi_swizzle.cu
Comment thread tests/pytorch/nvfp4/test_nvfp4_rht_quantize_swizzle_fusion.py
Comment thread transformer_engine/common/fused_attn/flash_attn.cu
Comment thread transformer_engine/pytorch/tensor/nvfp4_tensor.py
Comment thread transformer_engine/pytorch/tensor/storage/nvfp4_tensor_storage.py
Comment thread transformer_engine/pytorch/ops/basic/grouped_linear.py
Comment thread tests/pytorch/nvfp4/test_nvfp4_quantize_exact.py
Comment thread tests/pytorch/nvfp4/test_nvfp4_group_quantize.py
@github-actions

Copy link
Copy Markdown

Automated review summary

Scope reviewed: ROCm-side commits on top of the NVIDIA v2.17 upstream merge (commits fbcaca7a..6369177f, ~26 commits touching ~69 files). Upstream NVIDIA content brought in by the merge itself was not re-reviewed.

Code review: no additional correctness / safety findings beyond what @wenchenvincent has already flagged. The 21 existing inline questions on backend selection, ROCm-specific skips, fnuz handling, upstream API deltas, and dqkv layout still need author responses — this review does not duplicate them.

Copyright headers: 11 files need updates — see inline comments. All are ROCm-modified/added files whose headers still carry only the NVIDIA line. Per repo policy, an AMD line at the current year (2026) should sit above the NVIDIA line, and new files added by ROCm need both headers.

Comment thread build_tools/utils.py
Comment thread build_tools/VERSION.txt
Comment thread tests/cpp/operator/test_cast_nvfp4_transpose.cu
@@ -158,24 +190,26 @@ def _run_layer_with_overlap(
raise AssertionError(result.stderr.decode())


@pytest.mark.parametrize("use_cublasmp", (False, True))

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Should be only False for ROCm

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Gated use_cublasmp to (False,) if IS_HIP_EXTENSION else (False, True) on all six comm/gemm-overlap test parametrizations. On ROCm only the False cases are generated; CUDA still gets False+True.

def test_distributed_fuser_ops(world_size: int) -> None:
"""Launch parallel job that runs parallel tests"""
#TODO: find out why cannot align the following two lines with NV upstream
# TODO: find out why we cannot align the following two lines with NV upstream

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

TODO probably is not applicable - we know why the code is different. So it should rather be other comment, similar to test_cast_master_weights_to_fp8() which later solved the same problem


# Swizzled SF is emitted only by the Blackwell RHT cast-fusion kernel; no fusion path on ROCm.
if IS_HIP_EXTENSION and optimize_for_gemm:
pytest.skip("RHT cast-fusion swizzled SF is Blackwell-only")

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

The message should rather say it is not supported on ROCm

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Corrected. Applied the same correction to the sibling module-level skip in test_nvfp4_rht_quantize_swizzle_fusion.py for consistency.

Comment thread tests/pytorch/nvfp4/test_nvfp4_rht_quantize_swizzle_fusion.py
Comment thread tests/pytorch/references/ref_per_tensor_cs.py
Comment thread tests/pytorch/test_numerics.py
Comment thread tests/pytorch/test_recipe.py
Comment thread build_tools/jax.py
cxx_flags.append("-DNVTE_WITH_CUBLASMP")

# NCCL EP is a CUDA/NCCL-only feature; never enabled on ROCm.
# Disabled on ROCm

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

nit: usually for those nv side feature we didn't enabled during IFU, we put TODO so later grep can be easier. But if you keep a list of NV side feature then this is also fine

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

I enriched the jax.py comment to explain the gate, matching the fuller comment already in build_tools/pytorch.py, and kept # Disabled on ROCm as the grep marker.

Comment thread ci/pytorch.sh
NVTE_ROCM_ENABLE_MXFP8=1 run_default_fa 1 test_multi_tensor.py
run 1 test_numerics.py
NVTE_ROCM_ENABLE_MXFP8=1 run_default_fa_lbl "mxfp8" 1 test_numerics.py -k "recipe0 and 126m and not grouped"
NVTE_ROCM_ENABLE_MXFP8=1 run_default_fa_lbl "mxfp8" 1 test_numerics.py -k "MXFP8BlockScaling and 126m and not grouped"

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Why the model 126m cannot run mxfp8 recipe?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

126m does run mxfp8

The parametrize now uses ids=recipe_id which produces a readable name based id from the recipe class instead of recipe0, so the first recipe now shows up as MXFP8BlockScaling. I updated the filter from recipe0 to MXFP8BlockScaling so it keeps selecting the same MXFP8 cases.

}

void compareResults_mxfp4(const Tensor &test,
void compareResults_mxfp4(Tensor &test,

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Why previously we can use const Tensor& but now need Tensor&? do we edit some contents in test

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Upstream dropped the const qualifier from to_cpu() and the *_cpu_dptr() accessors in tests/cpp/test_common.h

compareResults_mxfp4 calls those on test, so once they became non-const the const Tensor& parameter stopped compiling and had to become Tensor&.

"requires using power of two scaling factors."
)

if recipe_name == "NVFP4BlockScaling" and async_save and IS_HIP_EXTENSION:

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Is this a new test from upstream? If so, can you file a ticket and put a TODO comment with this link here

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Working on it right now, will push an updated commit soon.

"""
_maybe_skip(recipe_name, quantized_model_init)

if (

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Need similar todo to github ticket for this fsdp2 issue

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Working on it right now, will push an updated commit soon.

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Why this file looks like a new file?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Row-scaled NVFP4 is new in upstream v2.17, so the IFU pulled it in and ROCm needed the row-wise amax to support it but upstream keeps that kernel inside the Blackwell-only quantize_transpose_nvfp4.cuh, which doesn't compile on ROCm, so instead of taking it as-is we extracted the portable amax into this new header.

if is_hip_extension():
# Fully-masked striped steps must contribute nothing; the CK kernel can
# emit non-finite output/LSE for them, so neutralize to avoid 0 * NaN.
softmax_aux_per_step = jnp.where(

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

I think this issue is worth a new git ticket for fully understanding.

If it's a CK issue in gfx942, we may need to report to CK team.

namespace jax {

// nvte_topk (common/util/topk.cu) is CUDA-only: standalone_topk.cuh relies on
// cub internals and cooperative_groups::reduce, which have no HIP equivalent, so

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Try hipcub

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Will implement as a feature after the IFU is merged.

src_device = tensor.device
if src_device.type != "cuda":
cuda_tensor = tensor.to(device=torch.device("cuda"))
cuda = torch.device("cuda")

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Is this a bug fix over NV side?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Yes, bug fix over NV. Upstream's tensor.to("cuda") on the NVFP4 subclass re-enters __torch_dispatch__ -> dequantize -> .to(), recursing infinitely when running the test_dcp_output_parity[NVFP4BlockScaling-sync] test.

We move the raw buffers via the metadata dict and rebuild NVFP4TensorStorage directly to avoid the recursion. Worth upstreaming.

tensor: QuantizedTensor,
) -> QuantizedTensor:
shape = tensor.shape
shape = tensor.size()

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Why the original .shape need to be changed to .size()?

@AllenFarcas AllenFarcas Jul 22, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

tensor is a quantized storage subclass (NVFP4/MXFP8) which overrides size() to infer the true logical shape from its raw packed buffers, while .shape isn't overridden, so it returns the raw packed storage dims and can be wrong or warn when only partial data is present.

_setup_conditional_transpose_storage needs the logical shape to size the transpose buffer, so we switched to the storage-safe .size() (and len(shape) instead of .ndim).

When running test_quantize in tests/pytorch/triton_kernels/test_cast.py I caught this.

AllenFarcas and others added 4 commits July 22, 2026 12:28
Use pathlib.Path(sys.executable) / pathlib.Path(__file__).resolve() like
upstream and test_cast_master_weights_to_fp8; drop the stale TODO. Verified
test_distributed_fuser_ops passes on ROCm (world_size 1 and 2).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
NVFP4 on gfx942 was simulation-only and required fnuz-specific workarounds
that were only needed because the fnuz e4m3 block scale (max 240) differs
from OCP (448). Restrict NVFP4 to gfx950 (gfx1250 remains behind the
existing TODO) and revert the fnuz workarounds:

- quantization.py: NVFP4 arch gate ((9,4),(9,5)) -> ((9,5),); message updated.
  FP8 / block-scaling gates keep (9,4) (unaffected).
- dequantize_nvfp4.cuh: drop the arch-aware factor_inv branch; use the
  upstream 1/(6*E4M3_MAX) form (no-op on gfx950/OCP).
- quantization_ref_nvfp4.py: simplify all is_fp8_fnuz() branches to 448 and
  drop the now-unused import (no-op on gfx950).

On gfx942 NVFP4 now reports unavailable, so all pytorch NVFP4 tests skip via
the centralized _compute_nvfp4_support gate; gfx950 behavior is unchanged.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Follow-up to bffa8f0, trimmed to the necessary changes:
- quantization_ref_nvfp4.py: keep only the be3feea revert (global scale);
  restore the incidental is_fp8_fnuz cleanup (no-op on gfx950, unreached on
  gfx942) to keep the diff minimal.
- test_dequantize_nvfp4.cu: skip on gfx942 (te_fp8_fnuz) — the dequant kernel
  now uses OCP 448 while its reference is fnuz-240, so it fails there.
- test_nvfp4_rht_quantize_exact.py: gate the ROCm-only test_hadamard_transform_amax
  on recipe_available (was IS_HIP_EXTENSION-only; uses the reverted reference).

Drops the unnecessary cast/grouped-gemm cpp skips (those tests pass on gfx942).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
68bc6e0 added the te_fp8_fnuz() skip to the CUDA-branch dequant bodies
(#else of __HIP_PLATFORM_AMD__), which don't run on ROCm — so gfx942 still
failed — and would break the CUDA build (te_fp8_fnuz is ROCm-only). Move the
skip to the ROCm-branch DequantizeNVFP4 test body (the one that actually runs
on gfx942) and drop it from the CUDA bodies.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
@AllenFarcas

AllenFarcas commented Jul 22, 2026

Copy link
Copy Markdown
Contributor Author

@ipanfilo @wangye805 related to your comment on tests/pytorch/distributed/test_fusible_ops.py line 1049 "TODO probably is not applicable - we know why the code is different. So it should rather be other comment, similar to test_cast_master_weights_to_fp8() which later solved the same problem"

I cannot comment there but ill add here my reply: I replaced the TODO with an explanatory comment in the # ROCm: style used in test_cast_master_weights_to_fp8.py: # ROCm: fnuz FP8 (gfx942) caps E4M3 at 240; OCP (gfx950+/CUDA) caps at 448., so the reference max must branch on is_fp8_fnuz(). Also removed an unused transformer_engine_torch import left behind.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.