Skip to content

Commit 7b77c52

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web'
2 parents 12f2a1c + 7ae7ca8 commit 7b77c52

File tree

8 files changed

+435
-68
lines changed

8 files changed

+435
-68
lines changed

‎sycl/doc/extensions/ComplexAlgorithms/ComplexAlgorithms.asciidoc

Lines changed: 129 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,129 @@
1+
= SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS
2+
:source-highlighter: coderay
3+
:coderay-linenums-mode: table
4+
5+
// This section needs to be after the document title.
6+
:doctype: book
7+
:toc2:
8+
:toc: left
9+
:encoding: utf-8
10+
:lang: en
11+
12+
:blank: pass:[ +]
13+
14+
// Set the default source code type in this document to C++,
15+
// for syntax highlighting purposes. This is needed because
16+
// docbook uses c++ and html5 uses cpp.
17+
:language: {basebackend@docbook:c++:cpp}
18+
19+
== Introduction
20+
IMPORTANT: This specification is a draft.
21+
22+
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.
23+
24+
This extension adds limited support for `std::complex<float>` and
25+
`std::complex<double>` to several SYCL group functions and algorithms.
26+
27+
== Notice
28+
29+
Copyright (c) 2021-2022 Intel Corporation. All rights reserved.
30+
31+
== Status
32+
33+
Working Draft
34+
35+
This is a proposed extension specification, intended to gather community
36+
feedback. Interfaces defined in this specification may not be implemented yet
37+
or may be in a preliminary state. The specification itself may also change in
38+
incompatible ways before it is finalized. Shipping software products should not
39+
rely on APIs defined in this specification.
40+
41+
== Version
42+
43+
Revision: 1
44+
45+
== Contacts
46+
47+
John Pennycook, Intel (john 'dot' pennycook 'at' intel 'dot' com)
48+
49+
== Dependencies
50+
51+
This extension is written against the SYCL 2020 specification, Revision 4.
52+
53+
== Feature Test Macro
54+
55+
This extension provides a feature-test macro as described in the core SYCL
56+
specification section 6.3.3 "Feature test macros". Therefore, an
57+
implementation supporting this extension must predefine the macro
58+
`SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS` to one of the values defined in the table
59+
below. Applications can test for the existence of this macro to determine if
60+
the implementation supports this feature, or applications can test the macro's
61+
value to determine which of the extension's APIs the implementation supports.
62+
63+
[%header,cols="1,5"]
64+
|===
65+
|Value |Description
66+
|1 |Initial extension version. Base features are supported.
67+
|===
68+
69+
== Overview
70+
71+
The types supported by some group functions and algorithms in SYCL 2020 are
72+
restricted to built-in scalar types and SYCL vector types. This extension
73+
relaxes these restrictions to permit `std::complex<float>` and
74+
`std::complex<double>` types.
75+
76+
Note that the following group functions and algorithms already accept
77+
`std::complex<float>` arguments and `std::complex<double>` because they
78+
are trivially copyable:
79+
80+
- `group_broadcast`
81+
- `group_shift_left`
82+
- `group_shift_right`
83+
- `permute_group_by_xor`
84+
- `select_from_group`
85+
86+
Usage of `std::complex<double>` requires support for double precision,
87+
and specifically the `sycl::aspect::fp64` device aspect.
88+
89+
== Extended Functions and Algorithms
90+
91+
The following group functions and algorithms accept `std::complex<float>`
92+
and `std::complex<double>` arguments only when used with a `BinaryOperation`
93+
argument of `sycl::plus`:
94+
95+
- `joint_reduce`
96+
- `reduce_over_group`
97+
- `joint_exclusive_scan`
98+
- `exclusive_scan_over_group`
99+
- `joint_inclusive_scan`
100+
- `inclusive_scan_over_group`
101+
102+
NOTE: `sycl::multiplies` is not currently supported because it cannot be
103+
implemented as an element-wise operation on the real and imaginary components.
104+
This restriction may be lifted in a future version of the extension.
105+
106+
NOTE: `sycl::bit_and`, `sycl::bit_or`, `sycl::bit_xor`, `sycl::logical_and`,
107+
`sycl::logical_or`, `sycl::minimum` and `sycl::maximum` are not supported
108+
because their behaviors are defined in terms of operators that `std::complex`
109+
does not implement.
110+
111+
== Issues
112+
113+
None.
114+
115+
//. asd
116+
//+
117+
//--
118+
//*RESOLUTION*: Not resolved.
119+
//--
120+
121+
== Revision History
122+
123+
[cols="5,15,15,70"]
124+
[grid="rows"]
125+
[options="header"]
126+
|========================================
127+
|Rev|Date|Author|Changes
128+
|1|2021-12-08|John Pennycook|*Initial public working draft*
129+
|========================================

‎sycl/include/CL/sycl/accessor.hpp

Lines changed: 50 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -1688,64 +1688,64 @@ class __SYCL_SPECIAL_CLASS accessor :
16881688

16891689
template <int Dims = Dimensions>
16901690
operator typename detail::enable_if_t<Dims == 0 &&
1691-
#ifdef __ENABLE_USM_ADDR_SPACE__
16921691
AccessMode == access::mode::atomic,
1693-
atomic<DataT>>() const {
1692+
#ifdef __ENABLE_USM_ADDR_SPACE__
1693+
atomic<DataT>
16941694
#else
1695-
AccessMode == access::mode::atomic,
1696-
atomic<DataT, AS>>() const {
1695+
atomic<DataT, AS>
16971696
#endif
1698-
const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1699-
return atomic<DataT, AS>(
1700-
multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1701-
}
1697+
>() const {
1698+
const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1699+
return atomic<DataT, AS>(
1700+
multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1701+
}
17021702

1703-
template <int Dims = Dimensions>
1704-
typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
1705-
atomic<DataT, AS>>
1706-
operator[](id<Dimensions> Index) const {
1707-
const size_t LinearIndex = getLinearIndex(Index);
1708-
return atomic<DataT, AS>(
1709-
multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1710-
}
1703+
template <int Dims = Dimensions>
1704+
typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
1705+
atomic<DataT, AS>>
1706+
operator[](id<Dimensions> Index) const {
1707+
const size_t LinearIndex = getLinearIndex(Index);
1708+
return atomic<DataT, AS>(
1709+
multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1710+
}
17111711

1712-
template <int Dims = Dimensions>
1713-
typename detail::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
1714-
atomic<DataT, AS>>
1715-
operator[](size_t Index) const {
1716-
const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
1717-
return atomic<DataT, AS>(
1718-
multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1719-
}
1720-
template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
1721-
typename AccessorCommonT::template AccessorSubscript<Dims - 1>
1722-
operator[](size_t Index) const {
1723-
return AccessorSubscript<Dims - 1>(*this, Index);
1724-
}
1712+
template <int Dims = Dimensions>
1713+
typename detail::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
1714+
atomic<DataT, AS>>
1715+
operator[](size_t Index) const {
1716+
const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
1717+
return atomic<DataT, AS>(
1718+
multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1719+
}
1720+
template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
1721+
typename AccessorCommonT::template AccessorSubscript<Dims - 1>
1722+
operator[](size_t Index) const {
1723+
return AccessorSubscript<Dims - 1>(*this, Index);
1724+
}
17251725

1726-
template <access::target AccessTarget_ = AccessTarget,
1727-
typename =
1728-
detail::enable_if_t<AccessTarget_ == access::target::host_buffer>>
1729-
DataT *get_pointer() const {
1730-
return getPointerAdjusted();
1731-
}
1726+
template <access::target AccessTarget_ = AccessTarget,
1727+
typename = detail::enable_if_t<AccessTarget_ ==
1728+
access::target::host_buffer>>
1729+
DataT *get_pointer() const {
1730+
return getPointerAdjusted();
1731+
}
17321732

1733-
template <
1734-
access::target AccessTarget_ = AccessTarget,
1735-
typename = detail::enable_if_t<AccessTarget_ == access::target::device>>
1736-
global_ptr<DataT> get_pointer() const {
1737-
return global_ptr<DataT>(getPointerAdjusted());
1738-
}
1733+
template <
1734+
access::target AccessTarget_ = AccessTarget,
1735+
typename = detail::enable_if_t<AccessTarget_ == access::target::device>>
1736+
global_ptr<DataT> get_pointer() const {
1737+
return global_ptr<DataT>(getPointerAdjusted());
1738+
}
17391739

1740-
template <access::target AccessTarget_ = AccessTarget,
1741-
typename = detail::enable_if_t<AccessTarget_ ==
1742-
access::target::constant_buffer>>
1743-
constant_ptr<DataT> get_pointer() const {
1744-
return constant_ptr<DataT>(getPointerAdjusted());
1745-
}
1740+
template <access::target AccessTarget_ = AccessTarget,
1741+
typename = detail::enable_if_t<AccessTarget_ ==
1742+
access::target::constant_buffer>>
1743+
constant_ptr<DataT> get_pointer() const {
1744+
return constant_ptr<DataT>(getPointerAdjusted());
1745+
}
17461746

1747-
bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
1748-
bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
1747+
bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
1748+
bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
17491749

17501750
private:
17511751
// supporting function for get_pointer()
@@ -1768,7 +1768,7 @@ bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
17681768
"buffer size must be greater than zero.",
17691769
PI_INVALID_VALUE);
17701770
}
1771-
}; // namespace sycl
1771+
};
17721772

17731773
#if __cplusplus > 201402L
17741774

0 commit comments

Comments
 (0)