From 51c07b83ab96cb8411d34fd7e98403a8712d7bb7 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Thu, 19 Dec 2024 13:06:49 -0500 Subject: [PATCH] Fix possible int overflow in compute_mixed_join_output_size --- cpp/src/join/mixed_join_size_kernel.cuh | 6 +++--- cpp/src/json/json_path.cu | 2 +- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/src/join/mixed_join_size_kernel.cuh b/cpp/src/join/mixed_join_size_kernel.cuh index 4049ccf35e1..98170ed719a 100644 --- a/cpp/src/join/mixed_join_size_kernel.cuh +++ b/cpp/src/join/mixed_join_size_kernel.cuh @@ -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); @@ -80,7 +80,7 @@ CUDF_KERNEL void __launch_bounds__(block_size) auto count_equality = pair_expression_equality{ 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) { diff --git a/cpp/src/json/json_path.cu b/cpp/src/json/json_path.cu index fd8629ed6f3..e6e01b9c9fe 100644 --- a/cpp/src/json/json_path.cu +++ b/cpp/src/json/json_path.cu @@ -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};