is NVFP4 not supported for rtx 50 series?

Hi
I locally compiled branch release 2.8.

when I tried to use nvfp4 on rtx 50 series it gave me error

/home/aza/workspace/projects/nvfp4/TransformerEngine/transformer_engine/common/util/nvfp4_transpose.cuh:234 in function mul_cvt_bf16_to_fp4_4x_with_rn (thread (95,0,0), block (2,2,0)): FP4 cvt PTX instructions are architecture-specific. Try recompiling with sm_XXXa instead of sm_XXX.

@yash3056

Hello @yash3056. The default NVFP4 recipe uses stochastic rounding, which is accelerated only on SM 100 and SM 103 (B100/B200/B300). You should be able to disable stochastic rounding and try to compile for the sm120a architecture as the round-to-nearest mode is supported there.
@Oleg-Goncharov we should add a fallback path for the stochastic rounding on non-10x cards.

Actually, looking at the error message the kernel which failed is the round to nearest kernel (the stochastic rounding will fail too but it did not get there yet). This is because we check for the SM 10X features, while we should be checking for 10X and 12X there. I can open a quick PR tomorrow to fix that.
In the meantime you can try removing the condition here: https://github.com/NVIDIA/TransformerEngine/blob/main/transformer_engine/common/util/nvfp4_transpose.cuh#L201 (just make sure you compile for SM120a).

@ptrendx

@ptrendx will disabling stochastic rounding and using round-to-nearest mode will not cause problem with model performance to degrade?

btw can you tell me how to compile for SM120a? I am trying to use this flag export CMAKE_CUDA_ARCHITECTURES="120a"
but it is not building for sm_120a

@yash3056

Hello @yash3056, stochastic rounding does impact the model performance, but its contribution depends on the model size, quantization granularity, and whether other techniques are used (e.g., the random Hadamard transform, which is currently the part of the NVFP4 recipe).
To build the TE for sm_120a, please try export NVTE_CUDA_ARCHS=120a.

I removed condition from
https://github.com/NVIDIA/TransformerEngine/blob/main/transformer_engine/common/util/nvfp4_transpose.cuh#L201
and compiled it using 120a cm flag but

I am getting this error

Error: Failed to set Shared Memory size.

most probable cause is difference between architecture as such it needs a fallback for rtx 50 series, just removing condition will not make it work.

to replicate it

CODE:

import torch
import transformer_engine.pytorch as te
from transformer_engine.common.recipe import Format, NVFP4BlockScaling

# Setup
device = torch.device("cuda")
fp4_format = Format.E2M1
fp4_recipe = NVFP4BlockScaling(fp4_format=fp4_format)

print("NVFP4 Matrix Multiplication Test")
print(f"Format: {fp4_format}")
print("Recipe: NVFP4BlockScaling\n")

# Create input
torch.manual_seed(42)
batch_size, seq_len, hidden_dim = 4, 128, 768
x = torch.randn(batch_size, seq_len, hidden_dim, dtype=torch.bfloat16).cuda()

# Create linear layer
linear = te.Linear(hidden_dim, hidden_dim, bias=True, params_dtype=torch.bfloat16).cuda()

print(f"Input shape: {x.shape}")
print(f"Input dtype: {x.dtype}")
print(f"Weight shape: {linear.weight.shape}\n")

# BF16 baseline
with torch.no_grad():
    out_bf16 = linear(x)

# NVFP4 forward pass
with te.fp8_autocast(enabled=True, fp8_recipe=fp4_recipe):
    out_fp4 = linear(x)

print(f"Output shape: {out_fp4.shape}")
print(f"BF16 output sample: {out_bf16[0, 0, :5]}")
print(f"FP4 output sample:  {out_fp4[0, 0, :5]}\n")

# Error metrics
rel_error = torch.abs(out_bf16 - out_fp4) / (torch.abs(out_bf16) + 1e-5)
print(f"Mean relative error: {rel_error.mean().item():.6f}")
print(f"Max relative error: {rel_error.max().item():.6f}")

Error:

NVFP4 Matrix Multiplication Test
Format: Format.E2M1
Recipe: NVFP4BlockScaling

Input shape: torch.Size([4, 128, 768])
Input dtype: torch.bfloat16
Weight shape: torch.Size([768, 768])

Error: Failed to set Shared Memory size.
Traceback (most recent call last):
  File "/home/aza/workspace/projects/nvfp4/TransformerEngine/../test.py", line 32, in <module>
    out_fp4 = linear(x)
  File "/home/aza/miniforge3/envs/transformer-engine/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1777, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "/home/aza/miniforge3/envs/transformer-engine/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1788, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/aza/miniforge3/envs/transformer-engine/lib/python3.13/site-packages/torch/_dynamo/eval_frame.py", line 1098, in _fn
    return fn(*args, **kwargs)
  File "/home/aza/miniforge3/envs/transformer-engine/lib/python3.13/site-packages/transformer_engine/pytorch/module/linear.py", line 1482, in forward
    out = linear_fn(*args)
  File "/home/aza/miniforge3/envs/transformer-engine/lib/python3.13/site-packages/torch/autograd/function.py", line 582, in apply
    return super().apply(*args, **kwargs)  # type: ignore[misc]
           ~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "/home/aza/miniforge3/envs/transformer-engine/lib/python3.13/site-packages/transformer_engine/pytorch/module/linear.py", line 254, in forward
    weightmat = module.get_weight_workspace(
        tensor=weight,
    ...<5 lines>...
        workspace_dtype=activation_dtype,
    )
  File "/home/aza/miniforge3/envs/transformer-engine/lib/python3.13/site-packages/transformer_engine/pytorch/module/base.py", line 1428, in get_weight_workspace
    out = quantizer.quantize(tensor, dtype=workspace_dtype)
  File "/home/aza/miniforge3/envs/transformer-engine/lib/python3.13/site-packages/transformer_engine/pytorch/tensor/quantized_tensor.py", line 203, in quantize
    return _QuantizeFunc.forward(None, tensor, self)
           ~~~~~~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^
  File "/home/aza/miniforge3/envs/transformer-engine/lib/python3.13/site-packages/transformer_engine/pytorch/tensor/quantized_tensor.py", line 282, in forward
    return tex.quantize(tensor, quantizer)
           ~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^
torch.AcceleratorError: CUDA error: invalid argument
Search for `cudaErrorInvalidValue' in https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html for more information.
CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.
For debugging consider passing CUDA_LAUNCH_BLOCKING=1
Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.

Ok, this is actually a slightly larger can of worms than I anticipated (in addition to the problem you flagged there are random things like CMake<4.0.2 not understanding "f" in the SM arch and erroring out 😞). Will continue looking at it on Monday.

Ok, this is actually a slightly larger can of worms than I anticipated (in addition to the problem you flagged there are random things like CMake<4.0.2 not understanding "f" in the SM arch and erroring out 😞). Will continue looking at it on Monday.

Some ptx is not supported on SM_120, such as 'cvt.rs.satfinite.e2m1x4.f32' 🥲
https://docs.nvidia.com/cuda/parallel-thread-execution/#:~:text=sm_103a-,cvt,-.rs%7B.e2m1x4/.e4m3x4

rtx 50 series instruction are closer to ada series gpu

@yash3056

The PR #2279 should help to move this forward a little bit - there is still a problem with the cvt.rs instructions (we need to add emulation for those unfortunately and that is not part of this PR), but this should hopefully enable you to run when setting the recipe to not do the stochastic rounding. Could you try @yash3056?

@ptrendx

@ptrendx the nvfp4 matmul is working now.

Here is the update code (the change in it is only torch.no_grad() added during inference)

import torch
import transformer_engine.pytorch as te
from transformer_engine.common.recipe import Format, NVFP4BlockScaling

# NVFP4 Matrix Multiplication Test
# This test demonstrates NVFP4 quantization for linear layers
# Key requirements:
# - Dimensions should be multiples of 16 for optimal NVFP4 block scaling (16x16 blocks)
# - fp8_autocast context must wrap all FP4 operations
# - Use torch.no_grad() when not training to avoid gradient computation

# Setup
device = torch.device("cuda")
fp4_format = Format.E2M1
fp4_recipe = NVFP4BlockScaling(fp4_format=fp4_format)

print("NVFP4 Matrix Multiplication Test")
print(f"Format: {fp4_format}")
print("Recipe: NVFP4BlockScaling\n")

# Create input - dimensions should be multiples of 16 for NVFP4 block scaling
torch.manual_seed(42)
batch_size, seq_len, hidden_dim = 4, 128, 768  # 768 is divisible by 16
x = torch.randn(batch_size, seq_len, hidden_dim, dtype=torch.bfloat16).cuda()

# Create linear layer
linear = te.Linear(hidden_dim, hidden_dim, bias=True, params_dtype=torch.bfloat16).cuda()

print(f"Input shape: {x.shape}")
print(f"Input dtype: {x.dtype}")
print(f"Weight shape: {linear.weight.shape}\n")

# BF16 baseline
with torch.no_grad():
    out_bf16 = linear(x)

# NVFP4 forward pass - wrap all FP4 operations in the context
print("Running NVFP4 forward pass...")
with torch.no_grad():
    with te.fp8_autocast(enabled=True, fp8_recipe=fp4_recipe):
        out_fp4 = linear(x)

print(f"Output shape: {out_fp4.shape}")
print(f"BF16 output sample: {out_bf16[0, 0, :5]}")
print(f"FP4 output sample:  {out_fp4[0, 0, :5]}\n")

# Error metrics
abs_error = torch.abs(out_bf16 - out_fp4)
rel_error = abs_error / (torch.abs(out_bf16) + 1e-5)

print(f"Mean absolute error: {abs_error.mean().item():.6f}")
print(f"Max absolute error: {abs_error.max().item():.6f}")
print(f"Mean relative error: {rel_error.mean().item():.6f}")
print(f"Median relative error: {rel_error.median().item():.6f}")

# Also compute RMSE for a more stable metric
rmse = torch.sqrt(torch.mean((out_bf16 - out_fp4) ** 2))
print(f"RMSE: {rmse.item():.6f}")

# Signal-to-noise ratio
signal_power = torch.mean(out_bf16 ** 2)
noise_power = torch.mean((out_bf16 - out_fp4) ** 2)
snr = 10 * torch.log10(signal_power / noise_power)
print(f"SNR: {snr.item():.2f} dB")

Output:

NVFP4 Matrix Multiplication Test
Format: Format.E2M1
Recipe: NVFP4BlockScaling

Input shape: torch.Size([4, 128, 768])
Input dtype: torch.bfloat16
Weight shape: torch.Size([768, 768])

Running NVFP4 forward pass...
Output shape: torch.Size([4, 128, 768])
BF16 output sample: tensor([-0.3867, -0.0835,  0.3047, -0.9023, -0.4121], device='cuda:0',
       dtype=torch.bfloat16)
FP4 output sample:  tensor([-0.3223, -0.1543,  0.3105, -0.8945, -0.4609], device='cuda:0',
       dtype=torch.bfloat16)

Mean absolute error: 0.073730
Max absolute error: 0.478516
Mean relative error: 1.031250
Median relative error: 0.146484
RMSE: 0.092285
SNR: 16.75 dB

in short it is working

@ptrendx nvfp4 matmul 现已开始工作。

这是更新代码（其中的变化仅在推理期间添加了 torch.no_grad()）

import torch
import transformer_engine.pytorch as te
from transformer_engine.common.recipe import Format, NVFP4BlockScaling

# NVFP4 Matrix Multiplication Test
# This test demonstrates NVFP4 quantization for linear layers
# Key requirements:
# - Dimensions should be multiples of 16 for optimal NVFP4 block scaling (16x16 blocks)
# - fp8_autocast context must wrap all FP4 operations
# - Use torch.no_grad() when not training to avoid gradient computation

# Setup
device = torch.device("cuda")
fp4_format = Format.E2M1
fp4_recipe = NVFP4BlockScaling(fp4_format=fp4_format)

print("NVFP4 Matrix Multiplication Test")
print(f"Format: {fp4_format}")
print("Recipe: NVFP4BlockScaling\n")

# Create input - dimensions should be multiples of 16 for NVFP4 block scaling
torch.manual_seed(42)
batch_size, seq_len, hidden_dim = 4, 128, 768  # 768 is divisible by 16
x = torch.randn(batch_size, seq_len, hidden_dim, dtype=torch.bfloat16).cuda()

# Create linear layer
linear = te.Linear(hidden_dim, hidden_dim, bias=True, params_dtype=torch.bfloat16).cuda()

print(f"Input shape: {x.shape}")
print(f"Input dtype: {x.dtype}")
print(f"Weight shape: {linear.weight.shape}\n")

# BF16 baseline
with torch.no_grad():
    out_bf16 = linear(x)

# NVFP4 forward pass - wrap all FP4 operations in the context
print("Running NVFP4 forward pass...")
with torch.no_grad():
    with te.fp8_autocast(enabled=True, fp8_recipe=fp4_recipe):
        out_fp4 = linear(x)

print(f"Output shape: {out_fp4.shape}")
print(f"BF16 output sample: {out_bf16[0, 0, :5]}")
print(f"FP4 output sample:  {out_fp4[0, 0, :5]}\n")

# Error metrics
abs_error = torch.abs(out_bf16 - out_fp4)
rel_error = abs_error / (torch.abs(out_bf16) + 1e-5)

print(f"Mean absolute error: {abs_error.mean().item():.6f}")
print(f"Max absolute error: {abs_error.max().item():.6f}")
print(f"Mean relative error: {rel_error.mean().item():.6f}")
print(f"Median relative error: {rel_error.median().item():.6f}")

# Also compute RMSE for a more stable metric
rmse = torch.sqrt(torch.mean((out_bf16 - out_fp4) ** 2))
print(f"RMSE: {rmse.item():.6f}")

# Signal-to-noise ratio
signal_power = torch.mean(out_bf16 ** 2)
noise_power = torch.mean((out_bf16 - out_fp4) ** 2)
snr = 10 * torch.log10(signal_power / noise_power)
print(f"SNR: {snr.item():.2f} dB")

输出：

NVFP4 Matrix Multiplication Test
Format: Format.E2M1
Recipe: NVFP4BlockScaling

Input shape: torch.Size([4, 128, 768])
Input dtype: torch.bfloat16
Weight shape: torch.Size([768, 768])

Running NVFP4 forward pass...
Output shape: torch.Size([4, 128, 768])
BF16 output sample: tensor([-0.3867, -0.0835,  0.3047, -0.9023, -0.4121], device='cuda:0',
       dtype=torch.bfloat16)
FP4 output sample:  tensor([-0.3223, -0.1543,  0.3105, -0.8945, -0.4609], device='cuda:0',
       dtype=torch.bfloat16)

Mean absolute error: 0.073730
Max absolute error: 0.478516
Mean relative error: 1.031250
Median relative error: 0.146484
RMSE: 0.092285
SNR: 16.75 dB

简而言之，它正在发挥作用

well,what about fp4 fuse attention？or te.fp8_model_init()?

@yash3056

@yash3056 Were you able to get this working for the backward pass also, or only the forward pass?

I've been able to build off the main branch and get the forward pass working but the backward pass is showing an unsupported GEMM.

Stacktrace in details.

Details

Traceback (most recent call last):
  File "/app/minimal_example.py", line 121, in <module>
    main()
  File "/app/minimal_example.py", line 107, in main
    loss.backward()
  File "/usr/local/lib/python3.12/dist-packages/torch/_tensor.py", line 625, in backward
    torch.autograd.backward(
  File "/usr/local/lib/python3.12/dist-packages/torch/autograd/__init__.py", line 354, in backward
    _engine_run_backward(
  File "/usr/local/lib/python3.12/dist-packages/torch/autograd/graph.py", line 829, in _engine_run_backward
    return Variable._execution_engine.run_backward(  # Calls into the C++ engine to run the backward pass
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/usr/local/lib/python3.12/dist-packages/torch/autograd/function.py", line 315, in apply
    return user_fn(self, *args)
           ^^^^^^^^^^^^^^^^^^^^
  File "/usr/local/lib/python3.12/dist-packages/transformer_engine/pytorch/module/linear.py", line 895, in backward
    wgrad, grad_bias_ = wgrad_gemm(inputmat_total, grad_output)
                        ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/usr/local/lib/python3.12/dist-packages/transformer_engine/pytorch/module/linear.py", line 875, in wgrad_gemm
    dw, db, *_ = general_gemm(x, dy, **wgrad_gemm_kwargs)
                 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/usr/local/lib/python3.12/dist-packages/transformer_engine/pytorch/cpp_extensions/gemm.py", line 149, in general_gemm
    out, bias_grad, gelu_input, extra_output = tex.generic_gemm(*args, **kwargs)
                                               ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
RuntimeError: /app/TransformerEngine/transformer_engine/common/gemm/cublaslt_gemm.cu:744 in function cublas_gemm: Assertion failed: status != CUBLAS_STATUS_NOT_SUPPORTED. Unable to find suitable cuBLAS GEMM algorithm

@ptrendx

for anyone wondering, this is how you can get NVFP4BlockScaling to work on SM120:

recipe = NVFP4BlockScaling(disable_rht= True ,disable_stochastic_rounding= True )

@ptrendx any progress on this?

@alint77

@alint77 isn't stochastic rounding required for FP4 training?

"stochastic rounding ensures that gradients are rounded up or down randomly, with probabilities proportional to how close a number lies between two representable values. This step is essential for reducing rounding bias, maintaining gradient flow during training, and ultimately improving model accuracy."

https://developer.nvidia.com/blog/nvfp4-trains-with-precision-of-16-bit-and-speed-and-efficiency-of-4-bit/

Is the team not planning to implement stochastic rounding for our SM120 blackwell cards?

@vgoklani

@vgoklani yeah I also thought it's mandatory but turns out it isn't and the model actually does converge although I haven't tried with model_init set to nvfp4.

you're welcome to try it on my nanogpt-fp8 repo

@alint77

@alint77 thanks for the followup, Is there a timeline to implement stochastic rounding etc for the SM120 blackwell series? We are specifically interested in nvfp4 and will use this with 4x RTX Blackwell MaxQ cards. For your nanogpt-fp8-repo is there a baseline that we can compare to? When you say the model actually does "converge" - how does that compare to a similar model trained on an H200? We want to get these things sorted out before we start the training runs, as those are very time-consuming. Thanks!

@alint77 isn't stochastic rounding required for FP4 training?

"stochastic rounding ensures that gradients are rounded up or down randomly, with probabilities proportional to how close a number lies between two representable values. This step is essential for reducing rounding bias, maintaining gradient flow during training, and ultimately improving model accuracy."

https://developer.nvidia.com/blog/nvfp4-trains-with-precision-of-16-bit-and-speed-and-efficiency-of-4-bit/

Is the team not planning to implement stochastic rounding for our SM120 blackwell cards?

The stochastic rounding is a SM101/SM110 hardware/cuda feature, pytorch have discussed a extra step or similar to implement something equivalent but slower since it will inherently require a higher precision to determine the rounding fudge. Nvidia pretty obviously, semi-brilliantly, semi-evil, figured out a way to create a hampered inference card (RTX Pros) that don't do training well so as to not hurt the DC training products (beyond the HBM).

for anyone wondering, this is how you can get NVFP4BlockScaling to work on SM120:

recipe = NVFP4BlockScaling(disable_rht= True ,disable_stochastic_rounding= True )

@ptrendx any progress on this?

Just wanted to say that this was crucial to getting nvfp4 working on the Nvida DGX Spark device

@Mocuto

@Mocuto sure - but is training stable when you disable those flags? How are the gradients?

For anyone interested, I vibed a polyfill for cvt.rs.satfinite.e2m1x4.f32 on SM120 family (https://github.com/Harry-Chen/fp4_sm120) to mimic the instruction behavior, which adds ~100 PTX instructions for each FP4x4 rounding.

The code has some proof-of-concept self-tests. We are also trying to use it in our training code. Feel free to try it out but please bear in mind that it is not tested thoroughly.

@Harry-Chen

@Harry-Chen thanks for sharing - do you have any examples of using the polyfill for training, i'm missing a step. thanks!

@vgoklani

@vgoklani For example, when you use NVFP4 without disable_stochastic_rounding on SM120 family devices, you will hit the following error (and in other functions too):

TransformerEngine/transformer_engine/common/util/ptx.cuh

Lines 596 to 599 in e879bf8

    
             NVTE_DEVICE_ERROR( 
        
                 "FP4 cvt PTX instructions are architecture-specific. " 
        
                 "Try recompiling with sm_XXXa instead of sm_XXX."); 
        
           }

You could replace those with_stochastic_rounding functions in this file with the polyfill I provided and recompile TE. Then SR should work out-of-box. I'm still working RHT GEMM kernels.

Tested both upstream-suggested fixes on consumer Blackwell sm_120 (RTX 5080):

TE rebuild with TORCH_CUDA_ARCH_LIST=12.0a (closing the gap from cycle Check for cuDNN frontend API when building #307's NVTE_CUDA_ARCHS=120a which only covered TE's own NVCC paths):
- Build success; 61 NVCC compile rules emit --generate-code=arch=compute_120a,code=[compute_120a,sm_120a].
- NVFP4BlockScaling() defaults still RED with same error: RuntimeError: row_col_rht_gemm_ntt_w_sfc.cu:1200: CUDA Error: invalid argument.
- This confirms the layer-2 hardware-feature blocker (TMEM + UMMA absence on sm_120 silicon) is not patchable via TORCH_CUDA_ARCH_LIST=12.0a flag fix or CUTLASS DSL whitelist patch ([JAX] Warmup FFIs with "initialize" stage #2800 admissible_archs += sm_120a).
- Recipe-level workaround NVFP4BlockScaling(disable_rht=True, disable_stochastic_rounding=True) (PR Overhaul the compilation for the arch-specific features #2279) works at 1.77× bf16 (170.64 TFLOPS @ M=1024, K=N=4096).
- Harry-Chen polyfill (github.com/Harry-Chen/fp4_sm120, Apache 2.0, pinned SHA 36de459) restores SR via software emulation: 1.16× bf16 (111.14 TFLOPS) — first honest measurement (cycle Check for cuDNN frontend API when building #307's 1.83× claim was a false positive due to namespace-package collision masking the polyfill code path).
- Note: The 35% throughput cost vs no-SR workaround reflects ~100 PTX ops per FP4×4 conversion in Harry-Chen's polyfill emulating the missing cvt.rs.satfinite.e2m1x4.f32 instruction.
PyTorch 2.13 from-source build with TORCH_CUDA_ARCH_LIST=12.0a:
- torch.cuda.get_arch_list() returns ['sm_120a'] confirming gencode honored.
- Cast torch.bf16 → torch.float4_e2m1fn_x2 STILL fails identically.
- Source inspection of aten/src/ATen/native/cuda/Copy.cu:243-246 reveals the cast is deliberately stubbed for same-dtype-only: gpu_kernel_nocast + TORCH_CHECK(dtype == iter.dtype(1)) enforces float4→float4 cast only.
- Source inspection of aten/src/ATen/native/cuda/ScaledBlas.cpp:529-532 reveals the FP4 _scaled_mm GEMM dispatch is wrapped in #ifdef USE_ROCM + TORCH_CHECK(ROCM_VERSION >= 70000) — the CUDA-side FP4 GEMM dispatch path does not exist. AMD MI300X (ROCm 7.0+) has it; CUDA on any Blackwell variant does not.
- Two layers missing on CUDA (cast op + GEMM dispatch). The cutlass sm120_blockscaled_mma_* primitives ARE vendored at third_party/mslk/external/cutlass/include/cutlass/gemm/collective/ but not consumed by any CUDA dispatch path.

Combined evidence: NVIDIA's FP4 on consumer Blackwell sm_120 is a multi-vendor scaffolding gap.

TE NVFP4 has sm_100/B100 hardware-feature dependency (TMEM/UMMA absent on sm_120) — TE upstream needs sm_120-specific kernel port using HMMA/QMMA + cp.async + 99 KB shared memory budget instead of TMEM/UMMA + 232 KB.
CUTLASS [JAX] Warmup FFIs with "initialize" stage #2800/Refactor Amax Kernel ldmatrix loads, TMA/compute barriers, swizzle_idx #2820 needs sm_120a in admissible_archs — patch is 3 LOC (Python DSL whitelist) + sm_120-aware kernel impl beneath.
PyTorch needs cast op + GEMM dispatch wired for CUDA — Meta has the implementation for ROCm 7.0; CUDA wiring requires unfreezing the #ifdef USE_ROCM gate.
cuBLASLt 13.4.1.1 has no FP4 enums — CUDA toolkit 13.5+ rumored to add them.

Suggestion (concrete, actionable): an sm_120-targeted Hadamard kernel port in TransformerEngine using HMMA/QMMA + cp.async + 99 KB smem (instead of TMEM/UMMA + 232 KB), would close the consumer-Blackwell gap independently of upstream PyTorch + CUTLASS work landing. Alternatively, a near-term fix: gate is_nvfp4_available() to return False on sm_120 so users get a clean signal instead of the cryptic cudaErrorInvalidValue crashes from row_col_rht_gemm_ntt_w_sfc.cu:1200.

Empirical reproduction (RTX 5080, CUDA 13.0, WSL2, all environment details preserved across 8 logs at /tmp/cycle308*.log + finetune/_preflight_te_fp4/cycle308{a,b}_results.json):

# TE 2.14.1 with TORCH_CUDA_ARCH_LIST=12.0a + Harry-Chen polyfill
import torch, transformer_engine.pytorch as te
from transformer_engine.common.recipe import NVFP4BlockScaling
x = torch.randn(1024, 4096, dtype=torch.bfloat16, device="cuda")
layer = te.Linear(4096, 4096, params_dtype=torch.bfloat16, device="cuda")
recipe = NVFP4BlockScaling()  # defaults: RHT enabled, SR enabled
with te.fp8_autocast(enabled=True, fp8_recipe=recipe):
    y = layer(x)  # crashes with CUDA Error: invalid argument

Opened #3195 with a datatype and hardware support matrix. It documents that NVFP4 is reported as available on compute capability 10.0 and higher, but the default recipe uses stochastic rounding whose FP4 conversion instructions are architecture specific to 10.0 and 10.3. On 12.0 (RTX 50 series) the default recipe raises an architecture-specific error and round-to-nearest is required. That lines up with the behavior and workaround discussed here.

docs: add datatype and hardware support matrix

Following up from the docs side (#3195 now documents this): the practical snag for people landing here is that check_nvfp4_support reports NVFP4 as available on compute capability 10.0 and higher, so on sm_120 the default recipe only fails later, deep in the kernel, with the architecture-specific FP4 cvt assert. Round-to-nearest (disable_stochastic_rounding=True) is the working path today.

One improvement that does not need sm_120 hardware: a fast-fail check that detects stochastic rounding on compute capability 12.0 and raises a clear error pointing to that workaround, instead of surfacing the low-level PTX assert. I do not have an sm_120 card to validate the kernel side, but I am happy to put up a PR for the clearer error path if that would be useful.

docs: add datatype and hardware support matrix (#3195)

Data from SM110 (Jetson AGX Thor), adding two points to this issue. First, the same gap exists on SM110, not only on SM120. Second, in a release build this path does not raise. It prints from every thread, keeps going, and returns wrong FP4 data. The earlier comments describe it as an error, which it is not in release builds.

Environment: Jetson AGX Thor (CC 11.0), JetPack 7.2, CUDA 13.2 (nvcc 13.2.78), cuDNN 9.20. TE main @ 9d9bd9e5, built from source with NVTE_CUDA_ARCHS=110 (CMake expands it to 110 + 110f). torch 2.13.0+cu130.

Summary: on SM110, is_nvfp4_available() returns True, and NVFP4BlockScaling() applies stochastic rounding to gradients by default (disable_stochastic_rounding=False). The SR conversion cvt.rs.satfinite.e2m1x4.f32 does not exist on this target. ptxas on CUDA 13.2:

target	`cvt.rs.satfinite.e2m1x4.f32`
sm_100a, sm_103a	OK
sm_100f	`Feature '.rs' not supported on .target 'sm_100f'`
sm_110, sm_110a, sm_110f	`Feature '.rs' not supported on .target 'sm_110a'`
sm_120a	`Feature '.rs' not supported on .target 'sm_120a'`

ARCH_HAS_STOCHASTIC_ROUNDING (common/util/ptx.cuh) correctly excludes these archs. The fallback branches, however, call NVTE_DEVICE_ERROR("FP4 cvt PTX instructions are architecture-specific. Try recompiling with sm_XXXa instead of sm_XXX."). Examples are mul_cvt_fp32_to_fp4_4x_with_stochastic_rounding (common/util/ptx_arch_spec.cuh) and StochasticNumericConverterBase (hadamard_transform/*cast_fusion.cu). NVTE_DEVICE_ERROR is printf + assert(0) (common/utils.cuh:77), and release builds compile with -DNDEBUG. The kernel therefore prints one line per thread and keeps running, returning wrong FP4 data. No error ever reaches the host.

Observed consequences:

In tests/pytorch/nvfp4, 15 tests fail. 14 of them fail with "Stochastic rounding failed - error larger than the round to nearest", and test_nvfp4_stochastic_rounding_allocation_independence fails with "Stochastic rounding did not respond to different seeds". The run's stdout grew to 10.5 GB of these device printfs.

Minimal repro: te.Linear(2048, 2048), BF16, M=512, default NVFP4BlockScaling(), forward+backward, compared against torch.nn.Linear:

	default (SR on)	`NVTE_NVFP4_DISABLE_STOCHASTIC_ROUNDING=1`
device printf lines	3928	0
process exit code	0	0
y cosine vs BF16	0.989	0.989
dgrad cosine	0.0	0.989
wgrad cosine	0.0 (wgrad is all zeros: mean \|g\| = 0.0 vs 18.05 reference)	0.991

The forward is fine. With the default recipe wgrad is all zeros, and dgrad has zero cosine to the reference. The prints come from StochasticNumericConverterBase in hadamard_transform/row_cast_col_hadamard_transform_cast_fusion.cu, i.e. the RHT + SR gradient path. Training with the default NVFP4 recipe on SM110 therefore makes no progress, and nothing raises.

The hint "recompiling with sm_XXXa" doesn't help here, because sm_110a has no .rs either.

Suggestions (any of them would make this fail loudly instead of silently):

Host side: when the device lacks SR conversion, either reject stochastic_rounding=True in the NVFP4 quantizer, or fall back to round-to-nearest with a one-time warning (or an emulated SR). The recipe default would then be safe on SM110/SM120.
Make NVTE_DEVICE_ERROR trap in release builds (__trap()), so that a path that should be unreachable cannot return garbage.
Skip the SR tests on archs without SR conversion.

I'm happy to send a PR for whichever direction you prefer, and I can validate it on Thor.

Correction to my comment above (#2255 (comment)). Two statements in it were wrong:

mul_cvt_fp32_to_fp4_4x_with_stochastic_rounding and the other helpers in common/util/ptx_arch_spec.cuh are not affected. They have used a software stochastic rounder since [Common] Fix NVFP4 stochastic rounding on architectures without cvt.rs #3281, so they were the wrong example.
"not only on SM120" suggested SM120 still has the gap. SM120/121 do not reach the kernels below: NVFP4Quantizer::is_eligible_for_rht_cast_fusion limits the grouped kernels to SM100–110, and nvte_quantize_with_hadamard_transform routes SM120/121 to the no-TMEM quantize_transpose path, which [Common] Fix NVFP4 stochastic rounding on architectures without cvt.rs #3281 covers.

What remains is the RHT cast-fusion kernels: the five StochasticNumericConverterBase copies in hadamard_transform/*cast_fusion.cu. They still fall back to NVTE_DEVICE_ERROR when cvt.rs is unavailable, and on SM110 they are where the device printfs and the zero wgrad came from. The rest of that comment (default recipe, measurements) stands.

#3706 reuses the #3281 rounder in those five places. On Jetson AGX Thor with that change, the te.Linear repro above under the default NVFP4BlockScaling() gives dgrad cosine 0.985 and wgrad cosine 0.986, with no device printf. In tests/pytorch/nvfp4, 9314 passed and 0 failed; the 15 SR failures I reported are among the passes.

Navigation Menu

Repository navigation

is NVFP4 not supported for rtx 50 series? #2255

CODE:

Error:

Here is the update code (the change in it is only torch.no_grad() added during inference)

Output:

这是更新代码（其中的变化仅在推理期间添加了 torch.no_grad()）

输出：

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

	NVTE_DEVICE_ERROR(
	"FP4 cvt PTX instructions are architecture-specific. "
	"Try recompiling with sm_XXXa instead of sm_XXX.");
	}

is NVFP4 not supported for rtx 50 series? #2255

Description

Activity

ptrendx commented on Oct 9, 2025

ptrendx commented on Oct 9, 2025

yash3056 commented on Oct 10, 2025

yash3056 commented on Oct 10, 2025

Oleg-Goncharov commented on Oct 10, 2025

yash3056 commented on Oct 10, 2025

CODE:

Error:

ptrendx commented on Oct 10, 2025

lixlbuaa commented on Oct 13, 2025

yash3056 commented on Oct 15, 2025

ptrendx commented on Oct 17, 2025

yash3056 commented on Oct 19, 2025

Here is the update code (the change in it is only torch.no_grad() added during inference)

Output:

ultism commented on Oct 21, 2025

这是更新代码（其中的变化仅在推理期间添加了 torch.no_grad()）

输出：

JamesDConley commented on Nov 1, 2025

alint77 commented on Nov 15, 2025

vgoklani commented on Nov 15, 2025

alint77 commented on Nov 15, 2025

vgoklani commented on Nov 15, 2025

danielwoz commented on Nov 26, 2025

Mocuto commented on Jan 13, 2026

vgoklani commented on Jan 13, 2026

Harry-Chen commented on Mar 24, 2026

vgoklani commented on Mar 24, 2026

Harry-Chen commented on Mar 25, 2026

osubotin commented on May 3, 2026

CodersAcademy006 commented on Jul 9, 2026

CodersAcademy006 commented on Jul 9, 2026

0z5a commented on Oct 16, 2026

0z5a commented on Oct 17, 2026

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions