Skip to content
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

Fix possible overflow in WriteCoalescingCallbackWrapper::TearDown #17642

Merged
merged 7 commits into from
Jan 3, 2025
Merged

Fix possible overflow in WriteCoalescingCallbackWrapper::TearDown #17642

Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
4a29d82
Fix possible overflow in WriteCoalescingCallbackWrapper::TearDown
davidwendt Dec 20, 2024
04ac7d9
Merge branch 'branch-25.02' into fst-overflow
davidwendt Dec 21, 2024
0bbfee2
Merge branch 'branch-25.02' into fst-overflow
PointKernel Dec 30, 2024
c6e5444
fix docs-build experiment
davidwendt Jan 2, 2025
7f7f0dd
change copyright year for style check
davidwendt Jan 2, 2025
7fed58c
test cmake rmm-logger change too
davidwendt Jan 2, 2025
4ea06f4
Merge remote-tracking branch 'upstream/branch-25.02' into fst-overflow
vyasr Jan 3, 2025
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
8 changes: 5 additions & 3 deletions cpp/src/io/fst/agent_dfa.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2024, NVIDIA CORPORATION.
* Copyright (c) 2022-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -308,12 +308,14 @@ class WriteCoalescingCallbackWrapper {
{
__syncthreads();
if constexpr (!DiscardTranslatedOutput) {
for (uint32_t out_char = threadIdx.x; out_char < tile_out_count; out_char += blockDim.x) {
for (thread_index_type out_char = threadIdx.x; out_char < tile_out_count;
out_char += blockDim.x) {
out_it[tile_out_offset + out_char] = temp_storage.compacted_symbols[out_char];
}
}
if constexpr (!DiscardIndexOutput) {
for (uint32_t out_char = threadIdx.x; out_char < tile_out_count; out_char += blockDim.x) {
for (thread_index_type out_char = threadIdx.x; out_char < tile_out_count;
out_char += blockDim.x) {
out_idx_it[tile_out_offset + out_char] =
temp_storage.compacted_offset[out_char] + tile_in_offset;
}
Expand Down
Loading