Hacktoberfest 2026：メンテナが10月に向けて印を付けた、オープンで初心者向けの issue。 Hacktoberfest の issue を見る

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

オープン
  • 何か問題がありますか？
#2,810 コメント 0 件 リアクション 0 件 担当者 1 名 GitHub で見る

メンテナーはふだん 2 日以内に返信

@kinjalpatel27 がすでに取り組んでいます。

2026年10月15日 から。

評価

難易度
3/5
見積もり時間
半日
初心者へのやさしさ
72/100
issue の種類
バグ
明瞭さ
明確に書かれている
活発さ
活発
技術スタック
python
領域
machine-learning

調査の方向性

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.

索引モデルが issue の本文から書いたものです。

説明

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
主要言語
Python
スター
5.4k
フォーク
767
平均マージ
3日 11時間
マージ済み PR（30日）
203

環境構築

はじめの一歩

  1. issue を最後まで読み、次にプロジェクトのコントリビューションガイドを読みます。
  2. 着手することを issue にコメントします — 二人が同じ作業をするのを防げます。
  3. リポジトリをフォークし、ブランチを切って変更します。
  4. issue 番号を参照したプルリクエストを送ります。

NVIDIA/Model-Optimizer のほかの issue

NVIDIA/Model-Optimizer の issue をすべて見る

似ている issue

Python の issue をもっと見る

新しい issue をメールで受け取る

初心者向けの GitHub issue を短くまとめたダイジェスト。