8000 [CSV 1/9] Fix out-of-bounds reads, false NA matches and unchecked column indexes in the CSV reader by karthikeyann · Pull Request #16 · karthikeyann/cudf · GitHub
[go: up one dir, main page]
Skip to content

[CSV 1/9] Fix out-of-bounds reads, false NA matches and unchecked column indexes in the CSV reader #16

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

Draft
karthikeyann wants to merge 6 commits into
base: coral/csv-base
Choose a base branch
from
Draft

[CSV 1/9] Fix out-of-bounds reads, false NA matches and unchecked column indexes in the CSV reader #16

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
0271a46
Fix out-of-bounds reads when parsing malformed timestamps
karthikeyann Sep 24, 2026
83970af
Fix reads past the end of fields when parsing durations
karthikeyann Sep 24, 2026
6f29aa3
Fix false matches in serialized trie lookups of longer keys
karthikeyann Sep 24, 2026
2a900c9
Support non-ASCII keys in serialized tries
karthikeyann Sep 24, 2026
e12237e
Fix negative string length for a lone quote field in the CSV reader
karthikeyann Sep 24, 2026
66b7b12
Validate column indexes of CSV reader column options
karthikeyann Sep 24, 2026
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
23 changes: 23 additions & 0 deletions cpp/include/cudf/io/csv.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -564,6 +564,12 @@ class csv_reader_options {
/**
* @brief Sets indexes of columns to read.
*
* `read_csv` throws `std::out_of_range` if an index is negative or does not match a column of the
* input. If the input (or byte range) contains no row, its columns are those named by `names`:
* if there are as many names as selected columns, the result has these columns, without rows;
* other names name all columns, and an index must match one of them; without names, the result
* has no columns.
*
* @param col_indices Vector of column indices that are needed
*/
void set_use_cols_indexes(std::vector<int> col_indices)
Expand Down Expand Up @@ -747,6 +753,8 @@ class csv_reader_options {
/**
* @brief Sets indexes of columns to read as datetime.
*
* Indexes that do not match a column of the input are ignored.
*
* @param col_indices Vector of column indices to infer as datetime
*/
void set_parse_dates(std::vector<int> col_indices)
Expand All @@ -767,6 +775,8 @@ class csv_reader_options {
/**
* @brief Sets indexes of columns to parse as hexadecimal
*
* Indexes that do not match a column of the input are ignored.
*
* @param col_indices Vector of column indices to parse as hexadecimal
*/
void set_parse_hex(std::vector<int> col_indices) { _parse_hex_indexes = std::move(col_indices); }
Expand Down Expand Up @@ -959,6 +969,12 @@ class csv_reader_options_builder {
/**
* @brief Sets indexes of columns to read.
*
* `read_csv` throws `std::out_of_range` if an index is negative or does not match a column of the
* input. If the input (or byte range) contains no row, its columns are those named by `names`:
* if there are as many names as selected columns, the result has these columns, without rows;
* other names name all columns, and an index must match one of them; without names, the result
* has no columns.
*
* @param col_indices Vector of column indices that are needed
* @return this for chaining
*/
Expand Down Expand Up @@ -1192,6 +1208,8 @@ class csv_reader_options_builder {
/**
* @brief Sets indexes of columns to read as datetime.
*
* Indexes that do not match a column of the input are ignored.
*
* @param col_indices Vector of column indices to read as datetime
* @return this for chaining
*/
Expand All @@ -1216,6 +1234,8 @@ class csv_reader_options_builder {
/**
* @brief Sets indexes of columns to parse as hexadecimal.
*
* Indexes that do not match a column of the input are ignored.
*
* @param col_indices Vector of column indices to parse as hexadecimal
* @return this for chaining
*/
Expand Down Expand Up @@ -1358,6 +1378,9 @@ class csv_reader_options_builder {
* auto result = cudf::io::read_csv(options);
* @endcode
*
* @throw std::out_of_range if an index in `use_cols_indexes` is negative or does not match a column
* of the input
*
* @param options Settings for controlling reading behavior
* @param stream CUDA stream used for device memory operations and kernel launches
* @param mr Device memory resource used to allocate device memory of the table in the returned
Expand Down
8 changes: 6 additions & 2 deletions cpp/src/io/csv/csv_gpu.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -365,16 +365,20 @@ CUDF_KERNEL void __launch_bounds__(csvparse_block_dim)
auto end = next_delimiter;
bool was_quoted = false;
if (not options.keepquotes) {
// A quoted string needs an opening and a closing quote; a field consisting of a single
// quote character (an unterminated quote at the end of the input) is kept as is
if (not options.detect_whitespace_around_quotes) {
if ((*field_start == options.quotechar) && (*(end - 1) == options.quotechar)) {
if (end - field_start >= 2 && (*field_start == options.quotechar) &&
(*(end - 1) == options.quotechar)) {
++field_start;
--end;
was_quoted = true;
}
} else {
// If the string is quoted, whitespace around the quotes get removed as well
auto const trimmed_field = trim_whitespaces(field_start, end);
if ((*trimmed_field.first == options.quotechar) &&
if (trimmed_field.second - trimmed_field.first >= 2 &&
(*trimmed_field.first == options.quotechar) &&
(*(trimmed_field.second - 1) == options.quotechar)) {
field_start = trimmed_field.first + 1;
end = trimmed_field.second - 1;
Expand Down
58 changes: 37 additions & 21 deletions cpp/src/io/csv/datetime.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,22 @@ __inline__ __device__ T to_non_negative_integer(char const* begin, char const* e
return value;
}

/**
* @brief Returns the position that follows a separator found by searching a range ending at `end`.
*
* A search that finds no separator returns `end`; the position after it would lie outside the
* range, so `end` is returned instead. Every range that starts after a separator (to search for the
* next separator, or to parse the next component) is thus non-inverted and within the field.
*
* @param sep_pos Result of the separator search: the separator position, or `end` if not found
* @param end Pointer to the first element after the searched range
* @return Pointer to the first element after the separator, or `end` if no separator was found
*/
__inline__ __device__ char const* next_after(char const* sep_pos, char const* end)
{
return sep_pos < end ? sep_pos + 1 : end;
}

/**
* @brief Extracts the Day, Month, and Year from a string.
*
Expand Down Expand Up @@ -79,7 +95,7 @@ __inline__ __device__ cuda::std::chrono::year_month_day extract_date(char const*
y = year{to_non_negative_integer<int32_t>(begin, sep_pos)}; // year is signed

// Month
auto s2 = sep_pos + 1;
auto s2 = next_after(sep_pos, end);
sep_pos = thrust::find(thrust::seq, s2, end, sep);

if (sep_pos == end) {
Expand All @@ -89,24 +105,24 @@ __inline__ __device__ cuda::std::chrono::year_month_day extract_date(char const*

} else {
m = month{to_non_negative_integer<uint32_t>(s2, sep_pos)};
d = day{to_non_negative_integer<uint32_t>((sep_pos + 1), end)};
d = day{to_non_negative_integer<uint32_t>(next_after(sep_pos, end), end)};
}

} else {
//--- if the dayfirst flag is set, then restricts the format options
if (dayfirst) {
d = day{to_non_negative_integer<uint32_t>(begin, sep_pos)};

auto s2 = sep_pos + 1;
auto s2 = next_after(sep_pos, end);
sep_pos = thrust::find(thrust::seq, s2, end, sep);

m = month{to_non_negative_integer<uint32_t>(s2, sep_pos)};
y = year{to_non_negative_integer<int32_t>((sep_pos + 1), end)};
y = year{to_non_negative_integer<int32_t>(next_after(sep_pos, end), end)};

} else {
m = month{to_non_negative_integer<uint32_t>(begin, sep_pos)};

auto s2 = sep_pos + 1;
auto s2 = next_after(sep_pos, end);
sep_pos = thrust::find(thrust::seq, s2, end, sep);

if (sep_pos == end) {
Expand All @@ -116,7 +132,7 @@ __inline__ __device__ cuda::std::chrono::year_month_day extract_date(char const*

} else {
d = day{to_non_negative_integer<uint32_t>(s2, sep_pos)};
y = year{to_non_negative_integer<int32_t>((sep_pos + 1), end)};
y = year{to_non_negative_integer<int32_t>(next_after(sep_pos, end), end)};
}
}
}
Expand All @@ -143,17 +159,16 @@ __inline__ __device__ cuda::std::chrono::hh_mm_ss<duration_ms> extract_time_of_d
{
constexpr char sep = ':';

// Adjust for AM/PM and any whitespace before
// Adjust for AM/PM and any whitespace before. The field can be shorter than the suffix (e.g. a
// lone "M"), so the checks never look before `begin`.
duration_h d_h{0};
auto last = end - 1;
if (*last == 'M' || *last == 'm') {
if (*(last - 1) == 'P' || *(last - 1) == 'p') { d_h = duration_h{12}; }
last = last - 2;
while (*last == ' ') {
--last;
if (end > begin && (end[-1] == 'M' || end[-1] == 'm')) {
if (end - begin >= 2 && (end[-2] == 'P' || end[-2] == 'p')) { d_h = duration_h{12}; }
end = (end - begin >= 2) ? end - 2 : begin;
while (end > begin && end[-1] == ' ') {
--end;
}
}
end = last + 1;

// Find hour-minute separator
auto const hm_sep = thrust::find(thrust::seq, begin, end, sep);
Expand All @@ -165,11 +180,12 @@ __inline__ __device__ cuda::std::chrono::hh_mm_ss<duration_ms> extract_time_of_d
duration_ms d_ms{0};

// Find minute-second separator (if present)
auto const ms_sep = thrust::find(thrust::seq, hm_sep + 1, end, sep);
auto const minutes_begin = next_after(hm_sep, end);
auto const ms_sep = thrust::find(thrust::seq, minutes_begin, end, sep);
if (ms_sep == end) {
d_m = duration_m{to_non_negative_integer<int32_t>(hm_sep + 1, end)};
d_m = duration_m{to_non_negative_integer<int32_t>(minutes_begin, end)};
} else {
d_m = duration_m{to_non_negative_integer<int32_t>(hm_sep + 1, ms_sep)};
d_m = duration_m{to_non_negative_integer<int32_t>(minutes_begin, ms_sep)};

// Find second-millisecond separator (if present)
auto const sms_sep = thrust::find(thrust::seq, ms_sep + 1, end, '.');
Expand Down Expand Up @@ -257,7 +273,7 @@ __inline__ __device__ timestamp_type to_timestamp(char const* begin, char const*
template <typename T>
__inline__ __device__ T parse_integer(char const** begin, char const* end)
{
bool const is_negative = (**begin == '-');
bool const is_negative = (*begin < end && **begin == '-');
T value = 0;

auto cur = *begin + is_negative;
Expand Down Expand Up @@ -288,7 +304,7 @@ __inline__ __device__ T parse_integer(char const** begin, char const* end)
template <typename T>
__inline__ __device__ T parse_optional_integer(char const** begin, char const* end, char delimiter)
{
if (**begin != delimiter) { return 0; }
if (*begin >= end || **begin != delimiter) { return 0; }

++(*begin);
return parse_integer<T>(begin, end);
Expand Down Expand Up @@ -357,7 +373,7 @@ __inline__ __device__ duration_type to_duration(char const* begin, char const* e
auto const after_days_sep = skip_if_starts_with(cur, end, "days");
auto const has_days_seperator = (after_days_sep != cur);
cur = skip_spaces(after_days_sep, end);
cur += (*cur == '+');
cur += (cur < end && *cur == '+');

duration_D d_d{0};
duration_h d_h{0};
Expand All @@ -376,7 +392,7 @@ __inline__ __device__ duration_type to_duration(char const* begin, char const* e

if constexpr (std::is_same_v<duration_type, cudf::duration_s>) { return output_d; }

auto const d_ns = (*cur != '.') ? duration_ns{0} : [&]() {
auto const d_ns = (cur >= end || *cur != '.') ? duration_ns{0} : [&]() {
auto const start_subsecond = ++cur;
auto const unscaled_subseconds = parse_integer<int64_t>(&cur, end);
auto const scale = min(9L, cur - start_subsecond) - 9;
Expand Down
39 changes: 33 additions & 6 deletions cpp/src/io/csv/reader_impl.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -789,6 +789,9 @@ table_with_metadata read_csv(cudf::io::datasource* source,

auto const num_actual_columns = static_cast<int32_t>(column_names.size());
auto num_active_columns = num_actual_columns;
auto const is_column_index = [num_actual_columns](int index) {
return index >= 0 && index < num_actual_columns;
};
auto column_flags =
cudf::detail::make_host_vector<column_parse::flags>(num_actual_columns, stream);
std::fill(
Expand Down Expand Up @@ -880,12 +883,34 @@ table_with_metadata read_csv(cudf::io::datasource* source,
reader_opts.get_names().size() == unique_use_cols_indexes.size(),
"Specify names of all columns in the file, or names of all selected columns");

// The columns of the input are named by its header or first row. Input without rows (an empty
// input or a byte range without a row start) has no data to read, and its columns are those
// named by the user, if any: as many names as selected columns name the selected columns (were
// they the names of all columns, all columns would be selected), other names name all columns.
auto const are_selected_columns_named =
detected_column_names.empty() and
reader_opts.get_names().size() == unique_use_cols_indexes.size();
for (auto const index : unique_use_cols_indexes) {
column_flags[index] = column_parse::enabled | column_parse::inferred;
if (are_opts_col_names_used) {
column_names[index] = reader_opts.get_names()[num_active_columns];
CUDF_EXPECTS(index >= 0 and
(is_column_index(index) or are_selected_columns_named or column_names.empty()),
"Selected column index is out of range",
std::out_of_range);
}

if (are_selected_columns_named) {
std::fill(
column_flags.begin(), column_flags.end(), column_parse::enabled | column_parse::inferred);
num_active_columns = num_actual_columns;
} else {
for (auto const index : unique_use_cols_indexes) {
// Without rows or names, no column is known and none is selected
if (not is_column_index(index)) { continue; }
column_flags[index] = column_parse::enabled | column_parse::inferred;
if (are_opts_col_names_used) {
column_names[index] = reader_opts.get_names()[num_active_columns];
}
++num_active_columns;
}
++num_active_columns;
}
}

Expand All @@ -905,8 +930,9 @@ table_with_metadata read_csv(cudf::io::datasource* source,
// User can specify which columns should be read as datetime
if (!reader_opts.get_parse_dates_indexes().empty() ||
!reader_opts.get_parse_dates_names().empty()) {
// Like names that match no column, indexes that match no column are ignored
for (auto const index : reader_opts.get_parse_dates_indexes()) {
column_flags[index] |= column_parse::as_datetime;
if (is_column_index(index)) { column_flags[index] |= column_parse::as_datetime; }
}

for (auto const& name : reader_opts.get_parse_dates_names()) {
Expand All @@ -919,8 +945,9 @@ table_with_metadata read_csv(cudf::io::datasource* source,

// User can specify which columns should be parsed as hexadecimal
if (!reader_opts.get_parse_hex_indexes().empty() || !reader_opts.get_parse_hex_names().empty()) {
// Like names that match no column, indexes that match no column are ignored
for (auto const index : reader_opts.get_parse_hex_indexes()) {
column_flags[index] |= column_parse::as_hexadecimal;
if (is_column_index(index)) { column_flags[index] |= column_parse::as_hexadecimal; }
}

for (auto const& name : reader_opts.get_parse_hex_names()) {
Expand Down
33 changes: 15 additions & 18 deletions cpp/src/io/utilities/parsing_utils.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -97,7 +97,7 @@ CUDF_HOST_DEVICE constexpr char to_lower(char const c)
*/
CUDF_HOST_DEVICE constexpr bool is_infinity(char const* begin, char const* end)
{
if (*begin == '-' || *begin == '+') begin++;
if (begin < end && (*begin == '-' || *begin == '+')) begin++;
char const* cinf = "infinity";
auto index = begin;
while (index < end) {
Expand Down Expand Up @@ -128,13 +128,13 @@ CUDF_HOST_DEVICE cuda::std::optional<T> parse_numeric(char const* begin,
constexpr bool as_hex = (base == 16);

// Handle negative values if necessary
int32_t sign = (*begin == '-') ? -1 : 1;
int32_t sign = (begin < end && *begin == '-') ? -1 : 1;

// Handle infinity
if (cuda::std::is_floating_point_v<T> && is_infinity(begin, end)) {
return sign * cuda::std::numeric_limits<T>::infinity();
}
if (*begin == '-' || *begin == '+') begin++;
if (begin < end && (*begin == '-' || *begin == '+')) begin++;

// Skip over the "0x" prefix for hex notation
if (base == 16 && begin + 2 < end && *begin == '0' && *(begin + 1) == 'x') { begin += 2; }
Expand Down Expand Up @@ -359,18 +359,12 @@ __device__ __inline__ cudf::size_type* infer_integral_field_counter(char const*
*/
__inline__ __device__ bool is_whitespace(char ch) { return ch == '\t' || ch == ' '; }

/**
* @brief Skips past the current character if it matches the given value.
*/
template <typename It>
__inline__ __device__ It skip_character(It const& it, char ch)
{
return it + (*it == ch);
}

/**
* @brief Adjusts the range to ignore starting/trailing whitespace and quotation characters.
*
* At most one quotation character is removed from each end. Only characters within the range are
* inspected, so a lone quotation character yields an empty range rather than an inverted one.
*
* @param begin Pointer to the first character in the parsing range
* @param end Pointer to the first character after the parsing range
* @param quotechar The character used to denote quotes; '\0' if none
Expand All @@ -382,13 +376,16 @@ __inline__ __device__ cuda::std::pair<char const*, char const*> trim_whitespaces
{
auto not_whitespace = [] __device__(auto c) { return !is_whitespace(c); };

auto const trim_begin = thrust::find_if(thrust::seq, begin, end, not_whitespace);
auto const trim_end = thrust::find_if(thrust::seq,
cuda::std::make_reverse_iterator(end),
cuda::std::make_reverse_iterator(trim_begin),
not_whitespace);
auto trim_begin = thrust::find_if(thrust::seq, begin, end, not_whitespace);
auto trim_end = thrust::find_if(thrust::seq,
cuda::std::make_reverse_iterator(end),
cuda::std::make_reverse_iterator(trim_begin),
not_whitespace)
.base();

return {skip_character(trim_begin, quotechar), skip_character(trim_end, quotechar).base()};
if (trim_begin < trim_end && *trim_begin == quotechar) { ++trim_begin; }
if (trim_begin < trim_end && *(trim_end - 1) == quotechar) { --trim_end; }
return {trim_begin, trim_end};
}

/**
Expand Down
Loading
Loading
0