Repository navigation
[CSV 1/9] Fix out-of-bounds reads, false NA matches and unchecked column indexes in the CSV reader - #16
[CSV 1/9] Fix out-of-bounds reads, false NA matches and unchecked column indexes in the CSV reader#16karthikeyann wants to merge 6 commits into
Conversation
`extract_date` and `extract_time_of_day`, which the CSV and JSON readers use to parse timestamps, assume that each separator they look for is present. When it is not, `thrust::find` returns `end` and the next search starts at `end + 1`, an inverted range: the search then walks past the field (and possibly past the input buffer) until it happens to find a matching character. Inputs such as `2024T10:00` (no date separator), `2024-01-01T10` (no minutes), `12T10:00`, or a time-only JSON string such as "10:30" trigger this. The parsed values then depend on whatever data follows the field, and near the end of the buffer the read faults (the JSON reader failed with an illegal memory access on "10:30"). The AM/PM suffix check has the same problem for short time-of-day strings: `*(last - 1)` and the whitespace scan read before `begin` for a time part such as `M`, and the resulting `end` can precede `begin`. Add `next_after`, which returns the position after a separator or `end` when the separator was not found, and use it wherever a search or a component starts after a separator that may be missing. Bound the AM/PM checks and the whitespace scan by `begin`. For well-formed input every search finds its separator and the AM/PM suffix is fully inside the field, so the parsed values are unchanged; malformed fields are now parsed from their own characters only (missing components parse as zero, as empty components already did). No performance regression (CSV reader benchmarks on sm_120). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
`to_duration` and its helpers `parse_integer` and `parse_optional_integer` test the character at the current position without checking that the position is still inside the field: the sign check, the `+` after "days", the `:` before minutes and seconds, and the `.` before the fraction. Once the parse has consumed the field, each of these reads the character after it. For well-formed durations such as `1 days`, `00:00:01` or `1 days 00:00:01` that is the delimiter, or, for the last field of the input, the byte after the buffer (reported by compute-sanitizer with exact allocations). The value then depends on that byte whenever it happens to be one of the tested characters. Check the position against `end` before each of these reads. Characters inside the field are tested exactly as before, so every duration that does not end at one of these checks parses to the same value. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>8000
`serialized_trie_contains` moves to the children of the current node
before matching each character after the first. A node without children
keeps the default `children_offset` of -1, so for a key that continues
past such a node the lookup stepped back to the preceding node and
resumed the search there. When that node, or one of the nodes after it
in the same children list, held the next key character and ended a word,
the lookup reported a match for a key that is not in the trie.
With the default CSV NA values, fields such as `NAA`, `NA/A`, `NAULL`
and `#NAA` were read as nulls. With na_values {"a", "b"}, `bb` and `ba`
were nulls. The JSON reader uses the same lookup for its NA values:
na_values {"1", "2"} turned 21 and 22 into nulls.
A node without children ends every key that reaches it, so a lookup that
still has characters left cannot match; return false instead of
following the negative offset. Lookups of keys in the trie and of keys
that diverge at a node with children are unaffected.
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
`create_serialized_trie` indexes each node's array of 128 children with the key character as a (signed) `char`. Bytes 0x80-0xFF, i.e. any non-ASCII character such as UTF-8 "é", are negative on the platforms libcudf supports, so building a trie from such a key indexed the array out of bounds: undefined behavior that corrupts memory and, in practice, leaves the key out of the trie. Non-ASCII na_values, true_values and false_values of the CSV reader, and na_values of the JSON reader, did not match (for quoted JSON values, building the trie could crash the process). Index the children by the character's unsigned byte value, over all 256 values. The serializer visits the children in index order, so they are stored in ascending unsigned byte order; make the device lookup compare characters as unsigned bytes too, so that its early exit is consistent with that order (with a signed comparison, "é" would sort before "a" and the lookup would stop at "a"). For ASCII keys both orders agree, so existing tries and lookups are unchanged. The tree used to build the trie now has 256 instead of 128 child pointers per node. It only exists while the trie is built on the host. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
A quote character that is never closed extends its field to the end of the input, so the last field of the input can consist of a single quote character. The decode kernel treats a string field as quoted when its first and last characters are quotes; for a one-character field both are the same character, so stripping them produced `end = begin - 1` and a string length of -1 (as `size_t`, a huge length). Building the strings column from it then accessed memory far outside the input (reading "a,b\n1,\"" failed with an illegal memory access). The same happened with detect_whitespace_around_quotes for a field such as ` "`. Require at least two characters before stripping the enclosing quotes; the lone quote is kept as the field's value. `trim_whitespaces_quotes`, used for non-string columns, had the same inverted range for a lone quote, and for an empty field it inspected the characters just outside the field (`*end` and `begin[-1]`, which for the first field of the input is the byte before the buffer). Rewrite the quote stripping to only inspect characters within the trimmed range, so a lone quote yields an empty range. Fields of two or more characters are trimmed exactly as before. This removes the last use of `skip_character`, which is deleted. The empty range is then handled downstream: `parse_numeric` and `is_infinity`, used for integer, floating-point and boolean columns (and shared with the JSON reader), tested the sign of the first character without checking that the range is non-empty. For the lone quote that is the byte after the input. Guard these reads; non-empty ranges parse as before, and an empty field parses as zero. Previously an empty field took its sign from the byte after it, so with the delimiter '-' (and na_filter disabled) an empty floating-point field parsed as -0.0; it now parses as +0.0. Duration columns are covered by the preceding fix to `to_duration`. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
`read_csv` uses the indexes passed in `use_cols_indexes`, `parse_dates` and `parse_hex` to write into the per-column flags without checking them against the number of columns. A negative index, or one past the last column, wrote out of bounds of the host flags array: heap corruption that could crash the process or go unnoticed. This includes inputs without rows, such as an empty input or a byte range that contains no row start (e.g. a partition of a larger file), where no column is detected and the flags array only covers the user-provided names, if any. A column selection that does not exist is a caller error, and silently dropping it would change the shape of the result, so `use_cols_indexes` now throws `std::out_of_range` for a negative index or an index past the last column (pandas likewise rejects out-of-bounds `usecols`). The columns of an input are named by its header or first row. An input without rows has nothing to read, and its columns are those named by the user: - As many names as selected columns name the selected columns (had they named all columns, the selection would be all columns as well). The result has these columns, in the order of the names, without rows: the same column names as for an input with rows. Validating the indexes against the names instead would make the result depend on which indexes happen to be below the number of names, and could give a column the name meant for another one. - Other names name all columns, and the indexes are validated against them: an index at or past the number of names throws. - Without names no column is known: the result has no columns, as for such inputs without a column selection. `parse_dates` and `parse_hex` only modify how columns are parsed; their names are already ignored when they match no column, so indexes that match no column are ignored the same way. Document the behavior on the options and on `read_csv`. Co-Authored-By: Claude 8000 Opus 5.5 (1M context) <noreply@anthropic.com>
|
Note Re-measured cleanly (2026-09-25); these numbers replace the earlier ones in this comment. The earlier numbers were taken while other validation jobs (tests, sanitizers, fuzzers) ran on the same host, against a baseline measured hours earlier, so the pinned-host, pageable-host and file results in particular could be skewed. Methodology (clean re-measurement). All ten builds of the stack (base and PRs 1–9) were measured in one session on the same GPU (host GPU 1), in two interleaved rounds (base, PR 1, …, PR 9, then again), so drift affects every build equally. Each benchmark run started only when no other process was using the GPU and no build/test/sanitizer/fuzzer was running on the host in our environment, and would have been re-run if anything appeared during it; all 160 runs were clean on the first attempt. Each number below is the best of the two rounds per side. CSV reader benchmarks: PR 1/9 (correctness fixes)This PR Speedup vs base (cumulative), by input source
Sources: the Full results: every CSV reader benchmark config
|
| Benchmark | Config | Base | This PR | vs base |
|---|---|---|---|---|
| csv_read_column_selection | column_selection=ALL, num_chunks=1, row_selection=ALL | 2.63 GiB/s | 2.73 GiB/s | 1.04× |
| csv_read_column_selection | column_selection=ALTERNATE, num_chunks=1, row_selection=ALL | 1.45 GiB/s | 1.45 GiB/s | 1.00× |
| csv_read_column_selection | column_selection=FIRST_HALF, num_chunks=1, row_selection=ALL | 1.49 GiB/s | 1.49 GiB/s | 1.00× |
| csv_read_column_selection | column_selection=SECOND_HALF, num_chunks=1, row_selection=ALL | 1.49 GiB/s | 1.48 GiB/s | 1.00× |
| csv_read_data_type | data_type=DECIMAL, io=DEVICE_BUFFER | 15.33 GiB/s | 20.28 GiB/s | 1.32× |
| csv_read_data_type | data_type=DURATION, io=DEVICE_BUFFER | 12.78 GiB/s | 16.06 GiB/s | 1.26× |
| csv_read_data_type | data_type=FLOAT, io=DEVICE_BUFFER | 18.66 GiB/s | 19.91 GiB/s | 1.07× |
| csv_read_data_type | data_type=INTEGRAL, io=DEVICE_BUFFER | 17.43 GiB/s | 20.11 GiB/s | 1.15× |
| csv_read_data_type | data_type=STRING, io=DEVICE_BUFFER | 4.68 GiB/s | 5.15 GiB/s | 1.10× |
| csv_read_data_type | data_type=TIMESTAMP, io=DEVICE_BUFFER | 12.37 GiB/s | 15.48 GiB/s | 1.25× |
| csv_read_io | io=FILEPATH | 4.55 GiB/s | 4.61 GiB/s | 1.01× |
| csv_read_io | io=HOST_BUFFER | 3.38 GiB/s | 3.47 GiB/s | 1.03× |
| csv_read_row_selection | column_selection=ALL, num_chunks=1, row_selection=BYTE_RANGE | 2.56 GiB/s | 2.60 GiB/s | 1.02× |
| csv_read_row_selection | column_selection=ALL, num_chunks=1, row_selection=NROWS | 2.54 GiB/s | 2.59 GiB/s | 1.02× |
| csv_read_row_selection | column_selection=ALL, num_chunks=1, row_selection=SKIPFOOTER | 2.67 GiB/s | 2.68 GiB/s | 1.00× |
| csv_read_row_selection | column_selection=ALL, num_chunks=8, row_selection=BYTE_RANGE | 1.93 GiB/s | 1.91 GiB/s | 0.99× |
| csv_read_row_selection | column_selection=ALL, num_chunks=8, row_selection=NROWS | 0.71 GiB/s | 0.67 GiB/s | 0.94× |
| csv_read_row_selection | column_selection=ALL, num_chunks=8, row_selection=SKIPFOOTER | 0.45 GiB/s | 0.43 GiB/s | 0.95× |
CSV_READER_QUOTING_IO_NVBENCH (12 configs)
| Benchmark | Config | Base | This PR | vs base |
|---|---|---|---|---|
| csv_read_quoting_io | data_size_mb=256, io=DEVICE_BUFFER, quote_pct=QUOTE_0_PCT | 17.11 GiB/s | 20.32 GiB/s | 1.19× |
| csv_read_quoting_io | data_size_mb=256, io=DEVICE_BUFFER, quote_pct=QUOTE_100_PCT | 4.87 GiB/s | 4.71 GiB/s | 0.97× |
| csv_read_quoting_io | data_size_mb=256, io=DEVICE_BUFFER, quote_pct=QUOTE_25_PCT | 12.75 GiB/s | 13.46 GiB/s | 1.06× |
| csv_read_quoting_io | data_size_mb=256, io=FILEPATH, quote_pct=QUOTE_0_PCT | 5.25 GiB/s | 5.53 GiB/s | 1.05× |
| csv_read_quoting_io | data_size_mb=256, io=FILEPATH, quote_pct=QUOTE_100_PCT | 4.69 GiB/s | 4.64 GiB/s | 0.99× |
| csv_read_quoting_io | data_size_mb=256, io=FILEPATH, quote_pct=QUOTE_25_PCT | 6.47 GiB/s | 6.58 GiB/s | 1.02× |
| csv_read_quoting_io | data_size_mb=256, io=HOST_BUFFER, quote_pct=QUOTE_0_PCT | 3.80 GiB/s | 4.09 GiB/s | 1.07× |
| csv_read_quoting_io | data_size_mb=256, io=HOST_BUFFER, quote_pct=QUOTE_100_PCT | 4.18 GiB/s | 4.29 GiB/s | 1.03× |
| csv_read_quoting_io | data_size_mb=256, io=HOST_BUFFER, quote_pct=QUOTE_25_PCT | 4.99 GiB/s | 5.19 GiB/s | 1.04× |
| csv_read_quoting_io | data_size_mb=256, io=PINNED_BUFFER, quote_pct=QUOTE_0_PCT | 9.57 GiB/s | 10.62 GiB/s | 1.11× |
| csv_read_quoting_io | data_size_mb=256, io=PINNED_BUFFER, quote_pct=QUOTE_100_PCT | 5.40 GiB/s | 5.38 GiB/s | 1.00× |
| csv_read_quoting_io | data_size_mb=256, io=PINNED_BUFFER, quote_pct=QUOTE_25_PCT | 9.77 GiB/s | 9.75 GiB/s | 1.00× |
CSV_READER_QUOTING_NVBENCH (6 configs)
| Benchmark | Config | Base | This PR | vs base |
|---|---|---|---|---|
| csv_read_quoting | data_size_mb=256, quote_pct=QUOTE_0_PCT | 17.51 GiB/s | 20.70 GiB/s | 1.18× |
| csv_read_quoting | data_size_mb=256, quote_pct=QUOTE_100_PCT | 4.80 GiB/s | 4.70 GiB/s | 0.98× |
| csv_read_quoting | data_size_mb=256, quote_pct=QUOTE_25_PCT | 13.88 GiB/s | 13.32 GiB/s | 0.96× |
| csv_read_quoting | data_size_mb=64, quote_pct=QUOTE_0_PCT | 14.10 GiB/s | 15.86 GiB/s | 1.12× |
| csv_read_quoting | data_size_mb=64, quote_pct=QUOTE_100_PCT | 1.37 GiB/s | 1.40 GiB/s | 1.03× |
| csv_read_quoting | data_size_mb=64, quote_pct=QUOTE_25_PCT | 4.57 GiB/s | 4.39 GiB/s | 0.96× |
CSV_READER_REALISTIC_IO_NVBENCH (12 configs)
| Benchmark | Config | Base | This PR | vs base |
|---|---|---|---|---|
| csv_read_realistic_io | data_size_mb=256, io=DEVICE_BUFFER, profile=ANALYTICS | 25.28 GiB/s | 28.98 GiB/s | 1.15× |
| csv_read_realistic_io | data_size_mb=256, io=DEVICE_BUFFER, profile=LOGS | 27.19 GiB/s | 26.64 GiB/s | 0.98× |
| csv_read_realistic_io | data_size_mb=256, io=DEVICE_BUFFER, profile=TAXI | 18.91 GiB/s | 23.16 GiB/s | 1.22× |
| csv_read_realistic_io | data_size_mb=256, io=FILEPATH, profile=ANALYTICS | 6.63 GiB/s | 6.43 GiB/s | 0.97× |
| csv_read_realistic_io | data_size_mb=256, io=FILEPATH, profile=LOGS | 9.85 GiB/s | 9.27 GiB/s | 0.94× |
| csv_read_realistic_io | data_size_mb=256, io=FILEPATH, profile=TAXI | 6.37 GiB/s | 6.31 GiB/s | 0.99× |
| csv_read_realistic_io | data_size_mb=256, io=HOST_BUFFER, profile=ANALYTICS | 4.83 GiB/s | 4.76 GiB/s | 0.99× |
| csv_read_realistic_io | data_size_mb=256, io=HOST_BUFFER, profile=LOGS | 7.16 GiB/s | 7.29 GiB/s | 1.02× |
| csv_read_realistic_io | data_size_mb=256, io=HOST_BUFFER, profile=TAXI | 4.77 GiB/s | 4.76 GiB/s | 1.00× |
| csv_read_realistic_io | data_size_mb=256, io=PINNED_BUFFER, profile=ANALYTICS | 12.58 GiB/s | 13.61 GiB/s | 1.08× |
| csv_read_realistic_io | data_size_mb=256, io=PINNED_BUFFER, profile=LOGS | 17.05 GiB/s | 16.94 GiB/s | 0.99× |
| csv_read_realistic_io | data_size_mb=256, io=PINNED_BUFFER, profile=TAXI | 11.20 GiB/s | 12.45 GiB/s | 1.11× |
CSV_READER_REALISTIC_NVBENCH (9 configs)
| Benchmark | Config | Base | This PR | vs base |
|---|---|---|---|---|
| csv_read_realistic | data_size_mb=1024, profile=ANALYTICS | 25.11 GiB/s | 28.95 GiB/s | 1.15× |
| csv_read_realistic | data_size_mb=1024, profile=LOGS | 31.88 GiB/s | 31.80 GiB/s | 1.00× |
| csv_read_realistic | data_size_mb=1024, profile=TAXI | 23.43 GiB/s | 26.89 GiB/s | 1.15× |
| csv_read_realistic | data_size_mb=256, profile=ANALYTICS | 25.24 GiB/s | 28.87 GiB/s | 1.14× |
| csv_read_realistic | data_size_mb=256, profile=LOGS | 27.37 GiB/s | 26.61 GiB/s | 0.97× |
| csv_read_realistic | data_size_mb=256, profile=TAXI | 19.01 GiB/s | 23.02 GiB/s | 1.21× |
| csv_read_realistic | data_size_mb=512, profile=ANALYTICS | 25.64 GiB/s | 29.51 GiB/s | 1.15× |
| csv_read_realistic | data_size_mb=512, profile=LOGS | 31.05 GiB/s | 30.85 GiB/s | 0.99× |
| csv_read_realistic | data_size_mb=512, profile=TAXI | 22.88 GiB/s | 26.29 GiB/s | 1.15× |
CSV_READER_SCALE_NVBENCH (5 configs)
| Benchmark | Config | Base | This PR | vs base |
|---|---|---|---|---|
| csv_read_scale | data_size_mb=1024 | 3.54 GiB/s | 3.50 GiB/s | 0.99× |
| csv_read_scale | data_size_mb=2048 | 3.52 GiB/s | 3.45 GiB/s | 0.98× |
| csv_read_scale | data_size_mb=256 | 3.59 GiB/s | 3.56 GiB/s | 0.99× |
| csv_read_scale | data_size_mb=4096 | 3.38 GiB/s | 3.39 GiB/s | 1.00× |
| csv_read_scale | data_size_mb=512 | 3.60 GiB/s | 3.56 GiB/s | 0.99× |
CSV_READER_TYPE_INFERENCE_IO_NVBENCH (64 configs)
| Benchmark | Config | Base | This PR | vs base |
|---|---|---|---|---|
| csv_read_with_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_FLOAT | 10.11 GiB/s | 12.17 GiB/s | 1.20× |
| csv_read_with_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL | 9.02 GiB/s | 11.01 GiB/s | 1.22× |
| csv_read_with_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_STRING | 4.40 GiB/s | 4.24 GiB/s | 0.96× |
| csv_read_with_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=MIXED | 6.59 GiB/s | 7.37 GiB/s | 1.12× |
| csv_read_with_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_FLOAT | 12.31 GiB/s | 14.10 GiB/s | 1.15× |
| csv_read_with_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL | 21.14 GiB/s | 21.68 GiB/s | 1.03× |
| csv_read_with_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_STRING | 11.91 GiB/s | 11.95 GiB/s | 1.00× |
| csv_read_with_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=MIXED | 21.15 GiB/s | 21.65 GiB/s | 1.02× |
| csv_read_with_inference_io | data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_FLOAT | 4.03 GiB/s | 4.51 GiB/s | 1.12× |
| csv_read_with_inference_io | data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_INTEGRAL | 4.12 GiB/s | 4.29 GiB/s | 1.04× |
| csv_read_with_inference_io | data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_STRING | 3.94 GiB/s | 3.91 GiB/s | 0.99× |
| csv_read_with_inference_io | data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=MIXED | 3.16 GiB/s | 3.54 GiB/s | 1.12× |
| csv_read_with_inference_io | data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_FLOAT | 4.37 GiB/s | 4.81 GiB/s | 1.10× |
| csv_read_with_inference_io | data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_INTEGRAL | 5.98 GiB/s | 5.60 GiB/s | 0.94× |
| csv_read_with_inference_io | data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_STRING | 7.49 GiB/s | 7.79 GiB/s | 1.04× |
| csv_read_with_inference_io | data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=MIXED | 5.79 GiB/s | 5.82 GiB/s | 1.01× |
| csv_read_with_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_FLOAT | 3.07 GiB/s | 3.28 GiB/s | 1.07× |
| csv_read_with_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL | 3.21 GiB/s | 3.25 GiB/s | 1.01× |
| csv_read_with_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_STRING | 3.50 GiB/s | 3.64 GiB/s | 1.04× |
| csv_read_with_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=MIXED | 2.60 GiB/s | 2.81 GiB/s | 1.08× |
| csv_read_with_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_FLOAT | 3.30 GiB/s | 3.47 GiB/s | 1.05× |
| csv_read_with_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL | 4.04 GiB/s | 4.02 GiB/s | 1.00× |
| csv_read_with_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_STRING | 6.36 GiB/s | 6.46 GiB/s | 1.02× |
| csv_read_with_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=MIXED | 4.06 GiB/s | 4.28 GiB/s | 1.05× |
| csv_read_with_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_FLOAT | 6.80 GiB/s | 7.67 GiB/s | 1.13× |
| csv_read_with_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL | 6.43 GiB/s | 7.33 GiB/s | 1.14× |
| csv_read_with_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_STRING | 4.36 GiB/s | 4.44 GiB/s | 1.02× |
| csv_read_with_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=MIXED | 4.90 GiB/s | 5.18 GiB/s | 1.06× |
| csv_read_with_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_FLOAT | 7.73 GiB/s | 8.59 GiB/s | 1.11× |
| csv_read_with_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL | 10.97 GiB/s | 11.18 GiB/s | 1.02× |
| csv_read_with_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_STRING | 10.30 GiB/s | 10.48 GiB/s | 1.02× |
| csv_read_with_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=MIXED | 10.97 GiB/s | 11.27 GiB/s | 1.03× |
| csv_read_without_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_FLOAT | 18.66 GiB/s | 20.01 GiB/s | 1.07× |
| csv_read_without_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL | 17.43 GiB/s | 20.85 GiB/s | 1.20× |
| csv_read_without_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_STRING | 4.91 GiB/s | 4.87 GiB/s | 0.99× |
| csv_read_without_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=MIXED | 14.50 GiB/s | 16.86 GiB/s | 1.16× |
| csv_read_without_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_FLOAT | 21.45 GiB/s | 21.30 GiB/s | 0.99× |
| csv_read_without_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL | 28.99 GiB/s | 28.95 GiB/s | 1.00× |
| csv_read_without_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_STRING | 20.79 GiB/s | 21.07 GiB/s | 1.01× |
| csv_read_without_inference_io | data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=MIXED | 29.16 GiB/s | 29.03 GiB/s | 1.00× |
| csv_read_without_inference_io | data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_FLOAT | 5.12 GiB/s | 5.29 GiB/s | 1.03× |
| csv_read_without_inference_io | data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_INTEGRAL | 5.02 GiB/s | 5.51 GiB/s | 1.10× |
| csv_read_without_inference_io | data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_STRING | 4.78 GiB/s | 4.70 GiB/s | 0.98× |
| csv_read_without_inference_io | data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=MIXED | 4.60 GiB/s | 5.03 GiB/s | 1.09× |
| csv_read_without_inference_io | data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_FLOAT | 5.29 GiB/s | 5.25 GiB/s | 0.99× |
| csv_read_without_inference_io | data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_INTEGRAL | 6.01 GiB/s | 6.38 GiB/s | |
| csv_read_without_inference_io | data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_STRING | 11.19 GiB/s | 11.25 GiB/s | 1.01× |
| csv_read_without_inference_io | data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=MIXED | 6.39 GiB/s | 6.29 GiB/s | 0.98× |
| csv_read_without_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_FLOAT | 3.75 GiB/s | 3.84 GiB/s | 1.02× |
| csv_read_without_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL | 3.79 GiB/s | 4.10 GiB/s | 1.08× |
| csv_read_without_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_STRING | 4.25 GiB/s | 4.27 GiB/s | 1.00× |
| csv_read_without_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=MIXED | 3.33 GiB/s | 3.59 GiB/s | 1.08× |
| csv_read_without_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_FLOAT | 3.90 GiB/s | 3.88 GiB/s | 1.00× |
| csv_read_without_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL | 4.30 GiB/s | 4.54 GiB/s | 1.06× |
| csv_read_without_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_STRING | 8.50 GiB/s | 8.64 GiB/s | 1.02× |
| csv_read_without_inference_io | data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=MIXED | 4.49 GiB/s | 4.53 GiB/s | 1.01× |
| csv_read_without_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_FLOAT | 9.60 GiB/s | 10.15 GiB/s | 1.06× |
| csv_read_without_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL | 9.67 GiB/s | 10.60 GiB/s | 1.10× |
| csv_read_without_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_STRING | 5.47 GiB/s | 5.29 GiB/s | 0.97× |
| csv_read_without_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=MIXED | 7.96 GiB/s | 8.77 GiB/s | 1.10× |
| csv_read_without_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_FLOAT | 10.51 GiB/s | 10.56 GiB/s | 1.00× |
| csv_read_without_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL | 12.94 GiB/s | 12.93 GiB/s | 1.00× |
| csv_read_without_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_STRING | 17.06 GiB/s | 17.08 GiB/s | 1.00× |
| csv_read_without_inference_io | data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=MIXED | 12.86 GiB/s | 12.84 GiB/s | 1.00× |
CSV_READER_TYPE_INFERENCE_NVBENCH (32 configs)
| Benchmark | Config | Base | This PR | vs base |
|---|---|---|---|---|
| csv_read_with_inference | data_size_mb=256, num_cols=64, type_profile=ALL_FLOAT | 10.18 GiB/s | 12.16 GiB/s | 1.19× |
| csv_read_with_inference | data_size_mb=256, num_cols=64, type_profile=ALL_INTEGRAL | 9.10 GiB/s | 11.02 GiB/s | 1.21× |
| csv_read_with_inference | data_size_mb=256, num_cols=64, type_profile=ALL_STRING | 4.40 GiB/s | 4.30 GiB/s | 0.98× |
| csv_read_with_inference | data_size_mb=256, num_cols=64, type_profile=MIXED | 6.80 GiB/s | 7.53 GiB/s | 1.11× |
| csv_read_with_inference | data_size_mb=256, num_cols=8, type_profile=ALL_FLOAT | 12.32 GiB/s | 14.56 GiB/s | 1.18× |
| csv_read_with_inference | data_size_mb=256, num_cols=8, type_profile=ALL_INTEGRAL | 21.29 GiB/s | 21.76 GiB/s | 1.02× |
| csv_read_with_inference | data_size_mb=256, num_cols=8, type_profile=ALL_STRING | 12.00 GiB/s | 12.55 GiB/s | 1.05× |
| csv_read_with_inference | data_size_mb=256, num_cols=8, type_profile=MIXED | 21.64 GiB/s | 21.66 GiB/s | 1.00× |
| csv_read_with_inference | data_size_mb=64, num_cols=64, type_profile=ALL_FLOAT | 10.15 GiB/s | 9.31 GiB/s | 0.92× |
| csv_read_with_inference | data_size_mb=64, num_cols=64, type_profile=ALL_INTEGRAL | 8.33 GiB/s | 9.74 GiB/s | 1.17× |
| csv_read_with_inference | data_size_mb=64, num_cols=64, type_profile=ALL_STRING | 1.42 GiB/s | 1.41 GiB/s | 0.99× |
| csv_read_with_inference | data_size_mb=64, num_cols=64, type_profile=MIXED | 4.05 GiB/s | 3.87 GiB/s | 0.96× |
| csv_read_with_inference | data_size_mb=64, num_cols=8, type_profile=ALL_FLOAT | 11.10 GiB/s | 12.86 GiB/s | 1.16× |
| csv_read_with_inference | data_size_mb=64, num_cols=8, type_profile=ALL_INTEGRAL | 18.41 GiB/s | 18.92 GiB/s | 1.03× |
| csv_read_with_inference | data_size_mb=64, num_cols=8, type_profile=ALL_STRING | 5.52 GiB/s | 5.53 GiB/s | 1.00× |
| csv_read_with_inference | data_size_mb=64, num_cols=8, type_profile=MIXED | 19.86 GiB/s | 18.90 GiB/s | 0.95× |
| csv_read_without_inference | data_size_mb=256, num_cols=64, type_profile=ALL_FLOAT | 18.10 GiB/s | 19.91 GiB/s | 1.10× |
| csv_read_without_inference | data_size_mb=256, num_cols=64, type_profile=ALL_INTEGRAL | 17.51 GiB/s | 20.88 GiB/s | 1.19× |
| csv_read_without_inference | data_size_mb=256, num_cols=64, type_profile=ALL_STRING | 4.70 GiB/s | 4.78 GiB/s | 1.02× |
| csv_read_without_inference | data_size_mb=256, num_cols=64, type_profile=MIXED | 14.57 GiB/s | 16.99 GiB/s | 1.17× |
| csv_read_without_inference | data_size_mb=256, num_cols=8, type_profile=ALL_FLOAT | 21.55 GiB/s | 21.42 GiB/s | 0.99× |
| csv_read_without_inference | data_size_mb=256, num_cols=8, type_profile=ALL_INTEGRAL | 29.30 GiB/s | 29.10 GiB/s | 0.99× |
| csv_read_without_inference | data_size_mb=256, num_cols=8, type_profile=ALL_STRING | 21.92 GiB/s | 22.65 GiB/s | 1.03× |
| csv_read_without_inference | data_size_mb=256, num_cols=8, type_profile=MIXED | 29.24 GiB/s | 28.97 GiB/s | 0.99× |
| csv_read_without_inference | data_size_mb=64, num_cols=64, type_profile=ALL_FLOAT | 15.91 GiB/s | 13.64 GiB/s | 0.86× |
| csv_read_without_inference | data_size_mb=64, num_cols=64, type_profile=ALL_INTEGRAL | 14.41 GiB/s | 16.04 GiB/s | 1.11× |
| csv_read_without_inference | data_size_mb=64, num_cols=64, type_profile=ALL_STRING | 1.44 GiB/s | 1.45 GiB/s | 1.01× |
| csv_read_without_inference | data_size_mb=64, num_cols=64, type_profile=MIXED | 9.73 GiB/s | 9.78 GiB/s | 1.00× |
| csv_read_without_inference | data_size_mb=64, num_cols=8, type_profile=ALL_FLOAT | 18.58 GiB/s | 19.11 GiB/s | 1.03× |
| csv_read_without_inference | data_size_mb=64, num_cols=8, type_profile=ALL_INTEGRAL | 25.71 GiB/s | 25.79 GiB/s | 1.00× |
| csv_read_without_inference | data_size_mb=64, num_cols=8, type_profile=ALL_STRING | 7.32 GiB/s | 8.40 GiB/s | 1.15× |
| csv_read_without_inference | data_size_mb=64, num_cols=8, type_profile=MIXED | 25.23 GiB/s | 24.95 GiB/s | 0.99× |
Summary. Six independent correctness fixes, each with regression tests. They were found by differential fuzzing during an automated optimization run of the CSV reader, then re-implemented minimally against the original code and reviewed by a separate principal-engineer review pass.
2024T10:00,2024-01-01T10,12T10:00, JSON"10:30", which faulted the JSON reader). The AM/PM check no longer reads before the field.1 days,00:00:01).NAA,NA/A,NAULL,#NAAwere null) and customna_valuesin CSV and JSON.na_values/true_values/false_valuesand JSONna_valueswork (for quoted JSON NA values the trie builder could crash). The builder used to index an array with a negativechar(undefined behavior); the lookup now compares unsigned bytes, matching the serialization order.use_cols_indexesentry, or one past the last column, throwsstd::out_of_rangeinstead of writing out of bounds of a host array. Inputs without rows (empty input, byte ranges without a row start) return a result without rows whose columns are named by the user: with as manynamesas selected columns, those columns in the order of the names (the same column names as for input with rows); with other names, which name all columns, the selected ones (an index past the names throws); without names, no columns. Out-of-rangeparse_dates/parse_hexindexes are ignored, like unmatched names.Design. Each fix is a bounds check or guard at the point of the bad access. For well-formed input the results are unchanged. This was checked with bit-exact oracles (floats, timestamps in many formats, quoting, long rows, integers) and 7,800 fuzzed reads against a reference build.
Testing.
compute-sanitizer --tool memcheck --rmm_mode=cudareports 0 errors on the CSV reader and JSON reader tests.DurationsAtEndOfInput,LoneQuoteNonStringFieldandEmptyFirstFieldonly detect regressions under memcheck with exact allocations (--rmm_mode=cuda); the default pool allocator hides 1-byte out-of-bounds reads.Notes for reviewers.
";std::out_of_rangefor baduse_cols_indexes. pylibcudf maps it toIndexError, while pandas raisesParserError(aValueError) for out-of-boundsusecols. Python tests were not run.convert_csv_to_cudfinlinesto_timestamp, so the datetime change alters code generation for the whole kernel (64→72 registers). On sm_120 the device-buffer read got faster on TAXI and also on ANALYTICS, which has no timestamp columns. The trim/is_infinityguards raisedata_type_detectionfrom 54 to 59 registers (the next allocation bucket on sm_120); the type-inference benchmarks show no regression. Output is identical; later changes may shift this again.Follow-up issues (out of scope):
int16node index and offsets overflow above 32,767 nodes.'\n'in a key collides with the trie's end-of-list marker.12:xx PMparses as the next day's 00:xx, and12:xx AMas 12:xx.use_cols_namesthrowscudf::logic_error("Nonexistent column selected") whileuse_cols_indexesreturns an empty table, so Python users get different behavior forusecols=["a"]andusecols=[0].This is PR 1 of a stack of draft PRs extracted from an automated (CORAL) optimization run of the CSV reader; the raw, unreviewed result is on branch
coral/csv-autoresearch-rawfor reference.🤖 Generated with Claude Code