Open nvdbaranec opened 2 years ago
The fix in basically all of these cases is quite simple: just make the index a size_t
I'd love to just add an algorithm to do this.
Or even a simple range helper for for
loops: https://github.com/harrism/hemi#simple-grid-stride-loops
The fix in basically all of these cases is quite simple: just make the index a
size_t
I think the general approach should be:
This issue has been labeled inactive-30d
due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d
if there is no activity in the next 60 days.
Still relevant.
Created https://github.com/harrism/ranger as a solution to this. Needs to be moved into libcudf.
This issue has been labeled inactive-30d
due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d
if there is no activity in the next 60 days.
Still relevant.
This issue has been labeled inactive-30d
due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d
if there is no activity in the next 60 days.
Still relevant.
This issue has been labeled inactive-90d
due to no recent activity in the past 90 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed.
Thanks @harrism for creating the ranger repo! Do you think we are ready to kick off an integration with libcudf, or does ranger need more development first?
@GregoryKimball I have now created a PR to use ranger in libcuspatial. You guys could use this as an example if you want to do the same in libcudf. https://github.com/rapidsai/cuspatial/pull/1178
wrt attempting to find locations where this might be happening. In host code, clang and gcc will warn if you add -Wsign-conversion
(not covered by -Wall -Wextra
) under some circumstances. Unfortunately there is no such option for nvcc.
#include <cstdint>
int what(int upper)
{
int i = 0; // no warning if this is a std::int64_t
unsigned int stride = 10;
while (i < upper) {
i = i + stride; // clang warns for this, so does gcc
}
i = 0;
while (i < upper) {
i += stride; // gcc warns for this, clang does not.
}
return i;
}
(updated Aug 2023)
Background
We found a kernel indexing overflow issue, first discovered in the
fused_concatenate
kernels (https://github.com/rapidsai/cudf/issues/10333) and this issue is present in a number of our CUDA kernels that take the following form:If we have an output_size of say 1.2 billion and a grid size that's the same, the following happens: Some late thread id, say 1.19 billion attempts to add 1.2 billion (blockDim.x * gridDim.x) and overflows the size_type (signed 32 bits).
We made a round of fixes in #10448, and then later found another instance of this error in #13838. Our first pass of investigation was not adequate to contain the issue, so we need to take another close look.
Part 1 - First pass fix kernels with this issue
copying/concatenate.cu
fused_concatenate_kernel
valid_if.cuh
valid_if_kernel
scatter.cu
marking_bitmask_kernel
replace/nulls.cu
replace_nulls_strings
replace/nulls.cu
replace_nulls
rolling/rolling_detail.cuh
gpu_rolling
rolling/jit/kernel.cu
gpu_rolling_new
transform/compute_column.cu
compute_column_kernel
copying/concatenate.cu
fused_concatenate_string_offset_kernel
replace/replace.cu
replace_strings_first_pass
replace_strings_second_pass
replace_kernel
copying/concatenate.cu
concatenate_masks_kernel
fused_concatenate_string_offset_kernel
fused_concatenate_string_chars_kernel
fused_concatenate_kernel
(int64)hash/helper_functions.cuh
init_hashtbl
null_mask.cu
set_null_mask_kernel
copy_offset_bitmask
count_set_bits_kernel
transform/row_bit_count.cu
compute_row_sizes
multibyte_split.cu
multibyte_split_init_kernel
multibyte_split_seed_kernel
(auto??)multibyte_split_kernel
io/utilities/parsing_utils.cu
count_and_set_positions
(uint64_t)conditional_join_kernels.cuh
compute_conditional_join_output_size
conditional_join
merge.cu
materialize_merged_bitmask_kernel
partitioning.cu
compute_row_partition_numbers
compute_row_output_locations
copy_block_partitions
json_path.cu
get_json_object_kernel
tdigest
compute_percentiles_kernel
(int)strings/attributes.cu
count_characters_parallel_fn
strings/convert/convert_urls.cu
url_decode_char_counter
(int)url_decode_char_replacer
(int)text/subword/data_normalizer.cu
kernel_data_normalizer
(uint32_t)text/subword/subword_tokenize.cu
kernel_compute_tensor_metadata
(uint32_t)text/subword/wordpiece_tokenizer.cu
init_data_and_mark_word_start_and_ends
(uint32_t)mark_string_start_and_ends
(uint32_t)kernel_wordpiece_tokenizer
(uint32_t)Part 2 - Take another pass over more challenging kernels
gridDim.x
orblockDim.x
to find more examplesPart 3 - Use ranger to prevent grid stride loop overflow
Additional information
There are also a number of kernels that have this pattern but probably don't ever overflow because they are indexing by bitmask words. (Example) Additional, In this kernel,
source_idx
probably overflows, but harmlessly.A snippet of code to see this in action:
Note: rmm may mask out of bounds accesses in some cases, so it's helpful to run with the plain cuda allocator.