Fix the ORC decoding bug for the timestamp data #17570

kingcrimsontianyu · 2024-12-10T20:08:59Z

Description

This PR introduces a band-aid class run_cache_manager to handle an exceptional case in TIMESTAMP data type, where the DATA stream (seconds) is processed ahead of SECONDARY stream (nanoseconds) and the excess rows are lost. The fix uses run_cache_manager (and also cache_helper, which is an implementation detail) to cache the potentially missed data from the DATA stream and let them be used in the next decoding iteration, thus preventing data loss.

Closes #17155

Checklist

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

copy-pr-bot · 2024-12-10T20:09:03Z

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

kingcrimsontianyu · 2024-12-11T04:18:44Z

/ok to test

kingcrimsontianyu · 2024-12-11T18:18:09Z

/ok to test

kingcrimsontianyu · 2024-12-11T18:28:52Z

/ok to test

kingcrimsontianyu · 2024-12-16T20:41:17Z

/ok to test

kingcrimsontianyu · 2024-12-16T22:23:40Z

/ok to test

kingcrimsontianyu · 2024-12-16T22:38:28Z

/ok to test

vuule

some old comment, not sure if applicable still

cpp/src/io/orc/stripe_data.cu

Matt711

Just a couple small suggestions.

cpp/src/io/orc/stripe_data.cu

python/cudf/cudf/tests/test_orc.py

cpp/src/io/orc/stripe_data.cu

ttnghia · 2024-12-18T18:31:46Z

Should we run a benchmark on this patch to see how much performance impact it causes?

vuule · 2024-12-18T19:22:38Z

Should we run a benchmark on this patch to see how much performance impact it causes?

Are there Spark-RAPIDS benchmarks that we can (also) run to check the impact?

mhaseeb123

Some comments. Overall looks good.

vyasr

I'll be on vacation for the next week and I don't want to block this PR, so I'm just leaving comments without requesting blocking changes. Feel free to ping me if you have thoughts though!

cpp/src/io/orc/stripe_data.cu

ttnghia · 2025-01-02T01:08:36Z

cpp/src/io/orc/stripe_data.cu

+  __shared__ run_cache_manager run_cache_manager_inst;
+  cache_helper cache_helper_inst(run_cache_manager_inst);


Do we need to make any changes to the shared memory allocation upon launching the kernel?

The size of the shared memory needed won't change throughout the kernel execution, hence the static allocation. Does this answer your question?

Correct me if I'm wrong: I don't understand why the shared memory size of a kernel does not change when we add a new shared memory object?

It has changed, but we don't need to declare this explicitly; size of the shared memory is known at compile time, so CUDA takes care of this for us.

Yeah this makes sense. Thanks.

run_cache_manager increases the shared memory usage per block by 12 bytes. Quite negligible in comparison to the existing usage by orcdec_state_s, which takes 37,776 bytes.

@ttnghia Yeah that's a lot of shared memory to use. The static arrays in orcdec_state_s such as the intermediate "byte streams" and the intermediate decoded output are the biggest contributors. This is required by the current design where each block has a hardcoded number of 1024 threads to be able to consume two 512-length runs at a time. We may improve this part in the future when needed.

… selected snappy data with the correct uncompressed version

kingcrimsontianyu · 2025-01-07T19:29:05Z

/merge

mhaseeb123

LGTM!

kingcrimsontianyu added bug Something isn't working non-breaking Non-breaking change labels Dec 10, 2024

kingcrimsontianyu self-assigned this Dec 10, 2024

github-actions bot added the libcudf Affects libcudf (C++/CUDA) code. label Dec 10, 2024

kingcrimsontianyu force-pushed the fix-orc-decode-bug branch from 43fdca1 to 88cfa09 Compare December 11, 2024 15:01

kingcrimsontianyu force-pushed the fix-orc-decode-bug branch from 88cfa09 to 6a8280f Compare December 11, 2024 18:28

vuule self-requested a review December 13, 2024 22:03

kingcrimsontianyu force-pushed the fix-orc-decode-bug branch from 6a8280f to e9f4328 Compare December 16, 2024 19:07

github-actions bot added the Python Affects Python cuDF API. label Dec 17, 2024

kingcrimsontianyu marked this pull request as ready for review December 17, 2024 20:07

kingcrimsontianyu requested review from a team as code owners December 17, 2024 20:07

kingcrimsontianyu requested review from vyasr, Matt711, ttnghia and mhaseeb123 December 17, 2024 20:07

vuule reviewed Dec 17, 2024

View reviewed changes

cpp/src/io/orc/stripe_data.cu Show resolved Hide resolved

cpp/src/io/orc/stripe_data.cu Outdated Show resolved Hide resolved

cpp/src/io/orc/stripe_data.cu Outdated Show resolved Hide resolved

Matt711 approved these changes Dec 17, 2024

View reviewed changes

cpp/src/io/orc/stripe_data.cu Outdated Show resolved Hide resolved

python/cudf/cudf/tests/test_orc.py Outdated Show resolved Hide resolved

mhaseeb123 reviewed Dec 17, 2024

View reviewed changes

cpp/src/io/orc/stripe_data.cu Show resolved Hide resolved

mhaseeb123 reviewed Dec 17, 2024

View reviewed changes

cpp/src/io/orc/stripe_data.cu Show resolved Hide resolved

kingcrimsontianyu force-pushed the fix-orc-decode-bug branch from ab4221d to 0312873 Compare December 18, 2024 15:58

mhaseeb123 reviewed Dec 20, 2024

View reviewed changes

vyasr reviewed Dec 20, 2024

View reviewed changes

ttnghia reviewed Jan 2, 2025

View reviewed changes

cpp/src/io/orc/stripe_data.cu Outdated Show resolved Hide resolved

ttnghia reviewed Jan 2, 2025

View reviewed changes

kingcrimsontianyu added 19 commits January 6, 2025 10:38

Fix the orc decoding bug

335baf2

Cleanup

95f6ee1

Add comments

cd8e04e

Add comments

5b84d72

Improve implementation

53905f6

Improve implementation more

b7dc249

Simplify

079cc53

Improve the comment

663746a

Add python test

9ca9deb

Remove forceinline

bdbd348

Modify comments according to reviewer's comment

2d8e3ce

Add more comments. Rename test orc data name. Replace the incorrectly…

d0cfcaa

… selected snappy data with the correct uncompressed version

Address review comments

caaf866

Address reviewer comment. Add more clarifying comments

17fc285

Move sync around

68a736c

Add cache_helper class to simplify function call

d13be8c

Simplify implementation at the sacrifice of encapsulation

8aa90ae

Minor change

ee41f21

Update comments

9c87960

kingcrimsontianyu force-pushed the fix-orc-decode-bug branch from c2a2992 to 9c87960 Compare January 6, 2025 15:44

Update copyright date

ebc3860

vuule approved these changes Jan 7, 2025

View reviewed changes

rapids-bot bot merged commit 4e97cd4 into rapidsai:branch-25.02 Jan 7, 2025
116 of 117 checks passed

mhaseeb123 reviewed Jan 7, 2025

View reviewed changes

kingcrimsontianyu mentioned this pull request Jan 10, 2025

Add special orc test data: timestamp interspersed with null values #17713

Merged

3 tasks

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Fix the ORC decoding bug for the timestamp data #17570

Fix the ORC decoding bug for the timestamp data #17570

kingcrimsontianyu commented Dec 10, 2024 •

edited

Loading

copy-pr-bot bot commented Dec 10, 2024

kingcrimsontianyu commented Dec 11, 2024

kingcrimsontianyu commented Dec 11, 2024

kingcrimsontianyu commented Dec 11, 2024

kingcrimsontianyu commented Dec 16, 2024

kingcrimsontianyu commented Dec 16, 2024

kingcrimsontianyu commented Dec 16, 2024

vuule left a comment

Matt711 left a comment

ttnghia commented Dec 18, 2024

vuule commented Dec 18, 2024 •

edited

Loading

mhaseeb123 left a comment

vyasr left a comment

ttnghia Jan 2, 2025

kingcrimsontianyu Jan 6, 2025

ttnghia Jan 6, 2025

vuule Jan 6, 2025

ttnghia Jan 6, 2025

kingcrimsontianyu Jan 7, 2025

kingcrimsontianyu Jan 7, 2025

kingcrimsontianyu commented Jan 7, 2025

mhaseeb123 left a comment

		__shared__ run_cache_manager run_cache_manager_inst;
		cache_helper cache_helper_inst(run_cache_manager_inst);

Fix the ORC decoding bug for the timestamp data #17570

Fix the ORC decoding bug for the timestamp data #17570

Conversation

kingcrimsontianyu commented Dec 10, 2024 • edited Loading

Description

Checklist

copy-pr-bot bot commented Dec 10, 2024

kingcrimsontianyu commented Dec 11, 2024

kingcrimsontianyu commented Dec 11, 2024

kingcrimsontianyu commented Dec 11, 2024

kingcrimsontianyu commented Dec 16, 2024

kingcrimsontianyu commented Dec 16, 2024

kingcrimsontianyu commented Dec 16, 2024

vuule left a comment

Choose a reason for hiding this comment

Matt711 left a comment

Choose a reason for hiding this comment

ttnghia commented Dec 18, 2024

vuule commented Dec 18, 2024 • edited Loading

mhaseeb123 left a comment

Choose a reason for hiding this comment

vyasr left a comment

Choose a reason for hiding this comment

ttnghia Jan 2, 2025

Choose a reason for hiding this comment

kingcrimsontianyu Jan 6, 2025

Choose a reason for hiding this comment

ttnghia Jan 6, 2025

Choose a reason for hiding this comment

vuule Jan 6, 2025

Choose a reason for hiding this comment

ttnghia Jan 6, 2025

Choose a reason for hiding this comment

kingcrimsontianyu Jan 7, 2025

Choose a reason for hiding this comment

kingcrimsontianyu Jan 7, 2025

Choose a reason for hiding this comment

kingcrimsontianyu commented Jan 7, 2025

mhaseeb123 left a comment

Choose a reason for hiding this comment

kingcrimsontianyu commented Dec 10, 2024 •

edited

Loading

vuule commented Dec 18, 2024 •

edited

Loading