-
Notifications
You must be signed in to change notification settings - Fork 46
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Develop stream 2024-09-12 #462
Open
NB4444
wants to merge
45
commits into
ROCm:develop
Choose a base branch
from
StreamHPC:develop-upstream-12-9-2024
base: develop
Could not load branches
Branch not found: {{ refName }}
Loading
Could not load tags
Nothing to show
Loading
Are you sure you want to change the base?
Some commits from the old base branch may be removed from the timeline,
and old review comments may become outdated.
+13,550
−13,115
Open
Changes from 1 commit
Commits
Show all changes
45 commits
Select commit
Hold shift + click to select a range
9075421
Fixed overflow bug for large sizes in thrust::shuffle
Beanavil e9397cc
Added definitions of execution space macros
Beanavil a7a5d20
Add missing overloads for thrust::pow
Beanavil 7006599
Refactors thrust::unique_by_key to use cub::DeviceSelect::UniqueByKey
Beanavil e5dbdaa
Fix a typo in thrust-config.cmake
Beanavil bd43018
Check that thrust::pair is trivially copyable
Beanavil b7b785e
Remove double ignore in discard_iterator.h docs
Beanavil 93b72cd
Replace deprecated _VSTD macro with std
Beanavil f3e2676
Update mode example to use thrust::unique_count
Beanavil 44d7369
Ensure that thrust fancy iterators are trivially_copy_constructible w…
Beanavil a32a67c
Use checked allocators in CUB catch2 tests
Beanavil b741017
Refactors thrust::copy_if to use cub::DeviceSelect
Beanavil 158fa53
Refactor thrust::[stable_]partition[_copy] to use cub::DevicePartition
Beanavil bc6c83b
Fix include of <thrust/random.h> with NVC++
Beanavil 489c073
Cleanup diagnostic handling
Beanavil 9f5a3ba
Rework config.h
Beanavil 1020a11
Bump version to 2.4.0
Beanavil 917c255
Fix issues with ambiguous calls to addressof in thrust::optional
Beanavil 5af1ef7
Try harder to unwrap nested thrust::tuple_of_iterator_references, CUD…
Beanavil bd5228c
Added missing element from thrust's tuple implementation
Beanavil 099a901
Ensure that we can run reduce_by_key with const inputs
Beanavil 9508470
Leave definitions of __host__ and __device__
Beanavil 6791366
Patched up CI because of CCCL2.4.0 tests' build failure
Beanavil 9fe0b04
Updated tests and examples for __host__ __device__ use
Beanavil 15a07b0
Updated CHANGELOG
Beanavil 158a1e1
Added operator to transform_reduce benchmark
NB4444 d0bf50f
Added mem allocator in benchmarks
NB4444 aa64ae7
Changes for review
NB4444 0673125
ci: set up sccache
Snektron 75c44cf
Added helper functions for choosing between different custom reporter
NB4444 a36adac
Added json and csv custom reporter for benchmarks
NB4444 e00ad3a
Changes for review
NB4444 1cc4c8b
Added hipstdpar tests
Beanavil 568f6f9
Relocated our ParallelSTL additions
Beanavil f584551
Fixed several naming issues
Beanavil 387dbcf
Added missing unimplemented algorithms
Beanavil 8aff938
Split hipstdpar_lib.hpp
Beanavil e2d548f
Added relevant information to README and CHANGELOG regarding HIPSTDPAR
Beanavil 28da1b1
Clarified upstream LLVM offload support
Beanavil ed35b28
Emit error when HIPSTDPAR macros are not defined
Beanavil 5023555
Move forwarding calls to rocPRIM to thrust's stubs
Beanavil e340947
Fix path to hipstdpar impl headers
Beanavil 1f060d8
Prevent building hipstdpar tests when no compatible libstdc++ is present
Beanavil 5fdf870
Disable TBB tests build
Beanavil 5be02c2
Merge branch 'develop' into develop-upstream-12-9-2024
NB4444 File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
Large diffs are not rendered by default.
Oops, something went wrong.
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,107 @@ | ||
| /* | ||
| * Licensed under the Apache License, Version 2.0 (the "License"); | ||
| * you may not use this file except in compliance with the License. | ||
| * You may obtain a copy of the License at | ||
| * | ||
| * http://www.apache.org/licenses/LICENSE-2.0 | ||
| * | ||
| * Unless required by applicable law or agreed to in writing, software | ||
| * distributed under the License is distributed on an "AS IS" BASIS, | ||
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| * See the License for the specific language governing permissions and | ||
| * limitations under the License. | ||
| */ | ||
|
|
||
| /*! \file thrust/system/hip/hipstdpar/include/batch.hpp | ||
| * \brief <tt>Batch operations</tt> implementation detail header for HIPSTDPAR. | ||
| */ | ||
|
|
||
| #pragma once | ||
|
|
||
| #if defined(__HIPSTDPAR__) | ||
|
|
||
| #include "hipstd.hpp" | ||
|
|
||
| #include <thrust/execution_policy.h> | ||
| #include <thrust/for_each.h> | ||
|
|
||
| #include <algorithm> | ||
| #include <execution> | ||
| #include <utility> | ||
|
|
||
| namespace std | ||
| { | ||
| // BEGIN FOR_EACH | ||
| template< | ||
| typename I, | ||
| typename F, | ||
| enable_if_t< | ||
| ::hipstd::is_offloadable_iterator<I>() && | ||
| ::hipstd::is_offloadable_callable<F>()>* = nullptr> | ||
| inline | ||
| void for_each(execution::parallel_unsequenced_policy, I f, I l, F fn) | ||
| { | ||
| ::thrust::for_each(::thrust::device, f, l, ::std::move(fn)); | ||
| } | ||
|
|
||
| template< | ||
| typename I, | ||
| typename F, | ||
| enable_if_t< | ||
| !::hipstd::is_offloadable_iterator<I>() || | ||
| !::hipstd::is_offloadable_callable<F>()>* = nullptr> | ||
| inline | ||
| void for_each(execution::parallel_unsequenced_policy, I f, I l, F fn) | ||
| { | ||
| if constexpr (!::hipstd::is_offloadable_iterator<I>()) { | ||
| ::hipstd::unsupported_iterator_category< | ||
| typename iterator_traits<I>::iterator_category>(); | ||
| } | ||
| if constexpr (!::hipstd::is_offloadable_callable<F>()) { | ||
| ::hipstd::unsupported_callable_type<F>(); | ||
| } | ||
|
|
||
| return | ||
| ::std::for_each(::std::execution::par, f, l, ::std::move(fn)); | ||
| } | ||
| // END FOR_EACH | ||
|
|
||
| // BEGIN FOR_EACH_N | ||
| template< | ||
| typename I, | ||
| typename N, | ||
| typename F, | ||
| enable_if_t< | ||
| ::hipstd::is_offloadable_iterator<I>() && | ||
| ::hipstd::is_offloadable_callable<F>()>* = nullptr> | ||
| inline | ||
| I for_each_n(execution::parallel_unsequenced_policy, I f, N n, F fn) | ||
| { | ||
| return | ||
| ::thrust::for_each_n(::thrust::device, f, n, ::std::move(fn)); | ||
| } | ||
|
|
||
| template< | ||
| typename I, | ||
| typename N, | ||
| typename F, | ||
| enable_if_t< | ||
| !::hipstd::is_offloadable_iterator<I>() || | ||
| !::hipstd::is_offloadable_callable<F>()>* = nullptr> | ||
| inline | ||
| I for_each_n(execution::parallel_unsequenced_policy, I f, N n, F fn) | ||
| { | ||
| if constexpr (!::hipstd::is_offloadable_iterator<I>()) { | ||
| ::hipstd::unsupported_iterator_category< | ||
| typename iterator_traits<I>::iterator_category>(); | ||
| } | ||
| if constexpr (!::hipstd::is_offloadable_callable<F>()) { | ||
| ::hipstd::unsupported_callable_type<F>(); | ||
| } | ||
|
|
||
| return | ||
| ::std::for_each_n(::std::execution::par, f, n, ::std::move(fn)); | ||
| } | ||
| // END FOR_EACH_N | ||
| } | ||
| #endif // __HIPSTDPAR__ | ||
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,157 @@ | ||
| /* | ||
| * Licensed under the Apache License, Version 2.0 (the "License"); | ||
| * you may not use this file except in compliance with the License. | ||
| * You may obtain a copy of the License at | ||
| * | ||
| * http://www.apache.org/licenses/LICENSE-2.0 | ||
| * | ||
| * Unless required by applicable law or agreed to in writing, software | ||
| * distributed under the License is distributed on an "AS IS" BASIS, | ||
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| * See the License for the specific language governing permissions and | ||
| * limitations under the License. | ||
| */ | ||
|
|
||
| /*! \file thrust/system/hip/hipstdpar/copy.hpp | ||
| * \brief <tt>Copy operations</tt> implementation detail header for HIPSTDPAR. | ||
| */ | ||
|
|
||
| #pragma once | ||
|
|
||
| #if defined(__HIPSTDPAR__) | ||
|
|
||
| #include "hipstd.hpp" | ||
|
|
||
| #include <thrust/copy.h> | ||
| #include <thrust/execution_policy.h> | ||
|
|
||
| #include <algorithm> | ||
| #include <execution> | ||
| #include <utility> | ||
|
|
||
| namespace std | ||
| { | ||
| // BEGIN COPY | ||
| template< | ||
| typename I, | ||
| typename O, | ||
| enable_if_t<::hipstd::is_offloadable_iterator<I, O>()>* = nullptr> | ||
| inline | ||
| O copy(execution::parallel_unsequenced_policy, I fi, I li, O fo) | ||
| { | ||
| return ::thrust::copy(::thrust::device, fi, li, fo); | ||
| } | ||
|
|
||
| template< | ||
| typename I, | ||
| typename O, | ||
| enable_if_t<!::hipstd::is_offloadable_iterator<I, O>()>* = nullptr> | ||
| inline | ||
| O copy(execution::parallel_unsequenced_policy, I fi, I li, O fo) | ||
| { | ||
| ::hipstd::unsupported_iterator_category< | ||
| typename iterator_traits<I>::iterator_category, | ||
| typename iterator_traits<O>::iterator_category>(); | ||
|
|
||
| return ::std::copy(::std::execution::par, fi, li, fo); | ||
| } | ||
| // END COPY | ||
|
|
||
| // BEGIN COPY_IF | ||
| template< | ||
| typename I, | ||
| typename O, | ||
| typename P, | ||
| enable_if_t< | ||
| ::hipstd::is_offloadable_iterator<I, O>() && | ||
| ::hipstd::is_offloadable_callable<P>()>* = nullptr> | ||
| inline | ||
| O copy_if(execution::parallel_unsequenced_policy, I fi, I li, O fo, P p) | ||
| { | ||
| return | ||
| ::thrust::copy_if(::thrust::device, fi, li, fo, ::std::move(p)); | ||
| } | ||
|
|
||
| template< | ||
| typename I, | ||
| typename O, | ||
| typename P, | ||
| enable_if_t< | ||
| !::hipstd::is_offloadable_iterator<I, O>() || | ||
| !::hipstd::is_offloadable_callable<P>()>* = nullptr> | ||
| inline | ||
| O copy_if(execution::parallel_unsequenced_policy, I fi, I li, O fo, P p) | ||
| { | ||
| if constexpr (!::hipstd::is_offloadable_iterator<I, O>()) { | ||
| ::hipstd::unsupported_iterator_category< | ||
| typename iterator_traits<I>::iterator_category, | ||
| typename iterator_traits<O>::iterator_category>(); | ||
| } | ||
| if constexpr (!::hipstd::is_offloadable_callable<P>()) { | ||
| ::hipstd::unsupported_callable_type<P>(); | ||
| } | ||
|
|
||
| return ::std::copy_if( | ||
| ::std::execution::par, fi, li, fo, ::std::move(p)); | ||
| } | ||
| // END COPY_IF | ||
|
|
||
| // BEGIN COPY_N | ||
| template< | ||
| typename I, | ||
| typename N, | ||
| typename O, | ||
| enable_if_t<::hipstd::is_offloadable_iterator<I, O>()>* = nullptr> | ||
| inline | ||
| O copy_n(execution::parallel_unsequenced_policy, I fi, N n, O fo) | ||
| { | ||
| return ::thrust::copy_n(::thrust::device, fi, n, fo); | ||
| } | ||
|
|
||
| template< | ||
| typename I, | ||
| typename N, | ||
| typename O, | ||
| enable_if_t<!::hipstd::is_offloadable_iterator<I, O>()>* = nullptr> | ||
| inline | ||
| O copy_n(execution::parallel_unsequenced_policy, I fi, N n, O fo) | ||
| { | ||
| ::hipstd::unsupported_iterator_category< | ||
| typename iterator_traits<I>::iterator_category, | ||
| typename iterator_traits<O>::iterator_category>(); | ||
|
|
||
| return ::std::copy_n(::std::execution::par, fi, n, fo); | ||
| } | ||
| // END COPY_N | ||
|
|
||
| // BEGIN MOVE | ||
| template< | ||
| typename I, | ||
| typename O, | ||
| enable_if_t<::hipstd::is_offloadable_iterator<I, O>()>* = nullptr> | ||
| inline | ||
| O move(execution::parallel_unsequenced_policy, I fi, I li, O fo) | ||
| { | ||
| return ::thrust::copy( | ||
| ::thrust::device, | ||
| make_move_iterator(fi), | ||
| make_move_iterator(li), | ||
| fo); | ||
| } | ||
|
|
||
| template< | ||
| typename I, | ||
| typename O, | ||
| enable_if_t<!::hipstd::is_offloadable_iterator<I, O>()>* = nullptr> | ||
| inline | ||
| O move(execution::parallel_unsequenced_policy, I fi, I li, O fo) | ||
| { | ||
| ::hipstd::unsupported_iterator_category< | ||
| typename iterator_traits<I>::iterator_category, | ||
| typename iterator_traits<O>::iterator_category>(); | ||
|
|
||
| return ::std::move(::std::execution::par, fi, li, fo); | ||
| } | ||
| // END MOVE | ||
| } | ||
| #endif // __HIPSTDPAR__ |
Oops, something went wrong.
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.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Doing the split is fine, since it aligns with Thrust convention, I guess; however, IMHO it'd be better if the master header (this), which is the object of implicit inclusion / is part of the public interface, would do nothing but include the split headers, as opposed to having any actual functionality. We could also probably guard the
interpose_allocations.hppone with its respective macros. Finally, we probably want to error on the impl headers if they're directly included without__HIPSTDPAR__being defined, as that would indicate user error.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I have some questions about these suggestions:
Do you mean that we should further split the master headers (the ones named after the C++ Standard Library categories for algorithms)?
What macros do you refer to? IINM the
interpose_allocations.hppis already guarded by the__HIPSTDPAR__and__HIPSTDPAR_INTERPOSE_ALLOC__macros (as it was before splittinghipstdpar_lib.hpp).And for this, I don't fully understand why would we need to error here. Originally there was no error handling in case that macro was not defined in the
hipstdpar_lib.hpp, so I'm not sure why it would be done for the impl headers. Also, as I had understood it, the forwarding header was not intended to be included directly by users, but implicitly included by the compiler when specifying the--hipstdparoption (which also triggers the definition of the__HIPSTDPAR__macro). So if this is true (please correct me if not), then it wouldn't make sense to check for the definition of this macro IMHO.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Not quite, what I meant is that the master header would pretty much be a list of
#includes for the impl headers, with no declarations / definitions or, really, any substantive code living there. In essence, the most we'd do here is check for environmental macros / possibly declare some future config macros, if necessary. Does that make sense?What I was suggesting as a possible alternative is to check for
__HIPSTDPAR_INTERPOSE_ALLOC__in the masterhipstdpar_lib.hppheader and do conditional include of the impl header in there, rather than guarding within it. This ties both to the point above re: what to do with the master header and to the one below about considering error-ing out.Whilst the vote of confidence is highly appreciated, I'd not take the original header to necessarily be complete / the best way to do things:) To wit, whilst the forwarding header was / is not meant to be included by users, that does not mean users won't do it (accidentally / intentionally). Error-ing out if the macro is not defined i.e. compilation for
hipstdparwas not desired would guard against accidents, and will at least force intentional abusers to manually define the macro, which'd help with eventual debugging. I suggested it because I think it might be a time saver / inform users about things potentially going awry. Ultimately, it's up to you if you want to go in this direction.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Oh alright, I was confused because the comment was made on
impl/batch.hpp. But then, the master headerhipstdpar_lib.hppis already just a macro definition check (__HIPSTDPAR__) and a bunch of includes (for theimpl/headers)Regarding those two points, the final logic should be then:
hipstdpar_lib.hppif__HIPSTDPAR__is defined and include theimplheaders if so, exceptinterpose_allocations.hppwhich only will be included if__HIPSTDPAR_INTERPOSE_ALLOC__is also defined.__HIPSTDPAR__is defined and err if it's not.interpose_allocations.hpp, check if__HIPSTDPAR_INTERPOSE_ALLOC__is defined and err if not.If that sounds good I'll add a commit with the changes
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
This sounds great, thanks and apologies for adding work.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Also added the err logic discussed before