xserv-chat: support gpt-oss-20b with TP; fix GEMV precision bug

- Add ChatModel enum dispatching between Qwen3 and GptOss based on
  config.is_moe(), following the TP engine pattern.
- Add --tp N flag for tensor-parallel inference (required for 39GB
  gpt-oss-20b which doesn't fit on a single 32GB GPU).
- Add gpt-oss harmony chat template with channel/message format.
- Replace hardcoded is_stop_token() with tokenizer.is_eos() for
  multi-model EOS support.
- Restore gpt-oss hardcoded prompt template in server api.rs, lost
  during the Jinja template refactor.
- Fix GEMV race condition: the K-split kernel zeroed the FP32
  accumulator inside the kernel (block k=0) while other blocks
  atomicAdd'd concurrently. Pre-zero with cudaMemsetAsync instead.
- Update benchmark docs with post-fix results.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

csrc/gemm/gemv.cu

@@ -2,16 +2,14 @@
 #include <cuda_runtime.h>
 #include "../common.cuh"
 
-// K-split GEMV for M=1 BF16 decode, fully self-contained (single launch).
+// K-split GEMV for M=1 BF16 decode.
 //
 //   y[n] = sum_k x[k] * W[k * N + n]
 //
 // Grid: (N / TILE_N, K / TILE_K).
-// Block k=0 for each column group initializes the FP32 accumulator to 0.
-// All blocks atomicAdd their partial sums. Block k=last converts FP32→BF16.
-//
-// This replaces the old 3-launch pattern (cudaMemsetAsync + gemv + convert)
-// with a single kernel launch while preserving the K-split occupancy.
+// All blocks atomicAdd their partial sums into a pre-zeroed FP32 buffer.
+// A separate conversion kernel writes the final BF16 output.
+// Launch sequence: cudaMemsetAsync(fp32) → accumulation kernel → convert kernel.
 
 #define GEMV_TILE_N 128
 #define GEMV_TILE_K 256
@@ -32,11 +30,6 @@ __global__ void gemv_bf16_fused_kernel(
 
     if (col >= N) return;
 
-    // First K-block: zero the accumulator
-    if (block_k == 0) {
-        y_fp32[col] = 0.0f;
-    }
-
     const int k_start = block_k * GEMV_TILE_K;
     const int k_end = min(k_start + GEMV_TILE_K, K);
     const int k_len = k_end - k_start;
@@ -53,15 +46,6 @@ __global__ void gemv_bf16_fused_kernel(
     }
 
     atomicAdd(&y_fp32[col], sum);
-
-    // Last K-block: convert FP32 → BF16
-    // We need a grid-level sync between the accumulation and the conversion.
-    // Since blocks within a grid-y column don't synchronize, we use a
-    // completion counter per column group.
-    // Simpler approach: just let the host launch the conversion separately.
-    // ... Actually for correctness with atomicAdd we need ALL k-blocks to
-    // finish before converting. We can't know when that happens from within
-    // the kernel without cooperative groups. Fall back to 2-kernel approach.
 }
 
 // Conversion kernel: FP32 accumulator -> BF16 output
@@ -88,6 +72,11 @@ void launch_gemv_bf16(
 ) {
     cudaStream_t s = (cudaStream_t)stream;
 
+    // Zero the FP32 accumulator BEFORE the kernel — the kernel uses atomicAdd
+    // across K-blocks with no inter-block ordering, so the buffer must be
+    // pre-zeroed to avoid accumulating on stale data.
+    cudaMemsetAsync(y_fp32_buf, 0, (size_t)N * sizeof(float), s);
+
     int num_k_blocks = (K + GEMV_TILE_K - 1) / GEMV_TILE_K;
     dim3 grid((N + GEMV_TILE_N - 1) / GEMV_TILE_N, num_k_blocks);