Skip to content

Reduction functors (cuco::static_reduction_map refactoring 1/N) #187

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.

Sign up for GitHub

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

Closed
sleeepyjack wants to merge 12 commits into from
Closed

Reduction functors (cuco::static_reduction_map refactoring 1/N) #187

Changes from 1 commit
Commits
Show all changes
12 commits
Select commit Hold shift + click to select a range
2998d9f
Update Catch2 to v2.13.9.
sleeepyjack Jul 11, 2022
2a8a50f
Added reduction functors that can be used by static_reduction_map.
sleeepyjack Jul 11, 2022
e4adb05
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Jul 11, 2022
bb78c49
Prevent extended __host__ __device__/__device__ lambdas from being us…
sleeepyjack Jul 11, 2022
905edd8
Remove spurious include.
sleeepyjack Jul 11, 2022
7d3aff3
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Jul 11, 2022
9cbe890
Workaround for compiler warning where nvcc is not able to see the ret…
sleeepyjack Jul 12, 2022
cef8906
Fix includes und use type trait aliases.
sleeepyjack Jul 12, 2022
80c1544
Reduction ops should use relaxed memory order.
sleeepyjack Jul 12, 2022
6fc5ff2
Make identity_value ctor explicit.
sleeepyjack Jul 12, 2022
99945f4
Internally-synced reduction functors cannot have immutable target obj…
sleeepyjack Aug 1, 2022
4435d05
Use human-readable boolean operators.
sleeepyjack Aug 1, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 18 additions & 1 deletion include/cuco/reduction_functors.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -135,8 +135,25 @@ class reduction_functor : detail::reduction_functor_base {
Func,
cuda::atomic<value_type, cuda::thread_scope_thread>&,
value_type>;
static constexpr bool atomic_const_invocable_ =
cuda::std::is_invocable_r_v<value_type,
Func,
cuda::atomic<value_type, cuda::thread_scope_system> const&,
value_type> ||
cuda::std::is_invocable_r_v<value_type,
Func,
cuda::atomic<value_type, cuda::thread_scope_device> const&,
value_type> ||
cuda::std::is_invocable_r_v<value_type,
Func,
cuda::atomic<value_type, cuda::thread_scope_block> const&,
value_type> ||
cuda::std::is_invocable_r_v<value_type,
Func,
cuda::atomic<value_type, cuda::thread_scope_thread> const&,
value_type>;
Comment on lines +138 to +155

@sleeepyjack sleeepyjack Aug 1, 2022

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Listing all possible values for cuda::thread_scope is very ugly. Does anyone have a better solution?


static_assert(atomic_invocable_ || naive_invocable_,
static_assert((atomic_invocable_ && !atomic_const_invocable_) || naive_invocable_,
"Invalid operator signature. Valid signatures are "
"(T const&, T const&)->T and (cuda::atomic<T, Scope>&, T const&)->T.");
static_assert(!(__nv_is_extended_device_lambda_closure_type(Func) ||
Expand Down