build(BACKEND-ROCM): route -Wall -Wextra -Werror to the HIP compile language and clear the 19 diagnostics it raises

[localai-org-maint-bot]

Row

BACKEND-ROCM — one row per PR. Spec: .agents/specs/rocm-hip-werror.md.

Before starting

• Issue/PR search and existing claim: #2713, open, unassigned, no branch. gh pr list shows no open pull request touching cmake/CompilerWarnings.cmake or the HIP warning arm. The five open ROCm pull requests (perf(GFX1100-TG200): T2b flips ROCm support_static_graph_mode #2777, perf(ROCm): parallel random sample with shared primitives #2778, Port async device-mirror combine/scatter kernels to ROCm #2779, feat(KERNEL-QUANT-CIQ-GEMM-ROCM): land the W1 keep-quant providers on kROCM #2782, docs(BACKEND-ROCM): reconcile merged Q8_K records #2625) are kernel and record work and touch none of these lines.
• Pull request shape selected at row claim: one pull request, the repository default. Spec committed first (6f2ec89f5), implementation after.
• Roadmap or matrix row: BACKEND-ROCM, already ACTIVE. No lifecycle change, so nothing is owed to a public document.
• Exact current-code anchors inspected: cmake/CompilerWarnings.cmake:54-66 (the non-MSVC branch, no HIP arm), CMakeLists.txt:404 enable_language(HIP), :1484 vllm_cpp_set_warnings(vllm), :1711-1734 the 23 .hip sources on that same target, src/vt/rocm/rocm_backend.hip:390-404 (the comment that recorded this finding and never got an issue).

What changed

vllm_cpp_set_warnings gains a $<$<COMPILE_LANGUAGE:HIP>:-Wall -Wextra ${_vllm_cpp_werror}> arm beside the CXX, OBJCXX and CUDA arms it already had, and the 19 diagnostics that arm raises are cleared at their source. Sixteen are dead code and are deleted; the other three are a (void) on three hipError_t results the runtime discards on purpose, the parentheses C++ already applies to one decode predicate, and one explicit unsigned cast. No kernel computes anything different.

${_vllm_cpp_werror} rather than the literal -Werror the OBJCXX line uses, so a HIP build under a sanitizer lane keeps the diagnostics visible and non-fatal for the same reason the CXX arm does.

The count, which the issue asks for first

Measured on strix:gpu0 (gfx1151, ROCm/HIP 7.2.53211), rc job 60f7861d, two arms, each a fresh configure into a fresh build directory, each compiling all 23 HIP objects:

Arm	HIP flags	raw warning lines	unique file:line:col:flag
A	as-committed	6	3
B	-Wall -Wextra, no -Werror	46	19

Zero of the 19 are in a ROCm or system header. Asked of the build system rather than the cmake source: on arm A a HIP compile line carries -Wall 0 times and -Werror 0 times; on arm B, -Wall once and -Werror still 0 times. By class: -Wunused-function 10, -Wunused-value 3, -Wunused-const-variable 2, -Wlogical-op-parentheses 2, -Wunused-parameter 1, -Wsign-compare 1.

19 is small, so this pull request adds the flag and clears them rather than reporting the size and stopping.

The ten unused-function reports are real dead code, not a HIP host-pass artefact

__device__ functions are not emitted in the host pass, so "clang reports a used __device__ helper as unused" is the obvious hypothesis. The line numbers refute it. Each of rocm_gdn_conv.hip, rocm_gdn_fused.hip and rocm_gdn_postconv.hip declares three Ld and three St overloads on consecutive lines — float, __half, __hip_bfloat16 — and in all three files exactly one Ld line and one St line is reported, and it is always the __half one. The neighbours are silent because kernels instantiate them. A pass artefact would report all six. After deletion, grep -c __half returns 0 in all three files. The other four warned symbols appear in their file exactly once, at their own definition.

Evidence

rc job c74c0cdf-804f-4666-a155-00fa0f59b963 on strix:gpu0. Four arms, each a fresh configure into a fresh build directory, ninja -k 0 over the same 23 HIP objects.

Arm	Tree	-Werror on a HIP compile line	ninja rc	objects	distinct errors	Result
RED	flag added, nothing cleared	1	1	14/23	15	BUILD_FAILED
GREEN	this branch	1	0	23/23	0	BUILD_OK
MUT	this branch + a bare hipError_t	1	1	—	≥1	BUILD_FAILED
CTRL	base tree + the SAME bare hipError_t	0	0	23/23	0	BUILD_OK

The mutation is the exact construct #2713 is about, applied identically to both trees. On this branch it reads error: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Werror,-Wunused-value]. On the base tree the same line reads warning: and the build goes green. The control is the part that matters: without it a red MUT proves only that the build broke, not that the new flag broke it.

RED reports 15 rather than 19 because the four missing ones — rocm_paged_attn.hip:30, :335, :448, :451 — are end-of-translation-unit diagnostics in the single unit that hard-fails earlier on -Wlogical-op-parentheses, and clang stops before emitting them. GREEN is what shows all 19 were cleared.

The arms were archived from commits later reworded to repair an Assisted-by value. The trees are byte-identical (da7b2417… RED, 38d96fce… GREEN), verified with git rev-parse <sha>^{tree} on both the pre- and post-reword commits, so the result binds to the commits on this branch.

After merging origin/main

main moved under this branch, so the measurement above was re-checked rather than assumed still valid. git diff --name-only ca07f6e94 origin/main shows main touched none of the eleven files this branch edits. The only headers it changed are include/vllm/multimodal/{audio_processor,dots3_note_processor}.h, and grep -rl "vllm/multimodal" src/vt/ returns nothing — no vt translation unit reaches them — so the preprocessed input to all 23 HIP units is byte-unchanged by the merge. Re-read on the merged tree: 23 HIP translation units, the HIP genex present once, enable_language(HIP) still at CMakeLists.txt:404, vllm_cpp_set_warnings(vllm) still at :1484.

Two rebuilds are queued on strix:gpu0 behind other sessions' leases, and neither result is in this body:

• merged — this branch's GREEN arm on the merged tree, so the count is re-measured rather than argued forward.

• cross — this branch merged together with row/BACKEND-ROCM-SCRATCH-SYNC (Two ROCm scratch allocators are unsynchronised: EnsureQuantScratch locks the map not the entry, and LtWorkspace frees with stream work possibly in flight #2712, PR fix(BACKEND-ROCM): hold one lock across the scratch entry, and retire the block instead of freeing it #2847). That branch rewrites LtWorkspace in a file this branch also edits. git merge-tree --write-tree reports a clean merge with no conflict, but a clean merge is textual: after both land, that new code is compiled under the -Werror this branch arms, which it has never been compiled under. Neither pull request's own gate covers that combination, which is why the arm exists. Both arms require -Werror-on-HIP == 1 as well as BUILD_OK, so an arm that silently lost the flag reads FAIL rather than PASS.

• scripts/agent-preflight.sh passes

• tests that cover this change: the build itself, on the only toolchain that can run it. There is no unit test to name, and adding a checker that asserts the genex exists would be a transcription of the line it guards.

• public docs changed only when this PR changes their owned user-facing facts: none changed. This edits cmake/, src/ and one spec.

Speed claims

• This PR makes NO speed claim.
• the operator ran the numbers under ${GPU_LOCK}.

Honest gaps

• CI cannot rerun this gate. No CI lane has a HIP toolchain, so the flag is enforced only where somebody builds the ROCm backend. Deletion of the genex is a regression nothing here detects. Recorded in the spec under ## Owed; not repaired, because the derived checker that would catch it needs an exception list for C and ASM, and a list is the transcription this file's own rules warn about.
• One architecture, one ROCm release. gfx1151 and ROCm 7.2.5 only. -Wall -Wextra output is target-independent for these six classes, but that is reasoning, not a measurement, and no second gfx target was built.
• No kernel was executed. Every arm is a compile. The claim that no repair changes what a kernel computes rests on per-site reading — deletions of code with zero call sites, a parenthesisation matching C++'s existing precedence, and a cast on a value mindiv returns positive by construction — not on a numeric gate.
• The job's own "red classes" tally line printed nothing: its regex expects [-Wflag] and the -Werror spelling is [-Werror,-Wflag]. The per-error list above it is complete and is what the table reports; the tally was not relied on.

Closes #2713

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5-1m [claude-code]