Skip to content

Implement row-scaled NVFP4 fprop recipe - #2931

Merged
timmoon10 merged 47 commits into
NVIDIA:mainfrom
zianglih:fp4-per-token
May 8, 2026
Merged

Implement row-scaled NVFP4 fprop recipe#2931
timmoon10 merged 47 commits into
NVIDIA:mainfrom
zianglih:fp4-per-token

Conversation

@zianglih

@zianglihzianglih commented Apr 27, 2026

Copy link
Copy Markdown
Contributor

Description

@HumansAnd

Implement per-token row-scaled NVFP4 recipe with fprop only.
Currently, the row-scaled scaling is handled by separate pytorch code.
Quantization kernels are bitwise exact with existing TE reference implementation.

The following tests passed on B200:

python3 -m pytest --tb=auto tests/pytorch/nvfp4/test_nvfp4_quantize_exact.py
python3 -m pytest --tb=auto tests/pytorch/nvfp4/test_nvfp4_gemm_exact.py
python3 -m pytest --tb=auto tests/pytorch/test_backward_override.py
python3 -m pytest --tb=auto tests/pytorch/test_sanity.py
python3 -m pytest --tb=auto tests/pytorch/test_recipe.py
python3 -m pytest --tb=auto tests/pytorch/test_torch_compile.py
python3 -m pytest --tb=auto tests/pytorch/test_cpu_offloading.py
PYTORCH_JIT=0 NVTE_TORCH_COMPILE=0 NVTE_ALLOW_NONDETERMINISTIC_ALGO=0 NVTE_FUSED_ATTN=0 python3 -m pytest --tb=auto tests/pytorch/test_cuda_graphs.py
NVTE_TEST_NVINSPECT_ENABLED=1 NVTE_TEST_NVINSPECT_CONFIG_FILE=tests/pytorch/debug/test_configs/dummy_feature.yaml NVTE_TEST_NVINSPECT_FEATURE_DIRS=transformer_engine/debug/features PYTORCH_JIT=0 NVTE_TORCH_COMPILE=0 NVTE_ALLOW_NONDETERMINISTIC_ALGO=0 python3 -m pytest --tb=auto tests/pytorch/test_sanity.py

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:

  • Add a row_scaled_activation field in nvfp4 recipe, can be turned on by NVTE_NVFP4_ROW_SCALED_ACTIVATION
  • New per-token nvfp4 quantize kernels in transformer_engine/common/cast/nvfp4/quantize_pertoken_nvfp4.cuh, bitwise exact with existing TE pytorch reference implementation and per-tesor nvfp4 emulated implmentation. New quantization kernels folded into existing nvfp4 quantization kernels.
  • Expand dequant kernel transformer_engine/common/cast/nvfp4/dequantize_nvfp4.cuh to correctly handle this row-scaled nvfp4
  • In TransformerEngine/transformer_engine/pytorch/cpp_extensions/gemm.py, if row-scaled nvfp4 is enabled, it conducts separate per-token scaling using pytorch code, after cublas gemm
  • Broad test coverage by expanding 7 python and 2 cpp test files
  • Modify 1d quant reference implementation in tests/cpp/operator/test_cast_nvfp4_transpose.cu to align with pytorch reference numerics

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

@zianglih
zianglih marked this pull request as draft April 27, 2026 06:24
@greptile-apps

greptile-appsBot commented Apr 27, 2026

Copy link
Copy Markdown
Contributor

Greptile Summary

This PR implements a row-scaled (per-activation-row) NVFP4 recipe for fprop, controlled via the new NVTE_NVFP4_ROW_SCALED_ACTIVATION env-var. Instead of a single global FP32 amax per tensor, each activation row gets its own amax; the block-level FP8 scales are recomputed relative to that per-row value, and the matching global correction is applied in FP32 after the cuBLAS GEMM.

  • New compute_rowwise_amax CUDA kernel computes one FP32 max-abs per row before the main quantize kernel; quantize_transpose_nvfp4.cuh gains a ROW_SCALED_NVFP4 template branch that uses those per-row values when computing FP8 block scales.
  • general_gemm gains a post-GEMM scaling path that replaces both operands' global amaxes with 1.0 before cuBLAS, then multiplies the FP32 output by per_row_amax_B × scalar_amax_A; the grouped-GEMM variant loops over individual GEMMs to achieve the same effect.
  • All tensor storage classes (NVFP4TensorStorage, GroupedTensorStorage, NVFP4Tensor) and their C++ counterparts propagate the new row_scaled_nvfp4 flag through allocation, copy, and serialisation paths.

Confidence Score: 4/5

Safe to merge for fprop-only workloads on B200; remaining assert-based guards in the GEMM path can produce silent wrong results under optimised Python, but backward is already blocked by a RuntimeError.

The core quantisation math is verified bitwise-exact against the reference implementation and an extensive test suite passes on B200. Two contract checks in the row-scaled GEMM path use Python assert rather than RuntimeError; under -O/-OO the layout check is silently dropped and a wrong-transpose GEMM produces numerically incorrect output with no error.

transformer_engine/pytorch/cpp_extensions/gemm.py — the row-scaled GEMM helper _nvfp4_row_scaled_gemm_inputs and the layout guard inside general_gemm's row-scaled branch.

Important Files Changed

FilenameOverview
transformer_engine/pytorch/cpp_extensions/gemm.pyAdds row-scaled NVFP4 path to general_gemm and general_grouped_gemm; post-GEMM FP32 scale application logic is correct but several API-contract guards use assert (disabled by -O) instead of RuntimeError, including the critical layout check.
transformer_engine/common/recipe/init.pyAdds row_scaled_activation field and NVTE_NVFP4_ROW_SCALED_ACTIVATION env-var toggle to NVFP4BlockScaling recipe; change is straightforward and correct.
transformer_engine/pytorch/quantization.pyPropagates row_scaled_nvfp4 flag to forward quantizers via idx % 3 != 1 heuristic; backward quantizers correctly hardcode row_scaled_nvfp4=False.
transformer_engine/pytorch/tensor/nvfp4_tensor.pyNVFP4Quantizer gains row_scaled_nvfp4 attribute and allocates per-row amax buffer; is_quantizable override returns False with a misleading docstring that obscures the intentional distributed all-gather fallback.
transformer_engine/pytorch/csrc/quantizer.cppC++ NVFP4Quantizer correctly reads and propagates row_scaled_nvfp4; create_tensor, convert_and_update_tensor, and quantize_impl all validate constraints and allocate per-row amax buffers appropriately.
transformer_engine/common/cast/nvfp4/quantize_transpose_nvfp4.cuhAdds compute_rowwise_amax kernel and ROW_SCALED_NVFP4 template branch to the main quantize kernel; per-row encode/decode scale logic appears correct.
transformer_engine/common/cast/nvfp4/dequantize_nvfp4.cuhPasses row_scaled_nvfp4 flag through to kernel and selects tensor_amax[y] vs tensor_amax[0] for row indexing; boundary check on amax size is added correctly.
transformer_engine/pytorch/tensor/storage/grouped_tensor_storage.pyCorrectly computes total_amax_elements as sum of flat first dims per tensor for row-scaled case; amax offset slicing in split_into_quantized_tensors is consistent with the allocation.
transformer_engine/common/cast/dispatch/quantize.cuhCalls compute_rowwise_amax before the main quantize kernel for row-scaled path; backward quantize path hard-codes row_scaled_nvfp4=false with an additional NVTE_CHECK guard.
transformer_engine/pytorch/custom_recipes/quantization_nvfp4.pyReference quantizer gains row_scaled_nvfp4 support with per-row amax computation and zero-guard handling; gemm_ref correctly broadcasts partial_alpha when per-row amaxes are used.

Sequence Diagram

sequenceDiagram
participant FWD as Forward Pass
participant Q as NVFP4Quantizer (row_scaled)
participant K as compute_rowwise_amax kernel
participant QK as quantize_transpose kernel
participant G as general_gemm
participant C as cuBLAS GEMM
participant S as FP32 post-scale
FWD->>Q: quantize(activation)
Q->>K: compute per-row amax
K-->>Q: amax_rowwise[M]
Q->>QK: quantize with per-row encode scales
QK-->>Q: FP4 data + FP8 block scales
FWD->>G: general_gemm(weight_A, activation_B)
G->>G: "set amax_A=1 amax_B=1, save rowwise_scales"
G->>C: "GEMM with global scale=1"
C-->>G: raw_out (FP32)
G->>S: "raw_out *= rowwise_scales"
S-->>FWD: correctly scaled output
Loading

Reviews (12): Last reviewed commit: "Update tests/pytorch/utils.py" | Re-trigger Greptile

Comment threadtransformer_engine/pytorch/cpp_extensions/gemm.py Outdated
// Compute "correct" per-block encoding scaling factor
const float S_enc_b_fp8 = S_dec_b_fp32 == 0.f ? 0.f : S_enc / S_dec_b_fp32;
const float S_enc_b_fp8 = S_dec_b_fp32 == 0.f ? 0.f :
fminf(1.0f / (S_dec_b_fp32 * (1.0f / S_enc)), Numeric_Traits<float>::maxNorm);

Copy link
Copy Markdown
ContributorAuthor

Choose a reason for hiding this comment

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

We have to change here to stay aligned with pytorch reference.

@zianglih
zianglih marked this pull request as ready for review April 27, 2026 09:14
@zianglih
zianglih marked this pull request as draft May 2, 2026 18:22
zianglihand others added 14 commits May 2, 2026 11:27
Signed-off-by: Ziang Li <ziangli@umich.edu>
Co-authored-by: Yigong Qin <qqqyyy1233@outlook.com>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
@ziang-and
ziang-andforce-pushed the fp4-per-token branch 2 times, most recently from 6998f64 to 5b2f606CompareMay 2, 2026 19:10
zianglih added 5 commits May 2, 2026 16:33
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Ziang Li <ziangli@umich.edu>
@zianglih

Copy link
Copy Markdown
ContributorAuthor

The following extended tests all passed:

python3 -m pytest --tb=auto tests/pytorch/nvfp4/test_nvfp4_quantize_exact.py
python3 -m pytest --tb=auto tests/pytorch/nvfp4/test_nvfp4_gemm_exact.py
python3 -m pytest --tb=auto tests/pytorch/test_backward_override.py
python3 -m pytest --tb=auto tests/pytorch/test_sanity.py
python3 -m pytest --tb=auto tests/pytorch/test_recipe.py
python3 -m pytest --tb=auto tests/pytorch/test_torch_compile.py
python3 -m pytest --tb=auto tests/pytorch/test_cpu_offloading.py
PYTORCH_JIT=0 NVTE_TORCH_COMPILE=0 NVTE_ALLOW_NONDETERMINISTIC_ALGO=0 NVTE_FUSED_ATTN=0 python3 -m pytest --tb=auto tests/pytorch/test_cuda_graphs.py
NVTE_TEST_NVINSPECT_ENABLED=1 NVTE_TEST_NVINSPECT_CONFIG_FILE=tests/pytorch/debug/test_configs/dummy_feature.yaml NVTE_TEST_NVINSPECT_FEATURE_DIRS=transformer_engine/debug/features PYTORCH_JIT=0 NVTE_TORCH_COMPILE=0 NVTE_ALLOW_NONDETERMINISTIC_ALGO=0 python3 -m pytest --tb=auto tests/pytorch/test_sanity.py
cd /root/TransformerEngine/tests/cpp
cmake --build build -j200
TEST_BIN="$(find build -type f -name test_operator -perm -u+x | head -n 1)"
"$TEST_BIN" --gtest_filter='*FusedCastTransposeNVFP4*:*DequantizeNVFP4*'
EOF

Comment threadtransformer_engine/pytorch/cpp_extensions/gemm.py
@timmoon10

Copy link
Copy Markdown
Member

We see a test failure when running on A100: https://github.com/NVIDIA/TransformerEngine/actions/runs/25475147141/job/74746937329
The distributed test failures also show up in the main branch, so they are not blocking.

@zianglih

Copy link
Copy Markdown
ContributorAuthor

Previous L0_pytorch_unittest--A100_1GPU failed due to a minor test file issue. L1_pytorch_distributed_unittest--H100_4GPU failed due to numerics I did not touch. L1_pytorch_distributed_unittest--B200_8GPU reached time limit.

Signed-off-by: Ziang Li <ziangli@umich.edu>
@greptile-apps

This comment was marked as off-topic.

Comment threadtests/pytorch/utils.py Outdated
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: Ziang Li <ziangli@umich.edu>
@timmoon10

Copy link
Copy Markdown
Member

/te-ci L1

else:
if _is_nvfp4_row_scaled_tensor(A):
raise NotImplementedError("Row-scaled NVFP4 GEMM does not support row-scaled A.")
assert layout[1] == "N", "Row-scaled NVFP4 GEMM currently supports N-layout B only."

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.

P1Layout constraint silenced by -O produces wrong numerical results

assert layout[1] == "N" is the single most dangerous assertion in this path: if skipped under -O/-OO, the code continues with layout[1] == "T", meaning B is transposed. In that case B._amax_rowwise still contains per-row amaxes from the pre-transposition orientation, so rowwise_global_scales broadcasts the wrong scale over each output row — the GEMM produces silently incorrect values rather than crashing. Unlike the grad check that was already hardened to raise RuntimeError, this constraint should receive the same treatment since a wrong-layout GEMM is numerically undetectable without extra validation.

Comment on lines +86 to +103
A_metadata = A.get_metadata()
weight_amax = A._amax_rowwise if transa else A._amax_columnwise
assert weight_amax is not None and weight_amax.numel() == 1
A_metadata["amax_rowwise" if transa else "amax_columnwise"] = weight_amax.new_ones(1)
A_metadata["row_scaled_nvfp4"] = False

B_metadata = B.get_metadata()
rhs_rowwise_amax = B._amax_rowwise
assert rhs_rowwise_amax is not None
B_metadata["amax_rowwise"] = rhs_rowwise_amax.new_ones(1)
B_metadata["row_scaled_nvfp4"] = False

assert rhs_rowwise_amax.dtype == torch.float32 and weight_amax.dtype == torch.float32
return (
NVFP4TensorStorage(**A_metadata),
NVFP4TensorStorage(**B_metadata),
(rhs_rowwise_amax * weight_amax).view(-1, 1),
)

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.

P1Contract assertions in helper are disabled by -O, turning clean errors into cryptic AttributeErrors

_nvfp4_row_scaled_gemm_inputs contains three assert statements that enforce API contracts:

  • assert weight_amax is not None and weight_amax.numel() == 1 — if weight_amax is None (e.g., a columnwise-only weight tensor with transa=False) and the assert is stripped, weight_amax.new_ones(1) raises AttributeError: 'NoneType' object has no attribute 'new_ones' deep inside the function.
  • assert rhs_rowwise_amax is not None — same failure mode for B.
  • assert rhs_rowwise_amax.dtype == torch.float32 and weight_amax.dtype == torch.float32 — wrong dtype silently produces incorrect scaling arithmetic.

These three checks guard the function's entire scaling logic and should be RuntimeError raises just like the grad check already was.

@timmoon10timmoon10 left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

LGTM. The PyTorch test failures are unrelated. The JAX failure seems related to MoE kernels, so that's also unrelated.

@timmoon10
timmoon10 merged commit c74e5aa into NVIDIA:mainMay 8, 2026
41 of 53 checks passed
timmoon10 added a commit to timmoon10/TransformerEngine that referenced this pull request May 9, 2026
Also do some cleanup and improve documentation.
Signed-off-by: Tim Moon <tmoon@nvidia.com>
@timmoon10timmoon10 mentioned this pull request May 9, 2026
13 tasks
@zianglihzianglih mentioned this pull request May 9, 2026
13 tasks
Oleg-Goncharov pushed a commit that referenced this pull request May 11, 2026
* Fix bug in NVFP4 quantize test where we set scale instead of amax
Refactor test tensor wrapper by removing recipe-specific logic whenever possible.
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Only get fp32 scale when tensor is expected to have fp32 scale
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Create dedicated class for managing GPU/CPU buffers
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Fix bugs in C++ test tensor infrastructure
- Fix syntax error in switch case (:: -> :)
- Fix double-underscore typo in variable name
- Fix wrong buffer passed to set_amax_columnwise
- Fix unique_ptr assignment from raw pointer (use reset())
- Remove dead duplicate NVTE_MXFP8_1D_SCALING branch in get_scales()
- Rename cpu_data -> cpu_buffer to match Buffer class API
- Remove const from Tensor::to_cpu/from_cpu and their callers,
since both methods write to the CPU buffer
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Debug compilation errors
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Remove type check when accessing raw pointers
CPU and GPU types are inconsistent, so the type checks cause too many problems.
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Debug distributed C++ tests
Also adopt review suggestions from @greptile-apps.
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Remove unused header
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
* Copy-paste error
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
* Use shared buffer for FP8 row-wise scale-inv and col-wise scale-inv
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Typo
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
* Debug merge conflicts with #2931
Also do some cleanup and improve documentation.
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Address code review feedback
- Restore amax buffer size assertion in compare_rowwise_amax
- Remove set_tensor_amax alias in favor of set_amax
- Extract fill_uniform_buffer helper to anonymous namespace,
eliminating duplication in fill_uniform_{rowwise,columnwise}_scale_inv
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>
---------
Signed-off-by: Tim Moon <tmoon@nvidia.com>
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
faradawn pushed a commit to faradawn/TransformerEngine that referenced this pull request May 14, 2026
* Adapt initial implementation and make quantization bitwise exact
Signed-off-by: Ziang Li <ziangli@umich.edu>
Co-authored-by: Yigong Qin <qqqyyy1233@outlook.com>
* Add col
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Add fp32
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up tests
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up ref
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up gemm wrapper
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up test
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Rename and reformat
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Avoid partial amax folding in gemm
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Expand test coverage
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Expand more tests
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Turn on test for grouped linear sanity
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Rename pertoken to per_token
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Expand .cu test
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Format after rebase
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Fix test after rebase
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up cpp test
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Extend cpp dequantize test
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Only pass `per_token_activation` to forward activation quantizer and clean up
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Minor fix test
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Improve accuracy by unfolding weight per-tensor fp32
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Fold row-wise quantization
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Drop column wise
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up column wise
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Move shared test helpers
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Minor clean up test
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Readability
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Rename
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Further refactor
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up bias
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up cast
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Avoid silently disable column wise
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up
Signed-off-by: Ziang Li <ziangli@umich.edu>
* `is_quantizable` returns false
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Error out grouped gemm
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Tighten test
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Rename verbose rowwise_amax_is_row_scaled
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Clean up
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Explicitly handle both gemm input and error out
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Minor
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Nits and lint
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Minor fix A100 ci
Signed-off-by: Ziang Li <ziangli@umich.edu>
* Update tests/pytorch/utils.py
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: Ziang Li <ziangli@umich.edu>
---------
Signed-off-by: Ziang Li <ziangli@umich.edu>
Co-authored-by: Yigong Qin <qqqyyy1233@outlook.com>
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
faradawn pushed a commit to faradawn/TransformerEngine that referenced this pull request May 14, 2026
* Fix bug in NVFP4 quantize test where we set scale instead of amax
Refactor test tensor wrapper by removing recipe-specific logic whenever possible.
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Only get fp32 scale when tensor is expected to have fp32 scale
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Create dedicated class for managing GPU/CPU buffers
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Fix bugs in C++ test tensor infrastructure
- Fix syntax error in switch case (:: -> :)
- Fix double-underscore typo in variable name
- Fix wrong buffer passed to set_amax_columnwise
- Fix unique_ptr assignment from raw pointer (use reset())
- Remove dead duplicate NVTE_MXFP8_1D_SCALING branch in get_scales()
- Rename cpu_data -> cpu_buffer to match Buffer class API
- Remove const from Tensor::to_cpu/from_cpu and their callers,
since both methods write to the CPU buffer
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Debug compilation errors
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Remove type check when accessing raw pointers
CPU and GPU types are inconsistent, so the type checks cause too many problems.
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Debug distributed C++ tests
Also adopt review suggestions from @greptile-apps.
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Remove unused header
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
* Copy-paste error
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
* Use shared buffer for FP8 row-wise scale-inv and col-wise scale-inv
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Typo
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
* Debug merge conflicts with NVIDIA#2931
Also do some cleanup and improve documentation.
Signed-off-by: Tim Moon <tmoon@nvidia.com>
* Address code review feedback
- Restore amax buffer size assertion in compare_rowwise_amax
- Remove set_tensor_amax alias in favor of set_amax
- Extract fill_uniform_buffer helper to anonymous namespace,
eliminating duplication in fill_uniform_{rowwise,columnwise}_scale_inv
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>
---------
Signed-off-by: Tim Moon <tmoon@nvidia.com>
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
aleozlx pushed a commit to flashinfer-ai/flashinfer that referenced this pull request May 21, 2026
<!-- .github/pull_request_template.md -->
## 📌 Description
@HumansAnd
Implement 4over6 nvfp4 from:
- Paper: https://arxiv.org/abs/2512.02010
- Code: https://github.com/mit-han-lab/fouroversix
TE PR: - NVIDIA/TransformerEngine#2972
For original nvfp4, only `cutlass_fused_moe` is supported.
For per-token nvfp4, only `trtllm_fp4_block_scale_routed_moe` and
`trtllm_fp4_block_scale_moe` are supported.
The results is bitwise exact with reference implementation by enabling:
- `TRTLLM_DISABLE_FP4_QUANT_FAST_MATH=1`
Under strict no fast math mode, the quantizer is bitwise exact with
pytorch reference implementation.
By default, 448 scaling is used for E4M3. To enable the 256/384 E4M3
scaling in the original paper, enable
`FLASHINFER_NVFP4_4OVER6_E4M3_USE_256`.
Error mode can be either `MAE` or `MSE`, default to `MAE`.
Need to rebase after:
- #3237
- #3027
Future work:
- TE recipe implementation after
NVIDIA/TransformerEngine#2931 is merged
- Performance optimization
<!-- What does this PR do? Briefly describe the changes and why they’re
needed. -->
## 🔍 Related Issues
<!-- Link any related issues here -->
## 🚀 Pull Request Checklist
Thank you for contributing to FlashInfer! Before we review your pull
request, please make sure the following items are complete.
### ✅ Pre-commit Checks
- [x] I have installed `pre-commit` by running `pip install pre-commit`
(or used your preferred method).
- [x] I have installed the hooks with `pre-commit install`.
- [x] I have run the hooks manually with `pre-commit run --all-files`
and fixed any reported issues.
> If you are unsure about how to set up `pre-commit`, see [the
pre-commit documentation](https://pre-commit.com/).
## 🧪 Tests
- [x] Tests have been added or updated as needed.
- [x] All tests are passing (`unittest`, etc.).
## Reviewer Notes
<!-- Optional: anything you'd like reviewers to focus on, concerns, etc.
-->
<!-- This is an auto-generated comment: release notes by coderabbit.ai
-->
## Summary by CodeRabbit
## Release Notes
* **New Features**
* Introduced NVFP4 "4-over-6" quantization mode for improved FP4
precision, configurable via environment variables
* Added MSE-based scale candidate selection to enhance quantization
accuracy
* Implemented runtime toggles for FP4 fast-math and optimization control
* **Improvements**
* Enhanced FP4 quantization kernel dispatch for flexible runtime
configuration
[![Review Change
Stack](https://storage.googleapis.com/coderabbit_public_assets/review-stack-in-coderabbit-ui.svg)](https://app.coderabbit.ai/change-stack/flashinfer-ai/flashinfer/pull/3264)
<!-- end of auto-generated comment: release notes by coderabbit.ai -->
---------
Co-authored-by: Lain <siyuanf@nvidia.com>
@zianglihzianglih mentioned this pull request Jun 23, 2026
13 tasks
cael-ling added a commit to cael-ling/TransformerEngine that referenced this pull request Jul 13, 2026
…quantize/GEMM APIs
Introduce the NVFP4 row-scaled transpose quantization and CUTLASS GEMM kernels
at the common (C/CUDA) layer, reached only through TransformerEngine's existing
quantize/GEMM entry points (no new top-level C API), following how NVIDIA#2931
integrates row-scaled NVFP4:
- Row-scaled transpose cast kernels (single + grouped): per-row (rowwise) and
per-col (columnwise/transpose) NVFP4 scales in one pass. Selected inside
nvte_quantize_v2 when a row-scaled NVFP4 output also allocates a columnwise
buffer (output->row_scaled_nvfp4 && has_columnwise_data()); no dedicated
config attribute.
- CUTLASS NVFP4 row-scaled GEMM with per-row*per-col rescale fused into the EVT
epilogue: routed from cublas_gemm() (nvte_cublas_gemm_v2) when both operands
carry the row_scaled_nvfp4 flag, replacing the previous hard reject. The
standalone post-scale and grouped GEMM kernels are internalized (no C API),
to be wired into their dispatch paths by follow-up PRs.
Naming follows the merged row-scaled NVFP4 convention (cutlass_nvfp4_*,
row_scaled_transpose_*). C-only; the PyTorch recipe/quantizer/module wiring, and
the backward/SR/RHT/2D-weight paths, land in follow-up PRs. Covered by C++
gtests for the row-scaled transpose cast and the row-scaled GEMM (GEMM suite
GTEST_SKIPs on pre-Blackwell).
Signed-off-by: Cael Ling <caell@nvidia.com>
cael-ling added a commit to cael-ling/TransformerEngine that referenced this pull request Jul 22, 2026
… backward
The NVIDIA#2931 row-scaled NVFP4 path only produced the rowwise forward activation; its columnwise/transpose output was rejected. That
made the per-token activation unusable in the backward wgrad GEMM, so row-scaled training had to fall back to a dequantized/high-precision backward.
This change extends the existing row-scaled path to also emit the columnwise (transpose) direction, so a training Linear with
row_scaled_activation=True now quantizes the forward activation row-scaled in both directions and the wgrad GEMM consumes the row-scaled transpose directly. It is a minimal extension of NVIDIA#2931 (no new CUTLASS kernels, no grouped path, no RHT/4over6 transpose fusion).
Signed-off-by: Cael Ling <caell@nvidia.com>
cael-ling added a commit to cael-ling/TransformerEngine that referenced this pull request Jul 22, 2026
… backward
The NVIDIA#2931 row-scaled NVFP4 path only produced the rowwise forward activation; its columnwise/transpose output was rejected. That
made the per-token activation unusable in the backward wgrad GEMM, so row-scaled training had to fall back to a dequantized/high-precision backward.
This change extends the existing row-scaled path to also emit the columnwise (transpose) direction, so a training Linear with
row_scaled_activation=True now quantizes the forward activation row-scaled in both directions and the wgrad GEMM consumes the row-scaled transpose directly. It is a minimal extension of NVIDIA#2931 (no new CUTLASS kernels, no grouped path, no RHT/4over6 transpose fusion).
Signed-off-by: Cael Ling <caell@nvidia.com>
timmoon10 pushed a commit that referenced this pull request Jul 25, 2026
… backward (#3206)
* [common][PyTorch] NVFP4: enable row-scaled transpose quantization for backward
The #2931 row-scaled NVFP4 path only produced the rowwise forward activation; its columnwise/transpose output was rejected. That
made the per-token activation unusable in the backward wgrad GEMM, so row-scaled training had to fall back to a dequantized/high-precision backward.
This change extends the existing row-scaled path to also emit the columnwise (transpose) direction, so a training Linear with
row_scaled_activation=True now quantizes the forward activation row-scaled in both directions and the wgrad GEMM consumes the row-scaled transpose directly. It is a minimal extension of #2931 (no new CUTLASS kernels, no grouped path, no RHT/4over6 transpose fusion).
Signed-off-by: Cael Ling <caell@nvidia.com>
* [pre-commit.ci] auto fixes from pre-commit.com hooks
for more information, see https://pre-commit.ci
---------
Signed-off-by: Cael Ling <caell@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Baibaifan pushed a commit to Baibaifan/TransformerEngine that referenced this pull request Jul 25, 2026
… backward (NVIDIA#3206)
* [common][PyTorch] NVFP4: enable row-scaled transpose quantization for backward
The NVIDIA#2931 row-scaled NVFP4 path only produced the rowwise forward activation; its columnwise/transpose output was rejected. That
made the per-token activation unusable in the backward wgrad GEMM, so row-scaled training had to fall back to a dequantized/high-precision backward.
This change extends the existing row-scaled path to also emit the columnwise (transpose) direction, so a training Linear with
row_scaled_activation=True now quantizes the forward activation row-scaled in both directions and the wgrad GEMM consumes the row-scaled transpose directly. It is a minimal extension of NVIDIA#2931 (no new CUTLASS kernels, no grouped path, no RHT/4over6 transpose fusion).
Signed-off-by: Cael Ling <caell@nvidia.com>
* [pre-commit.ci] auto fixes from pre-commit.com hooks
for more information, see https://pre-commit.ci
---------
Signed-off-by: Cael Ling <caell@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Signed-off-by: yangfan.bai <yangfan.bai@shopee.com>
Sign up for freeto join this conversation on GitHub. Already have an account? Sign in to comment

Labels

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

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants

@zianglih@ptrendx@timmoon10