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

[SYCL] Implement reduction properties extension #15804

Merged
merged 16 commits into from
Oct 29, 2024

Conversation

Pennycook
Copy link
Contributor

@Pennycook Pennycook commented Oct 22, 2024

Adds support for initialize_to_identity and deterministic properties.

Since this extension is only experimental, the implementation here avoids making significant changes to reduction-related classes (e.g, reducer).

A more straightforward implementation that attaches a compile-time property list to these classes is possible, but may be considered an ABI break.


A note to reviewers: it occurred to me that IsDeterministicOperator<BinaryOperation> might actually be better as IsDeterministicOperator<T, BinaryOperation>, so that the implementation can infer that it's safe to use the faster reductions for certain types (e.g., int and sycl::plus<>() can safely use atomics) but I wanted to get your feedback on the proposed approach before going any further.

EDIT: To clarify, I think such changes to IsDeterministicOperator could be left until a future PR, since the behavior of the deterministic property as implemented here should match the legacy SYCL_DETERMINISTIC_REDUCTION macro. We could implement optimizations later.

Adds support for initialize_to_identity and deterministic properties.

Since this extension is only experimental, the implementation here avoids
making significant changes to reduction-related classes (e.g, reducer).
A more straightforward implementation that attaches a compile-time property
list to these classes is possible, but may be considered an ABI break.

Signed-off-by: John Pennycook <[email protected]>
Also fixes a few typos in the specification that were uncovered during
implementation of the extension.

Signed-off-by: John Pennycook <[email protected]>
@Pennycook Pennycook added the spec extension All issues/PRs related to extensions specifications label Oct 22, 2024
Signed-off-by: John Pennycook <[email protected]>
Signed-off-by: John Pennycook <[email protected]>
@aelovikov-intel
Copy link
Contributor

Any chance we can be "smart" and change

template <typename T, typename BinaryOperation>
auto reduction(T *Var, const T &Identity, BinaryOperation Combiner,
               const property_list &PropList = {}) 

to

template <typename T, typename BinaryOperation, typename PropertyListTy = property_list>
auto reduction(T *Var, const T &Identity, BinaryOperation Combiner,
               const PropertyListTy &PropList = {}) 

instead of adding a bunch of new overloads?

Bug was hidden during development due to testing on the CPU.

Signed-off-by: John Pennycook <[email protected]>
@Pennycook
Copy link
Contributor Author

Any chance we can be "smart" and change

template <typename T, typename BinaryOperation>
auto reduction(T *Var, const T &Identity, BinaryOperation Combiner,
               const property_list &PropList = {}) 

to

template <typename T, typename BinaryOperation, typename PropertyListTy = property_list>
auto reduction(T *Var, const T &Identity, BinaryOperation Combiner,
               const PropertyListTy &PropList = {}) 

instead of adding a bunch of new overloads?

I'd rather keep the overloads for now, if that's okay. Because this extension is only experimental, I don't want us to make changes to the "real" reduction interface if we don't have to. Plus, if we're successful with getting compile-time properties into a future version of the SYCL specification, I expect we'll want to generate a deprecation notice for the property_list overloads, and I can't think of a way to do that without splitting things out.

@Pennycook Pennycook marked this pull request as ready for review October 23, 2024 12:27
@Pennycook Pennycook requested review from a team as code owners October 23, 2024 12:27
@Pennycook
Copy link
Contributor Author

@gmlueck - Could you please take a quick look at the extension specification changes here? They're just typo fixes so it hopefully shouldn't take too long to review.

@Pennycook
Copy link
Contributor Author

@intel/llvm-gatekeepers, I think this can be merged now.

@sarnex sarnex merged commit 22e5ced into intel:sycl Oct 29, 2024
14 checks passed
YixingZhang007 pushed a commit to YixingZhang007/llvm that referenced this pull request Oct 30, 2024
Adds support for initialize_to_identity and deterministic properties.

Since this extension is only experimental, the implementation here
avoids making significant changes to reduction-related classes (e.g,
reducer).

A more straightforward implementation that attaches a compile-time
property list to these classes is possible, but may be considered an ABI
break.

---

A note to reviewers: it occurred to me that
`IsDeterministicOperator<BinaryOperation>` might actually be better as
`IsDeterministicOperator<T, BinaryOperation>`, so that the
implementation can infer that it's safe to use the faster reductions for
certain types (e.g., `int` and `sycl::plus<>()` can safely use atomics)
but I wanted to get your feedback on the proposed approach before going
any further.

**EDIT**: To clarify, I think such changes to `IsDeterministicOperator`
could be left until a future PR, since the behavior of the
`deterministic` property as implemented here should match the legacy
`SYCL_DETERMINISTIC_REDUCTION` macro. We could implement optimizations
later.

---------

Signed-off-by: John Pennycook <[email protected]>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
spec extension All issues/PRs related to extensions specifications
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants