From 111ec8ed65d4519116d8eb346586112adb2fc173 Mon Sep 17 00:00:00 2001 From: maruoheng Date: Fri, 16 Jan 2026 02:27:08 +0000 Subject: [PATCH] kernel fix --- .../src/kernel/kunlun3cpp/text_image_gather_scatter.xpu | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/custom_ops/xpu_ops/src/plugin/src/kernel/kunlun3cpp/text_image_gather_scatter.xpu b/custom_ops/xpu_ops/src/plugin/src/kernel/kunlun3cpp/text_image_gather_scatter.xpu index 777af5491c8..8de0e76b567 100644 --- a/custom_ops/xpu_ops/src/plugin/src/kernel/kunlun3cpp/text_image_gather_scatter.xpu +++ b/custom_ops/xpu_ops/src/plugin/src/kernel/kunlun3cpp/text_image_gather_scatter.xpu @@ -60,9 +60,11 @@ static __device__ inline void text_image_gather( if (token_type == 0) { text_image_input = text_input; text_image_index = text_index; - } else { + } else if (token_type == 1) { text_image_input = image_input; text_image_index = image_index; + } else { + continue; } GM2LM(text_image_index + i, &text_image_token_idx, sizeof(int)); int input_offset = i * hidden_size; @@ -132,9 +134,11 @@ static __device__ inline void text_image_scatter( if (token_type == 0) { text_image_input = text_input; text_image_index = text_index; - } else { + } else if (token_type == 1) { text_image_input = image_input; text_image_index = image_index; + } else { + continue; } GM2LM(text_image_index + i, &text_image_token_idx, sizeof(int)); int input_offset = i * hidden_size;