Improve performance of expression evaluation #9210
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
…e wherever possible.
…rather than the constructor and forward it along.
This reverts commit 7f3c775.
vyasr
added
3 - Ready for Review
|
Here are some benchmarks Conditional joins on a Tesla T4: Benchmarks
Before: After:
Benchmarks
Before: After: Conditional joins on an RTX 8000: Benchmarks
Before: After:
Benchmarks
Before: After: |
jrhemstad
reviewed
Sep 10, 2021
| __device__ possibly_null_value_t<Element, has_nulls> resolve_input( | ||
| detail::device_data_reference device_data_reference, | ||
| CUDA_DEVICE_CALLABLE possibly_null_value_t<Element, has_nulls> resolve_input( | ||
| detail::device_data_reference const& input_reference, |
There was a problem hiding this comment.
Out of curiosity, did you try seeing if passing by const value had the same impact?
There was a problem hiding this comment.
I did test that, and it seemed to work in some cases but not others. In particular the nullable template had different performance characteristics from the non-nullable one, w.r.t to both this change and a number of others. The obvious culprit is the fact that the nullable template simply has more local state, so I need to reduce it even further to see measurable improvements in occupancy, but I also observed more subtle differences that I'd probably attribute to the compiler just not doing as good a job at some point. I have some ideas for redesigning the data references that might hopefully make these issues moot, but I'll need to play with those a bit and I figured that this perf boost was a nice short-term win while I do that.
bdice
requested changes
Sep 10, 2021
bdice
approved these changes
Sep 13, 2021
Codecov Report
@@ Coverage Diff @@
## branch-21.10 #9210 +/- ##
===============================================
Coverage ? 10.84%
===============================================
Files ? 115
Lines ? 19173
Branches ? 0
===============================================
Hits ? 2080
Misses ? 17093
Partials ? 0 Continue to review full report at Codecov.
|
|
rerun tests |
davidwendt
reviewed
Sep 14, 2021
|
rerun tests |
1 similar comment
|
rerun tests |
davidwendt
approved these changes
Sep 16, 2021
|
@gpucibot merge |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This PR does some minor reworking of the internals of expression evaluation to improve performance. The largest performance improvements come from passing device data references down the call stack by reference rather than by value. The nullable kernel template experiences significantly higher register pressure and these changes do not seem to be as effective at increasing occupancy on the benchmarks with null data, but in general we see performance improvements across the board for non-null data and in some cases for nulllable data, with improvements ranging up to 40%. This PR also does some minor cleanup: removing some unused functions, replacing
__device__withCUDA_DEVICE_CALLABLEto ensure compatibility with host compilers, and fixing the templating of various functions to ensure proper usage of CRTP. These changes are intended to facilitate future redesigning of the internals of thedevice_data_referencesto reduce the depth of these call stacks, simplify the code, and reduce register pressure.