quantization: add FP8 E4M3 W8A16 for gpt-oss MoE expert weights

Store expert gate_up_proj and down_proj weights in FP8 E4M3 (1 byte/elem)
with per-expert FP32 scale factors. At inference, a fused CUDA kernel
dequantizes to BF16 before the existing cuBLAS batched GEMM.

Results on gpt-oss-20b (50-problem GSM8K subset):
  - FP8 TP=1: 47/50 = 94.0% (single RTX 5090, ~25 GB VRAM)
  - BF16 TP=2: 47/50 = 94.0% (requires 2× RTX 5090, ~39 GB total)

No measurable accuracy degradation. Model size: 41.8 GB → 22.7 GB (−46%).

New files:
  - tools/quantize_fp8.py: offline BF16→FP8 conversion script
  - csrc/quantization/dequant_fp8.cu: per-expert-scale dequant kernel
  - crates/xserv-kernels/src/quantization.rs: Rust FFI wrapper
  - tools/eval_gsm8k_batch.sh: GSM8K accuracy evaluation harness

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

csrc/quantization/dequant_fp8.cu

@@ -0,0 +1,51 @@
+#include <cuda_bf16.h>
+#include <cuda_fp8.h>
+#include "../common.cuh"
+
+// Dequantize FP8 E4M3 → BF16 with per-expert (per-batch-slice) FP32 scale.
+//
+// Input:  src    [num_experts, rows, cols]  FP8 E4M3 (1 byte each)
+//         scales [num_experts]              FP32
+// Output: dst    [num_experts, rows, cols]  BF16
+//
+// Each element: dst[e, r, c] = bf16( float(src[e, r, c]) * scales[e] )
+
+__global__ void dequant_fp8e4m3_to_bf16_kernel(
+    const __nv_fp8_e4m3* __restrict__ src,
+    const float* __restrict__ scales,
+    __nv_bfloat16* __restrict__ dst,
+    int num_experts, int rows, int cols
+) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    int total = num_experts * rows * cols;
+    if (idx >= total) return;
+
+    int expert_stride = rows * cols;
+    int expert = idx / expert_stride;
+    float scale = scales[expert];
+    float val = float(src[idx]) * scale;
+    dst[idx] = __float2bfloat16(val);
+}
+
+extern "C" {
+
+void launch_dequant_fp8e4m3_to_bf16(
+    const void* src,
+    const void* scales,
+    void* dst,
+    int num_experts, int rows, int cols,
+    void* stream
+) {
+    int total = num_experts * rows * cols;
+    int block = 256;
+    int grid = (total + block - 1) / block;
+    dequant_fp8e4m3_to_bf16_kernel<<<grid, block, 0, (cudaStream_t)stream>>>(
+        (const __nv_fp8_e4m3*)src,
+        (const float*)scales,
+        (__nv_bfloat16*)dst,
+        num_experts, rows, cols
+    );
+    CUDA_CHECK_LAST_ERROR();
+}
+
+}