Hacktoberfest 2026: los issues que los mantenedores marcaron para octubre, abiertos y aptos para principiantes. Explorar issues de Hacktoberfest

fp4_kernel_hopper.py uses tl.make_block_ptr, which is removed in Triton main (breaks NVFP4 fake quant)

Abierto
  • ¿Algo no está bien?
#2,810 0 comentarios 0 reacciones 1 asignado Ver en GitHub

Los mantenedores suelen responder en 2 días

@kinjalpatel27 ya está trabajando en esto.

Desde el 15/10/2026.

Evaluación

Dificultad
3/5
Tiempo estimado
Medio día
Aptitud para principiantes
72/100
Tipo de issue
Error
Claridad
Bien especificado
Estado de actividad
Activo
Stack tecnológico
python
Área
machine-learning

Línea de trabajo

Start in modelopt/torch/kernels/quantization/gemm/fp4_kernel_hopper.py at fp4_fake_quant_kernel, where the two tl.make_block_ptr calls for load and store are, and at fp4_fake_quant_block, the entry point the repro calls. The caller tensor_quant._dynamic_block_quantize_impl routes NVFP4 here when Triton is available and compute capability is at least 8.9, so DISABLE_TRITON_KERNEL is the only switch to note. Done when the repro call compiles under Triton 3.9 and the output matches the pre-change results on an earlier Triton.

Escrito por el modelo de indexación a partir del texto del issue.

Descripción

Describe the bug

modelopt/torch/kernels/quantization/gemm/fp4_kernel_hopper.py (fp4_fake_quant_kernel) uses tl.make_block_ptr for its load and store. Triton removed block pointers in triton-lang/triton#10833 ("[BC breaking][FRONTEND] Remove block pointer support"), after deprecating them in triton-lang/triton#9667. The Triton that ships with current PyTorch nightly (pytorch-triton 3.9.0+gitfc8d31a5) already includes this, so the kernel fails to compile:

triton.compiler.errors.CompilationError: at 24:18:
    x_block_ptr = tl.make_block_ptr(
                  ^
Block pointers have been removed in favor of the tensor descriptor API

tensor_quant._dynamic_block_quantize_impl routes NVFP4 fake quantization to this kernel whenever Triton is available and compute capability is >= 8.9, so NVFP4 quantization (e.g. mtq.NVFP4_DEFAULT_CFG) breaks on Hopper/Ada/Blackwell with Triton 3.9. The only bypass is the module constant DISABLE_TRITON_KERNEL = False, which isn't user-configurable. This is the only make_block_ptr usage in ModelOpt, and it's still present on main and 0.48.0rc0. Triton 3.8 (the latest stable release) is unaffected.

Steps/Code to reproduce bug
python -m venv venv && . venv/bin/activate
pip install --pre torch --index-url https://download.pytorch.org/whl/nightly/cu130   # brings triton 3.9.0+gitfc8d31a5
pip install "nvidia-modelopt[hf]==0.47.0"

import torch
from modelopt.torch.kernels.quantization.gemm.fp4_kernel_hopper import fp4_fake_quant_block

x = torch.randn(128, 64, device="cuda", dtype=torch.float16)
fp4_fake_quant_block(x, x.abs().amax().float())
# triton.compiler.errors.CompilationError: ... Block pointers have been removed in favor of the tensor descriptor API

(Side note: plain pip install nvidia-modelopt==0.47.0 fails at import modelopt.torch with No module named 'requests' and then 'huggingface_hub'; both are imported unconditionally but only provided by the [hf] extra.)

Expected behavior

The NVFP4 fake-quant kernel compiles with current Triton. Replacing the two block pointers with plain pointer arithmetic and masks keeps it working on all Triton versions:

offs_m = row_start + tl.arange(0, TILE_M)
offs_n = col_start + tl.arange(0, TILE_N)
mask = (offs_m[:, None] < M) & (offs_n[None, :] < N)
tile = tl.load(x_ptr + offs_m[:, None] * stride_xm + offs_n[None, :] * stride_xn, mask=mask, other=0.0).to(tl.float32)
...
tl.store(y_ptr + offs_m[:, None] * stride_ym + offs_n[None, :] * stride_yn, tile_quant.to(OUT_DTYPE), mask=mask)

(or port to tl.make_tensor_descriptor).

System information

  • Container used: none (fresh virtualenv)
  • OS: Ubuntu
  • CPU architecture: x86_64
  • GPU: reproduced on RTX 3080 Ti; the error is raised by Triton's frontend while compiling, so it's GPU-independent
  • Library versions:
    • Python: 3.13
    • ModelOpt: 0.47.0 (same code on main / 0.48.0rc0)
    • PyTorch: 2.16.0.dev20261014+cu130 (nightly)
    • Triton: 3.9.0+gitfc8d31a5 (bundled with the nightly)
    • CUDA: 13.0
Lenguaje dominante
Python
Estrellas
5.3k
Forks
745
Merge medio
4 d 2 h
PR fusionados (30 d)
160

Preparar el entorno

Primeros pasos

  1. Lee el issue completo y luego la guía de contribución del proyecto.
  2. Comenta en el issue que vas a ocuparte — evita que dos personas hagan lo mismo.
  3. Haz un fork del repositorio y trabaja en una rama.
  4. Abre un pull request que haga referencia al número del issue.

Más de NVIDIA/Model-Optimizer

Todos los issues de NVIDIA/Model-Optimizer

Issues similares

Más issues de Python

Recibe los nuevos issues en tu correo

Un resumen breve de issues de GitHub para principiantes.