fix(BACKEND-ROCM): gfx1200 build break — on_gfx1151() is __device__-only but called from host code
メンテナーはふだん 1 日以内に返信
まだ誰も着手していません。
評価
- 難易度
- 3/5
- 見積もり時間
- 1〜2日
- 初心者へのやさしさ
- 64/100
- issue の種類
- バグ
- 明瞭さ
- おおむね明確
- 活発さ
- 活発
- 技術スタック
- cpp
- 領域
- backend, build-system
調査の方向性
Start with src/vt/rocm/rocm_skinny_gemm.hip, especially on_gfx1151() at line 40 and SelectYtileUnrl() around lines 68-72. Reproduce the gfx1200 build with ROCm's bundled Clang, then verify the qualifier and preprocessor behavior for GFX11 and GFX12 targets. Done means the affected ROCm targets compile without introducing a device-side host-API call.
索引モデルが issue の本文から書いたものです。
説明
Row: -
What's broken
main (current head, via commit 316d74e78f "perf(BACKEND-ROCM): add
gfx1151 dispatch tuning to wvSplitK", authored 2026-09-15) fails to build for
gfx1200 (and, on inspection, would fail for any ROCm target since the
defect is a qualifier mismatch, not an architecture branch):
src/vt/rocm/rocm_skinny_gemm.hip:72:7: error: no matching function for call to 'on_gfx1151'
72 | if (on_gfx1151()) {
| ^~~~~~~~~~
src/vt/rocm/rocm_skinny_gemm.hip:40:33: note: candidate function not viable: call to __device__ function from __host__ function
40 | __device__ __forceinline__ bool on_gfx1151() {
| ^
on_gfx1151() (rocm_skinny_gemm.hip:40) is declared __device__-only, but
SelectYtileUnrl() (rocm_skinny_gemm.hip:68), which calls it at line 72, is
a plain host-side inline function with no __device__/__global__
qualifier — the host compilation pass of the translation unit cannot see a
device-only callee.
Compiler: ROCm's bundled Clang 22.0.0 (/opt/rocm/lib/llvm/bin/clang++),
target gfx1200.
How found
Same session as #3254: independently reproducing the operator's gfx1200
toggle matrix for PR #3036. Second, unrelated build break hit immediately
after clearing the first (#3254). Neither file is touched by #3036's own
diff.
git blame on rocm_skinny_gemm.hip:39-50 attributes on_gfx1151 to
316d74e78f alone.
Likely fix
Mark on_gfx1151() __host__ __device__ instead of __device__-only. The
body is guarded by #if defined(__GFX11__), which is not defined when
compiling for gfx1200 (GFX12, not GFX11), so the host pass takes the
#else return false; branch and never reaches the hipDeviceGetAttribute
call — that call is itself host-side API, so this guard needs re-checking
for any GFX11 target where __GFX11__ IS defined during device-side
compilation, where calling a host API from generated device code would be
its own error. Applied locally, uncommitted, only to unblock the toggle-matrix
build for #3036; not verified against a GFX11 device.
Owed
No row claims this yet. Filed Row: - per the same rule cited in #3254;
needs triage to pick an owner.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-sonnet-5 [Claude Code]
- 主要言語
- C++
- スター
- 423
- フォーク
- 53
- 平均マージ
- 1日 7時間
- マージ済み PR（30日）
- 380
環境構築
はじめの一歩
- issue を最後まで読み、次にプロジェクトのコントリビューションガイドを読みます。
- 着手することを issue にコメントします — 二人が同じ作業をするのを防げます。
- リポジトリをフォークし、ブランチを切って変更します。
- issue 番号を参照したプルリクエストを送ります。
mudler/vllm.cpp のほかの issue
-
難易度 2/5 1〜3時間 初心者へのやさしさ 84/100
メンテナーはふだん 1 日以内に返信
-
難易度 2/5 1〜3時間 初心者へのやさしさ 82/100
メンテナーはふだん 1 日以内に返信
-
難易度 2/5 1〜3時間 初心者へのやさしさ 70/100
メンテナーはふだん 1 日以内に返信
-
難易度 1/5 1時間未満 初心者へのやさしさ 88/100
メンテナーはふだん 1 日以内に返信
-
難易度 2/5 1〜3時間 初心者へのやさしさ 84/100
メンテナーはふだん 1 日以内に返信
mudler/vllm.cpp の issue をすべて見る
似ている issue
-
Unconfirmed bug
難易度 1/5 1時間未満 初心者へのやさしさ 88/100
luanti-org/luanti#17605 · コメント 1 件 ·
メンテナーはふだん 2 日以内に返信
-
area: config area: firmware priority: P2 - medium size: S type: bug
難易度 2/5 1〜3時間 初心者へのやさしさ 76/100
Mizithra/ActiveTerrain#16 ·
-
難易度 2/5 1〜3時間 初心者へのやさしさ 84/100
grumpycoders/pcsx-redux#2171 ·
メンテナーはふだん 1 日以内に返信
-
難易度 2/5 1〜3時間 初心者へのやさしさ 70/100
メンテナーはふだん 2 日以内に返信
-
難易度 2/5 1〜3時間 初心者へのやさしさ 88/100
bytedance/trae-agent#524 · コメント 1 件 ·
メンテナーはふだん 1 日以内に返信