From bc4d38dc370a3ebe743bcc5c17581aa5cd73de6b Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Tue, 24 Oct 2023 18:13:57 -0400 Subject: [PATCH] fixing thread index overflow issue (#14290) Ran across this issue during the review of https://github.com/NVIDIA/spark-rapids-jni/pull/1502 and since the code was modeled after this code, I am pushing the fix here as well. Authors: - Mike Wilson (https://github.com/hyperbolic2346) Approvers: - David Wendt (https://github.com/davidwendt) - Nghia Truong (https://github.com/ttnghia) URL: https://github.com/rapidsai/cudf/pull/14290 --- cpp/src/strings/convert/convert_urls.cu | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/cpp/src/strings/convert/convert_urls.cu b/cpp/src/strings/convert/convert_urls.cu index 9e847131be2..511acc38d75 100644 --- a/cpp/src/strings/convert/convert_urls.cu +++ b/cpp/src/strings/convert/convert_urls.cu @@ -212,7 +212,8 @@ __global__ void url_decode_char_counter(column_device_view const in_strings, char* in_chars_shared = temporary_buffer[local_warp_id]; // Loop through strings, and assign each string to a warp. - for (size_type row_idx = global_warp_id; row_idx < in_strings.size(); row_idx += nwarps) { + for (thread_index_type tidx = global_warp_id; tidx < in_strings.size(); tidx += nwarps) { + auto const row_idx = static_cast(tidx); if (in_strings.is_null(row_idx)) { out_counts[row_idx] = 0; continue; @@ -296,7 +297,8 @@ __global__ void url_decode_char_replacer(column_device_view const in_strings, char* in_chars_shared = temporary_buffer[local_warp_id]; // Loop through strings, and assign each string to a warp - for (size_type row_idx = global_warp_id; row_idx < in_strings.size(); row_idx += nwarps) { + for (thread_index_type tidx = global_warp_id; tidx < in_strings.size(); tidx += nwarps) { + auto const row_idx = static_cast(tidx); if (in_strings.is_null(row_idx)) continue; auto const in_string = in_strings.element(row_idx);