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

[joral]

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]