Skip to content
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

Use sub_groups in SYCL parallel_scan #4147

Merged
merged 1 commit into from
Jul 19, 2021

Conversation

masterleinad
Copy link
Contributor

One more step towards using subgroups where appropriate.

  • For using subgroups with shuffles, we need to restrict the algorithm to scalar values (which we already only allow for CUDA for example) so I dropped all ValueOps::copy.
  • We need to allocate slightly more global memory to make sure that we can store intermediate results used for the scan of subgroup results.

@masterleinad masterleinad added this to In progress in Developer: Daniel Arndt Jul 8, 2021
@masterleinad masterleinad moved this from In progress to Awaiting Feedback in Developer: Daniel Arndt Jul 8, 2021
@masterleinad masterleinad removed the [WIP] label Jul 8, 2021
@masterleinad masterleinad marked this pull request as ready for review July 8, 2021 18:47
@@ -86,9 +86,9 @@ class ParallelScanSYCLBase {
void scan_internal(sycl::queue& q, const Functor& functor,
pointer_type global_mem, std::size_t size) const {
// FIXME_SYCL optimize
constexpr size_t wgroup_size = 32;
constexpr size_t wgroup_size = 128;
Copy link
Contributor Author

Choose a reason for hiding this comment

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

This change is not really necessary of course but is more in line with what we are doing for reductions at the moment.

@masterleinad masterleinad requested a review from nliber July 8, 2021 18:48
@crtrott crtrott added this to In progress in Kokkos Release 3.5 Jul 14, 2021
@crtrott crtrott moved this from In progress to Awaiting Feedback in Kokkos Release 3.5 Jul 14, 2021
@crtrott crtrott requested a review from Rombur July 14, 2021 18:53
@crtrott crtrott merged commit 5befc7b into kokkos:develop Jul 19, 2021
Kokkos Release 3.5 automation moved this from Awaiting Feedback to Done Jul 19, 2021
@masterleinad masterleinad moved this from Awaiting Feedback to Done in Developer: Daniel Arndt Jul 21, 2021
@masterleinad masterleinad deleted the sycl_subgroup_scan branch December 2, 2021 19:29
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Development

Successfully merging this pull request may close these issues.

None yet

4 participants