fix(BACKEND-ROCM): gfx1200 build break — on_gfx1151() is __device__-only but called from host code
维护者通常 1 天内回复
还没有人认领这个 Issue。
评估
- 难度
- 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 天 4 小时
- 30 天内合并 PR
- 366
环境准备
从这里开始
- 先读完整个 Issue，再读项目的贡献指南。
- 在 Issue 下留言说明你要接手 —— 这能避免两个人做同样的事。
- Fork 仓库，在一个分支上完成修改。
- 提交 Pull Request，并在描述里引用这个 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 天内回复
相似的 Issue
-
难度 1/5 1 小时以内 新手友好度 92/100
维护者通常 1 天内回复
-
难度 2/5 1-3 小时 新手友好度 72/100
cp-algorithms/cp-algorithms#1715 ·
维护者通常 1 天内回复
-
难度 2/5 1-3 小时 新手友好度 78/100
Icinga/icinga2#11058 · 1 条评论 ·
维护者通常 1 天内回复
-
status:needs-triage
难度 2/5 1-3 小时 新手友好度 88/100
PX4/PX4-Autopilot#28924 ·
维护者通常 1 天内回复
-
component: split-view platform: windows
难度 2/5 1-3 小时 新手友好度 74/100
zen-browser/desktop#15616 · 1 个 reaction ·
维护者通常 1 天内回复