Skip to content

Commit

Permalink
Fix possible int overflow in compute_mixed_join_output_size (#17633)
Browse files Browse the repository at this point in the history
Fixes possible integer overflow condition when the number of rows is near max int32 in `compute_mixed_join_output_size` kernel function.
Reference #10368

Authors:
  - David Wendt (https://github.com/davidwendt)

Approvers:
  - Bradley Dice (https://github.com/bdice)
  - Vukasin Milovanovic (https://github.com/vuule)

URL: #17633
  • Loading branch information
davidwendt authored Dec 20, 2024
1 parent 0f1bae8 commit f9f5f7d
Show file tree
Hide file tree
Showing 2 changed files with 4 additions and 4 deletions.
6 changes: 3 additions & 3 deletions cpp/src/join/mixed_join_size_kernel.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -62,8 +62,8 @@ CUDF_KERNEL void __launch_bounds__(block_size)
intermediate_storage + (threadIdx.x * device_expression_data.num_intermediates);

std::size_t thread_counter{0};
cudf::size_type const start_idx = threadIdx.x + blockIdx.x * block_size;
cudf::size_type const stride = block_size * gridDim.x;
auto const start_idx = cudf::detail::grid_1d::global_thread_id();
auto const stride = cudf::detail::grid_1d::grid_stride();
cudf::size_type const left_num_rows = left_table.num_rows();
cudf::size_type const right_num_rows = right_table.num_rows();
auto const outer_num_rows = (swap_tables ? right_num_rows : left_num_rows);
Expand All @@ -80,7 +80,7 @@ CUDF_KERNEL void __launch_bounds__(block_size)
auto count_equality = pair_expression_equality<has_nulls>{
evaluator, thread_intermediate_storage, swap_tables, equality_probe};

for (cudf::size_type outer_row_index = start_idx; outer_row_index < outer_num_rows;
for (auto outer_row_index = start_idx; outer_row_index < outer_num_rows;
outer_row_index += stride) {
auto query_pair = pair_func(outer_row_index);
if (join_type == join_kind::LEFT_JOIN || join_type == join_kind::FULL_JOIN) {
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/json/json_path.cu
Original file line number Diff line number Diff line change
Expand Up @@ -928,7 +928,7 @@ __launch_bounds__(block_size) CUDF_KERNEL
get_json_object_options options)
{
auto tid = cudf::detail::grid_1d::global_thread_id();
auto const stride = cudf::thread_index_type{blockDim.x} * cudf::thread_index_type{gridDim.x};
auto const stride = cudf::detail::grid_1d::grid_stride();

size_type warp_valid_count{0};

Expand Down

0 comments on commit f9f5f7d

Please sign in to comment.