Skip to content

[SYCL][Doc] Add sycl complex to complex algorithms extension #6717

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

Merged
merged 2 commits into from
Sep 9, 2022
Merged

[SYCL][Doc] Add sycl complex to complex algorithms extension #6717

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
e63601e
Modify complex algorithms extension to support sycl complex
Jul 22, 2022
d137d40
Update based on feedback
Sep 8, 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
48 changes: 33 additions & 15 deletions sycl/doc/extensions/proposed/sycl_ext_oneapi_complex_algorithms.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,8 +21,10 @@ IMPORTANT: This specification is a draft.

NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by permission by Khronos.

This extension adds limited support for `std::complex<float>` and
`std::complex<double>` to several SYCL group functions and algorithms.
This extension adds SYCL group functions and algorithms to complex types in device code.
This includes `std::complex<float>`, `std::complex<double>`,
`sycl::complex<sycl::half>`, `sycl::complex<float>`, `sycl::complex<double>`
and their `marray` equivalents.

== Notice

Expand Down Expand Up @@ -50,6 +52,15 @@ John Pennycook, Intel (john 'dot' pennycook 'at' intel 'dot' com)

This extension is written against the SYCL 2020 specification, Revision 4.

This extension builds on the `sycl::ext::oneapi::complex` and
`marray<sycl::ext::oneapi::complex>` classes. Therefore this extension is
dependent on link:sycl_ext_oneapi_complex.asciidoc[sycl_ext_oneapi_complex]
and
link:sycl_ext_oneapi_complex_marray.asciidoc[sycl_ext_oneapi_complex_marray].
Comment on lines +55 to +59
Copy link
Contributor

Choose a reason for hiding this comment

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

Prior to this change, developers could use the std::complex features without requiring those other two extensions. Does anything in our std::complex support now depend on these new extensions? If not, I think we should be explicit that only those parts of this extension which use sycl::complex or marray are dependent.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Yes, nothing in std::complex is dependent on these new extensions. I have added a bit stating that all function overloads that use std::complex are not dependent upon these extensions.

Function overloads which use `std::complex` are not dependent upon the
above extensions.


== Feature Test Macro

This extension provides a feature-test macro as described in the core SYCL
Expand All @@ -64,33 +75,43 @@ value to determine which of the extension's APIs the implementation supports.
|===
|Value |Description
|1 |Initial extension version. Base features are supported.
|2 |Proposal is extended to other complex extensions. Multiplication binary operation is added.
|===

== Overview

To reduce repetition, within this extension, `gencomplex` will be redefined to
include `std::complex` floating point types. Therefore, `gencomplex` is defined
as types `std::complex<float>`, `std::complex<double>`,
`sycl::complex<sycl::half>`, `sycl::complex<float>`, and
`sycl::complex<double>`. The extension will also refer to `mgencomplex` which
are types `sycl::marray<sycl::complex<sycl::half>>`,
`sycl::marray<sycl::complex<float>>`, and `sycl::marray<sycl::complex<double>>`.

The types supported by some group functions and algorithms in SYCL 2020 are
restricted to built-in scalar types and SYCL vector types. This extension
relaxes these restrictions to permit `std::complex<float>` and
`std::complex<double>` types.
relaxes these restrictions to permit `gencomplex` and `mgencomplex` types.

Note that the following group functions and algorithms already accept
`std::complex<float>` arguments and `std::complex<double>` because they
are trivially copyable:
`gencomplex` arguments and `mgencomplex` because they are trivially copyable:

- `group_broadcast`
- `group_shift_left`
- `group_shift_right`
- `permute_group_by_xor`
- `select_from_group`

Usage of `std::complex<double>` requires support for double precision,
and specifically the `sycl::aspect::fp64` device aspect.
Each function is only supported if the underlying value type is also supported.
For example, `group_broadcast` only supports
`sycl::ext::oneapi::complex<sycl::half>` and
`marray<sycl::ext::oneapi::complex<sycl::half>>` if the device supports
`sycl::aspect::fp16`.

== Extended Functions and Algorithms

The following group functions and algorithms accept `std::complex<float>`
and `std::complex<double>` arguments only when used with a `BinaryOperation`
argument of `sycl::plus`:
The following group functions and algorithms accept `gencomplex` arguments
only when used with a `BinaryOperation` argument of `sycl::plus` and
`sycl::multiplies`. These operations do not support `mgencomplex`:
Comment on lines +112 to +114
Copy link
Contributor

Choose a reason for hiding this comment

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

This change seems unrelated to the sycl::complex and sycl::marray parts of things. Can you say more about why you're lifting this restriction, and what SPIR-V instruction(s) you expect it to generate?

Copy link
Contributor Author

@AidanBeltonS AidanBeltonS Sep 8, 2022 •
edited
Loading

Choose a reason for hiding this comment

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

Yes this change is somewhat unrelated to sycl::complex and sycl::marray<complex>, this restriction is being lifted to allow for a more fleshed out version of this extension. The multiplies operation is a valid binary operation for reductions and scans so we would like to support it. We currently do not have a valid SPIR-V instruction for the multiplies operation, so we would be adding a new one to support this op. I think we would name the instruction something along the lines of __spirv_GroupCMulKHR


- `joint_reduce`
- `reduce_over_group`
Expand All @@ -99,10 +120,6 @@ argument of `sycl::plus`:
- `joint_inclusive_scan`
- `inclusive_scan_over_group`

NOTE: `sycl::multiplies` is not currently supported because it cannot be
implemented as an element-wise operation on the real and imaginary components.
This restriction may be lifted in a future version of the extension.

NOTE: `sycl::bit_and`, `sycl::bit_or`, `sycl::bit_xor`, `sycl::logical_and`,
`sycl::logical_or`, `sycl::minimum` and `sycl::maximum` are not supported
because their behaviors are defined in terms of operators that `std::complex`
Expand All @@ -126,4 +143,5 @@ None.
|========================================
|Rev|Date|Author|Changes
|1|2021-12-08|John Pennycook|*Initial public working draft*
|2|2022-31-08|Aidan Belton|*Extend to support sycl complex extension and multiplication binary operation*
|========================================