perf(cuda): add warp-cooperative TurboQuant FWHT

[Amidwestnoob]

Replaces the CUDA scalar TurboQuant FWHT path with a warp-cooperative implementation. HIP remains unchanged.

Validated on dual RTX 3090 SM86:

• Scalar/warp correctness: 10/10 bit-identical
• Current-upstream A/B/A: +1.58% wall throughput
• All 15 outputs byte-identical
• No VRAM increase or fatal runtime errors

[cubic-dev-ai]

1 issue found across 3 files

Prompt for AI agents (unresolved issues)

Check if these issues are valid — if so, understand the root cause of each and fix them. If appropriate, use sub-agents to investigate and fix each issue separately.


<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/turbo-wht.cu">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/turbo-wht.cu:71">
P3: The new `k_turbo_wht_warp` duplicates the whole group-coordinate decode (`gid -> g/rem -> i01/i02` and the `src_base + i01*src_nb1 + i02*src_nb2 + g*QK_TQ3_0_GROUP` row-pointer math) plus the 12-argument signature and launch argument list that the scalar (`k_turbo_wht_scalar`) and fused-quantize (`k_turbo_wht_quantize`) kernels already own. Since the scalar and warp paths are chosen by `#if defined(GGML_USE_HIP)`, only one compiles on a given backend, so today this is behavior-neutral. The risk is maintenance drift: a future change to the group indexing or stride handling would have to be applied identically in three places or the HIP and CUDA paths would diverge silently. Consider factoring the common decode (params + gid->(g,i01,i02) + row pointers) into a shared helper/inline used by all three kernels to keep them in sync.</violation>
</file>

<sub>Reply with feedback, questions, or to request a fix.

Re-trigger cubic</sub>

[cubic-dev-ai]

P3: The new k_turbo_wht_warp duplicates the whole group-coordinate decode (gid -> g/rem -> i01/i02 and the src_base + i01*src_nb1 + i02*src_nb2 + g*QK_TQ3_0_GROUP row-pointer math) plus the 12-argument signature and launch argument list that the scalar (k_turbo_wht_scalar) and fused-quantize (k_turbo_wht_quantize) kernels already own. Since the scalar and warp paths are chosen by #if defined(GGML_USE_HIP), only one compiles on a given backend, so today this is behavior-neutral. The risk is maintenance drift: a future change to the group indexing or stride handling would have to be applied identically in three places or the HIP and CUDA paths would diverge silently. Consider factoring the common decode (params + gid->(g,i01,i02) + row pointers) into a shared helper/inline used by all three kernels to keep them in sync.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/src/ggml-cuda/turbo-wht.cu, line 71:

<comment>The new `k_turbo_wht_warp` duplicates the whole group-coordinate decode (`gid -> g/rem -> i01/i02` and the `src_base + i01*src_nb1 + i02*src_nb2 + g*QK_TQ3_0_GROUP` row-pointer math) plus the 12-argument signature and launch argument list that the scalar (`k_turbo_wht_scalar`) and fused-quantize (`k_turbo_wht_quantize`) kernels already own. Since the scalar and warp paths are chosen by `#if defined(GGML_USE_HIP)`, only one compiles on a given backend, so today this is behavior-neutral. The risk is maintenance drift: a future change to the group indexing or stride handling would have to be applied identically in three places or the HIP and CUDA paths would diverge silently. Consider factoring the common decode (params + gid->(g,i01,i02) + row pointers) into a shared helper/inline used by all three kernels to keep them in sync.</comment>

<file context>
@@ -41,6 +43,55 @@ static __global__ void k_turbo_wht(
+    const int64_t gid = (int64_t)blockIdx.x * (blockDim.x / warp_size_local) + warp;
+    if (gid >= total_groups) return;
+
+    const int64_t g   = gid % groups_per_row;
+    const int64_t rem = gid / groups_per_row;
+    const int64_t i01 = rem % ne01;
</file context>