Skip to content

Fix AppendNextTokensToSequences heap overflow #2111

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
Baiju Meswani (baijumeswani) merged 10 commits into from
Jul 16, 2026
Merged

Fix AppendNextTokensToSequences heap overflow #2111

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
38594fa
Fix AppendNextTokensToSequences heap overflow
apsonawane Apr 30, 2026
a34a7f9
address comments
apsonawane Apr 30, 2026
e8d3c72
Merge branch 'main' into asonawane/heap
apsonawane May 5, 2026
e571280
Merge branch 'main' into asonawane/heap
apsonawane May 20, 2026
465f01f
Cleanup
apsonawane May 20, 2026
e7cde34
Fix
apsonawane May 20, 2026
27fe4d7
Merge branch 'main' into asonawane/heap
apsonawane Jun 26, 2026
2a7bc97
Merge branch 'main' into asonawane/heap
apsonawane Jul 7, 2026
4b97bf4
Revert search.cpp and search_cuda.cpp changes; keep only generators.c…
apsonawane Jul 15, 2026
d3a3cd1
Merge branch 'main' into asonawane/heap
apsonawane Jul 16, 2026
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
16 changes: 10 additions & 6 deletions src/cuda/search_cuda.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -166,12 +166,12 @@ void GreedySearch_Cuda::SampleTopKTopP(int k, float p, float temperature) {

// Append tokens
CUDA_CHECK(cudaStreamSynchronize(GetStream()));
if (!*done_cpu_) {
if (!*done_cpu_ && sequences_.GetSequenceLength() < static_cast<size_t>(params_->search.max_length)) {
cuda::Launch_AppendNextTokensToSequences(next_tokens_buffer_.Span(), sequences_.GetSequences().Span(), params_->BatchBeamSize(), sequences_.GetSequenceLength(), sequences_.max_length_, GetStream());
sequences_.AfterAppendNextTokens(next_tokens_buffer_, params_->BatchBeamSize());
}

if (sequences_.GetSequenceLength() == params_->search.max_length) {
if (sequences_.GetSequenceLength() >= params_->search.max_length) {
Comment thread
apsonawane marked this conversation as resolved.
Outdated
if (GetLogItems().enabled && GetLogItems().hit_max_length)
Log("hit_max_length", "greedy cuda hit");
*done_cpu_ = true;
Expand Down Expand Up @@ -223,17 +223,21 @@ void GreedySearch_Cuda::AppendTokens(DeviceSpan<int32_t>& next_tokens) {
ResetDone();

auto next_tokens_gpu = next_tokens.Span();
cuda::Launch_AppendNextTokensToSequences(next_tokens_gpu, sequences_.GetSequences().Span(), params_->BatchBeamSize(), sequences_.GetSequenceLength(), sequences_.max_length_, GetStream());
sequences_.AfterAppendNextTokens(next_tokens, params_->BatchBeamSize());
if (sequences_.GetSequenceLength() < static_cast<size_t>(params_->search.max_length)) {
cuda::Launch_AppendNextTokensToSequences(next_tokens_gpu, sequences_.GetSequences().Span(), params_->BatchBeamSize(), sequences_.GetSequenceLength(), sequences_.max_length_, GetStream());
sequences_.AfterAppendNextTokens(next_tokens, params_->BatchBeamSize());
}

if (sequences_.GetSequenceLength() >= params_->search.max_length) {
if (sequences_.GetSequenceLength() >= static_cast<size_t>(params_->search.max_length)) {
if (GetLogItems().enabled && GetLogItems().hit_max_length)
Log("hit_max_length", "greedy cuda hit");
*done_cpu_ = true;
return;
}

ResetDone();
// Only reset done_ if buffer is not full
if (sequences_.GetSequenceLength() < static_cast<size_t>(params_->search.max_length))
ResetDone();
}

void BeamSearch_Cuda::AppendTokens(DeviceSpan<int32_t>& next_tokens) {
Expand Down
5 changes: 5 additions & 0 deletions src/generators.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -596,6 +596,11 @@ void Generator::GenerateNextToken() {

ThrowErrorIfSessionTerminated(state_->session_terminated_);

if (search_->GetSequenceLength() >= state_->params_->search.max_length)
throw std::runtime_error(
"GenerateNextToken called with sequence length already at max_length (" +
std::to_string(state_->params_->search.max_length) + ")");

Comment thread
apsonawane marked this conversation as resolved.
Outdated
// RNNT models: yield one token per call from the decoder state machine
if (is_nemotron_speech_model_) {
state_->SetExtraInputs(extra_inputs_);
Expand Down
30 changes: 26 additions & 4 deletions src/search.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -408,8 +408,17 @@ void GreedySearch_Cpu::SetNextToken(size_t batch_id, int32_t token) {

void GreedySearch_Cpu::AppendNextTokensToSequences() {
// Append next token to each sequence.
auto sequences_span = sequences_.GetSequences().CpuSpan();
auto current_length = sequences_.GetSequenceLength();

// Bounds check: prevent writing past the allocated sequences buffer.
if (current_length >= params_->search.max_length) {
if (g_log.enabled && g_log.hit_max_length)
Log("hit_max_length", "greedy cpu hit");
done_ = true;
return;
}

auto sequences_span = sequences_.GetSequences().CpuSpan();
auto next_tokens = next_tokens_ptr_.Span(); // always on cpu
auto batch_beam_size = params_->BatchBeamSize();
for (int i = 0; i < batch_beam_size; i++) {
Expand All @@ -436,9 +445,14 @@ void GreedySearch_Cpu::AppendTokens(DeviceSpan<int32_t>& next_tokens) {
SetNextToken(i, next_tokens_cpu[i * tokens_count_per_batch + j]);
}
AppendNextTokensToSequences();
if (done_)
break;
Comment thread
apsonawane marked this conversation as resolved.
Outdated
}

ResetDone();
// Preserve done_=true if we have filled the buffer; resetting it would allow
// a subsequent GenerateNextToken to bypass the !done_ guard and OOB-write.
if (sequences_.GetSequenceLength() < sequences_.max_length_)
ResetDone();
}

void GreedySearch_Cpu::RewindTo(size_t index) {
Expand Down Expand Up @@ -483,10 +497,18 @@ bool BeamSearch_Cpu::IsDone() const {
}

void BeamSearch_Cpu::AppendNextTokensToSequences() {
auto sequences_span = sequences_.GetSequences().CpuSpan();
auto sequences_next_span = sequences_.GetNextSequences().CpuSpan();
auto max_length = sequences_.max_length_;
auto current_length = sequences_.GetSequenceLength();

// Bounds check: prevent writing past the allocated sequences buffer.
if (current_length >= static_cast<size_t>(params_->search.max_length)) {
if (g_log.enabled && g_log.hit_max_length)
Log("hit_max_length", "beam cpu hit");
return;
}

auto sequences_span = sequences_.GetSequences().CpuSpan();
auto sequences_next_span = sequences_.GetNextSequences().CpuSpan();
auto batch_beam_next_tokens = beam_scorer_->GetNextTokens().Span();
auto batch_beam_indices = beam_scorer_->GetNextIndices().Span();
auto batch_beam_size = params_->BatchBeamSize();
Expand Down
Loading