modelscope / modelscope/ms-swift
参考Qwen3.5最佳实践GKD训练,vLLM启动模型 FlashInfer JIT 构建失败
Nobody has claimed this yet.
- Dominant language
- Python
- Stars
- 15.7k
- Forks
- 1.7k
- Avg merge
- 1d 16h
- Merged PRs (30d)
- 136
Description
Checklist / 检查清单
- I have searched existing issues, and this is a new bug report. / 我已经搜索过现有的 issues,确认这是一个新的 bug report。
Bug Description / Bug 描述
参考Qwen3.5最佳实践构建环境,按照GKD训练的例子用视频数据训练qwen3.5-4B, vllm在运行 Qwen3.5时,FlashInfer JIT 构建失败。(H100 gpu)
报错为:
ERROR 03-16 08:35:10 [dump_input.py:79] Dumping scheduler output for model execution: SchedulerOutput(scheduled_new_reqs=[NewRequestData(req_id=0-b692e5ec,prompt_token_ids_len=1538,prefill_token_ids_len=None,mm_features=[MultiModalFeat
ureSpec(data={'video_grid_thw': MultiModalFieldElem(data=tensor([21, 14, 18]), field=MultiModalBatchedField(keep_on_cpu=True)), 'pixel_values_videos': MultiModalFieldElem(data=tensor([[ 1.0000, 1.0000, 1.0000, ..., 1.0000, 1.0000,
1.0000],
ERROR 03-16 08:35:10 [dump_input.py:79] [ 1.0000, 1.0000, 1.0000, ..., 1.0000, 1.0000, 1.0000],
ERROR 03-16 08:35:10 [dump_input.py:79] [ 0.9844, 0.9844, 0.9844, ..., 0.9844, 0.9844, 0.9844],
ERROR 03-16 08:35:10 [dump_input.py:79] ...,
ERROR 03-16 08:35:10 [dump_input.py:79] [-0.9219, -0.9219, -0.9219, ..., 1.0000, 1.0000, 1.0000],
ERROR 03-16 08:35:10 [dump_input.py:79] [-0.9219, -0.9219, -0.9219, ..., -1.0000, -1.0000, -1.0000],
ERROR 03-16 08:35:10 [dump_input.py:79] [-0.9219, -0.9219, -0.9219, ..., -1.0000, -1.0000, -1.0000]],
ERROR 03-16 08:35:10 [dump_input.py:79] dtype=torch.bfloat16), field=MultiModalFlatField(keep_on_cpu=False, slices=[(slice(0, tensor(5292), None),)], dim=0)), 'timestamps': MultiModalFieldElem(data=[0.25, 1.2833333333333334, 2.3
16666666666667, 3.35, 4.383333333333334, 5.416666666666667, 6.45, 7.466666666666667, 8.5, 9.533333333333335, 10.566666666666666, 11.600000000000001, 12.633333333333333, 13.666666666666668, 14.683333333333334, 15.716666666666667, 16.75,
17.783333333333335, 18.816666666666666, 19.85, 20.883333333333333], field=MultiModalBatchedField(keep_on_cpu=True))}, modality='video', identifier='78d10e5d945460e1164abb06b46f1b03a4e3f62e8904eeb90d98ebbbcd39b3ed', mm_position=Placeho
lderRange(offset=3, length=1502, is_embed=tensor([False, False, False, ..., True, True, False])), mm_hash='78d10e5d945460e1164abb06b46f1b03a4e3f62e8904eeb90d98ebbbcd39b3ed')],sampling_params=SamplingParams(n=1, presence_penalty=0.0,
frequency_penalty=0.0, repetition_penalty=1.0, temperature=0.9, top_p=0.9, top_k=50, min_p=0.0, seed=None, stop=['<|endoftext|>', '<|im_end|>'], stop_token_ids=[248044, 248046], bad_words=[], include_stop_str_in_output=False, ignore_e
os=False, max_tokens=31961, min_tokens=0, logprobs=0, prompt_logprobs=None, skip_special_tokens=True, spaces_between_special_tokens=True, structured_outputs=None, extra_args=None),block_ids=([0, 1], [0, 2], [0, 3], [4, 5]),num_computed
_tokens=0,lora_request=None,prompt_embeds_shape=None), NewRequestData(req_id=1-aaa222bb,prompt_token_ids_len=10562,prefill_token_ids_len=None,mm_features=[MultiModalFeatureSpec(data={'video_grid_thw': MultiModalFieldElem(data=tensor([1
40, 22, 12]), field=MultiModalBatchedField(keep_on_cpu=True)), 'pixel_values_videos': MultiModalFieldElem(data=tensor([[ 0.3574, 0.3574, 0.3574, ..., 0.3496, 0.3496, 0.3496],
7 errors detected in the compilation of "/root/.cache/flashinfer/0.6.4/90a/generated/gdn_prefill_sm90/gdn_prefill_kernel_bf16_g0b1a0i1.cu".
[33/34] /usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /root/.cache/flashinfer/0.6.4/90a/cached_ops/gdn_prefill_sm90/gdn_prefill_kernel_bf16_g0b0a1i0.cuda.o.d -DPy_LIMITED_API=0x03090000 -D_GLIBCXX_US
E_CXX11_ABI=1 -isystem /root/miniforge3/envs/qwen3_5/include/python3.10 -isystem /usr/local/cuda/include -isystem /usr/local/cuda/include/cccl -isystem /root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/tvm_ffi/include -isystem
/root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/tvm_ffi/include -isystem /root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/include -isystem /root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages
/flashinfer/data/csrc -isystem /root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/cutlass/include -isystem /root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/cutlass/tools/util/include -i
system /root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/spdlog/include --compiler-options=-fPIC --expt-relaxed-constexpr -DFLASHINFER_ENABLE_FP8_E8M0 -DFLASHINFER_ENABLE_FP4_E2M1 --threads=1 -use_fast_math -DF
LASHINFER_ENABLE_F16 -DFLASHINFER_ENABLE_BF16 -DFLASHINFER_ENABLE_FP8_E4M3 -DFLASHINFER_ENABLE_FP8_E5M2 -DNDEBUG -O3 -gencode=arch=compute_90a,code=sm_90a -DCUTE_SM90_EXTENDED_MMA_SHAPES_ENABLED -DFLASHINFER_ENABLE_FP8_E8M0 -DFLASHINFE
R_ENABLE_FP4_E2M1 -DFLAT_SM90A_ENABLED -std=c++20 -c /root/.cache/flashinfer/0.6.4/90a/generated/gdn_prefill_sm90/gdn_prefill_kernel_bf16_g0b0a1i0.cu -o /root/.cache/flashinfer/0.6.4/90a/cached_ops/gdn_prefill_sm90/gdn_prefill_kernel_b
f16_g0b0a1i0.cuda.o
FAILED: [code=2] /root/.cache/flashinfer/0.6.4/90a/cached_ops/gdn_prefill_sm90/gdn_prefill_kernel_bf16_g0b0a1i0.cuda.o
/usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /root/.cache/flashinfer/0.6.4/90a/cached_ops/gdn_prefill_sm90/gdn_prefill_kernel_bf16_g0b0a1i0.cuda.o.d -DPy_LIMITED_API=0x03090000 -D_GLIBCXX_USE_CXX11_
ABI=1 -isystem /root/miniforge3/envs/qwen3_5/include/python3.10 -isystem /usr/local/cuda/include -isystem /usr/local/cuda/include/cccl -isystem /root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/tvm_ffi/include -isystem /root/m
iniforge3/envs/qwen3_5/lib/python3.10/site-packages/tvm_ffi/include -isystem /root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/include -isystem /root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashin
fer/data/csrc -isystem /root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/cutlass/include -isystem /root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/cutlass/tools/util/include -isystem /
root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/spdlog/include --compiler-options=-fPIC --expt-relaxed-constexpr -DFLASHINFER_ENABLE_FP8_E8M0 -DFLASHINFER_ENABLE_FP4_E2M1 --threads=1 -use_fast_math -DFLASHINFE
R_ENABLE_F16 -DFLASHINFER_ENABLE_BF16 -DFLASHINFER_ENABLE_FP8_E4M3 -DFLASHINFER_ENABLE_FP8_E5M2 -DNDEBUG -O3 -gencode=arch=compute_90a,code=sm_90a -DCUTE_SM90_EXTENDED_MMA_SHAPES_ENABLED -DFLASHINFER_ENABLE_FP8_E8M0 -DFLASHINFER_ENABLE
_FP4_E2M1 -DFLAT_SM90A_ENABLED -std=c++20 -c /root/.cache/flashinfer/0.6.4/90a/generated/gdn_prefill_sm90/gdn_prefill_kernel_bf16_g0b0a1i0.cu -o /root/.cache/flashinfer/0.6.4/90a/cached_ops/gdn_prefill_sm90/gdn_prefill_kernel_bf16_g0b0
a1i0.cuda.o
/root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/include/flashinfer/flat/hopper/collective/flat_collective_store.hpp(251): error: namespace "cuda::ptx" has no member "tensormap_replace_global_dim"
ptx::tensormap_replace_global_dim(ptx::space_global, tensormap, ptx::n32_t<1>{},
^
/root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/include/flashinfer/flat/hopper/collective/flat_collective_store.hpp(251): error: namespace "cuda::ptx" has no member "n32_t"
ptx::tensormap_replace_global_dim(ptx::space_global, tensormap, ptx::n32_t<1>{},
^
/root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/include/flashinfer/flat/hopper/collective/flat_collective_store.hpp(251): error: expected an expression
ptx::tensormap_replace_global_dim(ptx::space_global, tensormap, ptx::n32_t<1>{},
^
/root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/include/flashinfer/flat/hopper/collective/flat_collective_store.hpp(256): error: namespace "cuda::ptx" has no member "fence_proxy_tensormap_generic"
ptx::fence_proxy_tensormap_generic(ptx::sem_release, ptx::scope_cta);
^
/root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/include/flashinfer/flat/hopper/collective/flat_collective_store.hpp(262): error: namespace "cuda::ptx" has no member "fence_proxy_tensormap_generic"
ptx::fence_proxy_tensormap_generic(ptx::sem_acquire, ptx::scope_cta, tensormap,
^
/root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/include/flashinfer/flat/hopper/collective/flat_collective_store.hpp(263): error: namespace "cuda::ptx" has no member "n32_t"
ptx::n32_t<128>{});
^
/root/miniforge3/envs/qwen3_5/lib/python3.10/site-packages/flashinfer/data/include/flashinfer/flat/hopper/collective/flat_collective_store.hpp(263): error: expected an expression
ptx::n32_t<128>{});
^
7 errors detected in the compilation of "/root/.cache/flashinfer/0.6.4/90a/generated/gdn_prefill_sm90/gdn_prefill_kernel_bf16_g0b0a1i0.cu".
ninja: build stopped: subcommand failed.
How to Reproduce / 如何复现
PYTORCH_CUDA_ALLOC_CONF='expandable_segments:True'
IMAGE_MAX_TOKEN_NUM=8100
VIDEO_MIN_TOKEN_NUM=64
VIDEO_MAX_TOKEN_NUM=64
VIDEO_MAX_PIXELS=50176
MAX_PIXELS=1003520
FPS_MAX_FRAMES=360
MASTER_PORT=29501
DECORD_EOF_RETRY_MAX=20480
NPROC_PER_NODE=6
CUDA_VISIBLE_DEVICES=0,1,2,3,4,5
swift rlhf
--rlhf_type gkd
--split_dataset_ratio 0.01
--model /model_zoo/Qwen3.5-4B/
--teacher_model /model_zoo/Qwen3.5-9B/
--gkd_logits_topk 64
--enable_thinking false
--tuner_type lora
--use_vllm true
--vllm_mode colocate
--vllm_gpu_memory_utilization 0.5
--vllm_tensor_parallel_size 1
--vllm_max_model_len 32000
--sleep_level 0
--lmbda 1
--seq_kd false
--beta 0.5
--torch_dtype bfloat16
--per_device_train_batch_size 2
--gradient_accumulation_steps 16
--learning_rate 5e-5
--logging_steps 1
--save_steps 100
--save_total_limit 10
--max_length 32000
--max_completion_length 32000
--warmup_ratio 0.1
--save_only_model true
--dataloader_num_workers 4 \
环境:
torch 2.10.0
transformers 5.3.0
vllm 0.17.1
flash_attn 2.8.3
flash-linear-attention 0.4.2
flashinfer-cubin 0.6.4
flashinfer-python 0.6.4
Additional Information / 补充信息
No response
Contributor guide
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Reproduce the GKD Qwen3.5 video-training example on an H100 and inspect the FlashInfer 0.6.4 JIT output under ~/.cache/flashinfer/0.6.4/90a/generated/gdn_prefill_sm90. Start with the reported nvcc command and flashinfer/flat/hopper/collective/flat_collective_store.hpp errors. Done means vLLM starts Qwen3.5 without the FlashInfer kernel compilation failure.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- python
- Domain
- build-system, machine-learning
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 35/100