From 310201ca58b4b56e9e756a1a28b3c5ec01e83a9b Mon Sep 17 00:00:00 2001 From: wanzihao <1060304770@qq.com> Date: Wed, 22 Jul 2026 21:13:59 +0800 Subject: [PATCH 1/2] Refactor layer handling in copy_to_cache function clear stale deepstack layers on slot reuse --- lightllm/server/embed_cache/copy_to_cache.py | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/lightllm/server/embed_cache/copy_to_cache.py b/lightllm/server/embed_cache/copy_to_cache.py index 9f275f79a2..1dac6e6d28 100644 --- a/lightllm/server/embed_cache/copy_to_cache.py +++ b/lightllm/server/embed_cache/copy_to_cache.py @@ -16,19 +16,23 @@ def _offload_embed_tensor_to_cache( cpu_stride1, cpu_stride2, start_index_in_cache, - layer_num, + source_layer_num: tl.constexpr, + cache_layer_num: tl.constexpr, hidden_size, BLOCK: tl.constexpr, ): token_index = tl.program_id(0).to(tl.int64) dest_index = (start_index_in_cache + token_index).to(tl.int64) - for layer_index in range(layer_num): + for layer_index in range(cache_layer_num): + layer_mask = layer_index < source_layer_num for block_index in range(tl.cdiv(hidden_size, BLOCK)): off = block_index * BLOCK + tl.arange(0, BLOCK) mask = off < hidden_size gpu_data = tl.load( - embed_tensor_ptr + token_index * gpu_stride0 + layer_index * gpu_stride1 + off * gpu_stride2, mask=mask + embed_tensor_ptr + token_index * gpu_stride0 + layer_index * gpu_stride1 + off * gpu_stride2, + mask=mask & layer_mask, + other=0.0, ) tl.store( cache_tensor_ptr + dest_index * cpu_stride0 + layer_index * cpu_stride1 + off * cpu_stride2, @@ -61,7 +65,8 @@ def offload_embed_tensor_to_cache( cpu_stride1=cache_tensor.stride(1), cpu_stride2=cache_tensor.stride(2), start_index_in_cache=start_index_in_cache, - layer_num=embed_tensor.shape[1], + source_layer_num=embed_tensor.shape[1], + cache_layer_num=cache_tensor.shape[1], hidden_size=embed_tensor.shape[2], BLOCK=256, num_warps=4, From cfded361c470eb08e1703402df96e2336779f46e Mon Sep 17 00:00:00 2001 From: wangzaijun Date: Mon, 27 Jul 2026 02:32:20 +0000 Subject: [PATCH 2/2] add log --- lightllm/models/qwen3_vl/layer_infer/pre_layer_infer.py | 4 ++++ .../qwen3_vl/triton_kernel/deepstack_multimodal_emb.py | 4 ++++ lightllm/server/embed_cache/copy_to_cache.py | 8 ++++++++ 3 files changed, 16 insertions(+) diff --git a/lightllm/models/qwen3_vl/layer_infer/pre_layer_infer.py b/lightllm/models/qwen3_vl/layer_infer/pre_layer_infer.py index 6be827ac0a..f6d0b29ab5 100644 --- a/lightllm/models/qwen3_vl/layer_infer/pre_layer_infer.py +++ b/lightllm/models/qwen3_vl/layer_infer/pre_layer_infer.py @@ -23,6 +23,10 @@ def context_forward( hidden_size = layer_weight.wte_weight_.weight.shape[1] for batch_id, p in enumerate(infer_state.multimodal_params): + # images 与 audios 共用后续 multimodal_emb / apply_deepstack 的定位信息。 + # embedding 只取 cache[:, 0, :]，音频写入后这里是正确的；deepstack 阶段会 + # 再按同一批 loc 叠加 cache[:, 1:, :]，因此音频占用的 slot 在写入时必须把 + # 无用的 deepstack 层清零（见 offload_embed_tensor_to_cache），否则会叠加上脏数据。 for img in p["images"] + p["audios"]: # skip the same image if img["token_id"] in img_start_token_ids: diff --git a/lightllm/models/qwen3_vl/triton_kernel/deepstack_multimodal_emb.py b/lightllm/models/qwen3_vl/triton_kernel/deepstack_multimodal_emb.py index 94e46c035f..fb7e98677f 100644 --- a/lightllm/models/qwen3_vl/triton_kernel/deepstack_multimodal_emb.py +++ b/lightllm/models/qwen3_vl/triton_kernel/deepstack_multimodal_emb.py @@ -101,6 +101,10 @@ def apply_deepstack_features( """ apply deepstack features for all images in qwen3-vl/qwen3-vl-moe """ + # deepstack 层数取自 cache.shape[1]-1（Qwen3 为 3），并按 pre_layer 收集的 + # img_*（实际含 images+audios）对对应 token 叠加 cache[:, layer_num+1, :]。 + # 音频本身没有 deepstack，语义上叠加量应为 0；依赖 copy_to_cache 在音频写入 + # 时将 cache[:, 1:, :] 置零，使误走此路径时数值上等价于不叠加。 deepstack_num_layers = infer_state.cpu_embed_cache_tensor.shape[1] - 1 diff --git a/lightllm/server/embed_cache/copy_to_cache.py b/lightllm/server/embed_cache/copy_to_cache.py index 1dac6e6d28..9e8839fbca 100644 --- a/lightllm/server/embed_cache/copy_to_cache.py +++ b/lightllm/server/embed_cache/copy_to_cache.py @@ -24,6 +24,13 @@ def _offload_embed_tensor_to_cache( token_index = tl.program_id(0).to(tl.int64) dest_index = (start_index_in_cache + token_index).to(tl.int64) + # Qwen3-VL/Omni 的 embed cache 为 [token, layer, H]，layer=4： + # cache[:, 0, :] 为主 embedding，cache[:, 1:4, :] 为 vision deepstack。 + # 音频只有主 embedding（source_layer_num=1），embedding 阶段只读 + # cache[:, 0, :]，本身不受影响；但 transformer 里 apply_deepstack 会对 + # images+audios 的 cache 位置统一叠加 cache[:, 1:4, :]。若 slot 曾被图像 + # 占用，只覆盖第 0 层会留下旧 deepstack，音频 token 就会被错误叠加。 + # 因此按 cache_layer_num 写满：有效源层正常拷贝，超出 source 的层用 0 填充。 for layer_index in range(cache_layer_num): layer_mask = layer_index < source_layer_num for block_index in range(tl.cdiv(hidden_size, BLOCK)): @@ -65,6 +72,7 @@ def offload_embed_tensor_to_cache( cpu_stride1=cache_tensor.stride(1), cpu_stride2=cache_tensor.stride(2), start_index_in_cache=start_index_in_cache, + # 音频 source=1、vision+deepstack cache=4；kernel 将多出的层写 0。 source_layer_num=embed_tensor.shape[1], cache_layer_num=cache_tensor.shape[1], hidden_size=embed_tensor.shape[2],