Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 11 additions & 7 deletions custom_ops/gpu_ops/noauxtc_kernel.h
Original file line number Diff line number Diff line change
Expand Up @@ -547,7 +547,7 @@ __global__ void group_idx_and_topk_idx_kernel(

if (case_id < num_tokens) {
// calculate group_idx
int32_t target_num_min = WARP_SIZE - n_group + topk_group;
int32_t want_neg_inf_num = WARP_SIZE - n_group + topk_group;
if (lane_id < n_group &&
(isfinite(cuda_cast<float, T>(
group_scores[lane_id])))) // The check is necessary to avoid
Expand All @@ -556,20 +556,24 @@ __global__ void group_idx_and_topk_idx_kernel(
value = group_scores[lane_id];
}

int count_equal_to_top_value = WARP_SIZE - n_group;
int pre_count_equal_to_top_value = 0;
int neg_inf_num = WARP_SIZE - n_group;
int last_neg_inf_num = 0;
// Use loop to find the largset top_group
while (count_equal_to_top_value < target_num_min) {
while (neg_inf_num < want_neg_inf_num) {
__syncwarp(); // Ensure all threads have valid data before reduction
topk_group_value = cg::reduce(tile, value, cg::greater<T>());
if (value == topk_group_value) {
value = neg_inf<T>();
}
pre_count_equal_to_top_value = count_equal_to_top_value;
count_equal_to_top_value =
last_neg_inf_num = neg_inf_num;

neg_inf_num =
__popc(__ballot_sync(FULL_WARP_MASK, (value == neg_inf<T>())));
}
num_equalto_topkth_group = target_num_min - pre_count_equal_to_top_value;
// There is a possible case:
// may have many different group holding the same score!
// but we only accept some of them!
num_equalto_topkth_group = want_neg_inf_num - last_neg_inf_num;
}
__syncthreads();

Expand Down
4 changes: 2 additions & 2 deletions tests/ci_use/EB_VL_Lite/test_EB_VL_Lite_serving.py
Original file line number Diff line number Diff line change
Expand Up @@ -205,9 +205,9 @@ def test_consistency_between_runs(api_url, headers, consistent_payload):
# base result
base_path = os.getenv("MODEL_PATH")
if base_path:
base_file = os.path.join(base_path, "ernie-4_5-vl-base-tp2-dev-1215")
base_file = os.path.join(base_path, "ernie-4_5-vl-base-tp2-dev")
else:
base_file = "ernie-4_5-vl-base-tp2-dev-1215"
base_file = "ernie-4_5-vl-base-tp2-dev"
with open(base_file, "r") as f:
content2 = f.read()

Expand Down
Loading
Loading