HIP: limit RDNA3.5 Q8_0 X-tile prefetch to J=128

The prefetch pipeline was gated on J == 96 || J == 128, but the batched WMMA
vec_dot it feeds is gated on J == 128 only, and rows_per_warp() returns 16
(ntx=1) only at J == 128 -- at J == 96 it returns 32 (ntx=2). So J == 96 paired
the prefetch with the generic two-minitile kernel. Q8_0 is also the only type
caching 2*(I/nwarps) = 32 ints where the others cache 16, on top of two A tiles
and 48 accumulators, so that combination spills.

mmq_rdna35_tuned_J maps Q8_0 J_occupancy == 64 to 96 to avoid the J == 64
pathology, which made every 64-wide batch land on the broken width. MoE Q8_0
(J > 32 -> 96) had the same exposure.

Cosmos-Reason2-8B Q8_0 prefill on gfx1151, llama-bench -p 512 -r 3 against
rocm/gfx11 03d2068:

  ubatch  64:  111.4 -> 178.7 t/s  (was -37.7% vs gfx11, now -0.1%)
  ubatch  96:    n/a -> 267.6 t/s  (now -0.4% vs gfx11)
  ubatch 128: 1426.8 -> 1430.9 t/s (+3.0% vs gfx11, unchanged)
  ubatch 512: 1581.5 -> 1581.9 t/s (+2.8% vs gfx11, unchanged)

The J == 128 win is retained in full. test-backend-ops MUL_MAT q8_0: 64/64.

Co-authored-by: Cursor <cursoragent@cursor.com>

Author: liangliangchang

ggml/src/ggml-cuda/mmq.cuh

@@ -1206,7 +1206,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile(
             }
             __syncthreads();
         }
-    } else if constexpr (type == GGML_TYPE_Q8_0 && (J == 96 || J == 128)) {
+    } else if constexpr (type == GGML_TYPE_Q8_0 && J == 128) {
         constexpr int qs_cache_size = 2*(I/nwarps);
         constexpr int d_cache_size  = 4;