cuda: port TQ kernel optimizations from TheTom/llama-cpp-turboquant

mmvq-tq.cu: complete rewrite synced from turboquant PR ggml-org#57
- AMD/ROCm: scalar half path (dp4a regresses on RDNA4) — replaces compute-bound V12 WHT shmem kernel
- NVIDIA: dp4a multi-token path with register-based __byte_perm centroid LUT
- Multi-token support up to ne[1]=8 for speculative decoding (was ne[1]=1 only)
- ggml_cuda_mul_mat_tq4_1s_cublas: large prefill via runtime fp16 dequant + cuBLAS
- Load-time q8_0 conversion now opt-in (GGML_TQ_CONVERT_Q8=1), default off

ggml-cuda.cu dispatch:
- Use ggml_cuda_mul_mat_tq for ne[1]<=8 (was ne[1]==1)
- Add cuBLAS path for large TQ4_1S prefill

fattn-common.cuh: ROCm FA pool OOM bypass
- HIP: raw cudaMalloc/Free for f16 temp buffers instead of pool
- Prevents permanent VRAM retention from peak-size pool allocations
- Fixes OOM when using quantized KV cache on ROCm

fattn.cu: force VEC FA path for quantized KV on HIP
- Prevents TILE/MMA paths from allocating temp f16 buffers for quant KV
- Also enables mixed f16/turbo KV type combinations

turbo-quant.cuh, turbo-innerq.cuh: minor correctness fixes
- Wrap bare cuda*Symbol calls in CUDA_CHECK
- Fix TURBO_IQ_API DLL export macros for static builds

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

Author: kmbandy

ggml/src/ggml-cuda/fattn-common.cuh

@@ -1297,8 +1297,33 @@ void launch_fattn(
     const int cc  = ggml_cuda_info().devices[id].cc;
     const int nsm = ggml_cuda_info().devices[id].nsm;
 
+#ifdef GGML_USE_HIP
+    // HIP/ROCm: bypass the memory pool for f16 temp buffers.
+    // The legacy pool (ggml_cuda_pool_leg) retains peak-sized allocations permanently.
+    // For quantized KV dequant, this means the f16 temp buffer stays allocated,
+    // consuming more VRAM than the quantized KV compression saves — causing OOM.
+    // Using raw alloc+free ensures the memory is released after the kernel completes.
+    struct hip_f16_alloc {
+        half * ptr = nullptr;
+        cudaStream_t stream;
+        hip_f16_alloc(cudaStream_t s) : stream(s) {}
+        ~hip_f16_alloc() {
+            if (ptr) {
+                cudaStreamSynchronize(stream);
+                cudaFree(ptr);
+            }
+        }
+        void alloc(size_t nelements) {
+            CUDA_CHECK(cudaMalloc(&ptr, nelements * sizeof(half)));
+        }
+        half * get() { return ptr; }
+    };
+    hip_f16_alloc K_f16(main_stream);
+    hip_f16_alloc V_f16(main_stream);
+#else
     ggml_cuda_pool_alloc<half>   K_f16(pool);
     ggml_cuda_pool_alloc<half>   V_f16(pool);
+#endif
     ggml_cuda_pool_alloc<int>    KV_max(pool);
     ggml_cuda_pool_alloc<float>  dst_tmp(pool);
     ggml_cuda_pool_alloc<float2> dst_tmp_meta(pool);

ggml/src/ggml-cuda/fattn.cu

@@ -472,6 +472,19 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
     // For small batch sizes the vector kernel may be preferable over the kernels optimized for large batch sizes:
     const bool can_use_vector_kernel = Q->ne[0] <= 256 && Q->ne[0] % 64 == 0 && K->ne[1] % FATTN_KQ_STRIDE == 0;
 
+#ifdef GGML_USE_HIP
+    // HIP/ROCm: the TILE/MMA/WMMA FA paths allocate unbounded f16 temp buffers
+    // for quantized KV types (K_f16, V_f16 in launch_fattn). The pool retains
+    // peak allocation size, so the temp buffer VRAM exceeds KV compression savings.
+    // This causes quantized KV to OOM before f16 on the same context length.
+    // Force VEC path which does inline dequant with zero temp buffer overhead.
+    // Trade-off: prefill is slower (sequential query processing).
+    // Limitation: head_dim > 256 cannot use VEC (falls through to TILE).
+    if ((ggml_is_quantized(K->type) || ggml_is_quantized(V->type)) && can_use_vector_kernel) {
+        return BEST_FATTN_KERNEL_VEC;
+    }
+#endif // GGML_USE_HIP
+
     // If Turing tensor cores are available, use them:
     if (turing_mma_available(cc) && Q->ne[0] != 40 && Q->ne[0] != 72) {
         if (can_use_vector_kernel) {

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -2469,9 +2469,12 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
         ggml_cuda_op_mul_mat(ctx, src0, src1, dst, ggml_cuda_op_mul_mat_vec_q, quantize_row_q8_1_cuda);
     } else if (use_mul_mat_q) {
         ggml_cuda_op_mul_mat(ctx, src0, src1, dst, ggml_cuda_op_mul_mat_q, quantize_mmq_q8_1_cuda);
-    } else if (!split && is_tq_weight && src1->ne[1] == 1) {
-        // Fused TQ weight mul_mat_vec with pre-rotated activations via warp shuffle WHT
-        ggml_cuda_mul_mat_vec_tq(ctx, src0, src1, dst);
+    } else if (!split && is_tq_weight && src1->ne[1] <= MMVQ_MAX_BATCH_SIZE) {
+        // Fused TQ weight mul_mat: handles decode (ne[1]=1) and speculative (ne[1]<=8)
+        ggml_cuda_mul_mat_tq(ctx, src0, src1, dst);
+    } else if (!split && src0->type == GGML_TYPE_TQ4_1S) {
+        // Large prefill: runtime TQ4_1S → fp16 dequant + cuBLAS tensor cores
+        ggml_cuda_mul_mat_tq4_1s_cublas(ctx, src0, src1, dst);
     } else {
         ggml_cuda_op_mul_mat(ctx, src0, src1, dst, ggml_cuda_op_mul_mat_cublas, nullptr);
     }