Skip to content

Add a PyTorch EP=1 backend for MoE dispatch and combine - #7

Draft
0z5a wants to merge 119 commits into
vthumbe1503:nccl_ep_opsfrom
0z5a:ep1-pytorch-dispatch-combine
Draft

0z5a wants to merge 119 commits into
vthumbe1503:nccl_ep_opsfrom
0z5a:ep1-pytorch-dispatch-combine

Conversation

@0z5a

@0z5a 0z5a commented Sep 28, 2026 •

Copy link
Copy Markdown

Summary

Add a bufferless PyTorch backend for MoeDispatch and MoeCombine when the expert-parallel group has one rank. Dispatch orders local routes by expert and combine sums expert outputs back into source tokens. Routing metadata belongs to each invocation, so two outstanding forwards do not overwrite each other. The local backend rejects remote expert IDs and communication settings it cannot honor.

This is the EP=1 follow-up requested on NVIDIA/TransformerEngine#3503. This draft is stacked on vthumbe1503:nccl_ep_ops at 261a81c; its signed-off implementation commit and pre-commit.ci formatting commit change only the EP=1 path, tests, and API documentation. The main-targeting upstream draft is open.

Validation

Check Result
CUDA_VISIBLE_DEVICES=3 PYTHONPATH=$PWD python -m pytest tests/pytorch/test_moe_ep_local.py -q 13 passed in 11.57 s on RTX 5090, PyTorch 2.13.0+cu130, at head 6ce761b
Five-op MoeDispatch → GroupedLinear → ScaledSwiGLU → GroupedLinear → MoeCombine BF16 forward and backward match an independent dense PyTorch MoE within the test's rtol=atol=2e-2, including token, gate-weight, and expert-weight gradients
Twelve-step training E2E with changing inputs and top-k=2 routes Full five-op forward, backward, SGD updates, and post-update inference matched the dense PyTorch reference; 8 parameter tensors changed, maximum output absolute difference 0.00178438 and maximum gradient absolute difference 0.00277808
Routing cases top-k=1 and 2, repeated and dropped routes, empty experts, two outstanding forwards, and unsupported settings covered
Repository license checker Passed

The test imported the PR source tree, ran without an initialized torch.distributed process group, and the loaded native extension exposed no EP entry points. The NCCL EP distributed suite was not run on this machine. The full L0 launcher was not run because it installs dependencies; the focused tests used the existing environment without package changes.

Route tokens locally in expert-major order, combine their outputs without an EpBuffer or NCCL EP, and validate unsupported communication settings. Cover the five-op MoE forward and backward path against a dense reference.

Signed-off-by: 0z5a <192209249+0z5a@users.noreply.github.com>
vthumbe1503 pushed a commit that referenced this pull request Sep 29, 2026
…IA#3137)

* start draft

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

* remove benchmark scripts

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>

* add license

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

* fix

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

* make cutlass dsl required

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

* fix linting errors

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>

* warn failed compilation

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>

* make tvm-ffi common dependency

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

* use a higher version for cutlass-dsl

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

* less comment

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

* skip zeroing buffer if noop flag set

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

* refactor activations

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

* fix

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

* lint

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

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

* nit

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

* maybe it's better to prepare tvm-ffi & cutedsl kernels in common/__init__.py

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>

* nit

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>

* 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

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

* fix

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

* fix

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

* fix

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

* also flush colwise scale in SMEM

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

* use explicit fp8 types

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>

* nit

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

* add checks

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

* refactor

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>

* nit

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

* fi

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>

* nit

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

* nit

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

* skip comparision test if cutedsl disabled

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

* fix

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

* Merge pull request #3 from janekb04/cutedsl_mxfp8_common

Perform noop tensor check on device

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

* add test to the qa script

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

* warn only once if dlopen fails

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

* fix

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

* let tests cover swizzling

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

* warn not chosen cutedsl kernel during testing

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

* improve lazyload so now the cache is uint32

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

* let tests be verbose about warnings so we know if cutedsl kernel is not dispatched

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

* support swizzled scale for specialized kernel but disable it

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>

* fix inline ptx

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

* more fix

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

* revert to mul

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

* nit

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

* add jax tests

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

* also ignore noop when dbias

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

* relax dim check for the kernels

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

* reject 2D quant

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

* run more tests with CuTeDSL path

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

* make rowwise only handle divisible cases

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

* fallback to CUDA if not TMA aligned

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

* specialize skip masking version

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

* nit

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

* add type annotations

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

* optimize masking

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

* skip masking for some acts

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

* allow user to opt out from building with CuTeDSL

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

* better warning

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

* doc

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

* add jax test coverage for cutedsl

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

* fix device sm query

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

* catch more python error

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

* nit

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>

* fix

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

* make tvm-ffi required for python users

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

* optimize dbias

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

* skip scale bound check whenever possible

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

* allow colwise to fuse relu too

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

* dispatch to rowwise specialized with swizzled scale

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

* don't link tvmffi if building a standalong C++ lib without framework

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

* change some logging

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

* use the right fast math flag

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

* also let test script use NVTE_DEBUG

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

* also let mutex cover the loading of tvm-ffi func

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

* locate tvm-ffi include dir from setup_common_extension

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

* fix how tests run with cutedsl enabled

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

* nit

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

* reduce framework scope tests run with cutedsl backend

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

* allow C++ tests to run with cutdsl

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

* remove tvmffi and cutlass dependencies from setup.py due to being optional

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

* encode sm numbers in tvm_ffi function name & cache key

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

* make C++ tests run in sequential

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

* make CuTeDSL default and cover C++ in L1

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

* Improvements and fix to CuTe DSL kernel launches (#7)

* [Common] Bypass Python wrapper overhead with __tvm_ffi_object__

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* [Common] Avoid heap allocations in DLTensorWrapper

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* [Common] Avoid querying device in DLTensorWrapper

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* [Common] Fix packed dtype handling in DLTensorWrapper

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Apply review suggestions

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

---------

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

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

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

* fix dispatching so it never uses specialized if grid won't fit

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

* also add tests

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

* Minor fixes to CuTe DSL infrastructure (#8)

* Improvements and fix to CuTe DSL kernel launches (#7)

* [Common] Bypass Python wrapper overhead with __tvm_ffi_object__

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* [Common] Avoid heap allocations in DLTensorWrapper

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* [Common] Avoid querying device in DLTensorWrapper

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* [Common] Fix packed dtype handling in DLTensorWrapper

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Apply review suggestions

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

---------

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

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

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

* [Common] Add missing dtypes to CuTe DSL dispatch

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* [Common] Add line information to CuTe DSL inline PTX helpers

Signed-off-by: Jan Bielak <jbielak@nvidia.com>

---------

Signed-off-by: Jan Bielak <jbielak@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

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

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

* add missing ip & loc

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

* refactor so we don't need #pragma to avoid collision in tests

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

* whoops

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

* whoops again

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

* rename NVTE_ENABLE_CUTEDSL_QUANT_BACKEND to NVTE_ENABLE_CUTEDSL_BACKEND

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

* switch back to NVTE_WARN_IF_CUTEDSL_BACKEND_NOT_CHOSEN

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

* use TE's to_string to convert dtype to str

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

* remove MXFP8QuantFused

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

* move python init from tests to TE lib

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

* software path for e8m0 conversion

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

* nit

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

* remove unreachable branch

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

* fold quantize_mxfp8_cutedsl_config.h back to quantize_mxfp8_cutedsl.h

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

* nit

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

* nit for some C++

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

* rewrite dsl_user_op to cute.jit if possible and let compiler optimize

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

* fix python init

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

* prefere inline_ptx instead of llvm.inline_asm

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

* nit

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

* fix activation encoding

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

* remove unused device_is_blackwell function

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

* link C++ test executable to libtvm_ffi.so

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

* move cutedsl utils from math.h to a separate file

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

* recover test_cast_mxfp8.cu

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

* no need to init extension now

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

* hoist importorskip in cutedsl python tests

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

* fix activation_func_to_enum so it works with CUDA 12.8

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

* make NVTE_WITH_CUTEDSL both a env var and a cmake option

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

* fix TE path searching for python bootstrap

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

* when enabling cutedsl manually, also initialize python if it's not done before

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

* use dlsym

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>

* let main TE lib link to Python::Module

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

* add a test to cover untested manually enable cutedsl path

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

* nit

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

* update min ver workflow

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

* rename L1 CUDA C++ UT

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

* fix build issues for CuTeDSL backend (NVIDIA#10)

* disable CuTeDSL support and other fix

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

* remove docs that mention CuTeDSL since they are off by default now

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

* fix

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

* remove test_cutedsl_backend.cpp

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

* make envvar.rst more human

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

---------

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

* fix

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

* fix python tests

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

* skip L1 C++ UT with CUDA when CuTeDSL is not enabled globally

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

* fix

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

* fix C++ UT python discovery

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

* fix deadlock for multigpu

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

* fix python init

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

* fix python init again thanks to codex

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

* refactor the serialized initial launch

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

* nit

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

---------

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>
Signed-off-by: Jan Bielak <jbielak@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: Jan Bielak <janekb04@icloud.com>
Co-authored-by: Jan Bielak <jbielak@nvidia.com>
cael-ling and others added 25 commits September 30, 2026 11:54
…VIDIA#3467)

* [Common] Add single-launch grouped fused amax for row-scaled NVFP4

Compute per-expert rowwise/columnwise amax over a packed (sum_M, K) grouped input in one kernel launch instead of one launch per expert, exposed via nvte_group_nvfp4_compute_amax and gated by NVTE_NVFP4_FUSED_AMAX. Adds a gtest that checks the result against a CPU reference and across launches.

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

* [Common] Enforce K alignment and per-expert amax buffers in grouped NVFP4 amax

The public entry point calls the launcher directly, bypassing the eligibility
check, so misaligned K silently dropped trailing columns and heterogeneous
per-expert amax buffers could dereference null. Add a column/split alignment
check in the launcher, take do_row/do_col as the union over experts with a
per-expert null guard in the kernel, and cover it with HeterogeneousColumnwise.

Signed-off-by: Cael Ling <caell@nvidia.com>

* [Common] nvfp4: drop unused grouped fused-amax gate and env switch

group_fused_amax_supported was never called (the C API calls the launcher directly and alignment checks already live there), and it carried the NVTE_NVFP4_FUSED_AMAX switch. Remove the dead function and its now-unused <cstdlib>.

Signed-off-by: Cael Ling <caell@nvidia.com>

---------

Signed-off-by: Cael Ling <caell@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
NVIDIA#3553)

* Isolate FlashAttention 2/3 imports so a broken install cannot break TE

The version gate only inspects distribution metadata, so a flash-attn whose
extension fails to load -- missing CUDA dependency, ABI mismatch -- raised
straight out of backends.py and made `import transformer_engine.pytorch` fail
outright, even for callers that never ask for FlashAttention.

FlashAttention 4 already guards its import this way, so give FA2 and FA3 the same
treatment: on ImportError leave the module aliases at None, mark the backend
unavailable, and warn with the version and the original error. A genuine TE
failure still raises; only the optional dependency is contained.

The regression test injects a metadata-only flash-attn distribution in a
subprocess, because the fault has to be in place before transformer_engine is
imported. It fails for both FA2 and FA3 without this change and passes with it.

Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

* Fix optional FlashAttention import regression coverage

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

---------

Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Co-authored-by: Przemek Tredak <ptredak@nvidia.com>
…A#3565)

* [common] Reject mismatched bias/out dtype for non-FP8/FP4 GEMM

cuBLASLt assumes bias dtype equals D when BIAS_DATA_TYPE is unset. TE only
sets that attribute on the FP8/FP4 path, so BF16 bias with FP32 out was
reinterpreted and read past the bias allocation. Match cuBLAS support tables
and Megatron-LM#6000: require bias dtype == D on the non-FP8/FP4 path.

Fixes NVIDIA#3562

Signed-off-by: Dundy Pasupuleti <dundysm@gmail.com>

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

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

* [qa] Wire test_gemm_bias_dtype into L0 pytorch unittest

L0 enumerates explicit pytest paths, so the new bias dtype regression
module was never invoked by community CI. Add it next to the other
top-level pytorch tests.

Signed-off-by: Dundy Pasupuleti <dundysm@gmail.com>

* [qa] Move non-fp8 bias dtype check into test_sanity

Signed-off-by: Dundy Pasupuleti <dundysm@gmail.com>

---------

Signed-off-by: Dundy Pasupuleti <dundysm@gmail.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: Santosh Bhavani <santosh.bhavani@live.com>
)

* [PyTorch] MoE Sequential block with Dispatch and Combine as basic ops

Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>

* minor bug fix in test

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

---------

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
…iled kernel (NVIDIA#3454)

* [Common] nvfp4: fuse row/col amax into a single TMA-tiled kernel

The row-scaled path ran two amax kernels; the columnwise one read global
memory column-major (uncoalesced) and dominated runtime. Compute both
directions in one kernel that streams 128x128 chunks through shared memory
via TMA and reduces columns from SMEM.

Gated by fused_amax_supported (BF16, 128-aligned dims); other cases keep the
two-kernel path. Set NVTE_NVFP4_FUSED_AMAX=0 to force fallback.

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

* [Common] Preserve NVFP4 fused amax buffers on no-op quantize

The fused wrapper zeroed both amax buffers with an unconditional memset while
the kernel returns early on noop[0]==1, so a skipped (graph-replay) call cleared
the previously published amax. Replace the memset with a noop-aware zero kernel
that returns early on the same flag, matching the standalone kernels' contract.

Signed-off-by: Cael Ling <caell@nvidia.com>

* [Common] nvfp4: hide row-scaled amax fused/unfused behind one entry

Drop the NVTE_NVFP4_FUSED_AMAX switch and make fused-vs-unfused an internal detail; callers use nvfp4::row_scaled::compute_amaxes, unfused stays as fallback.

Signed-off-by: Cael Ling <caell@nvidia.com>

* Assume fused amax kernel computes both row-wise and col-wise amaxes

Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: Tim Moon <4406448+timmoon10@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: Cael Ling <caell@nvidia.com>
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: 0z5a <Dezhen.lu@student.uni-tuebingen.de>
Signed-off-by: 0z5a <Dezhen.lu@student.uni-tuebingen.de>
…ackwell (NVIDIA#3313)

* Prefer Flash Attn 2 over Fused Attn for THD dropout on Blackwell

cuDNN's dropout kernels for THD attention are much slower than FlashAttention 2's
on SM100/103, and the cuDNN team confirmed a fix is not on their roadmap. The
generic Hopper+ rule prefers FusedAttention, so this combination silently took the
slow path with nothing reporting it.

Prefer FA2 for THD training with dropout on SM100/103 when FA2 is confirmed
usable, and leave every other configuration on the existing rule. The earlier
advisory is dropped: it ran before backend selection was final and could name
FusedAttention when a later filter selected something else.

Signed-off-by: Minho Ryu <ryumin93@gmail.com>

* Keep the THD dropout condition under the pylint bool limit

The six-term condition tripped R0916 (6/5), which the repo does not disable, so
qa/L0_pytorch_lint would have failed. Naming the hardware and workload half also
says what the condition is for.

Pin NVTE_FLASH_ATTN_V2 in the test as well. get_attention_backend reads it apart
from NVTE_FLASH_ATTN, so an environment that disables FA2 on purpose would fail
the availability assertion instead of skipping.

Signed-off-by: Minho Ryu <ryumin93@gmail.com>

* Remove the dedicated THD dropout backend test

Signed-off-by: Minho Ryu <ryumin93@gmail.com>

---------

Signed-off-by: Minho Ryu <ryumin93@gmail.com>
Co-authored-by: Kshitij Lakhani <33047503+KshitijLakhani@users.noreply.github.com>
Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>
Co-authored-by: Kshitij Lakhani <33047503+KshitijLakhani@users.noreply.github.com>
…30 GB -> 1.75 GB) (NVIDIA#3609)

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>
[PyTorch] Zero padded THD rows for FlashAttention and CP

Signed-off-by: xiaoyao0115 <1804647152@qq.com>
…irst load fallback (NVIDIA#3602)

The 128-bit fallback of ptx::ld_global_nc_evict_first_b32x8 emitted
ld.global.nc.L2::evict_first.v4.b32, but the .L2::evict_first load
qualifier is only valid on 256-bit accesses (CUDA 12.8's ptxas does not
accept it on ld at all). The fallback is compiled for Blackwell targets
with CUDA < 12.9, so building those targets with CUDA 12.8 failed.

Express the same hint through createpolicy.fractional.L2::evict_first
(fraction 1.0) and ld.global.nc.L2::cache_hint.v4.b32, as the helper's
comment already describes and as ld_global_nc_b32x8 with a policy does.
The 256-bit path is unchanged.

Fixes NVIDIA#3601

Signed-off-by: Ravi M <ravitejamajeti@gmail.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
* Fix JAX distributed unit test failure

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

* Scope change to 8 GPU config only

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

---------

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>
* Produce BOLT-compatible libraries

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

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

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

* Fix target arch detection

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

* fail fast when failing to detect arch

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

* reword env var description in docs

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

* add guard for when `CMAKE_SYSTEM_PROCESSOR` does not match CXX target arch

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

* Rename `build_target_arch` to `build_target_cpu_arch`

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

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

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

* Gate `-fno-reorder-blocks-and-partitions` on compiler identity

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

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

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

* Enable BOLT flags for all builds

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

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

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

* Enforce default compiler consistency for NCCL EP (needed to keep flags consistent)

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

* fix greptile PR comments about compiler probe failures and cross build arch mismatch

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>

---------

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Add MoE activation checkpoint names

Signed-off-by: Jeremy Berchtold <jberchtold@nvidia.com>
Initialize softmax custom calls before execution

Signed-off-by: Kshitij Lakhani <klakhani@nvidia.com>
Co-authored-by: Kshitij Lakhani <klakhani@nvidia.com>
* Prototype NVFP4 with UE5M3 scales

Co-authored-by: Teddy Do <tdophung@nvidia.com>
Co-authored-by: Varun Thumbe <vthumbe@nvidia.com>
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

* [PyTorch] Enable e5m3 fused GEMM kernels from cuDNN (NVIDIA#2)

* [PyTorch] Enable e5m3 fused GEMM kernels from cuDNN

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

* have to pad to 256 to use cuDNN

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

* fix: need to pass scale_dtype

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

* route wgrad to cuDNN's wgrad API

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

* Support grouped linear with NVFP4-UE5M3

NVFP4-UE5M3 grouped GEMM falls back to dense GEMMs. Generalize usage of wgrad kernel and use when tensors sizes are not 256-aligned. Fix inconsistent m,n,k GEMM notation. Remove ue5m3 hacks in op fuser tests. Add ue5m3 to grouped MLP tests.

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

* Fix typos

Co-authored-by: Codex <noreply@openai.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>

---------

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>
Co-authored-by: Tim Moon <tmoon@nvidia.com>
Co-authored-by: Codex <noreply@openai.com>

* Use custom recipe for NVFP4-UE5M3 tests

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

* Add grouped MLP kernel for GGEMM+SwiGLU+RHT+quant

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

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

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

* Debug integration with GGEMM+GLU+RHT+quant

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

* Remove scale max helper functions from NVFP4 cast utils

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

* Fix compile error

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

* Fix NVFP4 scale dtype ABI

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

* remove redundant output alloc

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

* no need to pad N now

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

* fix linting errors

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

* Disable cuDNN GGEMM+GLU+RHT+quant kernel

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

* Guard NVFP4 alpha scaling by scaling mode

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

* Restore UE5M3 NVFP4 cast support

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

* Localize NVFP4 4over6 scale policy

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

* Tweak arg order in C API functions

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

* Rename cuDNN GGEMM helper functions for general_gemm

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

* Fix compilation error in C++ test

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

* Treat nvfp4_e4m3_max=0 as unset value

Co-authored-by: Codex <noreply@openai.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>

* Debug torch.compile test failure

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

* Debug test failures

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

* Revert accidental CUTLASS commit change

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

* Enable cuDNN GGEMM+GLU+RHT+quant kernel

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

* Remove incorrect scale_inv_dtype arg to NVFP4Tensor constructor

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

* Fix bug when selecting cuDNN GGEMM+GLU+RHT+quant kernel

Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com>
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>

* Restore GGEMM+GLU+RHT+amax kernel with RHT sign mask

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

* Debug minor test failures

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

* Fix incorrect scale dtypes in grouped tensor builder method

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

* Avoid redundant amax ptr check in row-scaled NVFP4 quantize

Review suggestion from @ptrendx

Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>

* Clean up C++ unit tests

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

* Address some review comments from @ptrendx

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

* Respect no-tensor-scaling in grouped tensor helper function

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

* Debug C++ test compilation error

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

* Make sure GEMM alpha/beta scales are on GPU, even without amaxes

Co-authored-by: Codex <noreply@openai.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>

* Create new Mcore DDP integration functions rather than breaking backward compatibility

Co-authored-by: Codex <noreply@openai.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>

* Remove unrelated tests from NVFP4 dequantize C++ unit tests

Co-authored-by: Codex <noreply@openai.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>

* Add reference impl and test for NVFP4-UE5M3

Co-authored-by: Codex <noreply@openai.com>
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

* Make NVFP4-UE5M3 grouped GEMM CUDA-graph safe (NVIDIA#4)

* Make NVFP4-UE5M3 grouped GEMM CUDA-graph safe

The UE5M3 grouped GEMM builds its one-element int32 offsets tensor inline
at two call sites:

    torch.tensor([tokens], dtype=torch.int32, device=a_tensor.device)  # wgrad
    torch.tensor([N], dtype=torch.int32, device=device)                # fprop

Each of these materializes an unpinned CPU tensor and copies it host-to-device.
CUDA rejects that during graph capture:

    RuntimeError: Cannot copy between CPU and CUDA tensors during CUDA graph
    capture unless the CPU tensor is pinned. Please use tensor.pin_memory() or
    allocate the tensor with pin_memory=True.

Any model capturing these GEMMs in a CUDA graph therefore fails at capture
time. This was hit end-to-end on DeepSeek-V3 671B, where every convergence
config captures full forward/backward graphs; it aborts during warmup capture
and no training step completes.

Route both sites through a cached helper, mirroring get_cached_ones_tensor in
this same file -- which both functions already use for their `ones` tensor a
few lines earlier, and whose docstring notes it keeps "stable data pointers
across CUDA graph replays". Caching by (value, device) moves the host-to-device
copy to warmup and keeps the pointer stable across replays.

The offsets are constant for a given shape, so this is numerically neutral.

Signed-off-by: Michal Futrega <mfutrega@nvidia.com>

* Review suggestions

Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>

---------

Signed-off-by: Michal Futrega <mfutrega@nvidia.com>
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>

* Fix JAX NVFP4 GEMM device scalar handling

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

* Replace NVFP4 amax heuristic with explicit alpha/beta residency flag

b3415c1 gated the NVFP4 per-tensor scale branch in cublas_gemm() on
have_nvfp4_amax, an OR over the two operands' amax pointers. That branch
is also what moves alpha/beta from host into device workspace, and
CUBLASLT_POINTER_MODE_DEVICE is set unconditionally for every NVFP4 GEMM.
So when both operands had a null amax the branch was skipped, leaving
alpha/beta as host stack pointers that cuBLAS then dereferenced as device
addresses -- garbage output, not a mis-scaled result.

That is exactly the shape of the CI failure: only the both_unit_global_scale
parametrization of test_gemm_with_missing_nvfp4_amax went red, while the
x_/w_ cases passed because the OR still held.

Note that nvte_nvfp4_compute_per_tensor_scale with a null amax is an exact
identity on alpha (it substitutes scale_max * fp4_max), so skipping it never
saved meaningful work -- its real job on that path is materializing the
scalars on device. Making the call is always safe; skipping it is not.

Replace the heuristic with kNVTEMatmulConfigAlphaBetaOnDevice, which states
pointer residency directly instead of inferring it. Default false, so
PyTorch (nvte_cublas_gemm/_scaled pass host pointers, and generic_gemm
reaches nvte_cublas_gemm_v2 without setting the flag) is restored to its
pre-b3415c19 behavior. JAX sets it when is_nvfp4_scaling(), matching the
device buffers it binds for alpha/beta, and folds the per-tensor scale into
alpha itself.

Also document in the JAX helper that its block-scale bound is hardcoded to
E4M3's 448 while the C++ get_nvfp4_scale_max() returns 114688 for UE5M3,
which will need reconciling once UE5M3 reaches the JAX path.

Verified on GB200: tests/pytorch/nvfp4 9310 passed / 0 failed / 6852 skipped
(CI at b3415c1 was 9309 passed / 1 failed / 6852 skipped), tests/pytorch/mxfp8
371 passed, and JAX nvfp4+gemm 382 passed.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Signed-off-by: tdophung <tdophung@nvidia.com>

* Document bugs with `alpha_beta_on_device` GEMM config

Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>

* [PyTorch] Do not swizzle operand scales in place in the NVFP4-UE5M3 GEMM (NVIDIA#5)

* Do not swizzle operand scales in place in the NVFP4-UE5M3 GEMM

The cuDNN UE5M3 GEMM wrapper swizzled A/B scales in place and marked
the tensors as swizzled. For weights that mark persisted on the
parameter. The master-weight cast (_cast_master_weights_to_nvfp4_2d,
used by Megatron's fp4 param gather after every optimizer step) then
rewrote the scales in the unswizzled layout without clearing the mark,
so every later GEMM and dequantize read unswizzled scales as swizzled.
Gaussian weights hide it (near-uniform scales, ~0.2 relative error);
real DeepSeek-V3 weights are destroyed (2.6 relative error, loss 12.98
vs 7.86 for E4M3 under identical routing).

Swizzle clones of the scale tensors instead and leave the operands
untouched, matching what the cuBLAS path does in C++.

* Test that the NVFP4-UE5M3 GEMM leaves operand scales untouched

* Avoid suppressing errors in swizzle function.

Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>

---------

Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>

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

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

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

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

* Return the swizzled scales the UE5M3 GEMM helper actually produced (NVIDIA#6)

tex.swizzle_scales_for_gemm_ assigns new tensors to the scale attributes
instead of writing into the buffers it is given, so returning the clones
made before the call hands the cuDNN GEMM unswizzled scales. On DeepSeek-V3
671B this turned the forward into garbage: routing collapsed onto expert 0
and the EP-rank-0 GPUs ran out of memory in the first expert forward.
Read the attributes back after the swizzle, and make the regression test
check the GEMM result against the dequantized reference.

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

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

* Debug compilation errors

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

* Debug compile error in NVRTC NVFP4 4over6 kernel

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

* Debug NVFP4 UE5M3 in grouped MLP tests

Use list of NVFP4 variant names rather than repeating at callsites. Exclude SReLU variants with UE5M3.

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

* Fix typo

Review suggestion from @greptile-apps.

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

* Debug NVFP4 4over6 tests

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

---------

Signed-off-by: Tim Moon <tmoon@nvidia.com>
Signed-off-by: Kaining Zhong <kainingz@nvidia.com>
Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: Michal Futrega <mfutrega@nvidia.com>
Signed-off-by: tdophung <tdophung@nvidia.com>
Co-authored-by: Teddy Do <tdophung@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>
Co-authored-by: Kaining Zhong <44538064+kainzhong@users.noreply.github.com>
Co-authored-by: Codex <noreply@openai.com>
Co-authored-by: Kaining Zhong <kainingz@nvidia.com>
Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com>
Co-authored-by: Michal Futrega <michal.futrega@gmail.com>
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
Bump nccl ep

Signed-off-by: Phuong Nguyen <phuonguyen@nvidia.com>
Signed-off-by: Alex Y. Chan <alechan@nvidia.com>
Add Triton dependency to JAX extension

Signed-off-by: Jeremy Berchtold <jberchtold@nvidia.com>
The multi-tensor swizzle and unswizzle kernels launch with TB_DIM x TB_DIM
(1024) threads but, unlike the other swizzle kernels, had no
__launch_bounds__. ptxas could then use more than 64 registers per thread
(up to 99 for the int4 variants), which makes the launch fail with "too
many resources requested for launch". Which variants fail depends on the
target arch and CUDA version.

Add __launch_bounds__(TB_DIM * TB_DIM) to the four kernels and add test
shapes that reach the rowwise int2 and columnwise int4 variants.

Fixes NVIDIA#3621

Signed-off-by: Ravi M <ravitejamajeti@gmail.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
* Skip FA4 for deterministic backward on SM120

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

* Respect autograd and device context in FA4 selection

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

* Simplify attention device capability lookup

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

---------

Signed-off-by: jbernloehr <jbernloehr@nvidia.com>
remove first launch protection

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>
cael-ling and others added 30 commits October 26, 2026 23:17
…ath (NVIDIA#3649)

* [Common][PyTorch] Add NVFP4 row-scaled grouped cast with graph-safe path

Single-launch grouped row-scaled NVFP4 cast + fused amax, plus graph-safe amax/cast variants driven by device first_dims/tensor_offsets for capture-once/replay MoE. Wired into PyTorch (ep.nvfp4_group_row_scaled_cast) with C++ and PyTorch capture/replay tests.

Signed-off-by: Cael Ling <caell@nvidia.com>

* Address review: guard grouped row-scaled cast dispatch and inputs

Restrict the FUSED_NVFP4_ROW_SCALED selection to shapes the grouped kernels accept (128-aligned last dim, <= 64 experts) so unsupported requests fall back to the per-expert path. Return early for empty routing, make the graph-safe binding's input contiguous, and reject with_rht/with_2d_quantization/with_amax_reduction on both row-scaled entries. Add regression tests for the four guarded paths.

Signed-off-by: Cael Ling <caell@nvidia.com>

* Guard row-scaled grouped cast against non-2D, partial, and strided routing

The grouped row-scaled cast treats its input as 2D [rows, last_dim] indexed by
split counts, so route higher-rank inputs and partial splits (split sum below
the row count) to the per-expert path instead of the grouped kernel. Make
first_dims and tensor_offsets contiguous in the graph-safe cast so a strided
routing view is not misread. Add fallback and strided-routing tests.

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

* Fall back to per-expert cast for mixed NVFP4 row/tensor-scaled splits

The bulk NVFP4 split_quantize allocation sizes its amax buffer from the
first quantizer's row-scale setting, so a group mixing row-scaled and
tensor-scaled quantizers writes per-row amax into a per-tensor buffer and
corrupts neighbors. Detect the mixed case and fall back to the unfused
per-split path.

Signed-off-by: Cael Ling <caell@nvidia.com>

* Route row-scaled NVFP4 grouped quantize through nvte_group_quantize

Add the kNVTEGroupedRowScaledNVFP4 grouped tensor param and dispatch
row-scaled NVFP4 outputs in nvte_group_quantize to the graph-safe amax
and cast kernels. Handle row-scaled quantizers in tex.group_quantize.

Remove the row-scaled-specific C APIs, the host-split row-scaled cast
kernel, tex.nvfp4_group_row_scaled_cast_graph_safe, the ep.py helper,
and the fused row-scaled split_quantize path.

Add a top-level static_assert for unregistered types in TypeInfo.

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

* Require 128-row experts for uniform row-scaled NVFP4 group_quantize

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

* Fold row-scaled NVFP4 amax into nvte_group_amax

Dispatch nvte_group_amax and nvte_group_amax_graph_safe on the output
tensor format so that row-scaled NVFP4 outputs get per-row amaxes, and
deprecate nvte_group_nvfp4_compute_amax in favor of nvte_group_amax.
Rename the row-scaled amax helpers to group_row_amax(_graph_safe).

In split_quantize, fall back to unfused allocation and quantization
whenever any NVFP4 quantizer is row-scaled.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
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

* Fix missing include

Review suggestion from @greptile-apps

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

* Add missing header include

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

---------

Signed-off-by: Cael Ling <caell@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>
Co-authored-by: Tim Moon <tmoon@nvidia.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
…IA#3596)

Support user-provided routing and return expert counts

Add RouterComputationInfo and RoutingMapInfo to the JAX MoE API. moe() takes
one router_info argument: RouterComputationInfo has TE compute the gate and
top-k, while RoutingMapInfo supplies caller-computed expert indices and
differentiable routing weights (e.g. DeepSeek-V4 hash routing) and skips
TE's router.

moe() and _MoEBlock also return global per-expert assignment counts as a
fourth output. collect_expert_counts=False skips the associated all-reduce.

Signed-off-by: Deep Patel <deepatel@nvidia.com>
* [Common] Keep cuda_driver::call's symbol cache alive through exit

rtc::KernelManager's static destructor unloads its CUDA modules through
cuda_driver::call. Each instantiation of call caches driver entry points in
function-local statics, and those caches are usually first constructed after
the KernelManager, so they are destroyed before it. ~Kernel then reads a
destroyed unordered_map: with GCC 16's libstdc++ the process segfaults at exit;
with GCC 13's it writes into freed memory instead.

Allocate the cache and its mutex with new and never destroy them, and add a
test that the process exits cleanly after running NVRTC kernels.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Connor Baker <connorb@nvidia.com>

* [Common] Never destroy rtc::KernelManager or the fused-attention graph caches

Both were function-local statics whose destructors call into CUDA after main
returns: ~Kernel unloads modules through the driver, and destroying a cached
cuDNN plan that a CUDA graph retains releases a CUDA user object. The CUDA
Programming Guide makes any use of CUDA during program termination undefined
behavior, and both crashed at exit depending on static destruction order.

Allocate both with new and never delete them, as TVMFFIConfigCache::create()
and the fused-attention debug registry already do. Both caches only ever
insert, so nothing changes while the process runs.

Replace tests/pytorch/test_process_exit.py, whose LayerNorm only runs NVRTC
kernels while the NVTE_NORM_FWD_USE_CUDNN and NVTE_NORM_BWD_USE_CUDNN switches
are off, with MALLOC_PERTURB_=165 for the C++ util tests. UtilTest.NVRTC drives
rtc::KernelManager directly, and ctest checks each test's exit status.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Connor Baker <connorb@nvidia.com>

---------

Signed-off-by: Connor Baker <connorb@nvidia.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
* Fix DistributedWeight handling under saved-tensor hooks

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

* Reduce CPU overhead of the DistributedWeight check in prepare_for_saving

prepare_for_saving and restore_from_saved run for every tensor every TE
module saves, so the DistributedWeight check sits on a hot path for all
users, not just GTP. Test None and torch.Tensor first, read the marker
with an inline getattr (DistributedWeight implementers must be Tensor
subclasses), and only query torch.compiler.is_compiling() for a flagged
weight. On restore, test the common None/Tensor case first.

On a CPU microbenchmark this brings the non-distributed cost back to
roughly main's, e.g. 4.5 -> 5.2 us (was 10.0 us) per prepare_for_saving
for a 32-expert bf16 GroupedLinear.

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* Add saved-tensor hook tests for DistributedWeight modules

Port the Linear and GroupedLinear saved-tensor hook tests from NVIDIA#3545 to
the prepare_for_saving fix, and extend them:

- Run each module case both under identity saved_tensors_hooks and under
  te.checkpoint(use_reentrant=False), whose pack hook discards saved
  tensors and returns recomputed ones in backward.
- Cover LayerNormLinear, which also gates its backward on the saved
  weight being a DistributedWeight.
- Add CPU-only checks that prepare_for_saving keeps a DistributedWeight
  by reference and that backward gets the same object back with no
  hooks, identity hooks, and torch.utils.checkpoint.

Co-authored-by: Xiaowei Ren <xren@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* Still use the autograd checker for the dist weight

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

* Error for distweight + torch.compile

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

* Rename _SavedDistributedWeight to _SavedTensorRef and drop premise test

Remove test_saved_tensor_hooks_unpack_subclass_as_plain_tensor, which
only checked PyTorch's saved-tensor-hook behavior rather than TE's.

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

---------

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Co-authored-by: Xiaowei Ren <xren@nvidia.com>
* Remove QB bounds device side validation

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

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

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

---------

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
* [PyTorch] Test grouped FP8 block-scaling GEMM against the split-tensor path

MXFP8 GEMM on Blackwell accumulates at higher precision than FP8 block-scaling GEMM on Hopper.
The grouped FP8 block-scaling GEMM tests compared against an FP32 reference, so they passed with
the MXFP8 emulation on Blackwell but failed on Hopper.

Compare against the legacy split-tensor grouped GEMM instead, consistent with the existing MXFP8
grouped GEMM tests. Also remove the FP8 block-scaling zero-size-expert and mixed-operand tests,
which no other recipe has.

Signed-off-by: Alp Dener <adener@nvidia.com>

* [PyTorch] Test grouped FP8 block-scaling GEMM with uneven and empty experts

Dropping the varying splits and the zero-size-expert test left gaps: the remaining tests no longer
run grouped FP8 block-scaling GEMM with uneven or empty experts when the weights or outputs are
per-expert lists, when accumulating, or when K is not a multiple of 128.

Parametrize the grouped FP8 block-scaling GEMM test over uniform, varying and zero-expert splits,
checked against the split-tensor reference.

Signed-off-by: Alp Dener <adener@nvidia.com>

---------

Signed-off-by: Alp Dener <adener@nvidia.com>
NVIDIA#3513)

* fix: patch cuda hardcodes for plugin compatibility (complete debug directory coverage)

This commit completes the CUDA plugin compatibility patches across the entire codebase:

1. Debug directory (transformer_engine/debug/):
   - fake_quant.py: Replace is_cuda checks with te_device_type()
   - log_fp8_tensor_stats.py: Replace torch.device('cuda') and torch.cuda.get_device_capability()
   - per_tensor_scaling.py: Replace is_cuda checks with te_device_type()
   - utils/stats_buffer.py: Replace .cuda() calls with .to(te_device_type())

2. Additional missing files from cuda_patch_rc2.14:
   - permutation.py: Replace all is_cuda checks with device.type comparisons
   - tensor/float8_tensor.py: Replace is_cuda checks with device.type comparisons
   - tensor/float8_blockwise_tensor.py: Replace is_cuda checks with device.type comparisons
   - tensor/mxfp8_tensor.py: Replace is_cuda checks with device.type comparisons

3. Core infrastructure:
   - transformer_engine/__init__.py: Export te_platform and te_device_type from pytorch module
   - transformer_engine/pytorch/__init__.py: Define te_platform() and te_device_type() functions

All runtime CUDA hardcodes have been replaced with plugin-compatible abstractions.
Total: 64 files modified (+405/-289 lines)

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* fix: remove explicit import to avoid black parsing error

The te_platform and te_device_type functions are dynamically assigned to
the transformer_engine module namespace via:
  transformer_engine.te_platform = te_platform
  transformer_engine.te_platform = te_platform

in transformer_engine/pytorch/__init__.py, so explicit imports in
transformer_engine/__init__.py cause black to fail parsing due to
circular import detection.

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* fix: correct malformed imports that prevent module loading

Move te_platform and te_device_type imports outside of parenthesized
import lists to fix SyntaxError. The imports were incorrectly placed
inside multi-line import statements, preventing the modules from loading.

Fixed files:
- transformer_engine/debug/features/log_fp8_tensor_stats.py
- transformer_engine/pytorch/attention/dot_product_attention/backends.py
- transformer_engine/pytorch/attention/dot_product_attention/context_parallel.py
- transformer_engine/pytorch/attention/dot_product_attention/dot_product_attention.py
- transformer_engine/pytorch/attention/inference.py
- transformer_engine/pytorch/graph.py
- transformer_engine/pytorch/quantization.py
- transformer_engine/pytorch/quantized_tensor.py
- transformer_engine/pytorch/transformer.py

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* fix: add missing te_device_type import to files using it

Several files imported only te_platform but also called te_device_type(),
causing NameError at runtime. Add te_device_type to imports.

Fixed files:
- transformer_engine/pytorch/attention/dot_product_attention/flex_attention.py
- transformer_engine/pytorch/attention/rope.py
- transformer_engine/pytorch/cpp_extensions/gemm.py
- transformer_engine/pytorch/optimizers/fused_adam.py
- transformer_engine/pytorch/utils.py

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* fix: correct malformed imports in additional files

Move te_platform and te_device_type imports outside of parenthesized
import lists to fix SyntaxError across 20 more files.

Fixed files:
- transformer_engine/pytorch/attention/dot_product_attention/utils.py
- transformer_engine/pytorch/attention/multi_head_attention.py
- transformer_engine/pytorch/cpu_offload.py
- transformer_engine/pytorch/module/base.py
- transformer_engine/pytorch/module/grouped_linear.py
- transformer_engine/pytorch/module/layernorm_linear.py
- transformer_engine/pytorch/module/layernorm_mlp.py
- transformer_engine/pytorch/module/linear.py
- transformer_engine/pytorch/ops/basic/basic_linear.py
- transformer_engine/pytorch/ops/basic/grouped_linear.py
- transformer_engine/pytorch/ops/basic/l2normalization.py
- transformer_engine/pytorch/ops/fused/forward_linear_scale_add.py
- transformer_engine/pytorch/ops/fused/grouped_mlp.py
- transformer_engine/pytorch/ops/fused/userbuffers_backward_linear.py
- transformer_engine/pytorch/ops/fused/userbuffers_forward_linear.py
- transformer_engine/pytorch/ops/op.py
- transformer_engine/pytorch/tensor/nvfp4_tensor.py
- transformer_engine/pytorch/tensor/storage/mxfp8_tensor_storage.py
- transformer_engine/pytorch/triton/mhc.py
- transformer_engine/pytorch/triton/permutation.py

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* fix: use routing_map device instead of hardcoded device type

In make_row_id_map, use routing_map.device for tensor allocation
instead of te_device_type() to respect the input tensor's device.

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* revert: remove te_platform() replacements, keep te_device_type()

Revert all te_platform() method calls back to torch.cuda.* calls.
Keep te_device_type() string replacements for device type parameters.
Preserve te_platform definition in transformer_engine/pytorch/__init__.py.

Changes:
- Reverted 146 te_platform() calls back to torch.cuda.*
- Removed 'from transformer_engine import te_platform' imports
- Kept all te_device_type() replacements intact
- Kept te_platform definition in __init__.py for future use

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* fix: remove remaining te_platform() calls in utils, fused_adam, and distributed

Replace remaining te_platform() calls with torch.cuda equivalents:
- torch.cuda.nvtx.range() / range_push() / range_pop() for NVTX profiling
- torch.cuda.amp.autocast for autocast context manager
- Update docstring reference in fused_adam.py

All te_platform() usage has been removed except for the definition
in transformer_engine/pytorch/__init__.py.

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* keep sync with torch.cuda api usage

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* format

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Fix missing te_device_type imports and import ordering

Add the missing `te_device_type` import to eight modules that called it
without importing it, which would raise NameError at runtime: jit,
attention/inference, attention/dot_product_attention/{backends,
context_parallel,utils}, ops/fused/{grouped_mlp,
userbuffers_backward_linear}, and triton/mhc.

Move the `te_device_type` import ahead of relative imports where pylint
classifies it as a first-party import that must precede local ones. In
jit and distributed, place it so it does not split the torch import
group. Order `transformer_engine_torch` before `transformer_engine` in
ops/_common.

Add docstrings to `te_device_type` and `te_platform`.

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Query autocast state and dtype on the same device

The device-type patch had changed the autocast dtype lookups to the
configured device while leaving the enabled checks on CUDA. With a
non-CUDA device type, vendor autocast enabled, and CUDA autocast
disabled, these sites took the non-autocast branch and kept the input
dtype instead of the requested autocast dtype.

Route the operation sites through `maybe_autocast_dtype`, which already
queries one device for both calls and carries the torch version guard.
This covers the activation, SwiGLU, grouped linear, basic linear, and
fused linear forward operations.

Add `torch_is_autocast_gpu_enabled` next to `torch_get_autocast_gpu_dtype`
for the two sites outside the operation layer that had the same
mismatch: `set_activation_dtype`, which sets the activation dtype for
every module, and the TransformerLayer forward cast.

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Gate per-document THD partitioning on CUDA

`tex.thd_get_partitioned_indices` allocates a CUDA output tensor and
launches a CUDA kernel on the current CUDA stream, so it runs only on
CUDA. The device-type patch had changed its dispatch condition to the
configured device type, which sent packed sequences into that kernel on
a non-CUDA device instead of the device-capable fallback, failing the
`is_cuda` assertion before the fallback could run.

Restore the CUDA condition in `build_rank_indices` and in
`_get_thd_partition_cu_seqlens`, which had the same problem when
normalizing the boundaries that feed the kernel, and record why both
are gated on CUDA rather than on the configured device.

Keep the configured device in the NO_LOAD_BALANCE branch, which builds
indices with plain tensor operations and runs no CUDA kernel. Restore
the assertion message, which named the device type that triggers it.

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Drop incidental changes from the device-type patch

Narrow the diff against main to the device-type substitutions themselves:

- context_parallel.py: remove the comments on the THD partitioning
  device gates.
- cpp_extensions/gemm.py: restore the single-line
  Float8BlockwiseQTensorStorage import, which did not need wrapping.
- graph.py: remove an added blank line, leaving the file unchanged from
  main.

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Cover device-type checks missed in recently merged code

Two paths reached from main used CUDA device checks where the work below
them is not CUDA-specific:

- tensor/nvfp4_tensor.py: `NVFP4Tensor` kept CUDA checks in its
  `record_stream` dispatch and in `copy_`, while the sibling
  `Float8Tensor`, `MXFP8Tensor`, and `Float8BlockwiseQTensor` classes use
  the configured device type for the same logic. On another device type
  `copy_` would move the operand to the destination device and
  `record_stream` would skip stream tracking.
- cross_entropy.py: the input validation rejected non-CUDA tensors even
  though the Triton backend derives every allocation from the input
  device and contains nothing CUDA-specific.

Leave the checks that guard compiled CUDA kernels, the `device_types`
dispatch keys on custom operations, and the `is_cuda` properties that
implement the PyTorch tensor contract.

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Apply plugin device patches after the _nv alias is registered

The top-level hook in transformer_engine/__init__.py imported
<plugin>.patches before load_framework_extension() registered
transformer_engine_torch_nv. A plugin that imports that alias at package
level - a case the loader explicitly supports - failed there, and the
enclosing except only warned, so apply_patches() never ran. The loader
later imported the plugin successfully but never retried the patches,
leaving device patches unapplied while vendor calls still went through
CUDA.

Move the call into load_framework_extension(), after the alias is
registered and before the PyTorch submodules capture references to
torch.cuda entry points. Probe the submodule with find_spec so a plugin
that needs no patches may omit it, without masking real import failures
inside an existing patches module.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Accept the configured device type in the shared forward input check

The modules whose `device` argument now defaults to `te_device_type()`
allocate their parameters on the configured device, but the
`prepare_forward()` input check still rejected anything whose `is_cuda`
was false. That property is a tensor attribute, so patching the
`torch.cuda` namespace does not change it, and forward raised before a
plugin's native operation could run.

Compare `device.type` against `te_device_type()` instead, and report the
configured device type in the error message. One check covers Linear,
LayerNormLinear, GroupedLinear, and LayerNormMLP, which share this base
method. Behavior on CUDA is unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Restore the original spacing in the top-level package

Removing the plugin hook left an extra blank line behind. Drop it so the
file matches main and no longer appears in the diff.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Skip patch discovery for single-file plugins

find_spec() raises ModuleNotFoundError instead of returning None when
the parent module is not a package, so a single-file plugin that
provides load_plugins() failed the optional-patches probe. The exception
handler then disabled the plugin and restored the original extension.

Check for __path__ before searching for the submodule. Genuine import
failures inside an existing patches module still propagate.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Use the configured device type in GroupedTensor record_stream

The merged upstream GroupedTensor stream-tracking handler checked
`tensor.is_cuda`, which pins the path to CUDA. Match the device
abstraction used by the other quantized-tensor `record_stream`
handlers so plugin device types propagate allocator stream tracking.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Accept the configured device type in the attention input checks

The attention entry points asserted `is_cuda` on queries, keys, and
values, so a non-CUDA configured device failed validation before any
attention implementation could run, even though the surrounding state
was already created on the configured device.

Check the configured device type instead, in DotProductAttention and in
the matching FlashAttention and FusedAttention assertions. Apply the
same change to the per-document THD partitioning check in context
parallelism, and drop the redundant CUDA half of the mixed-THD
sequence_ids check, where the comparison against the Q/K/V device
already enforces the documented requirement.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Keep the paged KV cache and EP buffer on CUDA

The paged KV cache page table and the EpBuffer default device fed
tensors into torch.library custom ops registered with
device_types="cuda". Allocating them on a configured non-CUDA device
left no matching registered implementation, so the dispatcher failed
the call before the wrapped tex function could run.

Restore the CUDA allocations so these tensors match their op
registrations, and drop the now-unused imports. Supporting another
device on these paths requires changing the op registrations, which
depends on how the device is registered with PyTorch and on when the
device type is configured relative to module import.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Apply plugin device patches after the plugin finishes loading

apply_patches() rewrites process-global state in place - torch.cuda entry
points, TE_DEVICE_TYPE and TE_PLATFORM - with no way to undo it. Running it
before load_plugins() meant a load_plugins() failure left those patches
active while the except block restored only sys.modules, so import continued
with the CUDA extension under a vendor device type and failed later, far from
the cause.

Run load_plugins() first so everything that can fail is rolled back
completely. The patches still precede the PyTorch submodules that capture
torch.cuda entry points, and still follow the _nv alias registration that a
plugin package may need to import itself.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>

* Fix regex in test_cross_entropy

Signed-off-by: Przemyslaw Tredak <ptrendx@gmail.com>

---------

Signed-off-by: Xianduo Li <lixianduo@mail.nankai.edu.cn>
Signed-off-by: Przemyslaw Tredak <ptrendx@gmail.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Co-authored-by: Przemyslaw Tredak <ptrendx@gmail.com>
* Guard fused MoE auxiliary loss integer offsets

Signed-off-by: Srijan Keshri <212402043+arcusbuilds@users.noreply.github.com>

* Update tests/pytorch/test_fused_router.py

Co-authored-by: Alp Dener <adener@nvidia.com>
Signed-off-by: Srijan Keshri <212402043+arcusbuilds@users.noreply.github.com>

* Use 64-bit indexing for fused MoE auxiliary loss

Signed-off-by: Srijan Keshri <212402043+arcusbuilds@users.noreply.github.com>

* Unroll forward row loop in fused MoE aux loss

Signed-off-by: arcusbuilds <srijankeshri007@gmail.com>

* Revert 64-bit indexing, restore INT_MAX guard in fused MoE aux loss

Match the other fused router kernels, which reject
num_tokens * num_experts > INT_MAX. Keep #pragma unroll 4 on the
forward row loops.

Signed-off-by: arcusbuilds <srijankeshri007@gmail.com>

* Use 64-bit loop counters in fused MoE aux loss kernels

Shapes at the INT_MAX guard (e.g. INT_MAX x 1) overflowed the int loop
counters on their last increment and the backward grid-size math.
Flattened offsets stay 32-bit, which the guard bounds.

Signed-off-by: arcusbuilds <srijankeshri007@gmail.com>

---------

Signed-off-by: Srijan Keshri <212402043+arcusbuilds@users.noreply.github.com>
Signed-off-by: arcusbuilds <srijankeshri007@gmail.com>
Co-authored-by: Alp Dener <adener@nvidia.com>
* rely on recipe object for fusion, reduce cpu overhedads for deciding the grouped tensor path

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

* Keep fusion decisions local to op fuser

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

* Use recipe representation for fusion cache key

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

* Include CUDA device in grouped tensor support cache

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

* Declare derived quantizer settings as recipe fields

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

* Update transformer_engine/pytorch/ops/fuser.py

Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>

* Update transformer_engine/pytorch/module/grouped_linear.py

Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>

* Update transformer_engine/pytorch/ops/basic/grouped_linear.py

Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>

---------

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>
Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
…caled kernels (NVIDIA#3682)

Optimize scale swizzle kernels

Use optimized rowwise full-width and columnwise chunk kernels for MXFP8 and NVFP4 scales, including multi-tensor swizzling. Keep the 16-byte scale-buffer alignment contract and cover its validation in C++ tests.

Choose the rowwise full-width kernel after the fallback grid exceeds one SM wave, based on paired GB200 and VR200 measurements.

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>
Fix GH CI docs

Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
…re (NVIDIA#3698)

* [CUDA] Handle THD boundary capacity across CP utility kernels

THD metadata can exceed per-block shared memory even when the token tensors are valid. Use device-derived default and opt-in capacities before reading the existing global boundary arrays, without allocating scratch buffers.

Centralize staging and normalization so indexing helpers consume consistent boundaries while retaining pre-scaled shared reads. Preserve the current old_lse output-accumulation behavior and exercise its rescaling in boundary regression tests.

Focused native validation passed 348 boundary/dtype cases, 60 output-rescaling cases including 24 exact upstream comparisons, 33 graph captures with three replays each, and 324 partition stress cases up to one million documents. SM86 and SM100 compilation passed.

Signed-off-by: Sudhakar Singh <sudhakars@nvidia.com>

* [PyTorch] Guard THD boundary tests on available device limits

Older PyTorch builds do not expose the shared-memory properties used to choose exact boundary-test sizes. Check both capabilities inside the two affected tests and skip only those tests when either field is missing, without guessing a version cutoff or querying CUDA during import.

Include the running PyTorch version in the skip reason to make compatibility skips identifiable. Production kernels and the remaining functional tests are unchanged. Missing/present-property checks and both targeted native boundary tests passed.

Signed-off-by: Sudhakar Singh <sudhakars@nvidia.com>

* Use runtime storage selection for THD boundary tables

Avoid duplicating each dtype, layout, and functor specialization solely for
boundary storage. Pass a block-uniform storage flag and plain pointers with
an explicit read divisor, preserving shared pre-scaling and immutable
global metadata.

Use the device-reported default shared-memory budget and the existing
global table as fallback, removing kernel-specific opt-in state and the
boundary-view type. Retain per-boundary integer division and CP validation,
and test the two remaining storage regimes.

Signed-off-by: Sudhakar Singh <sudhakars@nvidia.com>

* Restore compile-time storage selection for THD boundaries

Runtime storage selection obscures the shared address space and unit read
divisor, adding work to boundary searches. Restore shared/global kernel
specializations so the compiler can eliminate global-only normalization
from the shared path.

Keep plain pointer and explicit divisor arguments instead of the boundary
view. Retain the default shared-memory budget, immutable global fallback,
and the existing default-boundary regression tests without restoring
shared-memory opt-in.

Signed-off-by: Sudhakar Singh <sudhakars@nvidia.com>

---------

Signed-off-by: Sudhakar Singh <sudhakars@nvidia.com>
* Fail fast on Hopper FA4 backward with affected Quack versions

Signed-off-by: Kshitij Lakhani <klakhani@nvidia.com>

* Remove redundant Quack guard documentation

Signed-off-by: Kshitij Lakhani <klakhani@nvidia.com>

---------

Signed-off-by: Kshitij Lakhani <klakhani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
…DIA#3769)

Link tests executables against cpython when CuTeDSL is enabled

Signed-off-by: Fred Heinecke <fheinecke@nvidia.com>
* [Docs] Add MoE feature overview

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

* [Docs] Refine MoE overview, add supporting snippets/figures and JAX API entries

- Add code snippets and SVG figures referenced by mixture_of_experts.rst
  (moe_permute / moe_unpermute / grouped_linear tabbed examples for both
  PyTorch and JAX)
- Add JAX API reference entries for token_dispatch, token_combine and
  grouped_dense so the cross-references from the MoE page resolve
- Make wording framework-neutral where it was PyTorch-only (Grouped GEMM
  instead of GroupedLinear/grouped linear in shared sections, both
  m_splits and group_sizes mentioned, figure labels generalized)
- Tighten routing-kernel intro: consolidate the redundant "multiple
  variants exist / see API ref" notes into one paragraph next to the
  example, and explicitly state that the kernels are differentiable
- Sharpen merging_probs explanation (top-1 vs top-k) and explicitly
  describe what token_dispatch / token_combine return
- Snippet cleanups: define previously undefined symbols, drop the JAX
  probs= argument from the basic example and explain its purpose in a
  comment, document the ignored permuted_probs / pad_offsets outputs
- Reorder MoE entry in the docs/index.rst toctree

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
Co-authored-by: Cursor <cursoragent@cursor.com>

* [Docs] Extend MoE overview: router, end-to-end layer, padding, low-precision

Add Router (score function + top-k + load-balancing loss) and Putting-it-together
sections, plus token-probabilities / padding-and-alignment / chunk-sort subsections
and a fused-expert-MLP note. New SVG figures and PyTorch/JAX snippets. Add the
router and moe_permute_and_pad_with_probs API reference entries (PyTorch and JAX)
and sort_chunks_by_index (JAX).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>

* [Docs] MoE: fused grouped MLP and expert parallelism sections

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

* [Docs] MoE: split into per-topic pages, match current docs style, document EP APIs

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

* [Docs] MoE: single article with sections, redraw expert parallelism diagram

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

* [Docs] MoE: warn about uneven low-precision support

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

* [Docs] MoE: shorten low-precision warning

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

* [Docs] MoE: reword low-precision note

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

* [Docs] MoE: drop API-reference pointer from low-precision note

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

* [Docs] MoE: reorganize into introduction, routing kernels, grouped GEMM, grouped MLP, EP

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

* [Docs] MoE: move end-to-end example to the end of the article

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

* [Docs] MoE: add expert-parallel layer diagram to the end-to-end example

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

* [Docs] MoE: use expert-parallel layer diagram as Figure 1, mention EP in introduction

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

* [Docs] MoE: fold EP into the stage list, surface load balancing in the introduction

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

* [Docs] MoE: explain grouped routing, simplify router wording

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

* [Docs] MoE: deduplicate expert_bias description

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

* [Docs] MoE: move routing-kernels lead paragraphs above the router subsection

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

* [Docs] MoE: describe what the blocks do and how to use them; move framework specifics to snippets

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

* [Docs] MoE: describe grouped GEMM execution paths; shorten grouped MLP conditions

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

* [Docs] MoE: move grouped GEMM paths above snippet; simplify grouped MLP figure

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

* [Docs] MoE: EP dispatch/combine replace token dispatch/combine; redraw layer figure

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

* [Docs] MoE: split routing kernels into Router and Token permutation sections

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

* [Docs] MoE: fix EP figure dispatch arrows; animate dispatch and combine

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

* [Docs] MoE: drop EP figure animation

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

* [Docs] MoE: restructure expert parallelism (overview, common API, per-framework API)

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

* [Docs] MoE: EP back in tabs without C API; describe sort_chunks_by_index

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

* [Docs] MoE: EP placement figure matches flow figure; clarify capacity, MXFP8 dispatch, shared experts

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

* [Docs] MoE: break long paragraphs into lists

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

* [Docs] MoE: explain NCCL EP and zero-copy mode; link examples

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

* [Docs] MoE: EP wording (optimized ops, worst-case capacity, MXFP8 status)

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

* [Docs] MoE: describe EP buffer allocation and data flow (PyTorch)

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

* [Docs] MoE: clarify what EpBuffer holds and when to reuse it

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

* [Docs] MoE: spell out the JAX EP primitives

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

* [Docs] MoE: precise wording for MXFP8 dispatch output

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

* [Docs] MoE: restructure EP overview along one step (buffer, payload, transfer path)

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

* [Docs] MoE: EP overview headings and terser bullets

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

* [Docs] MoE figures: larger fonts, fix label anchors and clipping, drop placement footer

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

* [Docs] MoE: EP overview - one-step permutation, buffer sizing modes, zero-copy vs CUDA graphs

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

* [Docs] MoE: drop zero-copy block from EP overview

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

* [Docs] MoE: shorten EpBuffer bullet

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

* [Docs] MoE: merge PyTorch EP API and data-flow lists

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

* [Docs] MoE: describe EP backward (PyTorch)

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

* [Docs] MoE: one-step dispatch figure; capacity sizing guidance

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

* [Docs] MoE: single-device example before expert parallelism

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

* [Docs] MoE: receive buffer block as nested bullets with overflow behaviour

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

* [Docs] MoE: drop EP backward bullet

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

* [Docs] MoE: larger row labels in layer figure

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

* [Docs] Fix MoE routing guidance and examples

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

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

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

* [Docs] Fix RST in router and moe docstrings so the API pages build with -W

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

* [Docs] MoE: mask padded receive slots before combine in the PyTorch EP snippet

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

* [Docs] Fix MoE EP buffer handling and grouped MLP guidance

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

* [Docs] Update MoE EP operations and fix guide examples

Document MoeDispatch and MoeCombine composition with internal channels. Preserve actual routing decisions for auxiliary loss, use the fusion-compatible GLU layout, and distinguish PyTorch and JAX EP behavior.

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

* [Docs] Link grouped MLP fusion requirements to TE v2.19

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

* Document compact top-k routing indices

Explain the PyTorch topk_indices output buffer, the unchanged probability tensor shape, and frontend support in the MoE router guide.

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

* Clarify top-k index buffer requirements

Document the CUDA device, contiguous layout and supported integer dtypes, show an allocation example, and clarify the routing map format restriction.

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

* Stop documenting standalone JAX token dispatch and combine

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

* Clarify that integrated JAX MoE requires multiple GPUs

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

* Fix lint and documentation tooling compatibility

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

* Install wheel for Python lint setup analysis

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

* docs: update grouped MLP alignment and weighting examples

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

* docs: use eager receive sizing in PyTorch EP example

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

* docs: normalize JAX grouped linear split sizes to int32

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

* docs: use nonsquare expert weights in JAX MoE example

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

---------

Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
* Validate grouped split alignment on device

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

* Use versioned offset scan API for split validation

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

* Consolidate grouped split alignment replay test

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

* Simplify device alignment test and isolate L0 invocation

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

* Move split alignment test into grouped tensor suite

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

* Reduce grouped split alignment scan overhead

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

* Test grouped split alignment during graph replay

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

* Cover both grouped split scan sizes

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

* Remove CUDA context poisoning alignment test

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

* Apply suggestion from @timmoon10

Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>

* Apply suggestion from @timmoon10

Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>

---------

Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>
Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
* [PyTorch] Skip FA2 in L3 tests on archs without native FA2 support

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

* [PyTorch] Build FA2 for the family base arch instead of skipping

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

---------

Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>
[JAX] Disable CuTeDSL at import time

Signed-off-by: Jeremy Berchtold <jberchtold@nvidia.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Disable TE torch compilation in KV-cache CI tests to avoid numerical mismatches

Signed-off-by: Kshitij Lakhani <klakhani@nvidia.com>
… Logic (NVIDIA#3613)

* Reduce module GroupedLinear CPU preparation overhead

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

* reduce cpu overhead for op fuser

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

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

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

* some fixes

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

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

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

* Fix GroupedLinear bias gradients with delayed WGRAD

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

* Fix GroupedLinear bias gradients with frozen weights

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

* resolve comments

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

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

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

* resolve comments

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

* grep tile comments

Signed-off-by: Zhongbo Zhu <zhongboz@nvidia.com>

* Remove quantizer caching in grouped linear op

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

* Test grouped linear across training and inference

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

* Refactor grouped linear module implementation

Disentangle backend selection from quantizer extraction. Remove overly-specific helper function for backend selection. Remove test that calls internal helper functions.

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

* Make quantized single grouped weights own their quantizer

Make sure to set column-wise usage in quantized weight param quantizer during forward pass.

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

---------

Signed-off-by: zhongboz <zhongboz@nvidia.com>
Signed-off-by: Zhongbo Zhu <zhongboz@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>
Co-authored-by: Tim Moon <tmoon@nvidia.com>
…#3587)

* initial draft

Signed-off-by: Zhongbo Zhu <zhongboz@nvidia.com>
Signed-off-by: zhongboz <zhongboz@nvidia.com>

* Align deterministic MoE aux-loss workspace queries with the common API

Query workspace shape and dtype through the deterministic forward APIs, retain stream and CUDA Graph isolation, and cover metadata queries and invalid scratch buffers. Fix the stream cache-key lint warning and upstream import ordering.

Validated on GB200: 201 TE tests passed (12 invalid top-k cases skipped), 6 C API tests passed, and 40 MCore checks passed with 1 expected atomic negative-control XPASS. Full PyTorch QA lint, pre-commit, and license checks passed.

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

* Test deterministic MoE aux-loss integer boundaries

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

* Simplify deterministic MoE aux-loss workspace contract

Apply Varun Thumbe's workspace-contract changes from
zhongbozhu#1. Require the queried 256-element FP32
capacity for every input shape and update the C API workspace tests.
Retain the existing MLA import grouping required by Python lint.

Original author: Varun Thumbe <vthumbe@nvidia.com>
Source: zhongbozhu@9a3e764

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

---------

Signed-off-by: Zhongbo Zhu <zhongboz@nvidia.com>
Signed-off-by: zhongboz <zhongboz@nvidia.com>
Signed-off-by: Varun Thumbe <vthumbe@nvidia.com>
Co-authored-by: Varun Thumbe <vthumbe@nvidia.com>
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
…NVIDIA#3760)

* Use v2 EP ops with no handle cache in PyTorch and mark the cache and non-v2 ops for removal in TE 3.0

* Disable EP handle cache by default and rename topk_idx dtype helpers

Signed-off-by: Phuong Nguyen <phuonguyen@nvidia.com>

---------

Signed-off-by: Phuong Nguyen <phuonguyen@nvidia.com>
* [Common] Add env var to let C++ UT run in parallel to debug

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

* cat log if C++ UT fails so we can see the failed tests

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

* Update qa/L0_cppunittest/test.sh

Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* Update qa/L1_cppunittest_cutedsl_off/test.sh

Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

* Update tests/cpp/operator/CMakeLists.txt

Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Kaining Zhong <kainingz@nvidia.com>

---------

Signed-off-by: Kaining Zhong <kainingz@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Accumulate delayed wgrads into shared (tied) parameters

backward_dw() assigned the popped delayed wgrad to Parameter.grad. When two
modules share one Parameter, each pops its own contribution and the later one
silently replaced the earlier, so the shared gradient came out too small -- on
the issue's own 16x16 / batch=2 BF16 repro it was 0.6015x the regular path, with
256/256 elements differing.

The bias takes the same pop path and hits the same bug, with an extra twist: the
second module's bgrad is numerically ~0, so it wiped out a contribution the first
module had already written correctly (1.0984 delivered vs 33.0859 expected).

Initialise on the first contribution of the step and add on every later one, for
both the weight and the bias. dtype/shape/buffer-alias contracts are unchanged,
and the fused wgrad-accumulation branch is untouched.

tests/pytorch/test_tied_delayed_wgrad.py covers the matrix: tied and non-tied,
both backward_dw() orders, multiple microbatches, two optimizer steps with
set_to_none True/False, the bias case, and parity with a plain nn.Linear oracle.
It uses only the public te.Linear API so it cannot be satisfied by a shared
implementation detail of the fix. Without the fix it fails 5 and passes 3; with
it, 8 pass.

Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

* Clarify the accumulate path for zero_grad(set_to_none=False)

Review question: is the `grad is None` first-contribution test safe when the optimizer was
built with `zero_grad(set_to_none=False)`? It is: that setting leaves a zeroed tensor where
`set_to_none=True` leaves `None`, so the first contribution takes the add path and lands on
the same value, and nothing in the branch depends on `grad` being `None`. Say so where the
branch is, and pin it in the test.

`test_two_optimizer_steps_do_not_pollute_each_other` already ran both settings, but
`second != first` cannot show pollution - a first-step value carried into the second still
differs from the first. It now compares the second step's gradient against a fresh model on
the same weights and input, bit-exactly.

Verified with the new test over an out-of-tree copy of an installed TransformerEngine:
5 failed, 3 passed against the unfixed `backward_dw` (the same five this PR describes), and
8 passed with this PR's `backward_dw` in place.

Signed-off-by: 0z5a <192209249+0z5a@users.noreply.github.com>
Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

* Address review feedback on the tied-wgrad test file

- `_tie(model, *, bias)` and `_build(delay, *, bias=False, device="cuda")`: the boolean
  arguments are keyword-only, so a call site says what it means.
- `_build` only builds. Tying moved out to the call sites, which is where a test decides
  whether it wants a tied pair or two independent ones.
- `order` is now `backward_dw_order`, and the docstring says what it orders.
- The module-level `DEVICE` and tolerance globals are gone: the device is a default on the
  two helpers that need it, and the multi-microbatch tolerance is local to the test that
  needs it.

Signed-off-by: 0z5a <192209249+0z5a@users.noreply.github.com>
Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

* fix(pytorch): accumulate all delayed tied-parameter gradients

Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

* Own delayed gradients during CUDA Graph replay

Capture deferred contributions in the graph pool and accumulate them into caller-owned gradients before delayed hooks. Cover both capture orders, default and explicit gradient resets, tied weights, microbatches, bias, retained gradients, and fused main_grad with independent PyTorch references. Native GPU validation is pending on the coordinated test machine.

Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

* Keep unfilled delayed grouped wgrads away from autograd

With delay_wgrad_compute=True and fuse_wgrad_accumulation=False,
GroupedLinear returned the still-empty wgrad buffer to autograd from the
main backward pass. Autograd stored it as (or added it into) .grad before
backward_dw() had written it, so accumulating the delayed result on top
read uninitialized memory. Return None instead, as Linear does, and let
backward_dw() own .grad.

The tied-gradient tests compared multi-term BF16 sums from the delayed
and regular paths with a fixed 1e-2 tolerance. The two paths add the same
terms in a different order, so allow a few BF16 roundings at the
gradient's magnitude. The graphed microbatch test compared against a
PyTorch FP32 reference bitwise, but TE runs FP32 GEMMs in TF32; compare
graph replay against the same TE modules run eagerly instead.

Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

* Iterate the non-tied test's module pairs with zip

Pair each delayed Linear with its regular counterpart directly instead of
indexing both Sequentials, as suggested in review.

Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

* [PyTorch] Use 128-token splits in the tied te.ops GroupedLinear test

te.ops.GroupedLinear now validates on device that every split is a
multiple of split_alignment (default 128), so 16-token splits trap in
the forward pass.

Signed-off-by: Alp Dener <adener@nvidia.com>

---------

Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>
Signed-off-by: 0z5a <192209249+0z5a@users.noreply.github.com>
Signed-off-by: Alp Dener <adener@nvidia.com>
Co-authored-by: 0z5a <192209249+0z5a@users.noreply.github.com>
Co-authored-by: Alp Dener <adener@nvidia.com>
* Fix LayerNormLinear dummy gradients for sharded weights

Use the original parameter shape for both fused-DDP dummy-gradient branches. Distributed weights can expose a full GEMM gradient buffer even though autograd expects a gradient matching the local parameter shard.

Add standalone LayerNorm/RMSNorm regression coverage for shard and full-weight shapes, both dummy branches, and two-microbatch gradient accumulation. Include the regression in the PyTorch L0 launcher.

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

* Address LayerNormLinear dummy-gradient review feedback

Keep the dummy-gradient explanation focused on the original autograd input. Consolidate LayerNormLinear coverage in test_distributed_weight.py and retain the eager test settings on its existing L0 launcher entry.

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

---------

Signed-off-by: zhongboz <zhongboz@nvidia.com>
…ed MLP (NVIDIA#3641)

* [PyTorch] Pass canonical layouts to cuDNN grouped GEMM in fused grouped MLP

When the cuDNN frontend grouped GEMM GLU and quant wrappers advertise
supports_canonical_layouts, the MXFP8 forward passes A as (sum(m), k),
B as (num_groups, n, k), probabilities as 1-D, and the packed scale
buffers as-is, and consumes the 2-D outputs directly. This drops the
per-call view/permute chains on the host. NVFP4, SReLU, and older
frontends keep the existing layouts.

The FC2 quant call also skips the unused output amax when the wrapper
advertises supports_optional_amax.

Signed-off-by: Hwanseo Choi <hwanseoc@nvidia.com>

* [PyTorch] Test and document canonical cuDNN layouts in fused grouped MLP

Check that natural operand layouts reach the cuDNN frontend and give bitwise-identical outputs and gradients to the legacy views, and note the layout behavior in the fused op docstring.

Signed-off-by: Hwanseo Choi <hwanseoc@nvidia.com>

* [PyTorch] Pass canonical layouts to cuDNN grouped GEMM in fused grouped MLP backward

When the cuDNN frontend dGLU wrapper advertises supports_canonical_layouts,
the MXFP8 backward passes 2-D grad output, activation input, and FC1 dgrad
operands, 1-D probabilities, packed scale buffers as-is, and the column-wise
weights as (num_groups, k, n) with b_major="n". It consumes the 2-D dGLU
outputs directly. wgrad passes flat scale buffers when the wgrad wrapper
supports them. NVFP4, SReLU, and older frontends keep the existing layouts.

Signed-off-by: Hwanseo Choi <hwanseoc@nvidia.com>

* [PyTorch] Test and document canonical cuDNN layouts in fused grouped MLP backward

Extend the canonical-vs-legacy test to the dGLU and FC1 dgrad calls, and note the backward in the fused op docstring.

Signed-off-by: Hwanseo Choi <hwanseoc@nvidia.com>

* [PyTorch] Keep the typed cuDNN weight views on the grouped weight

The FC1 and FC2 weight buffers persist across steps, so the dtype and shape
views cuDNN takes are built once per buffer instead of with two torch ops
before each grouped GEMM launch (forward FC1 and FC2, backward FC2 dGLU and
FC1 dgrad). A changed buffer is detected by its data pointer.

Signed-off-by: Hwanseo Choi <hwanseoc@nvidia.com>

* [PyTorch] Simplify the cached cuDNN weight view helper

cudnn_weight_data takes the grouped weight and reads its row- or column-wise
buffer itself, so call sites pass only the dtype and shape.

Signed-off-by: Hwanseo Choi <hwanseoc@nvidia.com>

* [PyTorch] Drop the cached cuDNN weight views

Build the typed weight views on every call again instead of keeping them
on the grouped weight. The cache saved about 5 us per step and risked
stale state.

Signed-off-by: Hwanseo Choi <hwanseoc@nvidia.com>

* [PyTorch] Decide cuDNN canonical layouts after choosing the kernel

The forward chooses the FC1 kernel first and passes canonical layouts only
to the MXFP8 GEMM + activation kernel; NVFP4's GEMM + activation + RHT
kernels keep the kernel-facing views. The backward decides with the dGLU
wrapper it calls. Nothing is decided in __init__ anymore.

Signed-off-by: Hwanseo Choi <hwanseoc@nvidia.com>

* [PyTorch] Drop the canonical-vs-legacy layout test

Existing grouped MLP tests cover the numerics, and the legacy layouts are
cuDNN frontend's to test.

Signed-off-by: Hwanseo Choi <hwanseoc@nvidia.com>

---------

Signed-off-by: Hwanseo Choi <hwanseoc@nvidia.com>
…ps (NVIDIA#3785)

* Fix grouped linear deferred-init test split alignment

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

* Skip grouped split alignment check for unquantized compute

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

* Reject negative grouped splits at any alignment

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

* Update transformer_engine/common/util/splits_to_offsets.cu

Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: vthumbe1503 <vthumbe@nvidia.com>

* Resolve grouped linear split alignment by recipe

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

* Require aligned FP8 current splits on Hopper

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

* Check discrete grouped GEMM input alignment

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: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
* [PyTorch] Add BF16 fused grouped MLP support

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

* Pass packed BF16 weights to the dgrad GEMMs as dense N-major operands and allocate the wgrad buffer directly

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

* Extend the BF16 fused grouped MLP to GeGLU, distributed weights, and FC2 scale_bias

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

* Review feedback fixes

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

---------

Signed-off-by: tingyangk <tingyangk@nvidia.com>
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
…h-combine

Resolved docs/api/pytorch.rst: upstream added an "Expert parallelism"
section that already autodocs EpConfig/MoeDispatch/MoeCombine, so the
PR's EP=1 local-backend paragraph moves there and the duplicate
autoapiclass entries are dropped.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

This branch has not been deployed

No deployments
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.