[CSV 2/9] Unescape CSV doubled quotes in the decode kernel; skip zero-initializing fixed-width output by karthikeyann · Pull Request #17 · karthikeyann/cudf

karthikeyann · 2026-09-24T07:25:41Z

Summary

With doublequote on, the CSV decode kernel now collapses escaped quote pairs ("" into ") in place, in the device copy of the input that the reader owns. This removes:
- the per-row is_quoted arrays;
- the per-string-column thrust::count / cudf::strings::replace / copy_if_else step;
- the host worker pool and forked streams that ran that step.
Fixed-width output buffers are no longer zero-filled. String pair arrays and null masks still are.

Motivation
String-heavy files built every string column up to three times just to handle escaped quotes. Every fixed-width column also paid for a memset that the decode kernel then overwrote.

Design

Matching rule: pairs are matched left to right without overlap, which is exactly the old replace("\"\"", "\"") result ("""" becomes ""). Only fields whose enclosing quotes were stripped are touched, and fields without a pair are not written.
In-place safety: each row is decoded by one thread, and every access for row i stays inside [row_offsets[i], row_offsets[i+1]), so rows never share bytes. Within a row, a field is rewritten only after its end has been found, and the write position never passes the read position. This invariant is documented at the kernel; later changes that load bytes outside a row must not use those bytes.
Ownership: the reader always decodes its own copy of the input, so the caller's buffers are never modified. Decoding is the last use of that copy before the string columns are built.
Fixed-width output: a value is written exactly where the field is valid and its validity bit is set. Other slots are under a null bit, which the cudf column contract allows to be uninitialized.
String pairs stay zeroed: fields missing from the end of a short row rely on the zeroed pair reading as null. Handling missing fields in or after the field loop spilled registers (all-integer read 8.6 → 47 ms), and writing null pairs up front was no faster than the memset on string-heavy reads.

Testing

New gtests: EscapedQuotePairs, EscapedQuotePairsOptions (custom quote char, whitespace around quotes, CRLF pair before the closing quote, doublequote off, quoting NONE), EscapedQuotePairsManyRows, EscapedQuotePairsByteRanges (concatenated byte ranges equal the full read), EscapedQuotePairsDeviceSourceUnchanged (guards future zero-copy paths), ShortRowsMissingFieldsAreNull (1–1000 rows, INT32/STRING/FLOAT64/BOOL8, NA fields, use_cols, whitespace delimiters).
All CSV and stream tests pass at each commit. compute-sanitizer memcheck with exact allocations is clean; initcheck is clean for host and device buffer sources (reports with file sources are an artifact of the kvikio file read filling device memory, not out-of-bounds reads).
Bit-exact oracles match the parent, and 9,000 fuzzed reads (host and device buffer sources, byte ranges, doublequote on and off, quoting NONE, comments) show 0 differences.

Performance (RTX PRO 6000, 256 MB, GPU time; official best-of-2 numbers for all sources are in the benchmark comment below)

Device buffers: QUOTE_25 −50%, QUOTE_100 −79%, TAXI −19%, LOGS −44%; all-string 64 columns (no inference): 49.4 → 10.5 ms.
Host buffers: QUOTE_25 −23%, QUOTE_100 −51%.

Notes for reviewers

decode_row_column_data now takes a mutable device_span<char>. A future zero-copy path must unescape into a scratch buffer, never the caller's memory.
Pre-existing limitation, unchanged here: a byte range that starts after offset 0 treats a quoted newline as a row end (worth a follow-up issue).
The planned "write missing trailing fields explicitly" piece was dropped: it is only needed if string pairs are left uninitialized, and both in-kernel variants were slower or no faster.

This is PR 2 of a stack of draft PRs extracted from an automated (CORAL) optimization run of the CSV reader, reviewed by a separate principal-engineer review pass; it builds on #16.

🤖 Generated with Claude Code

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>

karthikeyann · 2026-09-24T07:42:55Z

Note

Re-measured cleanly (2026-09-25); these numbers replace the earlier ones in this comment. The earlier numbers were taken while other validation jobs (tests, sanitizers, fuzzers) ran on the same host, against a baseline measured hours earlier, so the pinned-host, pageable-host and file results in particular could be skewed.

Methodology (clean re-measurement). All ten builds of the stack (base and PRs 1–9) were measured in one session on the same GPU (host GPU 1), in two interleaved rounds (base, PR 1, …, PR 9, then again), so drift affects every build equally. Each benchmark run started only when no other process was using the GPU and no build/test/sanitizer/fuzzer was running on the host in our environment, and would have been re-run if anything appeared during it; all 160 runs were clean on the first attempt. Each number below is the best of the two rounds per side.

CSV reader benchmarks: PR 2/9 (fused unescape, no fixed-width zero-init)

This PR 979961268c vs base coral/csv-base 63712844b0 and vs parent PR 66b7b12fda. 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	1.53× _{28.9 GiB/s}	1.25× _{14.0 GiB/s}	1.04× _{5.0 GiB/s}	1.15× _{7.3 GiB/s}
realistic LOGS 256MB	1.77× _{48.0 GiB/s}	1.39× _{23.7 GiB/s}	1.14× _{8.1 GiB/s}	1.20× _{11.9 GiB/s}
realistic ANALYTICS 256MB	1.17× _{29.5 GiB/s}	1.09× _{13.8 GiB/s}	1.01× _{4.9 GiB/s}	1.03× _{6.8 GiB/s}
inference MIXED 8 cols	1.06× _{22.4 GiB/s}	1.05× _{11.5 GiB/s}	1.00× _{4.1 GiB/s}	0.99× _{5.7 GiB/s}
inference MIXED 64 cols	1.40× _{9.2 GiB/s}	1.30× _{6.4 GiB/s}	1.18× _{3.1 GiB/s}	1.21× _{3.8 GiB/s}
no inference MIXED 8 cols	1.02× _{29.8 GiB/s}	1.02× _{13.1 GiB/s}	0.96× _{4.3 GiB/s}	1.01× _{6.5 GiB/s}
25% quoted 256MB	2.33× _{29.8 GiB/s}	1.75× _{17.1 GiB/s}	1.39× _{6.9 GiB/s}	1.46× _{9.4 GiB/s}
100% quoted 256MB	4.87× _{23.7 GiB/s}	3.38× _{18.3 GiB/s}	2.19× _{9.1 GiB/s}	2.50× _{11.7 GiB/s}
Weighted geomean	1.62×	1.39×	1.18×	1.24×

Speedup vs parent PR (this PR's own contribution)

Config	Device buffer	Pinned host	Pageable host	File (kvikio)
realistic TAXI 256MB	1.25× _{28.9 GiB/s}	1.13× _{14.0 GiB/s}	1.04× _{5.0 GiB/s}	1.16× _{7.3 GiB/s}
realistic LOGS 256MB	1.80× _{48.0 GiB/s}	1.40× _{23.7 GiB/s}	1.12× _{8.1 GiB/s}	1.28× _{11.9 GiB/s}
realistic ANALYTICS 256MB	1.02× _{29.5 GiB/s}	1.01× _{13.8 GiB/s}	1.02× _{4.9 GiB/s}	1.06× _{6.8 GiB/s}
inference MIXED 8 cols	1.03× _{22.4 GiB/s}	1.02× _{11.5 GiB/s}	0.95× _{4.1 GiB/s}	0.99× _{5.7 GiB/s}
inference MIXED 64 cols	1.25× _{9.2 GiB/s}	1.23× _{6.4 GiB/s}	1.09× _{3.1 GiB/s}	1.08× _{3.8 GiB/s}
no inference MIXED 8 cols	1.03× _{29.8 GiB/s}	1.02× _{13.1 GiB/s}	0.95× _{4.3 GiB/s}	1.03× _{6.5 GiB/s}
25% quoted 256MB	2.21× _{29.8 GiB/s}	1.76× _{17.1 GiB/s}	1.34× _{6.9 GiB/s}	1.43× _{9.4 GiB/s}
100% quoted 256MB	5.03× _{23.7 GiB/s}	3.40× _{18.3 GiB/s}	2.13× _{9.1 GiB/s}	2.52× _{11.7 GiB/s}
Weighted geomean	1.51×	1.34×	1.15×	1.23×

_{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	Parent	This PR	vs base	vs parent
csv_read_column_selection	column_selection=ALL, num_chunks=1, row_selection=ALL	2.63 GiB/s	2.73 GiB/s	2.97 GiB/s	1.13×	1.09×
csv_read_column_selection	column_selection=ALTERNATE, num_chunks=1, row_selection=ALL	1.45 GiB/s	1.45 GiB/s	1.56 GiB/s	1.07×	1.07×
csv_read_column_selection	column_selection=FIRST_HALF, num_chunks=1, row_selection=ALL	1.49 GiB/s	1.49 GiB/s	1.62 GiB/s	1.09×	1.09×
csv_read_column_selection	column_selection=SECOND_HALF, num_chunks=1, row_selection=ALL	1.49 GiB/s	1.48 GiB/s	1.69 GiB/s	1.13×	1.14×
csv_read_data_type	data_type=DECIMAL, io=DEVICE_BUFFER	15.33 GiB/s	20.28 GiB/s	20.19 GiB/s	1.32×	1.00×
csv_read_data_type	data_type=DURATION, io=DEVICE_BUFFER	12.78 GiB/s	16.06 GiB/s	16.09 GiB/s	1.26×	1.00×
csv_read_data_type	data_type=FLOAT, io=DEVICE_BUFFER	18.66 GiB/s	19.91 GiB/s	20.29 GiB/s	1.09×	1.02×
csv_read_data_type	data_type=INTEGRAL, io=DEVICE_BUFFER	17.43 GiB/s	20.11 GiB/s	20.06 GiB/s	1.15×	1.00×
csv_read_data_type	data_type=STRING, io=DEVICE_BUFFER	4.68 GiB/s	5.15 GiB/s	23.84 GiB/s	5.09×	4.63×
csv_read_data_type	data_type=TIMESTAMP, io=DEVICE_BUFFER	12.37 GiB/s	15.48 GiB/s	15.52 GiB/s	1.25×	1.00×
csv_read_io	io=FILEPATH	4.55 GiB/s	4.61 GiB/s	4.97 GiB/s	1.09×	1.08×
csv_read_io	io=HOST_BUFFER	3.38 GiB/s	3.47 GiB/s	3.69 GiB/s	1.09×	1.06×
csv_read_row_selection	column_selection=ALL, num_chunks=1, row_selection=BYTE_RANGE	2.56 GiB/s	2.60 GiB/s	2.88 GiB/s	1.13×	1.11×
csv_read_row_selection	column_selection=ALL, num_chunks=1, row_selection=NROWS	2.54 GiB/s	2.59 GiB/s	2.88 GiB/s	1.13×	1.11×
csv_read_row_selection	column_selection=ALL, num_chunks=1, row_selection=SKIPFOOTER	2.67 GiB/s	2.68 GiB/s	3.08 GiB/s	1.15×	1.15×
csv_read_row_selection	column_selection=ALL, num_chunks=8, row_selection=BYTE_RANGE	1.93 GiB/s	1.91 GiB/s	2.62 GiB/s	1.36×	1.37×
csv_read_row_selection	column_selection=ALL, num_chunks=8, row_selection=NROWS	0.71 GiB/s	0.67 GiB/s	0.80 GiB/s	1.13×	1.20×
csv_read_row_selection	column_selection=ALL, num_chunks=8, row_selection=SKIPFOOTER	0.45 GiB/s	0.43 GiB/s	0.49 GiB/s	1.09×	1.15×

`CSV_READER_QUOTING_IO_NVBENCH` (12 configs)

Benchmark	Config	Base	Parent	This PR	vs base	vs parent
csv_read_quoting_io	data_size_mb=256, io=DEVICE_BUFFER, quote_pct=QUOTE_0_PCT	17.11 GiB/s	20.32 GiB/s	20.50 GiB/s	1.20×	1.01×
csv_read_quoting_io	data_size_mb=256, io=DEVICE_BUFFER, quote_pct=QUOTE_100_PCT	4.87 GiB/s	4.71 GiB/s	23.70 GiB/s	4.87×	5.03×
csv_read_quoting_io	data_size_mb=256, io=DEVICE_BUFFER, quote_pct=QUOTE_25_PCT	12.75 GiB/s	13.46 GiB/s	29.76 GiB/s	2.33×	2.21×
csv_read_quoting_io	data_size_mb=256, io=FILEPATH, quote_pct=QUOTE_0_PCT	5.25 GiB/s	5.53 GiB/s	5.59 GiB/s	1.07×	1.01×
csv_read_quoting_io	data_size_mb=256, io=FILEPATH, quote_pct=QUOTE_100_PCT	4.69 GiB/s	4.64 GiB/s	11.71 GiB/s	2.50×	2.52×
csv_read_quoting_io	data_size_mb=256, io=FILEPATH, quote_pct=QUOTE_25_PCT	6.47 GiB/s	6.58 GiB/s	9.41 GiB/s	1.46×	1.43×
csv_read_quoting_io	data_size_mb=256, io=HOST_BUFFER, quote_pct=QUOTE_0_PCT	3.80 GiB/s	4.09 GiB/s	3.92 GiB/s	1.03×	0.96×
csv_read_quoting_io	data_size_mb=256, io=HOST_BUFFER, quote_pct=QUOTE_100_PCT	4.18 GiB/s	4.29 GiB/s	9.14 GiB/s	2.19×	2.13×
csv_read_quoting_io	data_size_mb=256, io=HOST_BUFFER, quote_pct=QUOTE_25_PCT	4.99 GiB/s	5.19 GiB/s	6.94 GiB/s	1.39×	1.34×
csv_read_quoting_io	data_size_mb=256, io=PINNED_BUFFER, quote_pct=QUOTE_0_PCT	9.57 GiB/s	10.62 GiB/s	10.80 GiB/s	1.13×	1.02×
csv_read_quoting_io	data_size_mb=256, io=PINNED_BUFFER, quote_pct=QUOTE_100_PCT	5.40 GiB/s	5.38 GiB/s	18.27 GiB/s	3.38×	3.40×
csv_read_quoting_io	data_size_mb=256, io=PINNED_BUFFER, quote_pct=QUOTE_25_PCT	9.77 GiB/s	9.75 GiB/s	17.14 GiB/s	1.75×	1.76×

`CSV_READER_QUOTING_NVBENCH` (6 configs)

Benchmark	Config	Base	Parent	This PR	vs base	vs parent
csv_read_quoting	data_size_mb=256, quote_pct=QUOTE_0_PCT	17.51 GiB/s	20.70 GiB/s	21.34 GiB/s	1.22×	1.03×
csv_read_quoting	data_size_mb=256, quote_pct=QUOTE_100_PCT	4.80 GiB/s	4.70 GiB/s	24.24 GiB/s	5.05×	5.16×
csv_read_quoting	data_size_mb=256, quote_pct=QUOTE_25_PCT	13.88 GiB/s	13.32 GiB/s	30.30 GiB/s	2.18×	2.27×
csv_read_quoting	data_size_mb=64, quote_pct=QUOTE_0_PCT	14.10 GiB/s	15.86 GiB/s	16.47 GiB/s	1.17×	1.04×
csv_read_quoting	data_size_mb=64, quote_pct=QUOTE_100_PCT	1.37 GiB/s	1.40 GiB/s	8.18 GiB/s	5.99×	5.83×
csv_read_quoting	data_size_mb=64, quote_pct=QUOTE_25_PCT	4.57 GiB/s	4.39 GiB/s	17.74 GiB/s	3.88×	4.04×

`CSV_READER_REALISTIC_IO_NVBENCH` (12 configs)

Benchmark	Config	Base	Parent	This PR	vs base	vs parent
csv_read_realistic_io	data_s 8000 ize_mb=256, io=DEVICE_BUFFER, profile=ANALYTICS	25.28 GiB/s	28.98 GiB/s	29.51 GiB/s	1.17×	1.02×
csv_read_realistic_io	data_size_mb=256, io=DEVICE_BUFFER, profile=LOGS	27.19 GiB/s	26.64 GiB/s	48.03 GiB/s	1.77×	1.80×
csv_read_realistic_io	data_size_mb=256, io=DEVICE_BUFFER, profile=TAXI	18.91 GiB/s	23.16 GiB/s	28.91 GiB/s	1.53×	1.25×
csv_read_realistic_io	data_size_mb=256, io=FILEPATH, profile=ANALYTICS	6.63 GiB/s	6.43 GiB/s	6.84 GiB/s	1.03×	1.06×
csv_read_realistic_io	data_size_mb=256, io=FILEPATH, profile=LOGS	9.85 GiB/s	9.27 GiB/s	11.85 GiB/s	1.20×	1.28×
csv_read_realistic_io	data_size_mb=256, io=FILEPATH, profile=TAXI	6.37 GiB/s	6.31 GiB/s	7.31 GiB/s	1.15×	1.16×
csv_read_realistic_io	data_size_mb=256, io=HOST_BUFFER, profile=ANALYTICS	4.83 GiB/s	4.76 GiB/s	4.85 GiB/s	1.01×	1.02×
csv_read_realistic_io	data_size_mb=256, io=HOST_BUFFER, profile=LOGS	7.16 GiB/s	7.29 GiB/s	8.15 GiB/s	1.14×	1.12×
csv_read_realistic_io	data_size_mb=256, io=HOST_BUFFER, profile=TAXI	4.77 GiB/s	4.76 GiB/s	4.97 GiB/s	1.04×	1.04×
csv_read_realistic_io	data_size_mb=256, io=PINNED_BUFFER, profile=ANALYTICS	12.58 GiB/s	13.61 GiB/s	13.76 GiB/s	1.09×	1.01×
csv_read_realistic_io	data_size_mb=256, io=PINNED_BUFFER, profile=LOGS	17.05 GiB/s	16.94 GiB/s	23.66 GiB/s	1.39×	1.40×
csv_read_realistic_io	data_size_mb=256, io=PINNED_BUFFER, profile=TAXI	11.20 GiB/s	12.45 GiB/s	14.04 GiB/s	1.25×	1.13×

`CSV_READER_REALISTIC_NVBENCH` (9 configs)

Benchmark	Config	Base	Parent	This PR	vs base	vs parent
csv_read_realistic	data_size_mb=1024, profile=ANALYTICS	25.11 GiB/s	28.95 GiB/s	29.72 GiB/s	1.18×	1.03×
csv_read_realistic	data_size_mb=1024, profile=LOGS	31.88 GiB/s	31.80 GiB/s	50.68 GiB/s	1.59×	1.59×
csv_read_realistic	data_size_mb=1024, profile=TAXI	23.43 GiB/s	26.89 GiB/s	30.81 GiB/s	1.31×	1.15×
csv_read_realistic	data_size_mb=256, profile=ANALYTICS	25.24 GiB/s	28.87 GiB/s	29.39 GiB/s	1.16×	1.02×
csv_read_realistic	data_size_mb=256, profile=LOGS	27.37 GiB/s	26.61 GiB/s	47.45 GiB/s	1.73×	1.78×
csv_read_realistic	data_size_mb=256, profile=TAXI	19.01 GiB/s	23.02 GiB/s	29.28 GiB/s	1.54×	1.27×
csv_read_realistic	data_size_mb=512, profile=ANALYTICS	25.64 GiB/s	29.51 GiB/s	30.18 GiB/s	1.18×	1.02×
csv_read_realistic	data_size_mb=512, profile=LOGS	31.05 GiB/s	30.85 GiB/s	51.04 GiB/s	1.64×	1.65×
csv_read_realistic	data_size_mb=512, profile=TAXI	22.88 GiB/s	26.29 GiB/s	30.63 GiB/s	1.34×	1.17×

`CSV_READER_SCALE_NVBENCH` (5 configs)

Benchmark	Config	Base	Parent	This PR	vs base	vs parent
csv_read_scale	data_size_mb=1024	3.54 GiB/s	3.50 GiB/s	3.59 GiB/s	1.01×	1.03×
csv_read_scale	data_size_mb=2048	3.52 GiB/s	3.45 GiB/s	3.74 GiB/s	1.06×	1.08×
csv_read_scale	data_size_mb=256	3.59 GiB/s	3.56 GiB/s	3.60 GiB/s	1.00×	1.01×
csv_read_scale	data_size_mb=4096	3.38 GiB/s	3.39 GiB/s	3.64 GiB/s	1.08×	1.07×
csv_read_scale	data_size_mb=512	3.60 GiB/s	3.56 GiB/s	3.62 GiB/s	1.01×	1.02×

`CSV_READER_TYPE_INFERENCE_IO_NVBENCH` (64 configs)

Benchmark	Config	Base	Parent	This PR	vs base	vs parent
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_FLOAT	10.11 GiB/s	12.17 GiB/s	12.27 GiB/s	1.21×	1.01×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	9.02 GiB/s	11.01 GiB/s	11.08 GiB/s	1.23×	1.01×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_STRING	4.40 GiB/s	4.24 GiB/s	14.81 GiB/s	3.36×	3.50×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=MIXED	6.59 GiB/s	7.37 GiB/s	9.23 GiB/s	1.40×	1.25×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_FLOAT	12.31 GiB/s	14.10 GiB/s	14.07 GiB/s	1.14×	1.00×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	21.14 GiB/s	21.68 GiB/s	22.32 GiB/s	1.06×	1.03×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_STRING	11.91 GiB/s	11.95 GiB/s	18.29 GiB/s	1.54×	1.53×
csv_read_with_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=MIXED	21.15 GiB/s	21.65 GiB/s	22.38 GiB/s	1.06×	1.03×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_FLOAT	4.03 GiB/s	4.51 GiB/s	4.50 GiB/s	1.11×	1.00×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_INTEGRAL	4.12 GiB/s	4.29 GiB/s	4.31 GiB/s	1.05×	1.00×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_STRING	3.94 GiB/s	3.91 GiB/s	9.07 GiB/s	2.30×	2.32×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=MIXED	3.16 GiB/s	3.54 GiB/s	3.82 GiB/s	1.21×	1.08×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_FLOAT	4.37 GiB/s	4.81 GiB/s	4.55 GiB/s	1.04×	0.94×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_INTEGRAL	5.98 GiB/s	5.60 GiB/s	5.88 GiB/s	0.98×	1.05×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_STRING	7.49 GiB/s	7.79 GiB/s	9.99 GiB/s	1.33×	1.28×
csv_read_with_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=MIXED	5.79 GiB/s	5.82 GiB/s	5.73 GiB/s	0.99×	0.99×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_FLOAT	3.07 GiB/s	3.28 GiB/s	3.40 GiB/s	1.11×	1.04×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	3.21 GiB/s	3.25 GiB/s	3.30 GiB/s	1.03×	1.02×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_STRING	3.50 GiB/s	3.64 GiB/s	7.32 GiB/s	2.10×	2.01×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=MIXED	2.60 GiB/s	2.81 GiB/s	3.07 GiB/s	1.18×	1.09×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_FLOAT	3.30 GiB/s	3.47 GiB/s	3.48 GiB/s	1.05×	1.00×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	4.04 GiB/s	4.02 GiB/s	4.08 GiB/s	1.01×	1.01×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_STRING	6.36 GiB/s	6.46 GiB/s	8.29 GiB/s	1.30×	1.28×
csv_read_with_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=MIXED	4.06 GiB/s	4.28 GiB/s	4.07 GiB/s	1.00×	0.95×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_FLOAT	6.80 GiB/s	7.67 GiB/s	7.72 GiB/s	1.13×	1.01×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	6.43 GiB/s	7.33 GiB/s	7.39 GiB/s	1.15×	1.01×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_STRING	4.36 GiB/s	4.44 GiB/s	12.33 GiB/s	2.83×	2.78×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=MIXED	4.90 GiB/s	5.18 GiB/s	6.36 GiB/s	1.30×	1.23×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_FLOAT	7.73 GiB/s	8.59 GiB/s	8.62 GiB/s	1.11×	1.00×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	10.97 GiB/s	11.18 GiB/s	11.32 GiB/s	1.03×	1.01×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_STRING	10.30 GiB/s	10.48 GiB/s	14.49 GiB/s	1.41×	1.38×
csv_read_with_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=MIXED	10.97 GiB/s	11.27 GiB/s	11.52 GiB/s	1.05×	1.02×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_FLOAT	18.66 GiB/s	20.01 GiB/s	20.45 GiB/s	1.10×	1.02×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	17.43 GiB/s	20.85 GiB/s	21.36 GiB/s	1.23×	1.02×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=ALL_STRING	4.91 GiB/s	4.87 GiB/s	23.83 GiB/s	4.85×	4.89×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=64, type_profile=MIXED	14.50 GiB/s	16.86 GiB/s	20.32 GiB/s	1.40×	1.21×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_FLOAT	21.45 GiB/s	21.30 GiB/s	21.69 GiB/s	1.01×	1.02×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	28.99 GiB/s	28.95 GiB/s	29.92 GiB/s	1.03×	1.03×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=ALL_STRING	20.79 GiB/s	21.07 GiB/s	51.44 GiB/s	2.47×	2.44×
csv_read_without_inference_io	data_size_mb=256, io=DEVICE_BUFFER, num_cols=8, type_profile=MIXED	29.16 GiB/s	29.03 GiB/s	29.76 GiB/s	1.02×	1.03×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_FLOAT	5.12 GiB/s	5.29 GiB/s	5.32 GiB/s	1.04×	1.01×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_INTEGRAL	5.02 GiB/s	5.51 GiB/s	5.64 GiB/s	1.12×	1.02×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=ALL_STRING	4.78 GiB/s	4.70 GiB/s	11.81 GiB/s	2.47×	2.51×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=64, type_profile=MIXED	4.60 GiB/s	5.03 GiB/s	5.14 GiB/s	1.12×	1.02×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_FLOAT	5.29 GiB/s	5.25 GiB/s	5.37 GiB/s	1.02×	1.02×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_INTEGRAL	6.01 GiB/s	6.38 GiB/s	6.41 GiB/s	1.07×	1.01×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=ALL_STRING	11.19 GiB/s	11.25 GiB/s	15.72 GiB/s	1.41×	1.40×
csv_read_without_inference_io	data_size_mb=256, io=FILEPATH, num_cols=8, type_profile=MIXED	6.39 GiB/s	6.29 GiB/s	6.47 GiB/s	1.01×	1.03×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_FLOAT	3.75 GiB/s	3.84 GiB/s	3.67 GiB/s	0.98×	0.96×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	3.79 GiB/s	4.10 GiB/s	3.93 GiB/s	1.04×	0.96×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=ALL_STRING	4.25 GiB/s	4.27 GiB/s	9.16 GiB/s	2.15×	2.15×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=64, type_profile=MIXED	3.33 GiB/s	3.59 GiB/s	3.80 GiB/s	1.14×	1.06×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_FLOAT	3.90 GiB/s	3.88 GiB/s	3.74 GiB/s	0.96×	0.96×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	4.30 GiB/s	4.54 GiB/s	4.57 GiB/s	1.06×	1.01×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=ALL_STRING	8.50 GiB/s	8.64 GiB/s	10.84 GiB/s	1.28×	1.26×
csv_read_without_inference_io	data_size_mb=256, io=HOST_BUFFER, num_cols=8, type_profile=MIXED	4.49 GiB/s	4.53 GiB/s	4.32 GiB/s	0.96×	0.95×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_FLOAT	9.60 GiB/s	10.15 GiB/s	10.38 GiB/s	1.08×	1.02×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_INTEGRAL	9.67 GiB/s	10.60 GiB/s	11.00 GiB/s	1.14×	1.04×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=ALL_STRING	5.47 GiB/s	5.29 GiB/s	18.26 GiB/s	3.34×	3.45×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=64, type_profile=MIXED	7.96 GiB/s	8.77 GiB/s	10.19 GiB/s	1.28×	1.16×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_FLOAT	10.51 GiB/s	10.56 GiB/s	10.69 GiB/s	1.02×	1.01×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_INTEGRAL	12.94 GiB/s	12.93 GiB/s	13.10 GiB/s	1.01×	1.01×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=ALL_STRING	17.06 GiB/s	17.08 GiB/s	30.28 GiB/s	1.78×	1.77×
csv_read_without_inference_io	data_size_mb=256, io=PINNED_BUFFER, num_cols=8, type_profile=MIXED	12.86 GiB/s	12.84 GiB/s	13.09 GiB/s	1.02×	1.02×

`CSV_READER_TYPE_INFERENCE_NVBENCH` (32 configs)

1.03×

Benchmark	Config	Base	Parent	This PR	vs base	vs parent
csv_read_with_inference	data_size_mb=256, num_cols=64, type_profile=ALL_FLOAT	10.18 GiB/s	12.16 GiB/s	12.34 GiB/s	1.21×	1.01×
csv_read_with_inference	data_size_mb=256, num_cols=64, type_profile=ALL_INTEGRAL	9.10 GiB/s	11.02 GiB/s	11.16 GiB/s	1.23×	1.01×
csv_read_with_inference	data_size_mb=256, num_cols=64, type_profile=ALL_STRING	4.40 GiB/s	4.30 GiB/s	15.47 GiB/s	3.51×	3.60×
csv_read_with_inference	data_size_mb=256, num_cols=64, type_profile=MIXED	6.80 GiB/s	7.53 GiB/s	9.28 GiB/s	1.36×	1.23×
csv_read_with_inference	data_size_mb=256, num_cols=8, type_profile=ALL_FLOAT	12.32 GiB/s	14.56 GiB/s	14.81 GiB/s	1.20×	1.02×
csv_read_with_inference	data_size_mb=256, num_cols=8, type_profile=ALL_INTEGRAL	21.29 GiB/s	21.76 GiB/s	22.40 GiB/s	1.05×	1.03×
csv_read_with_inference	data_size_mb=256, num_cols=8, type_profile=ALL_STRING	12.00 GiB/s	12.55 GiB/s	19.14 GiB/s	1.60×	1.53×
csv_read_with_inference	data_size_mb=256, num_cols=8, type_profile=MIXED	21.64 GiB/s	21.66 GiB/s	22.46 GiB/s	1.04×	1.04×
csv_read_with_inference	data_size_mb=64, num_cols=64, type_profile=ALL_FLOAT	10.15 GiB/s	9.31 GiB/s	9.70 GiB/s	0.96×	1.04×
csv_read_with_inference	data_size_mb=64, num_cols=64, type_profile=ALL_INTEGRAL	8.33 GiB/s	9.74 GiB/s	9.78 GiB/s	1.17×	1.00×
csv_read_with_inference	data_size_mb=64, num_cols=64, type_profile=ALL_STRING	1.42 GiB/s	1.41 GiB/s	6.82 GiB/s	4.79×	4.83×
csv_read_with_inference	data_size_mb=64, num_cols=64, type_profile=MIXED	4.05 GiB/s	3.87 GiB/s	6.91 GiB/s	1.70×	1.78×
csv_read_with_inference	data_size_mb=64, num_cols=8, type_profile=ALL_FLOAT	11.10 GiB/s	12.86 GiB/s	13.03 GiB/s	1.17×	1.01×
csv_read_with_inference	data_size_mb=64, num_cols=8, type_profile=ALL_INTEGRAL	18.41 GiB/s	18.92 GiB/s	19.42 GiB/s	1.05×
csv_read_with_inference	data_size_mb=64, num_cols=8, type_profile=ALL_STRING	5.52 GiB/s	5.53 GiB/s	15.94 GiB/s	2.89×	2.88×
csv_read_with_inference	data_size_mb=64, num_cols=8, type_profile=MIXED	19.86 GiB/s	18.90 GiB/s	19.82 GiB/s	1.00×	1.05×
csv_read_without_inference	data_size_mb=256, num_cols=64, type_profile=ALL_FLOAT	18.10 GiB/s	19.91 GiB/s	20.31 GiB/s	1.12×	1.02×
csv_read_without_inference	data_size_mb=256, num_cols=64, type_profile=ALL_INTEGRAL	17.51 GiB/s	20.88 GiB/s	21.24 GiB/s	1.21×	1.02×
csv_read_without_inference	data_size_mb=256, num_cols=64, type_profile=ALL_STRING	4.70 GiB/s	4.78 GiB/s	24.27 GiB/s	5.16×	5.07×
csv_read_without_inference	data_size_mb=256, num_cols=64, type_profile=MIXED	14.57 GiB/s	16.99 GiB/s	20.21 GiB/s	1.39×	1.19×
csv_read_without_inference	data_size_mb=256, num_cols=8, type_profile=ALL_FLOAT	21.55 GiB/s	21.42 GiB/s	21.71 GiB/s	1.01×	1.01×
csv_read_without_inference	data_size_mb=256, num_cols=8, type_profile=ALL_INTEGRAL	29.30 GiB/s	29.10 GiB/s	29.94 GiB/s	1.02×	1.03×
csv_read_without_inference	data_size_mb=256, num_cols=8, type_profile=ALL_STRING	21.92 GiB/s	22.65 GiB/s	52.64 GiB/s	2.40×	2.32×
csv_read_without_inference	data_size_mb=256, num_cols=8, type_profile=MIXED	29.24 GiB/s	28.97 GiB/s	30.09 GiB/s	1.03×	1.04×
csv_read_without_inference	data_size_mb=64, num_cols=64, type_profile=ALL_FLOAT	15.91 GiB/s	13.64 GiB/s	14.29 GiB/s	0.90×	1.05×
csv_read_without_inference	data_size_mb=64, num_cols=64, type_profile=ALL_INTEGRAL	14.41 GiB/s	16.04 GiB/s	16.51 GiB/s	1.15×	1.03×
csv_read_without_inference	data_size_mb=64, num_cols=64, type_profile=ALL_STRING	1.44 GiB/s	1.45 GiB/s	8.18 GiB/s	5.69×	5.64×
csv_read_without_inference	data_size_mb=64, num_cols=64, type_profile=MIXED	9.73 GiB/s	9.78 GiB/s	15.43 GiB/s	1.59×	1.58×
csv_read_without_inference	data_size_mb=64, num_cols=8, type_profile=ALL_FLOAT	18.58 GiB/s	19.11 GiB/s	18.53 GiB/s	1.00×	0.97×
csv_read_without_inference	data_size_mb=64, num_cols=8, type_profile=ALL_INTEGRAL	25.71 GiB/s	25.79 GiB/s	26.08 GiB/s	1.01×	1.01×
csv_read_without_inference	data_size_mb=64, num_cols=8, type_profile=ALL_STRING	7.32 GiB/s	8.40 GiB/s	33.53 GiB/s	4.58×	3.99×
csv_read_without_inference	data_size_mb=64, num_cols=8, type_profile=MIXED	25.23 GiB/s	24.95 GiB/s	25.96 GiB/s	1.03×	1.04×

karthikeyann and others added 2 commits September 24, 2026 06:46

github-actions Bot assigned karthikeyann Sep 24, 2026

github-actions Bot added the libcudf label Sep 24, 2026

This was referenced Sep 24, 2026

[CSV 3/9] Single-pass row gathering for whole-file CSV reads, with bounded host reads #18

Draft

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

Draft

Navigation Menu

Repository navigation

[CSV 2/9] Unescape CSV doubled quotes in the decode kernel; skip zero-initializing fixed-width output - #17

[CSV 2/9] Unescape CSV doubled quotes in the decode kernel; skip zero-initializing fixed-width output#17
karthikeyann wants to merge 2 commits into
coral/csv-pr01-correctnessfrom
coral/csv-pr02-fused-unescape

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!

Uh oh!

CSV reader benchmarks: PR 2/9 (fused unescape, no fixed-width zero-init)

Speedup vs base (cumulative), by input source

Speedup vs parent PR (this PR's own contribution)

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)