Skip to content

Commit a64afb4

Browse files
committed
[None][test] DSv4 PR-6 coverage and import safety
Signed-off-by: Fanrong Li <lfr-0531@users.noreply.github.com>
1 parent 552f462 commit a64afb4

11 files changed

Lines changed: 516 additions & 5 deletions

File tree

tensorrt_llm/_torch/attention_backend/flashinfer.py

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -34,9 +34,12 @@
3434
except RuntimeError:
3535
# Override TORCH_CUDA_ARCH_LIST for JIT compilation of flashinfer kernels
3636
# since the existed TORCH_CUDA_ARCH_LIST may be too general and flashinfer requires sm75+.
37-
capability = torch.cuda.get_device_capability()
38-
arch_list = f"{capability[0]}.{capability[1]}"
39-
os.environ["TORCH_CUDA_ARCH_LIST"] = arch_list
37+
# Guard on a visible GPU: with CUDA_VISIBLE_DEVICES="" (pure client) the
38+
# capability query would force a CUDA context at import time.
39+
if torch.cuda.is_available():
40+
capability = torch.cuda.get_device_capability()
41+
arch_list = f"{capability[0]}.{capability[1]}"
42+
os.environ["TORCH_CUDA_ARCH_LIST"] = arch_list
4043

4144
from tensorrt_llm._utils import prefer_pinned
4245

tensorrt_llm/_torch/attention_backend/triton_prefill.py

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,9 @@
3333
import triton
3434
import triton.language as tl
3535

36-
CUDA_CAPABILITY = torch.cuda.get_device_capability()
36+
# Guard on a visible GPU so `import tensorrt_llm` stays GPU-free under
37+
# CUDA_VISIBLE_DEVICES="" (pure client); (0, 0) is a safe sentinel there.
38+
CUDA_CAPABILITY = torch.cuda.get_device_capability() if torch.cuda.is_available() else (0, 0)
3739

3840

3941
def _get_block_sizes(Lq: int, Lv: int):

tensorrt_llm/_torch/cuda_tile_utils.py

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,13 @@ def ceil_div(a, b):
5050
except ImportError:
5151
logger.warning("cuda-tile package not found, TileIR kernels will not be available")
5252
else:
53-
if (cc := torch.cuda.get_device_properties()) and (cc.major, cc.minor) < (10, 0):
53+
# Guard the device-properties probe: with no visible GPU (e.g.
54+
# CUDA_VISIBLE_DEVICES="" on a pure client) it would force a CUDA
55+
# context just from `import tensorrt_llm`. TileIR stays unavailable,
56+
# which is correct for a GPU-less process.
57+
if not torch.cuda.is_available():
58+
logger.warning("No CUDA device visible, TileIR kernels will not be available")
59+
elif (cc := torch.cuda.get_device_properties()) and (cc.major, cc.minor) < (10, 0):
5460
logger.warning(
5561
f"TileIR requires compute capability 10.0 or higher, but the current device has "
5662
f"{cc.major}.{cc.minor}. TileIR kernels will not be available"

tests/integration/defs/accuracy/accuracy_core.py

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -261,6 +261,7 @@ def evaluate(self,
261261
f"Hypothesis testing report:\n{hypothesis_testing_params.report(score)}"
262262
)
263263
hypothesis_testing_params.assert_passing(score)
264+
return score
264265

265266

266267
class VoxPopuli(AccuracyTask):

tests/integration/defs/accuracy/references/gsm8k.yaml

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -140,6 +140,16 @@ deepseek-ai/DeepSeek-V4-Flash:
140140
# 95.11 reference still holds for the hypothesis test.
141141
- quant_algo: FP8_BLOCK_SCALES
142142
accuracy: 95.11
143+
deepseek-ai/DeepSeek-V4-Pro:
144+
# Full GSM8K aggregate gate for the Pro deployment path: TP=8, EP=8,
145+
# attention DP, TRTLLM MoE, FP8 KV cache, MTP max_draft_len=1, padded CUDA
146+
# graphs, custom DeepSeek-V4 tokenizer, and the same system prompt used by
147+
# the Pro GSM8K bench script. Measured 96.32 on the full Pro aggregate run
148+
# (GSM8K, 1319 samples); floor set slightly below for run-to-run margin.
149+
- quant_algo: FP8_BLOCK_SCALES
150+
kv_cache_quant_algo: FP8
151+
spec_dec_algo: MTP
152+
accuracy: 96.0
143153
Qwen3/Qwen3-4B:
144154
- spec_dec_algo: Eagle
145155
accuracy: 85.823

tests/integration/defs/accuracy/test_disaggregated_serving.py

Lines changed: 182 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2324,3 +2324,185 @@ def test_auto_dtype(self, use_py_transceiver, mocker):
23242324
with launch_disaggregated_llm(disagg_cfg, ctx_cfg, gen_cfg,
23252325
self.MODEL_PATH) as llm:
23262326
run_accuracy_test(llm, self.MODEL_NAME, ["GSM8K"])
2327+
2328+
2329+
@pytest.mark.timeout(DEFAULT_TEST_TIMEOUT)
2330+
@skip_pre_blackwell
2331+
class TestDeepSeekV4Flash(LlmapiAccuracyTestHarness):
2332+
MODEL_NAME = "deepseek-ai/DeepSeek-V4-Flash"
2333+
MODEL_PATH = f"{llm_models_root()}/DeepSeek-V4-Flash"
2334+
2335+
@pytest.mark.skip_less_device(4)
2336+
def test_auto_dtype(self):
2337+
# Disagg smoke test: CTX TP=2 + GEN TP=2 = 4 GPUs.
2338+
# NVFP4 weights ~71 GB/rank at TP=2, leaving ~107 GB for KV on B200.
2339+
# TRTLLM backend required (WIDEEP lacks MXFP4 support for V4-Flash).
2340+
# V4 uses pure-Python KVCacheManagerV2; needs Python transceiver.
2341+
# NIXL (not DEFAULT) skips the TRTLLM_USE_UCX_KVCACHE=1 fallback.
2342+
cache_transceiver_config = {
2343+
"backend": "NIXL",
2344+
"transceiver_runtime": "PYTHON",
2345+
"max_tokens_in_buffer": 4096,
2346+
}
2347+
ctx_server_config = {
2348+
"tensor_parallel_size": 2,
2349+
"moe_expert_parallel_size": 2,
2350+
"disable_overlap_scheduler": True,
2351+
"max_seq_len": 4096,
2352+
"kv_cache_config": {
2353+
"free_gpu_memory_fraction": 0.5,
2354+
},
2355+
"cache_transceiver_config": cache_transceiver_config,
2356+
}
2357+
gen_server_config = {
2358+
"tensor_parallel_size": 2,
2359+
"moe_expert_parallel_size": 2,
2360+
"enable_attention_dp": True,
2361+
"disable_overlap_scheduler": True,
2362+
"max_seq_len": 4096,
2363+
"moe_config": {
2364+
"backend": "TRTLLM",
2365+
},
2366+
"kv_cache_config": {
2367+
"free_gpu_memory_fraction": 0.5,
2368+
},
2369+
"cache_transceiver_config": cache_transceiver_config,
2370+
}
2371+
disaggregated_server_config = {
2372+
"hostname": "localhost",
2373+
"backend": "pytorch",
2374+
"context_servers": {
2375+
"num_instances": 1
2376+
},
2377+
"generation_servers": {
2378+
"num_instances": 1
2379+
},
2380+
}
2381+
# V4-Flash 148GB weight prefetch + warmup needs >35 min, default wait timeout times out.
2382+
with launch_disaggregated_llm(disaggregated_server_config,
2383+
ctx_server_config,
2384+
gen_server_config,
2385+
self.MODEL_PATH,
2386+
server_waiting_timeout=3600) as llm:
2387+
task = MMLU(self.MODEL_NAME)
2388+
task.evaluate(llm, is_integration_test=True)
2389+
2390+
@pytest.mark.skip_less_device(4)
2391+
def test_gen_first(self):
2392+
"""Gen-first quick validation for DSv4-Flash on KVCacheManagerV2 + NIXL python."""
2393+
cache_transceiver_config = {
2394+
"backend": "NIXL",
2395+
"transceiver_runtime": "PYTHON",
2396+
"max_tokens_in_buffer": 4096,
2397+
}
2398+
ctx_server_config = {
2399+
"tensor_parallel_size": 2,
2400+
"moe_expert_parallel_size": 2,
2401+
"disable_overlap_scheduler": True,
2402+
"max_seq_len": 4096,
2403+
"kv_cache_config": {
2404+
"free_gpu_memory_fraction": 0.5,
2405+
},
2406+
"cache_transceiver_config": cache_transceiver_config,
2407+
}
2408+
gen_server_config = {
2409+
"tensor_parallel_size": 2,
2410+
"moe_expert_parallel_size": 2,
2411+
"enable_attention_dp": True,
2412+
"disable_overlap_scheduler": True,
2413+
"max_seq_len": 4096,
2414+
"moe_config": {
2415+
"backend": "TRTLLM",
2416+
},
2417+
"kv_cache_config": {
2418+
"free_gpu_memory_fraction": 0.5,
2419+
},
2420+
"cache_transceiver_config": cache_transceiver_config,
2421+
}
2422+
disaggregated_server_config = {
2423+
"hostname": "localhost",
2424+
"backend": "pytorch",
2425+
"context_servers": {
2426+
"num_instances": 1
2427+
},
2428+
"generation_servers": {
2429+
"num_instances": 1
2430+
},
2431+
"schedule_style": "generation_first",
2432+
}
2433+
with launch_disaggregated_llm(disaggregated_server_config,
2434+
ctx_server_config,
2435+
gen_server_config,
2436+
self.MODEL_PATH,
2437+
server_waiting_timeout=3600) as llm:
2438+
task = MMLU(self.MODEL_NAME)
2439+
task.evaluate(llm, is_integration_test=True)
2440+
2441+
2442+
@pytest.mark.timeout(14400)
2443+
@skip_pre_blackwell
2444+
@pytest.mark.skip_less_device_memory(140000)
2445+
class TestDeepSeekV4FlashBase(LlmapiAccuracyTestHarness):
2446+
MODEL_NAME = "deepseek-ai/DeepSeek-V4-Flash-Base"
2447+
MODEL_PATH = f"{llm_models_root()}/DeepSeek-V4-Flash-Base"
2448+
2449+
@pytest.mark.skip_less_device(4)
2450+
def test_auto_dtype(self):
2451+
# Disagg smoke test: CTX TP=2 + GEN TP=2 = 4 GPUs.
2452+
# FP8 weights ~71 GB/rank at TP=4 → ~142 GB/rank at TP=2; requires
2453+
# ≥140 GB per GPU (fits on B300 288 GB, tight on B200 178 GB).
2454+
# TRTLLM backend: WIDEEP's FP8 block-scale path is Hopper-only.
2455+
# Compact batching keeps KV cache ~1 GB/rank (default ~100 GB requires fully-clean GPU memory).
2456+
# V4 uses pure-Python KVCacheManagerV2; needs Python transceiver.
2457+
# NIXL (not DEFAULT) skips the TRTLLM_USE_UCX_KVCACHE=1 fallback.
2458+
cache_transceiver_config = {
2459+
"backend": "NIXL",
2460+
"transceiver_runtime": "PYTHON",
2461+
"max_tokens_in_buffer": 4096,
2462+
}
2463+
ctx_server_config = {
2464+
"tensor_parallel_size": 2,
2465+
"moe_expert_parallel_size": 2,
2466+
"disable_overlap_scheduler": True,
2467+
"max_batch_size": 16,
2468+
"max_num_tokens": 4096,
2469+
"max_seq_len": 4096,
2470+
"kv_cache_config": {
2471+
"free_gpu_memory_fraction": 0.5,
2472+
},
2473+
"cache_transceiver_config": cache_transceiver_config,
2474+
}
2475+
gen_server_config = {
2476+
"tensor_parallel_size": 2,
2477+
"moe_expert_parallel_size": 2,
2478+
"enable_attention_dp": True,
2479+
"disable_overlap_scheduler": True,
2480+
"max_batch_size": 16,
2481+
"max_num_tokens": 4096,
2482+
"max_seq_len": 4096,
2483+
"moe_config": {
2484+
"backend": "TRTLLM",
2485+
},
2486+
"kv_cache_config": {
2487+
"free_gpu_memory_fraction": 0.5,
2488+
},
2489+
"cache_transceiver_config": cache_transceiver_config,
2490+
}
2491+
disaggregated_server_config = {
2492+
"hostname": "localhost",
2493+
"backend": "pytorch",
2494+
"context_servers": {
2495+
"num_instances": 1
2496+
},
2497+
"generation_servers": {
2498+
"num_instances": 1
2499+
},
2500+
}
2501+
# Same long-init reason as TestDeepSeekV4Flash above.
2502+
with launch_disaggregated_llm(disaggregated_server_config,
2503+
ctx_server_config,
2504+
gen_server_config,
2505+
self.MODEL_PATH,
2506+
server_waiting_timeout=3600) as llm:
2507+
task = MMLU(self.MODEL_NAME)
2508+
task.evaluate(llm, is_integration_test=True)

tests/integration/defs/accuracy/test_llm_api_pytorch.py

Lines changed: 65 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3449,6 +3449,41 @@ def _make_deepseekv4_eplb_config(model_path, layer_updates_per_iter, ep_size=8):
34493449
layer_updates_per_iter=0)
34503450

34513451

3452+
_DEEPSEEK_V4_GSM8K_SYSTEM_PROMPT = (
3453+
"Solve the problem carefully. End your response with a final line exactly "
3454+
"in the form #### <answer>, using the simplest numeric form without units "
3455+
"or trailing zeros.")
3456+
3457+
3458+
def _deepseekv4_pro_agg_llm_kwargs(**overrides):
3459+
kwargs = dict(
3460+
tensor_parallel_size=8,
3461+
moe_expert_parallel_size=8,
3462+
moe_config=MoeConfig(backend="TRTLLM"),
3463+
enable_attention_dp=True,
3464+
max_seq_len=4096,
3465+
max_batch_size=16,
3466+
max_num_tokens=8192,
3467+
custom_tokenizer="deepseek_v4",
3468+
# Cap the KV cache so the padded CUDA-graph capture (batch sizes up to
3469+
# 1024) and other post-KV resources have headroom; without this the
3470+
# default fraction (~0.9) leaves too little and engine init OOMs on
3471+
# large-memory GPUs. These tests need very little KV.
3472+
kv_cache_config=KvCacheConfig(enable_block_reuse=False,
3473+
tokens_per_block=128,
3474+
dtype="fp8",
3475+
host_cache_size=0,
3476+
free_gpu_memory_fraction=0.6),
3477+
cuda_graph_config=CudaGraphConfig(batch_sizes=[
3478+
1, 2, 4, 8, 16, 32, 64, 128, 192, 256, 320, 384, 448, 512, 768, 1024
3479+
],
3480+
enable_padding=True),
3481+
speculative_config=MTPDecodingConfig(max_draft_len=1),
3482+
)
3483+
kwargs.update(overrides)
3484+
return kwargs
3485+
3486+
34523487
def _run_deepseekv4_eplb(model_name,
34533488
model_path,
34543489
moe_backend,
@@ -3551,6 +3586,36 @@ def test_nvfp4_4gpus_online_eplb(self, moe_backend, mtp_nextn):
35513586
mtp_nextn=mtp_nextn)
35523587

35533588

3589+
@pytest.mark.timeout(14400)
3590+
@pytest.mark.skip_less_device(8)
3591+
@pytest.mark.skip_less_device_memory(140000)
3592+
@skip_pre_blackwell
3593+
class TestDeepSeekV4Pro(LlmapiAccuracyTestHarness):
3594+
MODEL_NAME = "deepseek-ai/DeepSeek-V4-Pro"
3595+
MODEL_PATH = f"{llm_models_root()}/DeepSeek-V4-Pro"
3596+
EXTRA_EVALUATOR_KWARGS = dict(
3597+
apply_chat_template=True,
3598+
system_prompt=_DEEPSEEK_V4_GSM8K_SYSTEM_PROMPT,
3599+
)
3600+
3601+
@pytest.mark.skip_less_mpi_world_size(8)
3602+
def test_gsm8k_full_accuracy(self):
3603+
with LLM(self.MODEL_PATH, **_deepseekv4_pro_agg_llm_kwargs()) as llm:
3604+
task = GSM8K(self.MODEL_NAME)
3605+
acc_params = task.get_hypothesis_testing_params(
3606+
dtype=llm.args.dtype,
3607+
quant_algo=llm.args.quant_config.quant_algo,
3608+
kv_cache_quant_algo=llm.args.quant_config.kv_cache_quant_algo,
3609+
spec_dec_algo=llm.args.speculative_config.decoding_type)
3610+
assert acc_params.num_samples == GSM8K.NUM_SAMPLES
3611+
with mock.patch.dict(os.environ, {"INTEGRATION_TEST": "0"}):
3612+
score = task.evaluate(
3613+
llm, extra_evaluator_kwargs=self.EXTRA_EVALUATOR_KWARGS)
3614+
assert score >= acc_params.ref_accuracy, (
3615+
f"GSM8K accuracy {score:.3f} is below recorded reference "
3616+
f"{acc_params.ref_accuracy:.3f}")
3617+
3618+
35543619
@pytest.mark.timeout(14400)
35553620
@pytest.mark.skip_less_device_memory(140000)
35563621
@skip_pre_blackwell

0 commit comments

Comments
 (0)