[CSV all] Upstreamable CSV reader optimizations (combined #16–#24) by karthikeyann · Pull Request #25 · karthikeyann/cudf

karthikeyann · 2026-09-25T03:11:02Z

Summary
All upstreamable CSV reader changes from the stacked draft PRs #16–#24, as one branch on top of coral/csv-base (upstream 9a8858dd3d + the CSV benchmarks). This PR is for reviewing and benchmarking the combined result; the individual PRs are the units intended for review and merging.

What's included (in stack order)

PR	Change
#16	Correctness fixes found by fuzzing: timestamp/duration out-of-bounds reads, trie false NA matches, non-ASCII NA/true/false keys, lone-quote fields, column-index validation
#17	`""` unescape fused into the decode kernel; no zero-fill of fixed-width output
#18	Single-pass speculative row gathering for whole-file reads; bounded (32MB) host reads
#19	Word-at-a-time field seek; warp-aggregated validity bits and type-detection counts; `batch_null_count` beyond 65,535 masks
#20	Shared-memory row staging for decode and type detection (occupancy-neutral rule; validated on sm_120 only)
#21	In-place parsing of device buffers (new `datasource::supports_zero_copy_device_read()`); all-zero integer overread fix
#22	Bit-exact faster float parsing and ISO 8601 timestamp fast path
#23	Host-side latency: trie length check, sorted-key trie builder (fixes >32,767-node tries), batched null-mask memset, single-copy kernel-input uploads
#24	Two-kernel string-column builder with separate pointer/size arrays; oversized-field overflow check

What's intentionally excluded (from the raw automated result, branch coral/csv-autoresearch-raw): SWAR plain-number fast paths, trie-skip flags in parse_options, the cached word reader and its duplicated helpers, kernel variants chosen by register counts, integer pre-decode during type inference, and the per-read exp10 table (measured slower on this stack).

How it was produced
Extracted from a 10-hour CORAL multi-agent optimization run, then re-implemented one idea per PR to a principal-engineer standard, each reviewed by a separate reviewer agent at maximum rigor (2–3 rounds per PR), with findings addressed. Validation per PR: all CSV/JSON test suites at every commit, compute-sanitizer memcheck with exact allocations (plus racecheck/initcheck where kernels changed), mutation testing, bit-exact oracles, and differential fuzzing against the parent.

Known limitations / follow-ups are listed in each PR's description (e.g. sm_120-only validation of the staging rule and float changes, unaligned device views take a slower row pass, pre-existing reader quirks found during review).

Benchmarks for the combined result vs coral/csv-base across device, pinned-host, pageable-host and file sources are posted as a comment below.

🤖 Generated with Claude Code

`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>

`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 Opus 5.5 (1M context) <noreply@anthropic.com>

With `doublequote` enabled, the CSV reader turned escaped quote pairs ("" -> ") into single quotes after decoding: the decode kernel recorded a per-row "was quoted" flag for every string column, and the host then ran, per string column, a thrust::count over the flags, a cudf::strings::replace over the whole column, and a copy_if_else to keep the unquoted rows unreplaced. The columns were spread over a host worker pool and forked streams. For string-heavy files this pipeline built every string column up to three times and was a large share of the read time. The decode kernel now collapses the pairs itself, in place, in the device copy of the input that the reader owns, and the string (pointer, length) pairs describe the unescaped content. The string columns are then built once by the regular column_buffer path. The is_quoted arrays, the per-column post-processing, the worker pool use and the stream forking are removed. Correctness: - Pairs are matched left to right without overlap, which is exactly the result of the previous replace of "" with " (`""""` -> `""`), and only in fields that had their enclosing quotes stripped, as before; fields without a pair are not written. - The in-place update is race-free: each row is decoded by one thread, a field's bytes belong to one row, and the thread locates the end of the field (and so the start of the next one) before compacting it. The compaction never writes ahead of the read position. - The reader always decodes its own device copy of the input, so the caller's buffers are never modified. The unescaped data only has to live until the string columns are built, which copy the characters. The kernel documentation now states the invariants the in-place rewrite relies on: all accesses for a row stay within its byte range, row ranges do not overlap, and decoding is the last use of the data before the string columns are built. Later changes that load bytes outside a row (e.g. word loads) must not use them. Tests cover pairs at the start/end of the content, adjacent pairs, pairs before a delimiter or around a newline, a pair before the closing quote with CRLF line ends, pairs in unquoted fields, empty quoted fields, a custom quote character, detect_whitespace_around_quotes, doublequote disabled, quoting NONE, inferred string columns, 10k rows of mixed quoted and unquoted fields, byte ranges (concatenated ranges equal the full read), and a device source that must be left unchanged. Measured on an RTX PRO 6000 (nvbench, 256 MB, GPU time, minimum of two interleaved rounds against the parent): input config before after device buffer QUOTE_25_PCT 17.1 ms 8.6 ms (-50%) device buffer QUOTE_100_PCT 51.2 ms 10.9 ms (-79%) device buffer TAXI 10.6 ms 8.6 ms (-19%) device buffer LOGS 9.2 ms 5.2 ms (-44%) host buffer QUOTE_25_PCT 48.5 ms 37.4 ms (-23%) host buffer QUOTE_100_PCT 53.9 ms 26.5 ms (-51%) String-only reads gain the most (all-string, 64 columns, no inference: 49.4 -> 10.5 ms); reads without string columns are unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

decode_data created every output column with a zeroed data buffer: a memset of num_rows * sizeof(T) for each fixed-width column, ahead of a decode kernel that writes nearly all of it. The kernel writes a fixed-width slot exactly when the field is valid (ConvertFunctor stores the value only when it returns true), and that is exactly where it sets the validity bit. Every other slot is under a null bit, so its contents are unspecified per the cudf column contract. The fixed-width data buffers are therefore created without zeroing; the null masks are still zero-initialized. String (pointer, length) pairs stay zero-initialized. The kernel writes them for valid and NA fields, but not for fields missing from the end of a short row, which read as null through the zeroed pair. Writing those entries in the kernel instead, and dropping this memset too, was considered and rejected: - Handling the missing fields in or after the per-row field loop raised register pressure enough to spill the loop's state to local memory; an all-integer read (8 columns, 256 MB, no inference) went from 8.6 ms to 47 ms. - Setting every string pair of the row to null before the field loop avoided the spill, but stores each pair twice for valid fields and loops over all columns for every row; it was no faster than the memset on string-heavy reads (within noise, trending slower on 64 string columns). Zeroing only the pair arrays is the simplest form that keeps the win. A new test reads 1 to 1000 rows whose field counts cycle from 5 down to 1 across INT32, STRING, FLOAT64 and BOOL8 columns, with NA fields, plus use_cols and delim_whitespace variants. It fails if the missing string fields are left unwritten and uninitialized. This removes one num_rows * sizeof(T) memset per fixed-width column. On an RTX PRO 6000 (nvbench, 256 MB, device buffer, shared GPU so approximate), an 8-column all-integer read without inference went from 8.6-8.7 to 8.4 ms (about 3%); string columns do the same work as before. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

The CSV reader loads and parses its input in 64MB chunks so that it can stop early once a byte range or the requested number of rows has been covered. When the whole file is read, which is the common case, the chunking only costs time: every chunk synchronizes the stream to bring the block row contexts of the first gathering phase to the host, and the row offsets vector is regrown and copied for every chunk. Parse the whole file as a single chunk when no row selection option is set. Chunked parsing is unchanged for byte ranges, skiprows, nrows and skipfooter. The input used to be read with one `host_read` per chunk for sources that do not prefer device reads, so a single chunk would have put a copy of the whole file in host memory for sources that allocate on reads. Read such sources in chunks of at most 32MB instead, each copied to the device before it is released: at most two chunks are held at a time, the one being copied and the next one, whose read is requested with host_read_async before the copy so that it runs on the host worker pool meanwhile. If the copy throws, the pending read is waited for before the exception propagates, since it uses the source. Sources that prefer device reads are read into device memory as before. Staging the chunks through pinned buffers instead was measured to be slower for pageable host buffers (47ms vs 41ms to copy 528MB on the test machine): the driver's pageable copy path is faster than a single-threaded memcpy into pinned memory. Also move the per-block row context merge tree of gather_row_offsets from a global memory scratch buffer (8KB per 16KB of input, e.g. 128MB for a 256MB file) into shared memory. Measured end to end with the CSV reader benchmarks on 256MB inputs (RTX PRO 6000 Blackwell, minimum of 4 interleaved runs): reads from files, now one device read instead of one per chunk, take 16% less time for the TAXI profile and 13% less for 25% quoted fields; device buffer reads change by -6% to +3% and pinned host buffers by -1% to -2%. Pageable host buffers are flat (+1% and +2%, within noise): their reads are bound by the host-to-device copy. Add tests that read a 45MB input from a user source, checking that no read exceeds 32MB, and from pageable and pinned host buffers, and that a failed or short read from the source is reported. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

gather_row_offsets computes, for each 32-character slice of the input, the row start bitmap and output parser state for every possible input state, by merging the transitions of the 32 characters one at a time. Most slices hold no quote or comment character, and in those only the characters that follow a terminator matter. When a slice is fully in range and 16-byte aligned, load it with two 16-byte loads and locate its terminator, quote and comment characters with SWAR byte comparisons, then: - If the slice holds no quote or comment character, write the result directly: every character after a terminator starts a row in the NONE and COMMENT (pending quote exit) states and in none in the QUOTE state; NONE and COMMENT end in NONE (COMMENT exits at the first character) and QUOTE stays QUOTE. This is what the per-character merge computes, since each character is then either a row start that is neither a quote nor a comment character, or a regular character. - Otherwise, merge only the characters that follow a terminator and the quote characters with their full transitions. Every other character has the regular transition (NONE->NONE, QUOTE->QUOTE, COMMENT->NONE, no row start), which is idempotent, so each run of them is merged once. Both paths produce exactly the result of the per-character loop, which still handles unaligned slices and slices that cross the end of the chunk. The per-character transition is factored into char_row_context, and the characters it depends on are passed as one struct (dropping the unused escape character argument). The slices are aligned for whole-file reads and for the first chunk of nrows and skipfooter reads. Byte range reads, and skiprows reads after their first chunk, keep one character before the data they parse, so their slices are one byte off alignment and take the per-character loop as before; their results are unaffected. Aligning them is left for a follow-up. On 256MB inputs, gather_row_offsets_gpu takes 38% to 45% less time per read (nsys, both phases): 1490us -> 850us for the TAXI profile of the realistic CSV reader benchmark, and 1803us -> 997us, 1072us -> 614us and 684us -> 423us for 0%, 25% and 100% quoted fields. End to end, reads from device memory take 3% to 12% less time (minimum of 4 interleaved runs, RTX PRO 6000 Blackwell): TAXI 7.82ms -> 7.32ms, LOGS 4.52ms -> 4.15ms, ANALYTICS 7.80ms -> 7.03ms, and 11.38ms -> 10.04ms, 7.96ms -> 7.38ms and 9.77ms -> 9.50ms for 0%, 25% and 100% quoted fields. Add tests that sweep quoted fields with terminators, delimiters and escaped quotes, and a comment line, across 16KB block boundaries for '\n', CR-only and non-ASCII terminators, that check an unterminated quote at the end of the data, byte ranges against a full read, row selection with a quoted field across the 64MB chunk boundary, and a comment character at the end of an aligned slice (which the quote and comment free shortcut must exclude). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

Row offsets are gathered in two phases: the first counts the rows of each 16KB block for every possible parser state at its start, the host then resolves the actual state of each block and the row counts, and the second phase writes the offsets. For a whole file this costs a device-to-host copy of the block contexts, a host loop over all blocks, a host-to-device copy back and a second read of the data. Parsing the whole file needs no row selection, so gather its rows in a single pass over the data with no host round trip: 1. A speculative pass processes 32KB tiles in parallel, each assuming that it starts in the NONE state (outside of quotes). For each tile it writes the row start bitmap of every 32-character slice and the number of rows, and records the tile's state transition function (the output state for each input state, packed in a byte). Within a tile, a block scan over the threads' transitions gives each thread its input state. 2. An exclusive scan over the tile transitions, composed as functions, gives the exact start state of every tile. This does not depend on the speculation, since the transitions cover every input state. 3. A repair pass processes again, with one block per tile, the tiles whose start state is not NONE: those where a quoted field spans or ends at the start of the tile. Other blocks exit at once. 4. An exclusive scan of the tile row counts gives the first row of every tile and the total, which is copied to the host (the only synchronization) to allocate the offsets, and a last pass converts the bitmaps to offsets. The result is identical to the two-phase gathering: the bitmaps of each tile are those selected by its exact start state, from the repair pass if the speculation was wrong. Each tile is processed at most twice, so even when the speculation fails for every tile (e.g. a quoted field that spans the whole file) the work is at most twice that of one pass and stays linear in the input size. The bitmaps take one bit per input character. The fraction of repaired tiles follows the fraction of quoted characters in the data. On the 256MB CSV reader benchmark inputs it is 0% for the ANALYTICS profile and unquoted fields, 2% for TAXI, 14% for LOGS, and 11% and 26% for 25% and 100% quoted fields; with 100% quoted fields the repair pass takes 52us next to 229us for the speculative pass. Chunked gathering is unchanged for row selection (byte range, skiprows, nrows, skipfooter), which can stop early. On 256MB inputs, the row gathering kernels (including the removal of blank rows) take 25% to 31% less time per read (nsys): 1046us -> 774us for TAXI, and 1085us -> 749us, 645us -> 471us and 440us -> 328us for 0%, 25% and 100% quoted fields. End to end, reads from device memory take up to 8% less time (minimum of 4 interleaved runs, RTX PRO 6000 Blackwell): TAXI 7.32ms -> 6.77ms, LOGS 4.15ms -> 3.91ms, ANALYTICS 7.03ms -> 6.61ms, and 7.38ms -> 7.23ms and 9.50ms -> 9.27ms for 25% and 100% quoted fields; unquoted fields are flat (10.04ms -> 10.17ms). Add tests where every tile but the first starts inside a quoted field (with more tiles than the device runs blocks at once), where a quoted field spans many tiles, where the characters of short probes are placed in turn at tile starts so that tiles start in each parser state, and a seeded test that compares whole-file and chunked reads of random multi-tile documents with random quoting, comment and blank line options. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

After gathering the row offsets of a whole file, the reader removes the offsets of blank and comment rows with a thrust::remove_if over all offsets, which reads every offset, the first character of every row, and writes the offsets again. The single-pass gathering already visits every row start while it selects the row start bitmaps, with the character in cache. Clear the bits of the rows that remove_blank_rows would remove there, before the rows of each tile are counted: the exclusive scan of the counts then gives the positions of the remaining rows, and the offsets of blank rows are never written. The row that marks the end of the data is always kept, as in remove_blank_rows. The blank row predicate is shared by count_blank_rows, remove_blank_rows and the gathering, so that the chunked and the single-pass paths skip the same rows. On 256MB inputs this removes the remove_if (217us per read for the TAXI profile of the realistic CSV reader benchmark, and 99us, 38us and 17us for 0%, 25% and 100% quoted fields, nsys) while the gathering kernels take the same time within 1us. End to end, reads from device memory take up to 7% less time (minimum of 4 interleaved runs, RTX PRO 6000 Blackwell): TAXI 6.77ms -> 6.53ms, LOGS 3.91ms -> 3.63ms, 0% quoted fields 10.17ms -> 9.53ms; the others are flat. Add a test with blank, CRLF blank and comment lines, alone and in runs that span many tiles, at the start and at the end of the data, with the '\n' and a custom terminator, and with skip_blank_lines disabled, comparing the whole-file read with the chunked row selection path. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

The type detection and decode kernels process one row per thread, and `seek_field_end` reads each row one character at a time. The loads of a warp are therefore scattered over 32 rows, and each character costs a separate memory transaction. Add `seek_field_end_by_words`, which returns the same position as `seek_field_end` (without escape characters) but loads the row in aligned 8-byte words, and use it in both kernels: - In a field that does not start with a quote, only a delimiter, a terminator or a '\r' can end the field. The bytes of the word that may hold one of them are flagged with the SWAR zero-byte test on the word XOR the broadcast character. That test flags every occurrence (it may also flag other bytes), and the flags are checked in increasing order with the exact test of `seek_field_end` ('\r' only ends a field before '\n'), so the first flag that passes is the position `seek_field_end` returns, and false flags fail the test. - The characters of a quoted field are run through the same quote state machine as in `seek_field_end`, from the loaded word. - Multi-character delimiters (delim_whitespace) keep using `seek_field_end`. Only words that lie entirely within [data begin, row end) are loaded, whatever the alignment of the data, so no memory outside of the data is read even for unaligned or unpadded buffers; the characters of the row outside of such words are read one at a time. (The reader parses its own aligned copy of the data for now, so the lower bound matters for sources that would be parsed in place.) A word may start in the previous row, whose bytes the decode kernel may be rewriting (in place unescaping); those bytes are never used, as the documented invariant of the decode kernel requires. Median kernel times (nsys, RTX PRO 6000, 256MB benchmarks): benchmark kernel before after REALISTIC TAXI decode 5.64 ms 3.90 ms REALISTIC LOGS decode 1.62 ms 1.25 ms REALISTIC ANALYTICS decode 5.57 ms 4.12 ms QUOTING QUOTE_0_PCT decode 9.38 ms 6.53 ms QUOTING QUOTE_25_PCT decode 4.72 ms 3.02 ms QUOTING QUOTE_100_PCT decode 1.91 ms 1.22 ms INFERENCE ALL_INTEGRAL, 64 detection 11.36 ms 8.89 ms INFERENCE ALL_FLOAT, 8 detection 6.61 ms 5.38 ms Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

`batch_count_set_bits`, which backs `cudf::batch_null_count`, indexes the bitmasks with the second dimension of the grid, whose size is limited to 65535, so counting more bitmasks failed with cudaErrorInvalidValue. Launch the kernel once per batch of at most 65535 bitmasks. Nothing changes for fewer bitmasks. The CSV reader is about to count the nulls of its columns with this function, and CSV files can have more columns than that. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

The decode kernel set the validity bit of each valid field with its own atomicOr, and counted the valid fields of each column with an atomicAdd per field. The 32 lanes of a warp decode 32 consecutive rows that start at a multiple of 32, so their bits of a column all lie in the same mask word and their atomics serialize on it; the counter of a column is a single address that every warp of the grid updates for every field. Group the lanes that mark the same column with __match_any_sync and let the lowest of them set all of their bits with one atomicOr (the bit of a row is its lane index). The lanes that are active together normally mark the same column, but that is not guaranteed after divergence, so they are grouped by column. The atomic is still required because different groups of the same warp can update the word at different times. Stop counting in the kernel: after decoding, count the nulls of the non-string columns from their masks with cudf::batch_null_count. This is exact because the masks are zero-initialized, bits are only set for the valid fields of rows in [0, num_records), and exactly that range is counted. String columns get their null counts when they are built from the decoded string pairs, as before, so their masks are not counted. The count replaces the copy of the kernel's counters to the host, so the number of synchronizations does not change. A test reads a table with more columns than the batched count can index in one grid (65535), which the previous commit enables. Median time of the decode kernel (nsys, RTX PRO 6000, 256MB benchmarks): benchmark before after INFERENCE ALL_INTEGRAL, 8 4.24 ms 2.26 ms INFERENCE ALL_INTEGRAL, 64 7.51 ms 4.36 ms INFERENCE MIXED, 8 3.83 ms 2.26 ms INFERENCE MIXED, 64 3.72 ms 2.97 ms QUOTING QUOTE_0_PCT 6.53 ms 4.33 ms QUOTING QUOTE_25_PCT 3.02 ms 2.31 ms REALISTIC TAXI 3.90 ms 3.39 ms REALISTIC LOGS 1.25 ms 1.10 ms Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

The type detection kernel incremented a member of the column's column_type_histogram with a global atomicAdd in each branch of its classification. The counters of a column are a few addresses that every thread of the grid updates for each row, so these atomics serialize. Classify each field into a `field_class`, which indexes an array of `num_field_classes` counters per column, and increment the counter in one place. The lanes of a warp that increment the same counter at the same time are grouped with __match_any_sync, and the lowest of them adds their number with a single atomic. The counts are unchanged. Integers are classified by the new classify_integral_field, which holds the logic of infer_integral_field_counter; the latter is now a thin wrapper around it for the JSON type inference. The histograms are assembled on the host from the counters, so the device code no longer depends on the layout of column_type_histogram. A new test decides column types with single fields in later blocks of rows and with the number of NA fields, so that a lost or miscounted increment changes an inferred type. Median time of the type detection kernel (nsys, RTX PRO 6000, 256MB CSV_READER_TYPE_INFERENCE_NVBENCH with inference): profile, columns before after ALL_INTEGRAL, 8 2.77 ms 1.85 ms ALL_INTEGRAL, 64 8.00 ms 5.67 ms MIXED, 64 11.48 ms 5.92 ms ALL_STRING, 8 8.68 ms 1.15 ms ALL_STRING, 64 6.08 ms 2.03 ms ALL_FLOAT, 8 5.38 ms 5.73 ms (see the next commit) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

Narrow tables have few type detection counters, and every warp of the grid updates them, so even warp-aggregated global atomics contend on them. This is worst when the lanes of a warp reach their increments at different times, as with variable-length floating point fields. For tables of up to 32 inferred columns, each block counts in shared memory, zeroed on entry, and adds its non-zero counts to the global counters at the end; threads without a row now stay until then instead of returning early. The counts are exact either way. The kernel is a template on where the counters are, rather than taking a runtime flag. With the counters known to be in shared memory, the compiler emits shared memory atomics (ATOMS in the SASS of this version); in a build that selected the counters with a runtime flag, the increment compiled to a generic atomic (ATOM.E.ADD), since the counter pointer can be in either address space. The global memory instance also keeps its early return and has no block barriers. The threshold only affects speed, and was measured on one GPU (RTX PRO 6000 Blackwell). Larger shared memory allocations per block made the kernel slower than counting in global memory (64 columns: 8.9 ms instead of 5.7 ms; 32 columns with the allocation padded to 2KB: 8.3 ms instead of 3.6 ms), likely because the kernel, which reads rows through the L1 cache, then gets a larger shared memory carveout. Wider tables spread their updates over enough counters that global memory atomics do not contend much on them. Median time of the type detection kernel (nsys, 256MB CSV_READER_TYPE_INFERENCE_NVBENCH with inference, num_cols overridden for 16 and 32): profile, columns global shared ALL_FLOAT, 8 5.6 ms 2.8 ms ALL_FLOAT, 32 4.9 ms 3.5 ms ALL_INTEGRAL, 16 6.4 ms 3.3 ms ALL_INTEGRAL, 32 5.5 ms 3.6 ms MIXED, 16 4.9 ms 3.2 ms ALL_INTEGRAL, 8 1.76 ms 1.88 ms Tables with 8 integer columns get slightly slower (ALL_INTEGRAL and MIXED: 2-3% end to end), which the gains above outweigh. The tests that decide column types with single fields in later blocks now also cover 32 and 33 columns, on both sides of the threshold, and fail if a block's shared counts are not added to the global counters. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

`load_data_and_gather_row_offsets` removed the header rows and applied the nrows limit, and its caller then applied the skipfooter limit. The limits are now all applied in one place, `load_data_and_gather_row_offsets`, before the header offsets are copied to the host. The data and its selected row offsets are returned in a named struct instead of a pair. This prepares for work on the rows to read that has to be done on the device before that copy, so that its result can be copied to the host in the same synchronization (the next commit needs this). The behavior is unchanged: the header rows are removed only if the header row exists, as before, the nrows and then the skipfooter limits are applied in the same order to the same offsets, and the header is read from the same offsets. `skipfooter` cannot be combined with `nrows` or with a byte range (both are rejected earlier), so the only interactions are with the header rows, skiprows and blank rows. A differential run over 20,736 reads with header, skiprows, skipfooter, nrows, blank rows and byte ranges gives byte-identical results and exception messages. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

The decode kernel parses one row per thread, so the 32 threads of a warp read 32 different rows: each character (or 8-byte word) load of the warp is up to 32 separate memory transactions, and the sectors it fetches are reused only while they stay in the L1 cache. Each block can now first copy its rows, which are contiguous in the data, into dynamic shared memory with coalesced 16-byte loads, and its threads then parse the copy. The copy keeps the characters' address modulo 16, so the aligned 16-byte words within the rows are copied as a whole and the characters of the rows outside of such words one at a time, whatever the alignment of the data. Only the characters of the block's rows are read. Strings still point into the reader-owned data, never into shared memory: a quoted field is unescaped from the staged copy into `data`, whose bytes of the field are still those of the copy, since a block stages its rows before any of its threads rewrites a field and no other block touches them. `unescape_doublequotes` therefore takes a source and a destination that holds a copy of it (or is the source itself). Every read of a field stays within the field, so no staged read touches the buffer before the first row or after the last one: the word loads of `seek_field_end_by_words` are bounded by the staged characters, and `trim_whitespaces_quotes` does not read before the field. Staging size: the shared memory needed to stage the rows of every block is computed on the device, over the rows to read (after the header rows and the nrows/skipfooter limits are applied, see the previous commit), and copied to the host in the synchronization that already fetched the header offsets. The blocks of the kernels must be those of the computation: removing the header rows shifts the blocks, and a size computed before that can be too small. Each staged block checks that its rows fit in the launch's buffer, and the kernel stops with an error otherwise (one comparison per block). It cannot fall back to parsing global memory without losing the staged instance's shared memory loads (see below). `block_row_range` is the single definition of the rows of a block. When to stage (`is_row_staging_used`): - The size must fit in the dynamic shared memory that the kernel can use without opting in to more (48KB on all GPUs, less the kernel's static shared memory, from cudaFuncGetAttributes). - Staging must not reduce the occupancy of the launch (resident blocks per SM, from the occupancy API for the current GPU and kernel). On an RTX PRO 6000 (sm_120, 100KB of shared memory per SM, 7 resident unstaged blocks per SM), staging was faster at equal occupancy in every case measured, while with fewer blocks per SM the outcome depended on the data: rows of 134 to 174 characters decoded 14% to 30% faster with 5 or 4 staged blocks per SM, but TAXI was 17% slower with 3. So staging is only used where it costs no occupancy. The rule has no tuned constant, but it was only measured on sm_120: on GPUs with more shared memory per SM (sm_80: 164KB, sm_90: 228KB) it stages larger blocks (up to about 22KB and 31KB), e.g. TAXI's 28KB blocks on sm_90, which have not been measured. It also depends on the register counts of the two instances (56 staged against 72 unstaged on sm_120), which compiler or code changes can move. Before merging, LOGS, TAXI, ANALYTICS and the 8-column type inference benchmarks should be measured on an sm_80 and an sm_90 GPU; a loss there would call for a condition on the rule, not a tuned threshold. - The rows of all blocks are staged or none, by separate kernel instances. The shared memory size of a launch applies to all of its blocks, so the blocks that could not be staged would run at the lower occupancy of the staged launch (1.41 ms instead of 1.10 ms for rows of about 400 characters). And in a single kernel whose rows could be in either memory, every character load compiles to a generic load, which made the unstaged decode twice slower (QUOTE_0_PCT: 4.3 ms -> 8.4 ms). - The `LIBCUDF_CSV_ROW_STAGING` environment variable (AUTO by default, ALWAYS or NEVER) overrides the occupancy condition. Like LIBCUDF_NVCOMP_POLICY and LIBCUDF_HOST_COMPRESSION, it is read on every call, so tests can set it: the staging tests run both parsing paths whatever the GPU, and staging never changes the results. Measured on an RTX PRO 6000 (256MB inputs; decode kernel time, minimum of 20+ runs under nsys): LOGS 1.09 ms -> 0.87 ms (staged, 13.3KB per block, 7 blocks per SM as without staging), rows of 94 characters with 8 integer, 4 float and one string column 1.26 ms -> 0.96 ms. TAXI, ANALYTICS and QUOTE_0/25/100_PCT are not staged and unchanged. End to end (nvbench GPU time, minimum of 2 rounds on a shared GPU), LOGS 3.63 ms -> 3.43 ms; the other benchmarks are within noise. The size computation adds a reduction of 4 to 7 us. Tests (each run with staging forced and disabled): blocks of rows at the staging limit, starting at the last position of a 16-byte word, with a multiple of 16 characters, and one character over the limit; rows after a header whose first block is longer than any block that includes the header; long rows (1,000 rows of 530 characters, and a 60KB row among short rows); rows starting at every position of a 16-byte word; empty first and last fields of rows staged at the very start of the buffer, or parsed in global memory, where compute-sanitizer with exact allocations detects reads outside of the rows; an invalid staging policy. Computing the size before the header rows are removed, or without the 15 bytes before the rows, fails these tests. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

Type detection parses one row per thread like the decode kernel, and has the same scattered loads. It now stages the rows of its blocks in shared memory under the same conditions as the decode kernel (see `is_row_staging_used`), using the staging size that the reader already computes for the decode, whose blocks are the same. The per-block counters of narrow tables (up to 32 inferred columns) stay at the start of the dynamic shared memory, and the staged rows follow them. Their size, 32 bytes per column, keeps the staged rows 16-byte aligned (checked by a static_assert). Both staging conditions account for the counters: the counters and the staged rows together must fit in the dynamic shared memory the kernel can use without opting in (48KB less its static shared memory), and the occupancy of the staged launch, with both, is compared with that of the unstaged launch, with the counters only. The staged and unstaged launches are separate instances of the kernel, as for the decode kernel. Measured on an RTX PRO 6000 (sm_120; 256MB inputs, 8 columns; detection kernel time, minimum of 20+ runs under nsys): - ALL_INTEGRAL: 1.96 ms -> 1.59 ms; MIXED: 1.97 ms -> 1.58 ms (10KB per block, 9 staged blocks per SM against 7 without staging). End to end (nvbench GPU time, minimum of 2 rounds on a shared GPU), with the previous commit: 6.38 ms -> 5.33 ms and 6.36 ms -> 5.23 ms. - ALL_FLOAT and ALL_STRING (17KB and 19KB per block, 5 and 4 staged blocks per SM) are not staged and unchanged. Staged anyway, they took 1.77 ms instead of 2.61 ms and 1.35 ms instead of 1.17 ms: with lower occupancy staging can be slower, which the occupancy condition rules out. As for the decode kernel, the rule was only measured on sm_120. Test: type inference of rows with fields of 1 to 3 characters, with the counts in shared memory (1 to 32 columns) and in global memory (33 and 36 columns), with staging forced and disabled. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

`classify_integral_field` skips the leading zeros of integers of 19 digits or more before comparing them with the 64-bit limits. The loop dereferenced the next character before checking that it is still in the field, so an integer made of zeros only read the character after it. This is a pre-existing bug in upstream libcudf, in code that the CSV and JSON type inference share. The CSV reader's type inference reads one byte past the data when such a field ends the input without a final terminator: past the reader's copy of the input, or, for a staged block, past its rows in shared memory. In JSON, values are always followed by a structural character, so the read stays within the input there. It is fixed in this series because parsing device buffers in place (later in this series) makes that byte the one past the end of the caller's buffer; it was found while auditing every read of the CSV kernels against exact-size caller buffers. Check the bound first. The classification does not change: the byte past the digits only ended the loop, which also ends at the bound. Test: inferred reads of 19 to 32 zeros at the end of the data, with rows staged in shared memory and not. The test passes without the fix: it only detects the read under compute-sanitizer memcheck with exact allocations (--rmm_mode=cuda), when the rows are not staged: 21 invalid reads without the fix, none with it. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

`datasource::device_read(offset, size, stream)` returns a buffer that may own a copy of the data, as for kvikIO file sources, or be a view of device memory that holds the data already, as for device buffer sources. Readers cannot tell the two apart, so they copy the data of all sources into memory of their own, even when it is already in device memory. Add `datasource::supports_zero_copy_device_read()`, false by default. A source that returns true (and supports device reads) guarantees that the buffer returned by `device_read(offset, size, stream)` points to device memory that holds the data in stream order on `stream`, and that stays valid and unmodified until the buffer is destroyed and the reader's work on `stream` is complete. Readers can then process the data in place, without modifying it, whether or not the source prefers device reads, since no read takes place. It does not apply to reads into a preallocated buffer, which always copy. The memory must be device memory, not host memory that the device can access (pinned, or managed memory resident on the host): readers may read the data several times, and such sources should let them copy it once instead. The doc of `device_read(offset, size, stream)` now also allows calling it when the source supports zero-copy device reads. Device buffer sources return true; `datasource::create(device_span)` now documents that readers may process the buffer in place, and that it must hold the data when a read starts and must not be modified or freed until the read completes (host reads of the source are not ordered on the stream of the read). The wrapper of user-implemented sources forwards the query, so that user sources that hold their data in device memory can opt in. The name follows `supports_device_read()`, and Arrow's `RandomAccessFile::supports_zero_copy()`: it is a capability of the source, not a per-size preference like `is_device_read_preferred()`. Adding a virtual function changes the vtable of `datasource`, so subclasses outside of libcudf must be rebuilt. Test: the query for device buffer, host buffer, file and user sources, and device reads of a device buffer and of user sources through the wrapper, which return views exactly for zero-copy sources, in a new DATASOURCE_TEST. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

The CSV reader copies its whole input into device memory it owns before parsing it, because decoding unescapes quoted strings (`""` -> `"`) in place. For device buffer sources this is a device-to-device copy of the whole input (0.75 ms for the 554MB input of the TAXI benchmark), and device memory of the size of the input for the whole read. Whole-file reads (no byte range, skiprows, nrows or skipfooter, and no compression) of sources whose device reads return views of their data (`datasource::supports_zero_copy_device_read`, true for device buffers) now parse the view in place, whether or not the source prefers device reads, since a view involves no read. Reads with row selection still copy the chunks they read, and host, pinned and file sources are read as before: kvikIO file sources read into the reader's buffer, which is already a single read with no extra copy, and they keep unescaping in place. Empty ranges (a BOM-only input) are not requested from the source. The caller's buffer is never written: `decode_row_column_data` takes the data as read-only, and an `unescape_buffer` at the same offsets that receives the unescaped quoted fields. It is the data itself when the reader owns it (in place, as before), and otherwise uninitialized scratch memory of the size of the data, allocated when decoding starts and only if `doublequote` is set and a column is a string column. `unescape_doublequotes` now takes whether the output already holds the content: in place, from the data or from its staged copy, the characters before the first pair are not rewritten; into scratch memory, the whole unescaped content is written. Fields without escaped pairs are not written, and their strings keep pointing into the source's data. No kernel reads outside of the data, so the view can be exactly the caller's buffer, at any alignment: - The single-pass row gathering loads 16-byte words only for 32-byte slices that are aligned and lie entirely within the data, and one character at a time (bounded by the end) otherwise; the character before a slice is read only if the slice is not at the start. The blank row test reads only offsets below the size. - Row staging copies exactly the rows of a block, word loads of the field seek are bounded by the data start and the row end, and field parsing stays within the field (with the first commit's fix). Unaligned views (a UTF-8 BOM makes the data start 3 bytes into the buffer, and subspans of buffers can start anywhere) are parsed in place too, but the row gathering then loads every slice one character at a time: 783 us instead of 403 us for 512 MiB of data, which still costs less than the copy (681 us) plus the gathering of the aligned copy (438 us). Aligning the slices to absolute addresses would give them the word loads; that is left for a follow-up. Memory: the scratch memory replaces the copy, is only allocated for decoding and only when string columns can need it, so the peak device memory of a read is never larger than before. Measured on an RTX PRO 6000 (nvbench GPU time, 256MB configurations of the IO benchmarks, interleaved runs; the device is shared with other jobs, so gains vary by a few percent between sessions): - DEVICE_BUFFER, minimum / median of 5 runs: TAXI 5.49 -> 4.67 / 5.58 -> 4.79 ms (-15% / -14%), LOGS 3.17 -> 3.04 / 3.45 -> 3.05 ms (-4% / -12%), ANALYTICS 5.62 -> 4.92 / 5.67 -> 4.94 ms (-12% / -13%), QUOTE_0_PCT 6.63 -> 5.76 / 6.79 -> 5.82 ms (-13% / -14%), QUOTE_25_PCT 5.84 -> 5.58 / 5.99 -> 5.61 ms (-4% / -6%), QUOTE_100_PCT 9.31 -> 9.40 / 9.52 -> 9.53 ms (unchanged, see below). - Peak memory drops where no scratch memory is needed (no string column: ANALYTICS 849.5 -> 287.5 MiB, QUOTE_0_PCT 913.6 -> 265.1 MiB), and is unchanged otherwise (TAXI 903.7 MiB). - FILEPATH, HOST_BUFFER and PINNED_BUFFER (TAXI) and HOST_BUFFER (QUOTE_25_PCT): unchanged, within the noise of these runs (up to 5%). - Decoding data whose every string field has escaped quotes writes the unescaped strings to scratch memory rather than to cache lines just read: the decode kernel took 1645 us instead of 1486 us for 512 MiB (plus 0.7 ms saved on the copy). With the string columns of QUOTE_100_PCT, the slower decode takes about what the copy saves. Tests: device buffers at offsets 0 to 15 that end at the end of their allocation (BOM, comment and blank lines, escaped quotes at the start, middle and end of strings, quoted delimiters and terminators, CRLF, no final terminator), read with and without type inference, unescaping and quoting, with rows staged in shared memory and not staged (through LIBCUDF_CSV_ROW_STAGING), the same as from a host buffer and from a file, and left unchanged; empty, BOM-only and header-only device buffers and single fields at the end of the buffer; another quote character; device buffers of rows at every alignment; a user source with unaligned device data, parsed in place if it reports zero-copy reads (whether or not it prefers device reads) and copied otherwise, and copied with row selection. compute-sanitizer memcheck with exact allocations reports no error. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

Parsing a device buffer in place, the reader allocates scratch memory of the size of the input for the quoted strings it unescapes, whenever `doublequote` is set (the default) and a column is a string column, even if no field has an escaped quote pair. An escaped pair is two consecutive quote characters, and the single-pass row gathering already loads every character of the data. Its speculative pass now also records whether any quote character follows a quote character: in the word-loaded slices, from the quote bitmap it already computes (`quotes & (quotes << 1)`, and bit 0 if the character before the slice is a quote), in the others, character by character. The flag is copied to the host with the number of rows, in the synchronization that already sized the row offsets, and the reader only allocates the scratch memory if it is set. The flag ignores the parser state on purpose: it is a superset of the fields that decoding unescapes, whatever the options. A state-aware test could exclude empty quoted fields (`""`), but decoding also unescapes fields that the row state machine does not see as quoted (e.g. `, "a""b"` with detect_whitespace_around_quotes), so it could miss pairs that decoding writes. This helps reads of data with string columns and no two consecutive quote characters anywhere (no escaped quotes and no empty quoted fields). No CSV reader benchmark has such data: the TAXI, LOGS, QUOTE_25_PCT and QUOTE_100_PCT inputs have consecutive quotes, and QUOTE_0_PCT and ANALYTICS have no string column, so the previous commit already allocates no scratch memory for them. Measured on an RTX PRO 6000: - 256 MiB of rows `12345,678.5,s12345,t` read from a device buffer with inferred types: peak device memory 1439.5 -> 1183.5 MiB. With an empty quoted field in 10% of the rows, it is unchanged (1432.5 MiB). - Speculative row gathering pass over 512 MiB: 409 -> 412 us aligned, 873 -> 878 us unaligned; end-to-end times of the benchmarks are unchanged. Tests: the only escaped pair of the data at positions around the boundaries of 32-character slices, 16-byte words and 32KB tiles, in aligned and unaligned device buffers (an undetected pair makes the decode write to the missing scratch memory); a pair in a field that the row gathering does not consider quoted but decoding unescapes (detect_whitespace_around_quotes); the peak memory of a read with and without escaped pairs. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

parse_numeric computes the place value of each fraction digit of a floating-point field by dividing the previous one by 10: one double division per digit. A double division is a long instruction sequence (reciprocal approximation and refinement), and it is expensive on GPUs with low FP64 throughput. The CSV reader and the JSON reader's type casting (data_casting) both parse floats with this function. The place value of a digit only depends on its position, so it is now looked up in a table of the results of the repeated division. The table is generated at compile time (constexpr) by the same division, and the device reads a copy in constant memory. In the CSV reader, the threads of a warp parse the same column of different rows and step through the fraction digits together, so their lookups mostly read the same element, which constant memory broadcasts (the JSON reader maps threads to values differently; divergent indices only serialize the constant-cache reads). The table holds the 323 nonzero results; the next division underflows to zero (a static_assert checks that the last element is at most 5 * denorm_min), and so does every later one, so digits past the table get 0 as before. Base-16 parsing keeps the division; no floating-point parse uses it. Exactness: the device's double division is correctly rounded (round to nearest even) and keeps subnormals, which is exactly how the constexpr evaluation computes the table, so every place value, and every parsed value, is unchanged. Checked by comparing the table with the device's run-time repeated division (all 323 elements identical, the 324th and 325th divisions give 0), and by a differential run against the parent build: 10M random strings (0-25 whole and fraction digits, 300-400 fraction digits, exponents around the double limits, subnormals, printf-formatted random doubles and floats, malformed strings) read as FLOAT64, FLOAT32 and INT64 by the CSV reader with three option sets and as FLOAT64 and FLOAT32 by the JSON reader: identical bits and validity. Measured on an RTX PRO 6000 (FP64 at 1/64 of the FP32 rate), decode kernel time for about 256 MB of CSV read from a device buffer, median of 80 reads: - doubles near 1e+307, as the CSV writer writes the random values of the ANALYTICS benchmark (2 INT64 + 6 FLOAT64 columns): 2044 -> 1070 us - the same doubles among the TAXI benchmark's column types: 1789 -> 1401 us - FLOAT32 values near 1000 ("872.6951904"): 2828 -> 963 us - doubles near 1e8 without exponent ("87269516.69"): 1153 -> 677 us - INT64 only: unchanged (779 -> 780 us) These inputs are generated by a custom program, not the benchmarks. On the CSV_READER_REALISTIC_NVBENCH data (256 MB, device buffer, minimum GPU time), the ANALYTICS read goes from 4.5 to 3.2 ms end to end. All measurements are on this sm_120 GPU with low FP64 throughput; the gain is expected on every GPU (a double division is a long dependent sequence everywhere), but was not measured on full-rate FP64 GPUs. Test: FractionDigitPlaceValues reads fields with a single nonzero fraction digit at positions 1-330 as FLOAT64 and FLOAT32 and compares their bits with the repeated division computed on the host, and checks that thousands separators and '+' in a fraction take no position. JsonReaderTest.FloatingPointBits checks the bits of JSON FLOAT64 and FLOAT32 reads of single fraction digits at positions 1, 2 and 322-324. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

parse_numeric accumulates the whole part of a number as `value = value * 10 + digit` in the output type, so for FLOAT64 fields it spends one double multiply-add per digit, which is slow on GPUs with low FP64 throughput. For doubles in base 10, the first 15 digits (numeric_limits<double>:: digits10) of the whole part are now accumulated in a uint64_t, and the double accumulation continues from the converted value for any further digits. The loop stops at the same characters as the double loop (decimal point, 'e', 'E'), skips the same ones (thousands separator, '+'), and decodes digits with the same validity check, leaving the decimal point to the double loop. Exactness: while the double accumulation has consumed at most 15 digits, every intermediate value is an integer below 10^15 < 2^53, and `value * 10 + digit` of such integers is exact whether or not the multiply-add is fused. So after the same digits the double loop would hold exactly the integer, whose conversion to double is exact, and continuing from it performs the same operations on the same values. The bound is 15 and not 16 on purpose: the 16th step can round, and taking it in the integer would only match if the compiler fuses the multiply-add (it does by default, so 16 happens to give the same results here, but that is not a guarantee). Leading zeros count as digits, which only makes the integer part shorter. Other output types are unchanged: floats keep accumulating in float (single precision is fast) and integer types already accumulate in integers. Checked by the differential run of the previous commit (10M random strings, CSV with three option sets and JSON, FLOAT64/FLOAT32/INT64) against the parent build: identical bits and validity. Measured on an RTX PRO 6000, decode kernel time for about 256 MB of CSV read from a device buffer, median of 80 reads: doubles near 1e8 written without exponent ("87269516.69", up to 9 whole digits): 677 -> 558 us. Doubles written with an exponent ("-4.743861625e+307") have a single whole digit and are unchanged. The input is generated by a custom program, not the benchmarks. This GPU (sm_120) has FP64 at 1/64 of the FP32 rate; on GPUs with full-rate FP64 (A100, H100, B200) the 64-bit integer multiply-add chain replaces a single DFMA per digit, and the effect there is unmeasured (possibly neutral). Test: WholeDigitsOfDoubles pins 15, 16 and 17 whole digits, including 2^53 + 1 and 10 * (2^53 + 1), which round at the 16th digit (the latter gives 10 * 2^53, not the nearest double to the decimal value, so accumulating 17 digits in the integer would change it), leading zeros, "-0", skipped thousands separators and '+', an invalid digit, 2^24 + 1 as FLOAT32, and upper-case exponents after 1, 2, 18 and 19 whole digits (compared with the same fields with a lower-case 'e'). FloatingPointSyntaxEdgeCases covers whole parts that end at a decimal point or an exponent without digits ("1.", ".5", "1e", "1e+", "1e-") and exponents beyond the double range ("1e-400", "1e0400", "1e309", "0e400" -> NaN -> null). JsonReaderTest.FloatingPointBits adds 16 whole digits, 2^53 + 1 and upper-case exponents read by the JSON reader. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

to_timestamp, which the CSV reader and the JSON reader's type casting use to parse timestamps, accepts many date and time layouts: it scans for the end of the date, searches the date for '/' and then '-', and searches the time of day for ':' and '.', after checking for an AM/PM suffix. A common layout is ISO 8601 `YYYY-MM-DDTHH:MM:SS[.fraction]`, which is how the CSV writer writes timestamps ("...SS.fffZ", or "...SSZ" for seconds), and it has every separator at a fixed position. Strings of 19 or more characters with digits and separators at those positions ('T' or ' ' between date and time), that end after the seconds, or with a 'Z' right after them, or continue with '.' and do not end with 'M' or 'm', are now parsed directly from their fixed positions. Exactness: for such strings, the general parsing finds every separator at the same fixed position, and then parses the same characters into the same types and combines them the same way: - the search for the end of the date sees no 'T' before position 10, and counts the two '-' and the digit after the second one, so a ' ' at position 10 ends the date too; - extract_date finds no '/', so it splits at the '-' at positions 4 and 7; the 4-digit first component makes the year come first whatever `dayfirst` is; - extract_time_of_day sees no AM/PM suffix (the last character is a digit, 'Z', or not 'M'/'m'), finds ':' at 13 and 16, and '.' at 19 if present; otherwise it parses the seconds from position 17 to the end, where to_non_negative_integer ignores the 'Z'. All other strings, including every near miss, take the general parsing unchanged. The characters after the '.' are parsed as milliseconds, as before (so the fast path keeps the general parsing's handling of fractions and suffixes). Checked by a differential run against the parent build: 10M strings (the layout with random and out-of-range components, 'T' and ' ', 0-24 fraction digits including an empty fraction, suffixes such as Z, +05:30, AM/PM, 'M'; mutated, truncated and extended near misses; and general layouts), read by the CSV and JSON readers with dayfirst off and on, as TIMESTAMP_DAYS/SECONDS/MILLISECONDS/MICROSECONDS/ NANOSECONDS: identical values. Measured on an RTX PRO 6000, decode kernel time for about 256 MB of CSV with the TAXI benchmark's column types (two TIMESTAMP_MILLISECONDS columns written by the CSV writer) read from a device buffer, median of 80 reads: 1408 -> 1252 us. Tests: FixedLayoutIsoTimestamps (CSV and JSON) reads each string next to the same string with '/' date separators, which takes the general parsing and parses '/'-free strings the same way, and compares them for all timestamp types (JSON: seconds and nanoseconds) and both dayfirst settings. The strings (new io_test_utils.hpp helpers) are hand-picked strings of the layout (both separators, 'Z', 0-12 fraction digits, "SS.", suffixes, AM/PM, out-of-range components) and about 2,000 generated near misses: three strings of the layout truncated to every length, and with a character deleted, or one of "0 -.:MPTZmtxz" substituted or inserted, at every position, or with a suffix such as "Z5", "Z+05:30" or " PM" appended. The CSV test also pins the current behavior that a fraction's digits count milliseconds whatever their number (".5" is 5 ms, which the fast path keeps). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

A serialized trie lookup compares the key with the nodes one character at a time and stops at the first character without a matching child. The CSV reader looks up every field in the NA trie, and fields decoded or classified as numbers or booleans also in the true and false tries; type detection looks up every field in all three. Most fields are longer than every key (the longest default NA value is the 8 characters "#N/A N/A"), but many share a prefix with one: numbers that start with "1" or "-" walk into "1.#IND", "-1.#QNAN" or "-nan" before they fail to match. The root node's children offset is never used: its children are the nodes that follow it, where lookups start. Store the length of the longest key there when the trie is built, and reject longer keys before the walk. A key longer than every key cannot be equal to one, so every result is unchanged. If the length does not fit in the 16-bit offset, the root keeps its offset of -1 and lookups walk the trie as before. The JSON reader and JSON type inference use the same tries and lookup. This also removes an out-of-bounds read: the trie of the empty key alone is its root node, and a lookup of any other key read the node after it. The CSV and JSON readers never build such a trie (they add default keys), but other users of `create_serialized_trie` can. New tests: a CSV read of keys that are prefixes of each other, the empty key, a duplicate key and non-ASCII keys, against fields that are equal to them, longer, or prefixes of them, with explicit and inferred types, and with the rows staged in shared memory and not (LIBCUDF_CSV_ROW_STAGING); and a new TRIE_TEST of lookups on the device, including the trie of the empty key alone. The effect depends on how many fields share a prefix with a key, and the measurements are noisy. Kernel durations of the 256 MB benchmarks, as ranges over three sets of interleaved nsys runs on an RTX PRO 6000 shared with other jobs (two at this commit, one at the end of the series, whose later commits do not change these kernels' work): - decode: TAXI -4% to -5%, LOGS -5% to -7%, ANALYTICS -8% to +1%; with inferred types and 64 columns, ALL_INTEGRAL -6% to -9%, ALL_STRING -2% to -4%, MIXED -6% to +4%; - type detection: ALL_STRING -18% to -21%, MIXED -12% to -20%, ALL_INTEGRAL -2% to -11%. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

`create_serialized_trie` builds a tree of heap-allocated nodes, each with 256 child pointers, before serializing it breadth-first. For the CSV reader's default NA values this takes about 15 us on the host, and every read builds three tries (NA, true and false values). The builder also keeps node indexes, and the serialized trie its children offsets, in 16-bit integers. In a trie of more than 32767 nodes, a node can be more than 32767 nodes before its first child, and the children of such a node were silently lost: the keys below it never matched. For example, of the 40000 NA values "k00000" to "k39999", "k39999" did not match, in both the CSV and the JSON reader. Serialize the trie directly from the distinct keys in sorted order. The keys below a node are then consecutive, with the key that ends at the node first, and the others grouped by their next character in ascending order: each group is a child. A breadth-first walk over (node, key range) pairs emits the same nodes, in the same order and with the same offsets, as the tree walk did. `std::string_view` compares characters as unsigned char, which is the order of the children that `serialized_trie_contains` relies on for non-ASCII keys. Store the children offsets in 32-bit integers, so that tries of any practical size are correct; a trie too large even for that throws std::overflow_error instead of losing keys. The length of the longest key now always fits in the root, so lookups no longer handle a root without it. Nodes grow from 4 to 8 bytes: in the same kernel measurements as the previous commit, the decode and type detection kernels show no consistent change. The output was compared with the previous builder on 30000 random sets of up to 24 keys (with duplicates, the empty key, and every byte value but the terminating '\n'): the nodes are identical in every case. Building the default NA trie takes 0.6 us instead of 15.2 us, and the default true or false trie 0.13 us instead of 2.9 us. The new CSV, JSON and TRIE_TEST tests of a trie of 40000 keys, and the TRIE_TEST test of a key of 40000 characters, fail without this change. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

`make_parse_options` builds the tries of the true, false and NA values with a `create_serialized_trie` each, and each one uploads its nodes with a synchronous copy: three stream synchronizations at the start of every read_csv. Add `create_serialized_tries`, which serializes several sets of keys, uploads all their tries, and then synchronizes the stream once. The synchronization is still needed: the copies may read the host nodes only when the stream reaches them, so the nodes must outlive the copies. The CSV reader builds its three tries with it. A set without keys (the NA values with na_filter off) gives an empty trie, which matches nothing, like the empty trie the reader created before. A read makes two fewer stream synchronizations (nsys, 1 KiB file of 8 columns with inferred types: 30 instead of 32). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

`decode_data` allocates each output buffer with `column_buffer::create`, which allocates an all-null mask with a memset of its own, and zeroes the (pointer, length) pairs of each string column with another memset. The masks of string columns are never used: string columns get their null masks from the decoded pairs, when they are built. Allocate the masks of the other columns uninitialized, and zero them all with `cudf::detail::batched_memset`, as the Parquet reader does; give the buffers of string columns no mask. Each mask is zeroed over its whole allocation, as `create_null_mask` does, so the masks are the same, including the padding past the last row. Instead of one memset per column, a read makes one upload of the mask spans and the three kernel launches of `cub::DeviceCopy::Batched`, a fixed cost that pays off with more than a handful of columns (nsys, per read of a 1 KiB file: 8 columns, 2 of them strings, 8 memsets instead of 16, plus 1 copy and 3 launches; 64 columns, 16 of them strings, 35 memsets instead of 99). The new test reads 100 columns of every decoded type, strings among them, and checks every null mask up to the end of its allocation, with the rows staged in shared memory and not. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

Type detection uploads the column flags and zeroes its counters with a memset, then copies the counters back into a new pinned host vector. Decoding uploads the column flags, the column types, and the column data and mask pointers with four copies, from two new pinned host vectors. Allocating pinned host memory synchronizes the stream, so the counters' allocation waits for the detection kernel, and the read then waits again for their copy; `infer_column_types` also synchronizes once more after detection, for nothing. Add `hostdevice_arrays`, which places arrays of different element types in one `hostdevice_vector`, each at an offset aligned for its type, and returns them as `hostdevice_span`s. Like `hostdevice_vector`, it copies all arrays between the host and the device at once, and skips the copies when the device accesses the host memory directly (the integrated memory optimization). Detection fills its flags and zeroed counters on the host, uploads them with one copy, and copies them back to the same host memory, which is allocated before the kernel is launched. Decoding allocates its four arrays first, while the stream is idle, fills them as it allocates the output buffers, and uploads them with one copy. The column flags and types that `read_csv` keeps on the host are no longer uploaded directly, so they no longer need pinned memory, whose allocations synchronized the stream too. The new SPAN_TEST tests check the layout and the copies, with and without the integrated memory optimization. A read makes four fewer stream synchronizations, three fewer copies and one fewer memset (nsys, 1 KiB file of 8 columns with inferred types: 26 synchronizations instead of 30, 23 copies instead of 26). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

convert_csv_to_cudf read column_flags[col] and dtypes[actual_col] from global memory at every use, up to eight times per field. The kernel stores the decoded values through untyped column pointers, which the compiler must assume can alias the flags and types, so it cannot keep them in registers across those stores and reloads them after each one. Load the column's flags once per field, and its type once per enabled field (only enabled columns have a type: `actual_col` can run past the end of `dtypes` for disabled trailing columns). The values are the same as before, so the output is unchanged. This is a code clarity and code generation change with no measurable end-to-end effect: register and stack usage of both kernel instances is unchanged, and decode kernel time moves by -0.5% to -0.9% (TAXI, QUOTE_100, LOGS; nsys, 256 MB inputs, median of 20 reads, two interleaved runs), which is within noise. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

The CSV reader built each string column separately with make_strings_column: a scan for the offsets, a valid_if kernel whose null count synchronizes the stream, and cub::DeviceMemcpy::Batched for the characters. A read with 64 string columns (QUOTE_100) ran some 450 small kernels and 64 synchronizations, and the batched memcpy alone took 1.8 ms of its 4.6 ms of GPU time. Build all the string columns of a read together, in a new make_strings_columns (io/csv/string_columns.cu), called by decode_data, which now returns the finished columns: - compute_validity_and_sizes_kernel reads the (pointer, size) pairs of a block of 256 rows of one column per thread block. Each warp writes the mask word of its 32 rows, and the block adds its valid rows and characters to the column's counts and records its characters. - A scan by column of the blocks' characters gives each block's offset. - One copy of the counts to the host, the only synchronization of the string columns, sizes the offsets (64-bit from LIBCUDF_LARGE_STRINGS_THRESHOLD on, through create_offsets_child_column, which also raises std::overflow_error when large strings are disabled, as before) and the characters. The fixed-width columns' null counts still take their own synchronization (batch_null_count), so a decode has two. - write_strings_kernel scans the sizes within each block, writes the offsets, and copies the characters. The strings of a warp's 32 rows are contiguous in the output, and the warp writes them in windows of one aligned 4-byte word per lane, which it assembles from the bytes of the strings the word overlaps (found by a binary search over the rows' offsets in shared memory). Words that overlap other warps' characters are stored byte by byte, so warps never write each other's bytes. The blocks of the same rows of all columns are adjacent in the grid, so their strings, which are close in the source, share cached reads. Each block handles a single block of rows (the grid is 2D only for more blocks than its x dimension allows): a grid-stride loop costs the kernel six registers, and the copy needs the occupancy. The number of blocks of rows comes from the host, so it cannot overflow. The warp copy is fast for the short strings typical of CSV, but a warp copies about one window per memory latency, so long strings would keep single warps busy, and few rows of long strings leave too few warps to hide the latency (1 MiB strings: 220 ms instead of 0.35 ms). Strings of more than 1 KiB are therefore listed by the warps and copied afterwards with one batched memcpy; they split a warp's rows into runs whose strings are copied as above. With strings of a single size, the warp copy is faster up to about 1 KiB and the batched memcpy from about 2 KiB. This cutoff (and the 256-row blocks) were measured on one GPU (RTX PRO 6000 Blackwell, sm_120). The worst case is inputs made only of strings of exactly 1 KiB, the longest the warps copy: between +1% and +14% of GPU time per 256 MB read in separate runs (+11 to +190 us; the high end was measured with the grid-stride loop described above). Inputs of 2-4 KiB strings take up to 80 us (5%) more, for the two kernels that the batched memcpy now follows. Fields longer than a string can be (2^31 - 1 characters) were truncated to 32 bits before: fields of 2^31 to 2^32 - 1 characters got negative sizes, and the read failed with an allocation of about 16 EiB, while longer fields were silently read as wrong strings or nulls. The decode kernel now writes a distinct size for them, the validity kernel counts them (without adding them to the character counts), and make_strings_columns raises std::overflow_error on the host before any character is copied. Its pairs are now written as cuda::std::pair<char const*, size_type>, the type the string buffers hold, rather than as std::pair<char const*, size_t>. The characters are always copied, so the columns never alias the input, which can be the caller's device buffer. The builder only depends on the (pointer, size) pairs; it lives with the CSV reader because it relies on all columns having the same number of rows, and because make_strings_column_batch, which it could replace, has other users and columns of different lengths. Moving the warp copy to cudf::strings::detail (make_chars_buffer, make_strings_column_batch) is a possible follow-up. Tests (each under LIBCUDF_CSV_ROW_STAGING=ALWAYS and NEVER unless noted): - StringColumnsAcrossWarpsAndBlocks: 1 to 1000 rows around the warp and block sizes; all-null, empty, edge-null and escaped-quote columns. - ManyStringColumnsWithMissingFields: 70 columns, rows missing up to 69 trailing fields, selected columns. - MoreStringColumnsThanGridDimension: 70,000 columns (one policy; the decoding is not what it tests). - LongStringsAmongShortStrings: 1000 B to 1 MiB strings at warp and block edges. - StringColumnsWithLargeOffsets: 64-bit offsets from a low LIBCUDF_LARGE_STRINGS_THRESHOLD, std::overflow_error without large strings. - StringsOfDeviceBuffersAreCopied: the columns do not change when the source device buffer is overwritten after the read. - InferredStringColumnsOfRaggedRows, WhitespaceDelimitedStringColumnsOf- ShortRows, SelectedStringColumnsOfShortRows, LeadingStringColumnsOf- ShortRows: fields missing from short and blank rows with inferred types, delim_whitespace (including rows of whitespace only, the last one without a terminator) and use_cols (including 68 trailing columns that are not selected). The whitespace-delimited test pins the current strings, which keep all but one of the whitespace characters that follow a field, as before this change. - LongStringsAtEndOfUnalignedDeviceBuffers: short, long and escaped strings of device buffers at every alignment, ending at the end of their allocations. - StringColumnsOfByteRangesAndRowSelections, HeaderOnlyStringColumns. - LARGE_STRINGS_TEST CsvLargeReaderTest.StringColumnOverTwoGigabytes (a column of more than 2^31 characters, compared byte for byte) and DISABLED_FieldLongerThanStringSizeLimit (std::overflow_error; disabled because the one thread of its 2 GiB row takes about a minute; before this change the read failed with a 16 EiB allocation instead). The tests of fields missing from short rows and of long strings fill every device allocation with a byte pattern first, so that memory the reader reads without writing it changes their results deterministically. The csv_test.cpp environment variable guard of the row staging policy is generalized to set the large strings variables too. Measured before the rebase onto the host latency changes, on the float and timestamp parsing changes: GPU time per read of the 256 MB benchmark inputs (nsys, all kernels and memsets, median of 20 reads): QUOTE_100 4601 -> 2696 us, QUOTE_25 4058 -> 3623 us, LOGS 2206 -> 1955 us, TAXI 3220 -> 3190 us, ALL_STRING with type inference 6688 -> 4781 us. Reader benchmarks (nvbench, best of 3 interleaved runs): QUOTE_100 8.37 -> 3.34 ms (-60%), ALL_STRING 8.37 -> 3.22 ms (-61%) and 10.48 -> 5.27 ms with inference (-50%), QUOTE_25 5.18 -> 3.90 ms (-25%), LOGS 2.51 -> 2.16 ms (-14%); TAXI and ANALYTICS within noise. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

The decode kernel wrote a 16-byte (pointer, size) pair per row of each string column, into arrays that had to be zeroed beforehand: fields missing from the end of short rows were not written, and relied on the zeroed pair reading as null. The zeroing wrote as much memory as the kernel (140 us per QUOTE_100 read), and the string column kernels read all 16 bytes of every pair, although they only need the sizes to compute validity and offsets. Decode the strings of a string column to `decoded_strings`: the pointers of all rows, followed by their 4-byte sizes, with a size of -1 for a null row (whose pointer is not written) and -2 for a field too long for a string. The validity and sizes kernel now reads 4 bytes per row, and the write kernel reads pointers only for non-empty strings. The memory is not initialized: the decode kernel now also sets the strings of the fields missing from the end of a short row to null, after the row's last field, so every row of every string column is written. The tests of the previous commit on short, blank and whitespace-delimited rows and ShortRowsMissingFieldsAreNull cover this, and compute-sanitizer initcheck reports no read of the decoded strings: over the reader tests, the only reports in the string kernels are 1709 one-byte reads of source characters in write_strings_kernel, all in tests that read files, the same number for each test as the parent's reports in cub's batched memcpy. Where the sizes start depends on the number of rows, so the host passes the number of rows the memory is sized for to decode_row_column_data, which checks it against the row offsets, instead of both sides deriving it. The decoded strings of all string columns share one scratch allocation, and the string columns no longer get column buffers. Measured before the rebase onto the host latency changes, on the float and timestamp parsing changes: GPU time per read of the 256 MB benchmark inputs (nsys, all kernels and memsets, median of 20 reads): QUOTE_100 2696 -> 2415 us (-10%), LOGS 1955 -> 1786 us (-9%), QUOTE_25 3623 -> 3417 us (-6%), ALL_STRING with type inference 4781 -> 4475 us (-6%), TAXI 3190 -> 3141 us (-2%). Reader benchmarks (nvbench, best of 3 interleaved runs): QUOTE_100 3.34 -> 2.72 ms (-19%), ALL_STRING 3.22 -> 2.77 ms (-14%) and 5.27 -> 4.65 ms with inference (-12%), QUOTE_25 3.90 -> 3.63 ms (-7%), LOGS 2.16 -> 2.03 ms (-6%); ANALYTICS (no strings) unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

karthikeyann · 2026-09-25T06:34:21Z

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.

Per-PR contributions are in the benchmark comments on #16–#24. Note PR 6 (#21) trades ~4% on pageable-host and file reads for its device-buffer gain.

CSV reader benchmarks: combined stack (#16–#24) vs coral/csv-base

This PR 3ffd97beee vs base coral/csv-base 63712844b0. Measured on one NVIDIA RTX PRO 6000 Blackwell Server Edition, CUDA 13.3, Release build, with no builds running during the measurements. Throughput is NVBench bytes_per_second (uncompressed table bytes / mean GPU time); each side is the best of 2 run(s) (this PR) and 2 run(s) (base) per config. Run-to-run noise is a few percent, so treat changes under 3% as noise. Bold means at least 3% faster.

Speedup vs base (cumulative), by input source

Config	Device buffer	Pinned host	Pageable host	File (kvikio)
realistic TAXI 256MB	3.86× _{73.0 GiB/s}	1.66× _{18.6 GiB/s}	1.17× _{5.6 GiB/s}	1.60× _{10.2 GiB/s}
realistic LOGS 256MB	4.02× _{109.2 GiB/s}	1.75× _{29.9 GiB/s}	1.20× _{8.6 GiB/s}	1.58× _{15.6 GiB/s}
realistic ANALYTICS 256MB	2.94× _{74.4 GiB/s}	1.43× _{18.0 GiB/s}	1.06× _{5.1 GiB/s}	1.38× _{9.2 GiB/s}
inference MIXED 8 cols	2.78× _{58.9 GiB/s}	1.44× _{15.7 GiB/s}	1.15× _{4.7 GiB/s}	1.49× _{8.6 GiB/s}
inference MIXED 64 cols	3.86× _{25.4 GiB/s}	2.19× _{10.7 GiB/s}	1.45× _{3.8 GiB/s}	2.05× _{6.5 GiB/s}
no inference MIXED 8 cols	3.37× _{98.2 GiB/s}	1.34× _{17.3 GiB/s}	1.06× _{4.7 GiB/s}	1.38× _{8.8 GiB/s}
25% quoted 256MB	5.12× _{65.3 GiB/s}	2.39× _{23.4 GiB/s}	1.54× _{7.7 GiB/s}	2.04× _{13.2 GiB/s}
100% quoted 256MB	17.15× _{83.5 GiB/s}	6.78× _{36.6 GiB/s}	2.78× _{11.6 GiB/s}	4.38× _{20.5 GiB/s}
Weighted geomean	4.24×	1.96×	1.33×	1.80×

_{Sources: the *_IO_NVBENCH benchmarks (same data as the primary benchmarks, 256MB, with the input in device memory, pinned host memory, pageable host memory, or a file read via kvikio; file reads hit the OS page cache). Weighted geomean uses weight 2 for TAXI, MIXED-8 with inference and 25% quoted, weight 1 otherwise.}

Full results: every CSV reader benchmark config

`CSV_READER_NVBENCH` (18 configs)

Benchmark	Config	Base	This PR	vs base
csv_read_column_selection	column_selection=ALL, num_chunks=1, 10BC8 row_selection=ALL	2.63 GiB/s	3.50 GiB/s	1.33×
csv_read_column_selection	column_selection=ALTERNATE, num_chunks=1, row_selection=ALL	1.45 GiB/s	1.77 GiB/s	1.22×
csv_read_column_selection	column_selection=FIRST_HALF, num_chunks=1, row_selection=ALL	1.49 GiB/s	1.84 GiB/s	1.24×
csv_read_column_selection	column_selection=SECOND_HALF, num_chunks=1, row_selection=ALL	1.49 GiB/s	1.84 GiB/s	1.23×
csv_read_data_type	data_type=DECIMAL, io=DEVICE_BUFFER	15.33 GiB/s	44.78 GiB/s	2.92×
csv_read_data_type	data_type=DURATION, io=DEVICE_BUFFER	12.78 GiB/s	40.74 GiB/s	3.19×
csv_read_data_type	data_type=FLOAT, io=DEVICE_BUFFER	18.66 GiB/s	49.68 GiB/s	2.66×
csv_read_data_type	data_type=INTEGRAL, io=DEVICE_BUFFER	17.43 GiB/s	47.83 GiB/s	2.74×
csv_read_data_type	data_type=STRING, io=DEVICE_BUFFER	4.68 GiB/s	82.10 GiB/s	17.53×
csv_read_data_type	data_type=TIMESTAMP, io=DEVICE_BUFFER	12.37 GiB/s	48.06 GiB/s	3.88×
csv_read_io	io=FILEPATH	4.55 GiB/s	6.91 GiB/s	1.52×
csv_read_io	io=HOST_BUFFER	3.38 GiB/s	3.74 GiB/s	1.10×
csv_read_row_selection	column_selection=ALL, num_chunks=1, row_selection=BYTE_RANGE	2.56 GiB/s	3.23 GiB/s	1.26×
csv_read_row_selection	column_selection=ALL, num_chunks=1, row_selection=NROWS	2.54 GiB/s	3.22 GiB/s	1.27×
csv_read_row_selection	column_selection=ALL, num_chunks=1, row_selection=SKIPFOOTER	2.67 GiB/s	3.48 GiB/s	1.30×
csv_read_row_selection	column_selection=ALL, num_chunks=8, row_selection=BYTE_RANGE	1.93 GiB/s	3.16 GiB/s	1.64×
csv_read_row_selection	column_selection=ALL, num_chunks=8, row_selection=NROWS	0.71 GiB/s	0.75 GiB/s	1.06×
csv_read_row_selection	column_selection=ALL, num_chunks=8, row_selection=SKIPFOOTER	0.45 GiB/s	0.45 GiB/s	1.01×

`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	48.14 GiB/s	2.81×
csv_read_quoting_io	data_size_mb=256, io=DEVICE_BUFFER, quote_pct=QUOTE_100_PCT	4.87 GiB/s	83.50 GiB/s	17.15×
csv_read_quoting_io	data_size_mb=256, io=DEVICE_BUFFER, quote_pct=QUOTE_25_PCT	12.75 GiB/s	65.28 GiB/s	5.12×
csv_read_quoting_io	data_size_mb=256, io=FILEPATH, quote_pct=QUOTE_0_PCT	5.25 GiB/s	7.86 GiB/s	1.50×
csv_read_quoting_io	data_size_mb=256, io=FILEPATH, quote_pct=QUOTE_100_PCT	4.69 GiB/s	20.53 GiB/s	4.38×
csv_read_quoting_io	data_size_mb=256, io=FILEPATH, quote_pct=QUOTE_25_PCT	6.47 GiB/s	13.20 GiB/s	2.04×
csv_read_quoting_io	data_size_mb=256, io=HOST_BUFFER, quote_pct=QUOTE_0_PCT	3.80 GiB/s	4.52 GiB/s	1.19×
csv_read_quoting_io	data_size_mb=256, io=HOST_BUFFER, quote_pct=QUOTE_100_PCT	4.18 GiB/s	11.64 GiB/s	2.78×
csv_read_quoting_io	data_size_mb=256, io=HOST_BUFFER, quote_pct=QUOTE_25_PCT	4.99 GiB/s	7.68 GiB/s	1.54×
csv_read_quoting_io	data_size_mb=256, io=PINNED_BUFFER, quote_pct=QUOTE_0_PCT	9.57 GiB/s	14.21 GiB/s	1.49×
csv_read_quoting_io	data_size_mb=256, io=PINNED_BUFFER, quote_pct=QUOTE_100_PCT	5.40 GiB/s	36.61 GiB/s	6.78×
csv_read_quoting_io	data_size_mb=256, io=PINNED_BUFFER, quote_pct=QUOTE_25_PCT	9.77 GiB/s	23.38 GiB/s	2.39×

`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	48.46 GiB/s	2.77×
csv_read_quoting	data_size_mb=256, quote_pct=QUOTE_100_PCT	4.80 GiB/s	81.74 GiB/s	17.04×
csv_read_quoting	data_size_mb=256, quote_pct=QUOTE_25_PCT	13.88 GiB/s	65.53 GiB/s	4.72×
csv_read_quoting	data_size_mb=64, quote_pct=QUOTE_0_PCT	14.10 GiB/s	34.19 GiB/s	2.42×
csv_read_quoting	data_size_mb=64, quote_pct=QUOTE_100_PCT	1.37 GiB/s	46.36 GiB/s	33.91×
csv_read_quoting	data_size_mb=64, quote_pct=QUOTE_25_PCT	4.57 GiB/s	45.54 GiB/s	9.97×

`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	74.41 GiB/s	2.94×
csv_read_realistic_io	data_size_mb=256, io=DEVICE_BUFFER, profile=LOGS	27.19 GiB/s	109.22 GiB/s	4.02×
csv_read_realistic_io	data_size_mb=256, io=DEVICE_BUFFER, profile=TAXI	18.91 GiB/s	73.05 GiB/s	3.86×
csv_read_realistic_io	data_size_mb=256, io=FILEPATH, profile=ANALYTICS	6.63 GiB/s	9.19 GiB/s	1.38×
csv_read_realistic_io	data_size_mb=256, io=FILEPATH, profile=LOGS	9.85 GiB/s	15.58 GiB/s	1.58×
csv_read_realistic_io	data_size_mb=256, io=FILEPATH, profile=TAXI	6.37 GiB/s	10.20 GiB/s	1.60×
csv_read_realistic_io	data_size_mb=256, io=HOST_BUFFER, profile=ANALYTICS	4.83 GiB/s	5.11 GiB/s	1.06×
csv_read_realistic_io	data_size_mb=256, io=HOST_BUFFER, profile=LOGS	7.16 GiB/s	8.62 GiB/s	1.20×
csv_read_realistic_io	data_size_mb=256, io=HOST_BUFFER, profile=TAXI	4.77 GiB/s	5.60 GiB/s	1.17×
csv_read_realistic_io	data_size_mb=256, io=PINNED_BUFFER, profile=ANALYTICS	12.58 GiB/s	17.97 GiB/s	1.43×
csv_read_realistic_io	data_size_mb=256, io=PINNED_BUFFER, profile=LOGS	17.05 GiB/s	29.87 GiB/s	1.75×
csv_read_realistic_io	data_size_mb=256, io=PINNED_BUFFER, profile=TAXI	11.20 GiB/s	18.58 GiB/s	1.66×

`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	89.74 GiB/s	3.57×
csv_read_realistic	data_size_mb=1024, profile=LOGS	31.88 GiB/s	126.78 GiB/s	3.98×
csv_read_realistic	data_size_mb=1024, profile=TAXI	23.43 GiB/s	88.40 GiB/s	3.77×
csv_read_realistic	data_size_mb=256, profile=ANALYTICS	25.24 GiB/s	78.19 GiB/s	3.10×
csv_read_realistic	data_size_mb=256, profile=LOGS	27.37 GiB/s	112.29 GiB/s	4.10×
csv_read_realistic	data_size_mb=256, profile=TAXI	19.01 GiB/s	72.76 GiB/s	3.83×
csv_read_realistic	data_size_mb=512, profile=ANALYTICS	25.64 GiB/s	85.21 GiB/s	3.32×
csv_read_realistic	data_size_mb=512, profile=LOGS	31.05 GiB/s	129.74 GiB/s	4.18×
csv_read_realistic	data_size_mb=512, profile=TAXI	22.88 GiB/s	82.61 GiB/s	3.61×

`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.91 GiB/s	1.10×
csv_read_scale	data_size_mb=2048	3.52 GiB/s	3.91 GiB/s	1.11×
csv_read_scale	data_size_mb=256	3.59 GiB/s	3.89 GiB/s	1.08×
csv_read_scale	data_size_mb=4096	3.38 GiB/s	3.97 GiB/s	1.17×
csv_read_scale	data_size_mb=512	3.60 GiB/s	3.92 GiB/s	1.09×

`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	25.28 GiB/s	2.50×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	9.02 GiB/s	25.90 GiB/s	2.87×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_STRING	4.40 GiB/s	54.43 GiB/s	12.36×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=MIXED	6.59 GiB/s	25.42 GiB/s	3.86×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_FLOAT	12.31 GiB/s	40.34 GiB/s	3.28×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	21.14 GiB/s	58.63 GiB/s	2.77×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_STRING	11.91 GiB/s	71.39 GiB/s	5.99×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=MIXED	21.15 GiB/s	58.88 GiB/s	2.78×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_FLOAT	4.03 GiB/s	6.22 GiB/s	1.54×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_INTEGRAL	4.12 GiB/s	6.72 GiB/s	1.63×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_STRING	3.94 GiB/s	18.36 GiB/s	4.66×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=MIXED	3.16 GiB/s	6.49 GiB/s	2.05×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_FLOAT	4.37 GiB/s	7.04 GiB/s	1.61×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_INTEGRAL	5.98 GiB/s	7.97 GiB/s	1.33×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_STRING	7.49 GiB/s	18.70 GiB/s	2.50×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=MIXED	5.79 GiB/s	8.63 GiB/s	1.49×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_FLOAT	3.07 GiB/s	3.67 GiB/s	1.20×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	3.21 GiB/s	3.79 GiB/s	1.18×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_STRING	3.50 GiB/s	11.06 GiB/s	3.16×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=MIXED	2.60 GiB/s	3.78 GiB/s	1.45×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_FLOAT	3.30 GiB/s	3.75 GiB/s	1.14×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	4.04 GiB/s	4.53 GiB/s	1.12×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_STRING	6.36 GiB/s	11.43 GiB/s	1.80×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=MIXED	4.06 GiB/s	4.67 GiB/s	1.15×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_FLOAT	6.80 GiB/s	10.86 GiB/s	1.60×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	6.43 GiB/s	11.43 GiB/s	1.78×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_STRING	4.36 GiB/s	29.75 GiB/s	6.82×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=MIXED	4.90 GiB/s	10.74 GiB/s	2.19×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_FLOAT	7.73 GiB/s	12.97 GiB/s	1.68×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	10.97 GiB/s	15.69 GiB/s	1.43×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_STRING	10.30 GiB/s	33.90 GiB/s	3.29×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=MIXED	10.97 GiB/s	15.75 GiB/s	1.44×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_FLOAT	18.66 GiB/s	48.62 GiB/s	2.60×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	17.43 GiB/s	49.31 GiB/s	2.83×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_STRING	4.91 GiB/s	81.90 GiB/s	16.67×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=MIXED	14.50 GiB/s	50.83 GiB/s	3.50×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_FLOAT	21.45 GiB/s	57.74 GiB/s	2.69×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	28.99 GiB/s	92.27 GiB/s	3.18×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_STRING	20.79 GiB/s	101.96 GiB/s	4.91×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=MIXED	29.16 GiB/s	98.25 GiB/s	3.37×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_FLOAT	5.12 GiB/s	7.41 GiB/s	1.45×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_INTEGRAL	5.02 GiB/s	8.11 GiB/s	1.61×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_STRING	4.78 GiB/s	20.46 GiB/s	4.28×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=MIXED	4.60 GiB/s	7.23 GiB/s	1.57×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_FLOAT	5.29 GiB/s	7.69 GiB/s	1.45×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_INTEGRAL	6.01 GiB/s	9.00 GiB/s	1.50×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_STRING	11.19 GiB/s	21.69 GiB/s	1.94×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=MIXED	6.39 GiB/s	8.79 GiB/s	1.38×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_FLOAT	3.75 GiB/s	4.11 GiB/s	1.09×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	3.79 GiB/s	4.44 GiB/s	1.17×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_STRING	4.25 GiB/s	11.78 GiB/s	2.77×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=MIXED	3.33 GiB/s	3.96 GiB/s	1.19×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_FLOAT	3.90 GiB/s	4.02 GiB/s	1.03×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	4.30 GiB/s	4.79 GiB/s	1.11×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_STRING	8.50 GiB/s	11.67 GiB/s	1.37×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=MIXED	4.49 GiB/s	4.74 GiB/s	1.06×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_FLOAT	9.60 GiB/s	13.54 GiB/s	1.41×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	9.67 GiB/s	14.32 GiB/s	1.48×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_STRING	5.47 GiB/s	36.50 GiB/s	6.67×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=MIXED	7.96 GiB/s	13.50 GiB/s	1.70×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_FLOAT	10.51 GiB/s	14.34 GiB/s	1.36×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	12.94 GiB/s	17.28 GiB/s	1.34×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_STRING	17.06 GiB/s	39.10 GiB/s	2.29×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=MIXED	12.86 GiB/s	17.26 GiB/s	1.34×

`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	25.43 GiB/s	2.50×
csv_read_with_inference	data_size_mb=256, num_cols=64, type_profile=ALL_INTEGRAL	9.10 GiB/s	25.95 GiB/s	2.85×
csv_read_with_inference	data_size_mb=256, num_cols=64, type_profile=ALL_STRING	4.40 GiB/s	53.74 GiB/s	12.20×
csv_read_with_inference	data_size_mb=256, num_cols=64, type_profile=MIXED	6.80 GiB/s	26.12 GiB/s	3.84×
csv_read_with_inference	data_size_mb=256, num_cols=8, type_profile=ALL_FLOAT	12.32 GiB/s	40.42 GiB/s	3.28×
csv_read_with_inference	data_size_mb=256, num_cols=8, type_profile=ALL_INTEGRAL	21.29 GiB/s	59.48 GiB/s	2.79×
csv_read_with_inference	data_size_mb=256, num_cols=8, type_profile=ALL_STRING	12.00 GiB/s	70.79 GiB/s	5.90×
csv_read_with_inference	data_size_mb=256, num_cols=8, type_profile=MIXED	21.64 GiB/s	59.06 GiB/s	2.73×
csv_read_with_inference	data_size_mb=64, num_cols=64, type_profile=ALL_FLOAT	10.15 GiB/s	18.58 GiB/s	1.83×
csv_read_with_inference	data_size_mb=64, num_cols=64, type_profile=ALL_INTEGRAL	8.33 GiB/s	20.01 GiB/s	2.40×
csv_read_with_inference	data_size_mb=64, num_cols=64, type_profile=ALL_STRING	1.42 GiB/s	31.85 GiB/s	22.37×
csv_read_with_inference	data_size_mb=64, num_cols=64, type_profile=MIXED	4.05 GiB/s	22.13 GiB/s	5.46×
csv_read_with_inference	data_size_mb=64, num_cols=8, type_profile=ALL_FLOAT	11.10 GiB/s	30.44 GiB/s	2.74×
csv_read_with_inference	data_size_mb=64, num_cols=8, type_profile=ALL_INTEGRAL	18.41 GiB/s	43.05 GiB/s	2.34×
csv_read_with_inference	data_size_mb=64, num_cols=8, type_profile=ALL_STRING	5.52 GiB/s	46.74 GiB/s	8.47×
csv_read_with_inference	data_size_mb=64, num_cols=8, type_profile=MIXED	19.86 GiB/s	42.44 GiB/s	2.14×
csv_read_without_inference	data_size_mb=256, num_cols=64, type_profile=ALL_FLOAT	18.10 GiB/s	49.38 GiB/s	2.73×
csv_read_without_inference	data_size_mb=256, num_cols=64, type_profile=ALL_INTEGRAL	17.51 GiB/s	48.80 GiB/s	2.79×
csv_read_without_inference	data_size_mb=256, num_cols=64, type_profile=ALL_STRING	4.70 GiB/s	82.06 GiB/s	17.46×
csv_read_without_inference	data_size_mb=256, num_cols=64, type_profile=MIXED	14.57 GiB/s	50.66 GiB/s	3.48×
csv_read_without_inference	data_size_mb=256, num_cols=8, type_profile=ALL_FLOAT	21.55 GiB/s	57.18 GiB/s	2.65×
csv_read_without_inference	data_size_mb=256, num_cols=8, type_profile=ALL_INTEGRAL	29.30 GiB/s	95.63 GiB/s	3.26×
csv_read_without_inference	data_size_mb=256, num_cols=8, type_profile=ALL_STRING	21.92 GiB/s	104.41 GiB/s	4.76×
csv_read_without_inference	data_size_mb=256, num_cols=8, type_profile=MIXED	29.24 GiB/s	93.95 GiB/s	3.21×
csv_read_without_inference	data_size_mb=64, num_cols=64, type_profile=ALL_FLOAT	15.91 GiB/s	30.38 GiB/s	1.91×
csv_read_without_inference	data_size_mb=64, num_cols=64, type_profile=ALL_INTEGRAL	14.41 GiB/s	34.02 GiB/s	2.36×
csv_read_without_inference	data_size_mb=64, num_cols=64, type_profile=ALL_STRING	1.44 GiB/s	47.35 GiB/s	32.94×
csv_read_without_inference	data_size_mb=64, num_cols=64, type_profile=MIXED	9.73 GiB/s	33.36 GiB/s	3.43×
csv_read_without_inference	data_size_mb=64, num_cols=8, type_profile=ALL_FLOAT	18.58 GiB/s	41.80 GiB/s	2.25×
csv_read_without_inference	data_size_mb=64, num_cols=8, type_profile=ALL_INTEGRAL	25.71 GiB/s	60.84 GiB/s	2.37×
csv_read_without_inference	data_size_mb=64, num_cols=8, type_profile=ALL_STRING	7.32 GiB/s	62.79 GiB/s	8.57×
csv_read_without_inference	data_size_mb=64, num_cols=8, type_profile=MIXED	25.23 GiB/s	60.14 GiB/s	2.38×

karthikeyann and others added 30 commits September 24, 2026 02:44

karthikeyann and others added 5 commits September 24, 2026 19:29

github-actions Bot assigned karthikeyann Sep 25, 2026

github-actions Bot added libcudf CMake labels Sep 25, 2026

karthikeyann mentioned this pull request Sep 26, 2026

[CSV merge] 2.8x faster CSV reader with +836/−328 library lines (device 3.3x, pinned 1.9x, pageable 3.8x, file 2.6x) #26

Draft

3 tasks done

Navigation Menu

Repository navigation

[CSV all] Upstreamable CSV reader optimizations (combined #16–#24) - #25

[CSV all] Upstreamable CSV reader optimizations (combined #16–#24)#25
karthikeyann wants to merge 35 commits into
coral/csv-basefrom
coral/csv-upstreamable

Uh oh!

`CSV_READER_NVBENCH` (18 configs)

`CSV_READER_QUOTING_IO_NVBENCH` (12 configs)

`CSV_READER_QUOTING_NVBENCH` (6 configs)

`CSV_READER_REALISTIC_IO_NVBENCH` (12 configs)

`CSV_READER_REALISTIC_NVBENCH` (9 configs)

`CSV_READER_SCALE_NVBENCH` (5 configs)

`CSV_READER_TYPE_INFERENCE_IO_NVBENCH` (64 configs)

`CSV_READER_TYPE_INFERENCE_NVBENCH` (32 configs)

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Conversation

Uh oh!

CSV reader benchmarks: combined stack (#16–#24) vs coral/csv-base

Speedup vs base (cumulative), by input source

CSV_READER_NVBENCH (18 configs)

CSV_READER_QUOTING_IO_NVBENCH (12 configs)

CSV_READER_QUOTING_NVBENCH (6 configs)

CSV_READER_REALISTIC_IO_NVBENCH (12 configs)

CSV_READER_REALISTIC_NVBENCH (9 configs)

CSV_READER_SCALE_NVBENCH (5 configs)

CSV_READER_TYPE_INFERENCE_IO_NVBENCH (64 configs)

CSV_READER_TYPE_INFERENCE_NVBENCH (32 configs)

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

`CSV_READER_NVBENCH` (18 configs)

`CSV_READER_QUOTING_IO_NVBENCH` (12 configs)

`CSV_READER_QUOTING_NVBENCH` (6 configs)

`CSV_READER_REALISTIC_IO_NVBENCH` (12 configs)

`CSV_READER_REALISTIC_NVBENCH` (9 configs)

`CSV_READER_SCALE_NVBENCH` (5 configs)

`CSV_READER_TYPE_INFERENCE_IO_NVBENCH` (64 configs)

`CSV_READER_TYPE_INFERENCE_NVBENCH` (32 configs)