nvidia_agent_tasks: qwen38_fp4 KDA solution

[Waterpine]

Hi BBuf,

Here is the KDA solution for qwen38_nvfp4__fp4_w4a4_skinny_gemm. Could you please take a look and let me know if you have any feedback? Thanks!

[BBuf]

Validated PR head 516c976 on an NVIDIA RTX PRO 6000 Blackwell Server Edition (SM120, GPU 0).

Benchmark setup:

• task-provided tools/bench_harness.py
• CUDA graph timing via triton.testing.do_bench
• cold L2, 200 iterations, 7 interleaved A/B trials
• all 16 workload rows passed the declared correctness gates

Summary:

• task geomean: 1.1086x–1.1197x across three complete sweeps; median 1.1136x
• real-call-weighted aggregate: 1.1224x (about 10.9% lower aggregate latency for the targeted kernels)
• a fully version-aligned FlashInfer 0.6.17 control run reproduced the result: 1.1175x geomean / 1.1225x weighted

Median FP4 GEMM speedups across the three sweeps:

GEMM	M=1 decode	M=9 DSpark verify	M=4369 prefill
MLP gate/up	1.185x	1.171x	1.012x
MLP down	1.360x	1.409x	1.089x
LM head	1.091x	1.086x	1.009x

The main gains are therefore on the skinny decode/verify shapes, especially MLP down at M=1/M=9. The standalone fp4_quantize rows are around 6 us and show timer/P-state quantization; M=8 fluctuated in complete sweeps, but five focused repeats were effectively neutral at about 1.00x, so I did not observe a stable regression there.

Overall: approximately 1.12x targeted-kernel speedup on RTX PRO 6000, with correctness passing on 16/16 rows.