Hacktoberfest 2026: le issue che i maintainer hanno segnato per ottobre, aperte e adatte ai principianti. Sfoglia le issue Hacktoberfest

fix(BACKEND-ROCM): gfx1200 build break — on_gfx1151() is __device__-only but called from host code

Aperta
  • Qualcosa non va?
#3,255 0 commenti 0 reazioni 0 assegnatari Vedi su GitHub

I maintainer di solito rispondono entro 1 giorno

Nessuno ha ancora preso questa issue.

Valutazione

Difficoltà
3/5
Tempo stimato
1-2 giorni
Idoneità per principianti
64/100
Tipo di issue
Bug
Chiarezza
Abbastanza chiara
Stato di attività
Attiva
Stack tecnologico
cpp
Ambito
backend, build-system

Direzione di ricerca

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.

Scritto dal modello di indicizzazione a partire dal testo della issue.

Descrizione

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]

Lingua principale
C++
Stelle
423
Fork
53
Merge medio
1g 7h
PR unite (30g)
380

Preparare l'ambiente

Come iniziare

  1. Leggi tutta la issue e poi la guida ai contributi del progetto.
  2. Commenta sulla issue per dire che te ne occupi tu — evita che due persone facciano lo stesso lavoro.
  3. Fai un fork del repository e lavora su un branch.
  4. Apri una pull request che faccia riferimento al numero della issue.

Altre issue di mudler/vllm.cpp

Tutte le issue di mudler/vllm.cpp

Issue simili

Altre issue su C++

Ricevi le nuove issue nella tua casella

Un breve riepilogo di issue GitHub adatte ai principianti.