rocm: fix build on gfx1201 (RDNA4, R9700) — gate the WMMA v1 kernel

[cm999club]

Fixes #521.

Problem

Building on an AMD Radeon AI PRO R9700 (gfx1201, RDNA4 discrete GPU) fails
with:

Cannot select: intrinsic %llvm.amdgcn.wmma.f32.16x16x16.f16

matmul_q8_0_f32_batch_wmma_4w_kernel in rocm/ds4_rocm_q8.cuh uses
__builtin_amdgcn_wmma_f32_16x16x16_f16_w32, which is the RDNA3 (WMMA v1)
builtin. RDNA4 (gfx12xx) uses a different WMMA v2 intrinsic, so the compiler
can't select the RDNA3 builtin for a gfx1201 target and the build aborts.

Fix

This is a compile/runtime workaround, not an RDNA4 WMMA v2
implementation:

• ds4_rocm_runtime.cuh: at ds4_gpu_init(), detect the arch from
  prop.gcnArchName and set a new g_wmma_v1 flag (0 on any gfx12*
  target, 1 otherwise).
• ds4_rocm_q8.cuh: guard the RDNA3 WMMA v1 builtin calls inside
  matmul_q8_0_f32_batch_wmma_4w_kernel with
  #if !defined(__gfx1200__) && !defined(__gfx1201__), so the kernel still
  compiles for gfx12xx targets (the guarded body becomes a no-op there)
  instead of failing at builtin selection.
• ds4_rocm_matmul.cuh: gate dispatch into that kernel on g_wmma_v1, so
  gfx12xx never calls into the no-op kernel at runtime — it falls back to
  the existing non-WMMA matmul path instead.

Net effect: the build succeeds on gfx1201, and inference runs correctly,
but the accelerated WMMA-4w kernel is not used there — gfx12xx runs on the
existing fallback path. A real RDNA4 WMMA v2 kernel would recover that
performance; this PR intentionally scopes to "unbreak the build and run
correctly" rather than attempting the v2 intrinsic without hardware to
validate the accumulator/layout semantics against.

Testing

Running in production on the reported hardware (R9700, 32GB VRAM, ROCm
7.2.2) for two weeks with --rocm --ssd-streaming, serving DeepSeek V4
Flash inference daily without correctness issues.

Happy to adjust the detection (e.g. a proper HIP_ARCH/target macro instead
of the gcnArchName substring check) if there's a preferred convention in
the codebase — I didn't see one, so I went with what's inspectable at
runtime.