mirror of
https://github.com/kvcache-ai/ktransformers.git
synced 2026-11-05 01:07:28 +00:00
[GH-ISSUE #2001] fixed: DeepSeek-V4-Flash outputs all-zero tokens on single RTX PRO 4000 Blackwell (SM_120, 24GB) #1230
Labels
No milestone
No project
No assignees
1 participant
Notifications
Due date
No due date set.
Dependencies
No dependencies set.
Reference
kvcache-ai/ktransformers#1230
Loading…
Add table
Add a link
Reference in a new issue
No description provided.
Delete branch "%!s()"
Deleting a branch is permanent. Although the deleted branch may continue to exist for a short time before it actually gets removed, it CANNOT be undone in most cases. Continue?
Originally created by @lizhucui02-arch on GitHub (May 13, 2026).
Original GitHub issue: https://github.com/kvcache-ai/ktransformers/issues/2001
Reminder
System Info
Hardware:
GPU: NVIDIA RTX PRO 4000 Blackwell, 24GB GDDR7 ECC, SM_120 (Compute Capability 12.0)
CPU: Intel Xeon, 64 physical cores / 128 threads, 4 NUMA nodes, AMX supported
RAM: 256GB+
Driver: 590.48.01, CUDA 13.1
Software:
kt-kernel: 0.6.2.post3
sglang-kt: 0.6.2.post3
transformers: 4.57.1
flashinfer-python: 0.6.11
tilelang: 0.1.9
torch: 2.9.1+cu128
Reproduction
Launch Command
export FLASHINFER_CUDA_ARCH_LIST=12.0a
export TORCH_CUDA_ARCH_LIST="12.0+PTX"
export SGLANG_DSV4_MODE=2604
export SGLANG_DSV4_2604_SUBMODE=2604B
numactl --interleave=all python -m sglang.launch_server
--host 0.0.0.0 --port 30000
--model /data/models/DeepSeek-V4-Flash
--kt-weight-path /data/models/DeepSeek-V4-Flash
--kt-method MXFP4
--kt-num-gpu-experts 10
--kt-cpuinfer 60
--kt-threadpool-count 2
--kt-gpu-prefill-token-threshold 4096
--kt-enable-dynamic-expert-update
--tensor-parallel-size 1
--context-length 16384
--attention-backend flashinfer
--mem-fraction-static 0.85
--chunked-prefill-size 2048
--max-prefill-tokens 2048
--max-running-requests 2
--watchdog-timeout 1200
--disable-shared-experts-fusion
--trust-remote-code
--cuda-graph-bs 1
--cuda-graph-max-bs 1
--disable-radix-cache
--skip-server-warmup
Behavior
Server starts successfully and accepts requests. With temperature=0, all output tokens are 0:
{"text":"","output_ids":[0,0,0,...,0]}
With temperature>0, the scheduler crashes with torch.multinomial device-side assert (NaN in logits).
Root Cause Analysis (Detailed Investigation)
After extensive debugging, the root cause is in the compressed attention backend's Lightning Indexer (fp8_paged_mqa_logits).
Step 1: Confirmed NaN is NOT from MoE experts
Setting --kt-num-gpu-experts 0 (all experts on CPU, no matmul_ogs called) still produces all-zero output. This rules out matmul_ogs / MXFP4 GPU expert path as the cause.
Step 2: Confirmed NaN is NOT from tilelang MHC kernel
Tested hc_split_sinkhorn directly — outputs correct values, no NaN.
Step 3: Identified root cause: fp8_paged_mqa_logits in compressed attention
The compressed attention backend's Lightning Indexer calls get_paged_mqa_logits_metadata / fp8_paged_mqa_logits. The kt-sglang fork has two fallback paths for non-DeepGEMM hardware:
tilelang path (SGLANG_OPT_USE_TILELANG_INDEXER=1 or auto on SM_120): tilelang_fp8_paged_mqa_logits — produces NaN on SM_120
torch reference path (SGLANG_FP8_PAGED_MQA_LOGITS_TORCH=1): fp8_paged_mqa_logits_torch — also produces NaN on SM_120
Both fallback paths produce incorrect results (all-zero output / NaN logits) on SM_120.
Step 4: Verified fp8_paged_mqa_logits_torch has a suspicious F.relu
In
indexer.py
:
score = F.linear(kvcache_value, q)
score = F.relu(score) # ← This relu should not be here for attention scores
score *= q_scale[None, :]
score = score.sum(dim=1)
score *= kvcache_scale
The F.relu truncates all negative attention scores to 0, causing the attention distribution to collapse. Removing it does not fix the NaN (the underlying computation is still incorrect), but it indicates the torch reference implementation may have logic errors.
Step 5: Confirmed _patch_strided_mxfp() is correctly triggered
_use_strided_layout(): True # SM_120 correctly detected
has_native_mxfp(): False # correctly patched
is_persistent constraint: False
_v4_assert_patched: True
All patches in v4_triton_kernels_moe.py are applied correctly. The issue is deeper in the compressed attention forward pass.
Related Issues
ktransformers #1988: RTX 5090D single card assert num_stages >= 1 (same root area)
sglang #23657: No SM120 fallback for Lightning Indexer (upstream sglang)
DeepGEMM #317: Missing SM120 kernels for paged_mqa_logits
Request
Fix fp8_paged_mqa_logits_torch reference implementation — the F.relu on attention scores appears incorrect
Fix tilelang_fp8_paged_mqa_logits to produce correct results on SM_120
Add end-to-end validation on SM_120 single card (RTX 5090 / RTX PRO 4000 / RTX PRO 6000)
The hardware (RTX PRO 4000 Blackwell, SM_120) is listed as "Validated" in the tutorial table, but produces all-zero output on every inference request.
Others
No response
@lhw365 commented on GitHub (May 13, 2026):
Have you tried removing these directories under ~/.cache/ before building from source?
ccache flashinfer torch_extensions tvm-ffi
@fairydreaming commented on GitHub (May 14, 2026):
I'm also trying to run DeepSeek V4 Flash on Blackwell (RTX PRO 6000 Max-Q on Epyc 9374F) but in my case (I followed the DeepSeek V4 Flash tutorial) sglang starts (
The server is fired up and ready to roll!) and queues requests (I can see"POST /v1/chat/completions HTTP/1.1" 200 OKfrom kt chat) but nothing happens apart fromAbort queued request. req.rid='8b5dde99f6fe4716aefd60ba95ffb71a'when I finally abort kt chat. :(I'm a bit lost how to debug this since there are multiple processes involved. Any advice? Tried what @lhw365 posted but didn't help. @lizhucui02-arch did you notice similar symptoms initially, maybe you already know the solution?
@yyj6666667 commented on GitHub (May 14, 2026):
github.com/kvcache-ai/sglang@ebaff7729brecently we update kt-sglang and that may help solving the launch hanging problem.
If you would like to give it a try, please run ./install.sh under ktransformers to install it or directly install the newest kt-sglang from source(notice it's not on the wheel for sm120)
@yyj6666667 commented on GitHub (May 14, 2026):
I'll try my best to reproduce and solve it.
@fairydreaming commented on GitHub (May 14, 2026):
I managed to solve my issue with hanging requests - it was because of
--max-running-requestsset to 1. When I set it to 2 like in the original example everything started working normally.By the way I tried DeepSeek V4 Pro and it works too!
@trilog-inc commented on GitHub (May 14, 2026):
hi @fairydreaming, what was the RAM and VRAM usage like?
@yyj6666667 commented on GitHub (May 14, 2026):
thanks! i already locate the --max-running-requests bug, it will be update with the next post, glad to see you succeed running it!
@fairydreaming commented on GitHub (May 14, 2026):
nvidia-smi shows 83261MiB / 97887MiB VRAM used after start, 85467MiB during request processing.
top shows 79.3% of memory used after start (I have 12 x 96GB), 80% during request processing.
llama-benchy results: