Add init element overload to insert_or_apply
#555
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
srinivasyadav18
requested review from
sleeepyjack and
PointKernel
as code owners
srinivasyadav18
changed the title
init element overload to insert_or_applyinit element overload to insert_or_apply
PointKernel
added
type: feature request
|
/ok to test |
PointKernel
reviewed
Jul 23, 2024
sleeepyjack
reviewed
Jul 24, 2024
| @@ -594,6 +713,8 @@ class operator_impl< | |||
| auto probing_iter = probing_scheme(key, storage_ref.window_extent()); | |||
| auto const empty_value = ref_.impl_.empty_slot_sentinel().second; | |||
|
|
|||
| auto constexpr use_insert = PayloadWrite and (sizeof(value_type) > 8); | |||
There was a problem hiding this comment.
The sizeof(value_type) > 8 just checks if we use packed_cas or not, right? This value might change based on the architecture (on sm_90 we have 16B CAS). It's fine for now but we have to find a better solution to the problem, e.g., based on a trait system or making it queryable from the slot_inserter (see comment).
PointKernel
reviewed
Jul 26, 2024
There was a problem hiding this comment.
@srinivasyadav18 can you please add missing docs for the new functions added?
sleeepyjack
reviewed
Jul 29, 2024
| @@ -76,33 +76,55 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void insert_or_assign(InputIt first, | |||
| * | |||
| * @note Callable object to perform binary operation should be able to invoke as | |||
| * Op(cuda::atomic_ref<T,Scope>, T>) | |||
| * @note If `HasInit` is `true` and if `init` value equals to the `sentinel value`, we directly | |||
There was a problem hiding this comment.
Suggested change
| * @note If `HasInit` is `true` and if `init` value equals to the `sentinel value`, we directly | |
| * @note If `HasInit` is `true` and `init == empty_value_sentinel`, we directly |
| @@ -76,33 +76,55 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void insert_or_assign(InputIt first, | |||
| * | |||
| * @note Callable object to perform binary operation should be able to invoke as | |||
| * Op(cuda::atomic_ref<T,Scope>, T>) | |||
| * @note If `HasInit` is `true` and if `init` value equals to the `sentinel value`, we directly | |||
| * `apply` the `op` instead of atomic store and then waiting for the payload to get materalized. | |||
| * This has potential speedups when sizeof(value_type) > 8. | |||
There was a problem hiding this comment.
Is the only relevant part here the size of the value or is it the size of the pair, i.e., when using packed CAS vs. CAS + dependent write?
There was a problem hiding this comment.
I think its size of value_type of the map, because that will determine on which type we are performing CAS operation.
There was a problem hiding this comment.
Right. The exact size value might change in the future as we're about to get 128b CAS on Hopper+ exposed through libcu++. That means we can use the packed CAS approach with up-to 16B value types
Benchmarks on H100We can see performance benifits on I64 types from 5-40% by using |
|
/ok to test |
PointKernel
requested changes
Jul 31, 2024
|
/ok to test |
sleeepyjack
approved these changes
Aug 6, 2024
|
/ok to test |
PointKernel
reviewed
Aug 7, 2024
PointKernel
reviewed
Aug 7, 2024
|
/ok to test |
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 adds a new overload to
insert_or_applywhich takes init parameter that represents the identity element of the binary operatoropused for apply operation. This allows to do optimization by skippingwait_for_payloadin case sentienl value of the map equals to the init.