[BUG] NVFP4 fails on RTX 5070 Laptop (SM120): CUDA Error: no kernel image is available for execution on the device

[Yash-456]

Environment

• GPU: NVIDIA GeForce RTX 5070 Laptop GPU
• Compute Capability: SM120 (12.0)
• CUDA Toolkit: 13.3
• Python: 3.13
• WSL2 Ubuntu 24.04
• TransformerEngine built from source
• Local commit: 6bf7790d
• Upstream main: 9d92fa0

Problem

TransformerEngine builds successfully, but NVFP4 training fails immediately with:

RuntimeError:
/home/.../transformer_engine/common/recipe/current_scaling.cu:77

launch_amax_kernel

CUDA Error: no kernel image is available for execution on the device

The failure happens immediately after:

zero_amax_kernel<<<1,1,0,stream>>>(...);
NVTE_CHECK_CUDA(cudaGetLastError());

[Yash-456]

Complete debugging history

This comment documents the complete debugging timeline leading to the current runtime failure.

---

Environment

• GPU: NVIDIA GeForce RTX 5070 Laptop GPU
• Compute Capability: SM120 (12.0)
• CUDA Toolkit: 13.3
• Python: 3.13
• Ubuntu 24.04 (WSL2)
• TransformerEngine built from source (editable install)

PyTorch correctly detects the GPU:

>>> import torch
>>> torch.cuda.get_device_name(0)
'NVIDIA GeForce RTX 5070 Laptop GPU'

>>> torch.cuda.get_device_capability(0)
(12, 0)

---

Git information

Current local branch:

local-build-fixes

Current local commit:

6bf7790d

Current upstream main:

9d92fa05

---

Previous successful editable build

I successfully built and installed TransformerEngine from source before investigating the runtime issue.

Successfully built transformer_engine
Successfully installed transformer_engine-2.18.0.dev0+b782ea6e

After updating my local branch, I rebuilt TransformerEngine again.

The rebuild also completed successfully:

Successfully built transformer_engine
Successfully installed transformer_engine-2.18.0.dev0+6bf7790d

This indicates that the project compiles and installs successfully.

---

Build modification

To obtain reliable incremental builds I modified build_tools/build_ext.py from

if found_ninja():
    configure_command.append("-GNinja")

to

if found_ninja() and False:
    configure_command.append("-GNinja")

forcing CMake to use gmake instead of Ninja.

This change only affects the build system and was made to simplify debugging.

---

CUDA architecture verification

python - <<'PY'
from build_tools.utils import cuda_archs
print(cuda_archs())
PY

returns

75;80;89;90;100;120

CMakeCache.txt contains

CMAKE_CUDA_ARCHITECTURES=75;80;89;90;100;120

Generated nvcc commands include

--generate-code=arch=compute_120f,code=sm_120f

---

cuobjdump investigation

current_scaling.cu.o contains

sm_75
sm_80
sm_89
sm_90
sm_100
sm_120

However,

hadamard_transform.cu.o

only reports

sm_75

when inspected using

cuobjdump --list-elf

despite its generated compile command including

--generate-code=arch=compute_120f,code=sm_120f

I do not know whether this difference is expected or indicative of the runtime issue.

---

Runtime progression

Initially NVFP4 failed inside

transformer_engine/common/hadamard_transform/hadamard_transform.cu

with

CUDA Error: no kernel image is available for execution on the device

I then changed the recipe to

NVFP4BlockScaling(
    disable_rht=True
)

Setting disable_rht=True changes the failure location from

transformer_engine/common/hadamard_transform/hadamard_transform.cu

to

transformer_engine/common/recipe/current_scaling.cu

indicating that execution progresses further into the NVFP4 pipeline.

The runtime now consistently fails inside

launch_amax_kernel()

immediately after

zero_amax_kernel<<<1,1,0,stream>>>(...);
NVTE_CHECK_CUDA(cudaGetLastError());

with

CUDA Error: no kernel image is available for execution on the device

---

Training command

python train_llm.py \
    --use_nvfp4 \
    --fresh \
    --no-grad_checkpoint

The failure occurs during the first forward pass.

---

Things already verified

• TransformerEngine builds successfully.
• Editable installation succeeds.
• PyTorch correctly detects the GPU as SM120.
• CUDA Toolkit 13.3 is being used.
• CMAKE_CUDA_ARCHITECTURES includes SM120.
• Generated nvcc commands include compute_120f.
• current_scaling.cu.o contains an sm_120 cubin.
• Setting disable_rht=True moves the failure from hadamard_transform.cu to current_scaling.cu, indicating that execution progresses further before failing.

---

Summary

At this point the issue no longer appears to be related to:

• compilation
• installation
• CUDA architecture detection
• CMake architecture generation

Instead, it appears to be a runtime issue in the NVFP4 execution path on a consumer Blackwell SM120 GPU after successful initialization.

If this is a known issue on consumer Blackwell (SM120), or if there is a newer branch or patch that addresses it, I'd be happy to test it and report the results.