Skip to content

Commit 6251100

Browse files
[Bugfix][SpecDecode] Fix rejection sampling spec length 1 in Triton (vllm-project#11458)
### What this PR does / why we need it? This PR optimizes the `rejection_greedy_sample_spec_len_1_triton` kernel by removing the loop and the helper function `bonus_renew_1`, replacing them with a vectorized comparison and masked store. It also squeezes `bonus_token_ids` in `rejection_greedy_sample_with_triton`. ### Does this PR introduce _any_ user-facing change? No. ### How was this patch tested? - vLLM version: v0.23.0 - vLLM main: vllm-project/vllm@1f486d9 --------- Signed-off-by: liyishi <1252651434@qq.com>
1 parent b417b95 commit 6251100

1 file changed

Lines changed: 4 additions & 23 deletions

File tree

vllm_ascend/ops/triton/reject_sample.py

Lines changed: 4 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -31,16 +31,6 @@ def cal_grid_and_block_size(batch_size: int):
3131
return grid, block_size
3232

3333

34-
@triton.jit(do_not_specialize=["max_spec_len"])
35-
def bonus_renew_1(
36-
bonus_token_ids_ptr,
37-
position,
38-
output_token_ids_ptr,
39-
):
40-
bonus_token_id = tl.load(bonus_token_ids_ptr + position)
41-
tl.store(output_token_ids_ptr + position * 2 + 1, bonus_token_id)
42-
43-
4434
@triton.jit(do_not_specialize=["max_spec_len"])
4535
def rejection_greedy_sample_spec_len_1_triton(
4636
output_token_ids_ptr, # [batch_size, 2]
@@ -56,21 +46,12 @@ def rejection_greedy_sample_spec_len_1_triton(
5646

5747
draft_token_id = tl.load(draft_token_ids_ptr + offset, mask)
5848
target_argmax_id = tl.load(target_argmax_ptr + offset, mask)
49+
bonus_token_id = tl.load(bonus_token_ids_ptr + offset, mask)
50+
5951
tl.store(output_token_ids_ptr + offset * 2, target_argmax_id, mask)
6052

61-
# Add validity check for pos within the loop
62-
for pos in tl.range(0, BLOCK_SIZE):
63-
# Calculate the global position of the current token
64-
global_pos = block_idx * BLOCK_SIZE + pos
65-
if global_pos < vec_len:
66-
draft_token_id1 = get_element(draft_token_id, (pos,))
67-
target_argmax1 = get_element(target_argmax_id, (pos,))
68-
if draft_token_id1 == target_argmax1:
69-
bonus_renew_1(
70-
bonus_token_ids_ptr,
71-
global_pos,
72-
output_token_ids_ptr,
73-
)
53+
accept_mask = (draft_token_id == target_argmax_id) & mask
54+
tl.store(output_token_ids_ptr + offset * 2 + 1, bonus_token_id, accept_mask)
7455

7556

7657
@triton.jit(do_not_specialize=["max_spec_len"])

0 commit comments

Comments
 (0)