perf(GFX1100-TG200): put back the HW dot instruction T4a replaced with a scalar loop

T4a (51f5222, #2790) rewrote Dp4a in rocm_grouped_gemm.hip from __ockl_sdot4
to four int8 multiplies and four adds, and rewrote the comment above it to say
the hardware instruction is "a perf lever, not a correctness requirement". Dp4a
is the integer core of every K-quant GEMM on this backend.

scripts/check-rocm-dp4a-intrinsic.py exists for exactly this and says why in its
own docstring: the scalar form is bit-identical but ~1.4x slower on the
KQuantGemmK prefill path, a CPU-only ctest stays green either way because the
ROCm kernel is not compiled there, so only a source checker can hold the lever.
Its mutation test then failed in preflight with "mutation did not apply" --
there was no longer an intrinsic call to mutate.

Both descriptions agree the forms are bit-identical, so this changes no output.
It restores the instruction, the ~1.4x, and the gate: the checker reports OK and
tests/scripts/test_check_rocm_dp4a_intrinsic.py is 6/6.

Closes #2939.

FOLLOWING_AGENTS_PROTOCOL

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

Author: mudler

src/vt/rocm/rocm_grouped_gemm.hip

@@ -57,13 +57,17 @@ __device__ __forceinline__ int GetIntB2(const int8_t* qs, int i32) {
   return static_cast<int>(x16[2 * i32 + 0]) | (static_cast<int>(x16[2 * i32 + 1]) << 16);
 }
 
-// Signed 8-bit x4 dot-product-accumulate, bit-identical to __dp4a (integer
-// math is exact either way). The HW dot instruction (v_dot4_i32_i8 /
-// __ockl_sdot4) is a perf lever, not a correctness requirement.
+// Signed 8-bit x4 dot-product-accumulate. Uses the HW v_dot4_i32_i8
+// instruction (__ockl_sdot4) on gfx1100 — one instruction instead of 4
+// int8 multiplies + 4 adds. Bit-identical: signed int8×int8→int32 dot
+// product is exact either way (the HW instruction and the scalar expansion
+// compute the same integer result), which is why only a source checker can
+// hold it — scripts/check-rocm-dp4a-intrinsic.py, ~1.4x on KQuantGemmK.
 __device__ __forceinline__ int Dp4a(int a, int b, int acc) {
-  const int8_t* a8 = reinterpret_cast<const int8_t*>(&a);
-  const int8_t* b8 = reinterpret_cast<const int8_t*>(&b);
-  return acc + a8[0] * b8[0] + a8[1] * b8[1] + a8[2] * b8[2] + a8[3] * b8[3];
+  using char4_native = char __attribute__((ext_vector_type(4)));
+  char4_native va = *reinterpret_cast<const char4_native*>(&a);
+  char4_native vb = *reinterpret_cast<const char4_native*>(&b);
+  return __ockl_sdot4(va, vb, acc, false);
 }
 
 // ---- activation quantizers ----