modelscope / modelscope/ms-swift

参考Qwen3.5最佳实践GKD训练,vLLM启动模型 FlashInfer JIT 构建失败

Open
#8,349 4 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

bug stale
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

Open the contributing guide

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.