Deterministically compute parquet page fragments that first see dictionary keys - #22323
Deterministically compute parquet page fragments that first see dictionary keys#22323mhaseeb123 wants to merge 35 commits into
Conversation
Repurpose the unused mapped_type slot in populate_chunk_hash_maps_kernel to carry the fragment index of the block that first inserts each value. Rewrite collect_map_entries_kernel to bucket dict_ids by that fragment index so pages that only reference earlier fragments' values see small max dict_index. No file-size delta in isolation; prerequisite for the per-page bit-width change landing next. Made-with: Cursor
…#13995) Each data page now RLE-encodes dictionary indices using ceil(log2(page_max_dict_index + 1)) bits instead of the chunk-wide maximum. Combined with the first-appearance dict_id ordering from the prior commit, this closes the ~30% file-size gap vs Spark on moderate-cardinality INT64 and STRING workloads. Page size estimation continues to use the chunk-wide bits as a conservative upper bound; the dictionary page itself is unaffected. Made-with: Cursor
| // Initialize all fragment offsets to 0 | ||
| for (auto idx = t; idx < num_frags; idx += block_size) { | ||
| fragment_offsets[idx] = 0; | ||
| } | ||
| __syncthreads(); | ||
|
|
||
| // Iterate over slots and count the number of dict values first seen page fragment | ||
| { | ||
| for (auto slot_idx = t; slot_idx < chunk.dict_map_size; slot_idx += block_size) { | ||
| auto const* slot = map_storage.data() + chunk.dict_map_offset + slot_idx; | ||
| if (slot->first != KEY_SENTINEL) { | ||
| auto const frag_loc = static_cast<size_type>(slot->second) - frag_start; | ||
| cudf_assert(frag_loc >= 0 && frag_loc < num_frags && | ||
| "fragment index in the slot is out of range of the chunk"); | ||
| atomicAdd(&fragment_offsets[frag_loc], 1); | ||
| } | ||
| } | ||
| __syncthreads(); | ||
| } |
There was a problem hiding this comment.
@pmattione-nvidia Part 1, phase 1 with this PR becomes this
There was a problem hiding this comment.
Is the time spent here on reading slot->second or on the add? If it's on reading slot->second you could keep counters while filling the map in the first place.
…22279) Contributes to #13995 This PR enables the Parquet writer to assign lower dictionary indices to elements appearing earlier in each column chunk thereby writing variable (reduced) number of dictionary key bits for earlier pages. New algorithm: 1. `map_insert_fn` now inserts this pair in the static map: `{row_idx, frag_idx}` instead of `{row_idx, row_idx}`. `frag_idx` is just `blockIdx.x` and indicates which page fragment actually inserted this entry in the static map (CAS race dependent but earlier thread blocks generally win this race) 2. `map_insert_fn` also writes the `PageFragment::num_dict_vals` field for each page fragment to keep a record of number of unique values inserted by it. 3. Then in `collect_map_entries_kernel` (where we assign keys to dict values), we optionally assign spatially-local keys to dictionary values inserted by the same page fragment. See algorithm details [here](#22279 (comment)) 4. Finally, we launch `compute_page_dict_bits_kernel` to compute the max bit width required to encode dict keys for each parquet page. See algorithm [here](#22279 (comment)) See follow up PR #22323 that deterministically writes the index of the first page fragment that sees each dictionary key Authors: - Muhammad Haseeb (https://github.com/mhaseeb123) Approvers: - Vukasin Milovanovic (https://github.com/vuule) - Paul Mattione (https://github.com/pmattione-nvidia) - Nghia Truong (https://github.com/ttnghia) URL: #22279
Performance Impact
SummaryParquet writer-speed regression reported across 58 configurations of the benchmark suite, typically +5–12% in Raw numbers (sorted by % diff)Click to expand
|
|
Make sure to modify this to reduce the # of atomic operations, per 2nd paragraph of #2) at #13995 (comment) |
Good idea. Let me first verify if the slowdown is in the insert or collect kernel. If it's in collect from the new loops, then this PR may be doomed anyway |
|
10% slowdown is not a big deal if the disk space savings are large. Do you have an idea what those are? Are these tests large enough to trigger a large disk savings? The perf numbers will probably be bad if you don't implement the block reordering as it will be too many atomic cas loops. |
I can certainly compare but I am not sure how much further disk saving by simply making this deterministic would we get as compared to
I analyzed both the populate and collect kernels using ncu and nsys and the main regression (roughly 20-30% -> e2e 5-10% in benchmarks at least) is inside the |
Performance analysis (part 2)Experimental SetupWrite 50 columns of mixed types — cycling through INT32, INT64, STRING, LIST, and STRUCT<INT32, STRING>, all dictionary-encoded at the leaf level (dictionary_policy::ALWAYS) for a more realistic workload. Results
|
Description
Follow up #22279. Closes #13995
This PR enables parquet writer to deterministically compute the fragment index that first inserts a key into the dictionary improving the spatial locality of keys across written pages.
Checklist