Fix writing of Parquet files with many fragments #11869

etseidl · 2022-10-06T01:18:45Z

Description

This PR fixes an error that can occur when very small page sizes are used when writing Parquet files. #11551 changed from fixed 5000 row page fragments to a scaled value based on the requested max page size. For small page sizes, the number of fragments to process can exceed 64k. The number of fragments is used as the y dimension when calling gpuInitPageFragments, and when it exceeds 64k the kernel fails to launch, ultimately leading to an invalid memory access.

Checklist

I am familiar with the Contributing Guidelines.
New or existing tests cover these changes.
The documentation is up to date with these changes.

etseidl · 2022-10-06T01:20:07Z

Alternative approaches to consider are either using a single dimension and calculating column and fragment indexes, or using a fixed y dimension and looping over fragments.

I added a test that fails when FRAGSWAP is set to 0.

nvbench and nsys profile don't show any performance degradation with the swap in place.

codecov · 2022-10-06T03:39:51Z

Codecov Report

Base: 87.40% // Head: 88.12% // Increases project coverage by +0.71% 🎉

Coverage data is based on head (56ebf7c) compared to base (f72c4ce).
Patch coverage: 88.55% of modified lines in pull request are covered.

❗ Current head 56ebf7c differs from pull request most recent head a072482. Consider uploading reports for the commit a072482 to get more accurate results

Additional details and impacted files

@@               Coverage Diff                @@
##           branch-22.12   #11869      +/-   ##
================================================
+ Coverage         87.40%   88.12%   +0.71%     
================================================
  Files               133      133              
  Lines             21833    21905      +72     
================================================
+ Hits              19084    19304     +220     
+ Misses             2749     2601     -148

Impacted Files	Coverage Δ
python/cudf/cudf/core/dataframe.py	`93.77% <ø> (ø)`
python/cudf/cudf/core/indexed_frame.py	`92.03% <ø> (ø)`
python/cudf/cudf/core/udf/__init__.py	`97.05% <ø> (+47.05%)`	⬆️
python/cudf/cudf/io/orc.py	`92.94% <ø> (-0.09%)`	⬇️
python/cudf/cudf/testing/dataset_generator.py	`72.83% <ø> (-0.42%)`	⬇️
...thon/dask_cudf/dask_cudf/tests/test_distributed.py	`18.86% <ø> (+4.94%)`	⬆️
python/cudf/cudf/core/_base_index.py	`82.20% <43.75%> (-3.35%)`	⬇️
python/cudf/cudf/io/text.py	`91.66% <66.66%> (-8.34%)`	⬇️
python/strings_udf/strings_udf/__init__.py	`84.31% <76.00%> (-12.57%)`	⬇️
python/cudf/cudf/core/index.py	`92.96% <95.71%> (+0.33%)`	⬆️
... and 20 more

Help us with your feedback. Take ten seconds to tell us how you rate us. Have a feature suggestion? Share it here.

☔ View full report at Codecov.
📢 Do you have feedback about the report comment? Let us know in this issue.

etseidl · 2022-10-06T17:54:58Z

A downside to simply swapping x and y is that now the number of columns is limited to 64k. A 1D grid would allow either number of fragments or number of columns to exceed the 64k limit, but not both obviously. Maybe a fixed y with looping (as suggested offline by @vuule) is the best fix?

cpp/src/io/parquet/page_enc.cu

vuule

I like this! Got nothing to contribute :)

bdice · 2022-10-17T23:08:31Z

cpp/src/io/parquet/page_enc.cu

-    g->col       = &col_desc[column_id];
-    g->start_row = fragments[column_id][frag_id].start_value_idx;
-    g->num_rows  = fragments[column_id][frag_id].num_leaf_values;
+  uint32_t const lane_id                  = threadIdx.x & 0x1f;


There are some rather innocuous-seeming magic values that all related to cudf::detail::warp_size in this function. I'll point them out, but I am fine with doing nothing if we feel the current code is better not to change.

Suggested change

uint32_t const lane_id = threadIdx.x & 0x1f;

uint32_t const lane_id = threadIdx.x % cudf::detail::warp_size;

sounds good (although using the mod operator makes my teeth itch 🤣). Does anyone happen to know if there are constants anywhere for the max threadblock dimensions? Or are those per-card values?

I believe there are no constants for that, and that's why we defined cudf::detail::warp_size. It is a constant for all NVIDIA GPUs as far as I am aware.

These two snippets should compile out roughly the same. Compilers can recognize that unsigned modulo by $2^N$ is equivalent to bitwise-and with $2^N - 1$. Evidence: https://godbolt.org/z/r4c41va5P

I went ahead and added a constexpr for the warp mask (before I read your reply)...there are several other instances of 0x1f sprinkled about in this file that can be replaced later.

Thanks for the link @bdice! Should I get rid of my mask constexpr and just use cudf::detail::warp_size everywhere?

To be a bit more precise here, CUDA does provide warpSize, which is available inside device code, and the getDeviceProperties host function, which returns a struct containing the warp size. However, neither of them is a constant and therefore cannot be used in constant expressions (e.g. for declaring a C-style or std::array). The warp size is indeed constant across all current compute capabilities. In theory that's not something that we promise, so the technically correct answer is that we can't use a compile-time constant because in theory someone could run on a new architecture with a different answer. In practice, NVIDIA has no plans to change the warp size AFAIK and many examples of GPU code (even lots of code written by NVIDIA) define a warp_size constant. Lots of places use it assuming that it is in fact a compile-time constant and would have to be rewritten if we ever had any cards with a different warp size, so that's a much bigger problem to deal with another day if that ever changes :)

bdice · 2022-10-17T23:08:49Z

cpp/src/io/parquet/page_enc.cu

+  uint32_t const column_id                = blockIdx.x;
+  uint32_t const num_fragments_per_column = fragments.size().second;
+
+  uint32_t frag_id = blockIdx.y * 4 + (threadIdx.x >> 5);


Suggested change

uint32_t frag_id = blockIdx.y * 4 + (threadIdx.x >> 5);

uint32_t frag_id = blockIdx.y * 4 + (threadIdx.x * cudf::detail::warp_size);

bdice

A couple of magic values that don't need to be magical -- otherwise LGTM.

etseidl · 2022-10-19T17:04:32Z

Tests seem to be failing on the mimesis stuff now. Should I merge with 22.12 to pull in #11906?

bdice · 2022-10-19T17:07:37Z

@etseidl Merging the upstream or commenting “rerun tests” should work.

bdice · 2022-10-19T17:07:49Z

rerun tests

vuule · 2022-10-20T19:18:53Z

@gpucibot merge

etseidl and others added 5 commits October 5, 2022 17:15

swap columns/fragments in gpuInitPageFragments

2d0ec6d

get rid of some print statements

f148442

add test for too many fragments

1e9039e

formatting

690bd65

Merge branch 'rapidsai:branch-22.12' into feature/fragments_fix

e14cd7c

github-actions bot added the libcudf Affects libcudf (C++/CUDA) code. label Oct 6, 2022

etseidl added 6 commits October 6, 2022 13:34

switch to loop driven version

571d00c

formatting

94e3e02

remove leftover newline

09851df

add consts

d69fc2d

rework gpuInitFragmentStats() to move sync out of the loop.

5dff33a

get rid of shared variables and synchronization

3488d25

vuule added bug Something isn't working non-breaking Non-breaking change labels Oct 11, 2022

vuule reviewed Oct 11, 2022

View reviewed changes

cpp/src/io/parquet/page_enc.cu Outdated Show resolved Hide resolved

vuule approved these changes Oct 17, 2022

View reviewed changes

etseidl marked this pull request as ready for review October 17, 2022 21:08

etseidl requested a review from a team as a code owner October 17, 2022 21:08

etseidl requested review from bdice and karthikeyann October 17, 2022 21:08

bdice reviewed Oct 17, 2022

View reviewed changes

bdice approved these changes Oct 17, 2022

View reviewed changes

karthikeyann approved these changes Oct 18, 2022

View reviewed changes

reduce magic numbers per review suggestion

a072482

jbrennan333 mentioned this pull request Oct 18, 2022

[FEA] Support ZSTD compression with Parquet and Orc NVIDIA/spark-rapids#3037

Closed

rapids-bot bot merged commit 98185fe into rapidsai:branch-22.12 Oct 20, 2022

etseidl deleted the feature/fragments_fix branch October 20, 2022 22:23

bdice mentioned this pull request Apr 26, 2023

Clean up and simplify gpuDecideCompression #13202

Merged

3 tasks

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Fix writing of Parquet files with many fragments #11869

Fix writing of Parquet files with many fragments #11869

etseidl commented Oct 6, 2022 •

edited

Loading

etseidl commented Oct 6, 2022 •

edited

Loading

codecov bot commented Oct 6, 2022 •

edited

Loading

etseidl commented Oct 6, 2022

vuule left a comment

bdice Oct 17, 2022 •

edited

Loading

etseidl Oct 18, 2022

bdice Oct 18, 2022 •

edited

Loading

etseidl Oct 18, 2022 •

edited

Loading

vyasr Oct 18, 2022

bdice Oct 17, 2022

bdice left a comment

etseidl commented Oct 19, 2022

bdice commented Oct 19, 2022

bdice commented Oct 19, 2022

vuule commented Oct 20, 2022

	uint32_t const lane_id = threadIdx.x & 0x1f;
	uint32_t const lane_id = threadIdx.x % cudf::detail::warp_size;

	uint32_t frag_id = blockIdx.y * 4 + (threadIdx.x >> 5);
	uint32_t frag_id = blockIdx.y * 4 + (threadIdx.x * cudf::detail::warp_size);

Fix writing of Parquet files with many fragments #11869

Fix writing of Parquet files with many fragments #11869

Conversation

etseidl commented Oct 6, 2022 • edited Loading

Description

Checklist

etseidl commented Oct 6, 2022 • edited Loading

codecov bot commented Oct 6, 2022 • edited Loading

Codecov Report

etseidl commented Oct 6, 2022

vuule left a comment

Choose a reason for hiding this comment

bdice Oct 17, 2022 • edited Loading

Choose a reason for hiding this comment

etseidl Oct 18, 2022

Choose a reason for hiding this comment

bdice Oct 18, 2022 • edited Loading

Choose a reason for hiding this comment

etseidl Oct 18, 2022 • edited Loading

Choose a reason for hiding this comment

vyasr Oct 18, 2022

Choose a reason for hiding this comment

bdice Oct 17, 2022

Choose a reason for hiding this comment

bdice left a comment

Choose a reason for hiding this comment

etseidl commented Oct 19, 2022

bdice commented Oct 19, 2022

bdice commented Oct 19, 2022

vuule commented Oct 20, 2022

etseidl commented Oct 6, 2022 •

edited

Loading

etseidl commented Oct 6, 2022 •

edited

Loading

codecov bot commented Oct 6, 2022 •

edited

Loading

bdice Oct 17, 2022 •

edited

Loading

bdice Oct 18, 2022 •

edited

Loading

etseidl Oct 18, 2022 •

edited

Loading