Skip to content

Add DeepSeek V4 Pro B200 workload#45

Open
khluu wants to merge 4 commits into
mainfrom
codex/add-deepseek-v4-pro-b200
Open

Add DeepSeek V4 Pro B200 workload#45
khluu wants to merge 4 commits into
mainfrom
codex/add-deepseek-v4-pro-b200

Conversation

@khluu

@khluu khluu commented Jul 22, 2026

Copy link
Copy Markdown
Member

This PR was authored with assistance from Codex.

Summary

  • add an 8x B200 DeepSeek V4 Pro NVFP4 workload
  • follow the official Blackwell TP8 + expert-parallel topology with FP8 KV cache, FP4 indexer cache, and full/piecewise CUDA graphs
  • retain the existing DeepSeek GSM8K, BFCL, and random serving benchmark coverage
  • pin TileLang to the image's system CUDA toolkit so its JIT compiler and headers stay version-matched
  • skip the pinned image's generic FA4/MLA JIT warmup, which is incompatible with DeepSeek V4's CSA/HCA configuration; model-specific warmup and FlashInfer autotuning remain enabled

Recipe: https://recipes.vllm.ai/deepseek-ai/DeepSeek-V4-Pro

Local validation

  • parser smoke test with a stubbed lm_eval registry
  • bash -n lib/run.sh lib/server.sh lib/run_lm_eval.sh lib/run_vllm_bench.sh
  • python3 .buildkite/test_generate_pipeline.py (6/6 passed)
  • explicit WORKLOADS=deepseek_v4_pro_b200 pipeline generation
  • git diff --check

GPU validation

  • Buildkite #313: acquired 8x B200, downloaded and loaded the 64-shard NVFP4 checkpoint, then failed when TileLang selected the pip CUDA compiler package and hit incompatible CUDA compiler/toolkit headers while compiling mhc_pre_big_fuse_broadcast_with_norm_tilelang
  • commit 7e9eb94 sets CUDA_HOME=/usr/local/cuda, selecting the CUDA 13.0 compiler and headers installed together in the vLLM image
  • Buildkite #318: all TileLang MHC kernels compiled successfully, validating the CUDA fix; startup then failed in generic fa4_cutedsl_warmup() because DeepSeek V4 intentionally has no legacy MLA qk_nope_head_dim or v_head_dim
  • Buildkite #320: DP8 + EP reproduced the identical generic warmup failure, ruling out parallel topology as the cause
  • commit eb441cc restores the recipe-default TP8 + EP topology and sets kernel_config.enable_jit_warmup=false, skipping only the incompatible generic FA4/MLA warmup
  • Buildkite #321: passed in 34m00s on 8x B200; model-specific warmup, FlashInfer autotuning, CUDA graph capture, server health, the 512-request serving benchmark, GSM8K, and all five BFCL categories completed successfully

khluu and others added 2 commits July 22, 2026 01:39
Follow the official NVFP4 Blackwell recipe and retain the existing accuracy, BFCL, and serving benchmark coverage.

Co-Authored-By: Codex <noreply@anthropic.com>
Signed-off-by: khluu <khluu000@gmail.com>
Point TileLang at the CUDA 13.0 toolkit installed in the vLLM image so its NVCC compiler and CCCL headers remain version-matched.\n\nCo-Authored-By: Codex <noreply@anthropic.com>

Signed-off-by: khluu <khluu000@gmail.com>

khluu commented Jul 22, 2026

Copy link
Copy Markdown
Member Author

Buildkite #313 reached the DeepSeek V4 MHC kernel warmup after loading the 64-shard checkpoint, then failed consistently on all ranks while compiling mhc_pre_big_fuse_broadcast_with_norm_tilelang:

CUDA compiler and CUDA toolkit headers are incompatible
Command: /usr/local/lib/python3.12/dist-packages/nvidia/cu13/bin/nvcc ...

TileLang's discovery order selected the pip nvidia-cuda-nvcc package before the CUDA toolkit already installed by the vLLM image. Commit 7e9eb94 sets CUDA_HOME=/usr/local/cuda so TileLang uses the image's matching system NVCC and headers.

khluu commented Jul 22, 2026

Copy link
Copy Markdown
Member Author

Buildkite #318 has now crossed the exact #313 failure boundary. All TileLang MHC kernels, including mhc_pre_big_fuse_broadcast_with_norm_tilelang, completed compilation across the 8 ranks with CUDA_HOME=/usr/local/cuda; the run is continuing through DeepSeek V4 sparse-MLA warmup. This validates the compiler/toolkit-path fix while full workload validation continues.

Co-Authored-By: Codex <noreply@anthropic.com>
Signed-off-by: khluu <khluu000@gmail.com>

khluu commented Jul 22, 2026

Copy link
Copy Markdown
Member Author

Buildkite #318 validated the CUDA-toolkit fix, then exposed a separate topology issue: TP8 reduced DeepSeek V4's MLA dimensions to qk_nope_head_dim=0 and v_head_dim=0, so backend selection correctly failed. Commit aaab28e switches to the official DP8 + EP deployment topology, which leaves the attention dimensions intact and distributes the experts across all 8 B200s. Local parser/generator/shell checks remain green; replacement GPU validation is next.

khluu commented Jul 22, 2026

Copy link
Copy Markdown
Member Author

Replacement DP8 + EP validation is running in Buildkite #320 on aaab28e.

Co-Authored-By: Codex <noreply@anthropic.com>
Signed-off-by: khluu <khluu000@gmail.com>

khluu commented Jul 22, 2026

Copy link
Copy Markdown
Member Author

Buildkite #320 reproduced the same zero-dimension error with TP=1 / DP=8, ruling out topology. The exact stack is kernel_warmup -> fa4_cutedsl_warmup -> get_mla_prefill_backend; DeepSeek V4 replaces legacy MLA dimensions with head_dim and CSA/HCA, so this generic warmup sees intentional missing fields as zeros. Commit eb441cc restores recipe-default TP8 + EP and sets kernel_config.enable_jit_warmup=false, skipping the incompatible FA4/MLA warmup while retaining the model-specific sparse-attention warmup and FlashInfer autotuning.

khluu commented Jul 22, 2026

Copy link
Copy Markdown
Member Author

Replacement validation is running in Buildkite #321 on eb441cc.

khluu commented Jul 22, 2026

Copy link
Copy Markdown
Member Author

Buildkite #321 passed in 34m00s on eb441cc using 8x B200 with TP8 + EP. The run completed model load, DeepSeek V4 sparse-attention warmup, FlashInfer autotuning, mixed/full CUDA graph capture, and server health. Workload stages all passed: the 512-request 8K/1K serving benchmark (8m45s), GSM8K (2m35s), and BFCL simple_python, multiple, parallel, parallel_multiple, and multi_turn.

@khluu
khluu marked this pull request as ready for review July 22, 2026 10:49
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant