zero deepstack layers on slot when audio reuse

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:

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
 

lightllm/server/embed_cache/copy_to_cache.py

@@ -16,19 +16,30 @@ 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):
+    # 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)):
             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 +72,9 @@ 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=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],
         BLOCK=256,
         num_warps=4,