test(deepseek4): CPU-parity test for HC pre-mix kernels

[cheese-cakee]

What

Adds a CPU-parity unit test for the DeepSeek-V4 Hierarchical-Controller pre-mix kernels (server/src/deepseek4/deepseek4_hc_cuda.cu), registered as the deepseek4_hc_cuda ctest. HIP-only, no runtime code touched.

Why

The HC pre step is hand-written CUDA/HIP with block reductions whose width follows the wavefront, and it had no correctness backstop — the sibling test_rms_norm_hip covers only RMSNorm. This guards the HC path (mix matvec + NHC=4 Sinkhorn + gated combine) against reduction and wave-width regressions on AMD the same way the RMSNorm test does: it compares the GPU kernels against a CPU reference on the device's actual wavefront width, so wave32 cards (gfx1100 / gfx1151) and any future wave64 card each validate the mode they run.

How

• server/test/test_deepseek4_hc_cuda.cpp drives deepseek4_cuda_hc_pre with the DS4-Flash config (n_hc = 4, so mix_dim = 24; n_embd = 512 so hc_dim = 2048 spans several strided load iterations in both reductions).
• The CPU reference reproduces both device kernels exactly:
  • hc_mix_norm_kernel — per-row RMS-normalized dot fn[row] . state.
  • hc_finish_kernel — sigmoid pre/post gates, per-dst row-softmax followed by the sinkhorn_iters normalization sweeps, then the gated combine working[d] = sum_h pre[h] * state[h][d].
• fn is round-tripped through __half on the host so both sides consume identical weights; the check covers all three outputs (working, post, comb).
• The CMake target mirrors the existing test_rms_norm_hip block (LANGUAGE HIP, includes src + hip_compat, links dflash_common) and prints the device warpSize so the run doubles as a wave-width report.

Verification

Compiled the test together with its kernel translation unit directly (hipcc --offload-arch=gfx1151) and ran on gfx1151 (Radeon 8060S, wave32):

[ds4-hc-test] device=Radeon 8060S Graphics warpSize=32 n_hc=4 n_embd=512 sinkhorn_iters=4
[ds4-hc-test]   working  max_abs=1.490e-07
[ds4-hc-test]   post     max_abs=1.192e-07
[ds4-hc-test]   comb     max_abs=8.941e-08
[ds4-hc-test] max_abs_diff=1.490e-07 max_rel_diff=1.717e-05
[ds4-hc-test] PASS

Tolerance is 1e-4 (matching test_rms_norm_hip), ~3 orders of magnitude above the observed margin, leaving headroom for wave64 reduction ordering. The deepseek4_hc_cuda CMake target is wired verbatim from the merged test_rms_norm_hip target; a full CMake build of the test was not run in this pass.

[cubic-dev-ai]

No issues found across 2 files

_{Re-trigger cubic}

[cubic-dev-ai]

All reported issues were addressed across 1 file (changes from recent commits).

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

Re-trigger cubic</sub>