|
| 1 | += sycl_ext_oneapi_complex |
| 2 | + |
| 3 | +:source-highlighter: coderay |
| 4 | +:coderay-linenums-mode: table |
| 5 | + |
| 6 | +// This section needs to be after the document title. |
| 7 | +:doctype: book |
| 8 | +:toc2: |
| 9 | +:toc: left |
| 10 | +:encoding: utf-8 |
| 11 | +:lang: en |
| 12 | +:dpcpp: pass:[DPC++] |
| 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 | + |
| 20 | +== Notice |
| 21 | + |
| 22 | +[%hardbreaks] |
| 23 | +Copyright (C) 2022-2023 Codeplay Ltd. All rights reserved. |
| 24 | + |
| 25 | +Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks |
| 26 | +of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by |
| 27 | +permission by Khronos. |
| 28 | + |
| 29 | + |
| 30 | +== Contact |
| 31 | + |
| 32 | +To report problems with this extension, please open a new issue at: |
| 33 | + |
| 34 | +https://github.com/intel/llvm/issues |
| 35 | + |
| 36 | + |
| 37 | +== Dependencies |
| 38 | + |
| 39 | +This extension is written against the SYCL 2020 revision 5 specification. All |
| 40 | +references below to the "core SYCL specification" or to section numbers in the |
| 41 | +SYCL specification refer to that revision. |
| 42 | + |
| 43 | +== Status |
| 44 | + |
| 45 | +This is an experimental extension specification, intended to provide early |
| 46 | +access to features and gather community feedback. Interfaces defined in this |
| 47 | +specification are implemented in {dpcpp}, but they are not finalized and may |
| 48 | +change incompatibly in future versions of {dpcpp} without prior notice. |
| 49 | +*Shipping software products should not rely on APIs defined in this |
| 50 | +specification.* |
| 51 | + |
| 52 | +== Overview |
| 53 | + |
| 54 | +While {dpcpp} has support for `std::complex` in device code, it limits the |
| 55 | +complex interface and operations to the existing C++ standard. This proposal |
| 56 | +defines a SYCL complex extension based on but independent of the `std::complex` |
| 57 | +interface. This framework would allow for further development of complex math |
| 58 | +within oneAPI. Possible areas for deviation with `std::complex` include adding |
| 59 | +complex support for `marray` and `vec` and overloading mathematical |
| 60 | +functions to handle the element-wise operations. |
| 61 | + |
| 62 | +== Specification |
| 63 | + |
| 64 | +=== Feature test macro |
| 65 | + |
| 66 | +This extension provides a feature-test macro as described in the core SYCL |
| 67 | +specification. An implementation supporting this extension must predefine the |
| 68 | +macro `SYCL_EXT_ONEAPI_COMPLEX` to one of the values defined in the table |
| 69 | +below. Applications can test for the existence of this macro to determine if |
| 70 | +the implementation supports this feature, or applications can test the macro's |
| 71 | +value to determine which of the extension's features the implementation |
| 72 | +supports. |
| 73 | + |
| 74 | +[%header,cols="1,5"] |
| 75 | +|=== |
| 76 | +|Value |
| 77 | +|Description |
| 78 | + |
| 79 | +|1 |
| 80 | +|The APIs of this experimental extension are not versioned, so the feature-test macro always has this value. |
| 81 | +|=== |
| 82 | + |
| 83 | +=== Complex Class |
| 84 | + |
| 85 | +The core of this extension is the complex math class. This class contains a real |
| 86 | +and imaginary component and enables mathematical operations between complex |
| 87 | +numbers and decimals. The complex class interface and operations (shown below) |
| 88 | +are available in both host code and device code. Some operations, however, are |
| 89 | +available only in host code as noted below. |
| 90 | + |
| 91 | +The complex type is trivially copyable and type trait `is_device_copyable` |
| 92 | +should resolve to `std::true_type`. |
| 93 | + |
| 94 | +The `T` template parameter must be one of the types float, double, or |
| 95 | +sycl::half. |
| 96 | + |
| 97 | +Note: When performing operations between complex numbers and decimals, |
| 98 | +the decimal is treated as a complex number with a real component equal to |
| 99 | +the decimal and an imaginary component equal to 0. |
| 100 | + |
| 101 | +```C++ |
| 102 | +namespace sycl::ext::oneapi::experimental { |
| 103 | + |
| 104 | + template <typename T> |
| 105 | + class complex { |
| 106 | + public: |
| 107 | + using value_type = T; |
| 108 | + |
| 109 | + /// Constructs the complex number from real and imaginary parts. |
| 110 | + constexpr complex(value_type __re = value_type(), value_type __im = value_type()); |
| 111 | + |
| 112 | + /// Converting constructor. Constructs the object from a complex number of a different type. |
| 113 | + template <typename X> |
| 114 | + constexpr complex(const complex<X> &); |
| 115 | + |
| 116 | + /// Converting constructor. Constructs the object from a std::complex number. |
| 117 | + template <class X> |
| 118 | + constexpr complex(const std::complex<X> &); |
| 119 | + |
| 120 | + /// Constructs a std::complex number from a sycl::complex. |
| 121 | + template <class X> |
| 122 | + constexpr operator std::complex<X>() const; |
| 123 | + |
| 124 | + /// Returns the real part. |
| 125 | + constexpr value_type real() const; |
| 126 | + /// Returns the imaginary part. |
| 127 | + constexpr value_type imag() const; |
| 128 | + |
| 129 | + /// Sets the real part from value. |
| 130 | + void real(value_type value); |
| 131 | + /// Sets the imaginary part from value. |
| 132 | + void imag(value_type value); |
| 133 | + |
| 134 | + /// Assigns x to the real part of the complex number. Imaginary part is set to zero. |
| 135 | + complex<value_type> &operator=(value_type x); |
| 136 | + /// Adds and assigns real number y to complex number z. |
| 137 | + friend complex<value_type> &operator+=(complex<value_type> &z, value_type y); |
| 138 | + /// Subtracts and assigns real number y to complex number z. |
| 139 | + friend complex<value_type> &operator-=(complex<value_type> &z, value_type y); |
| 140 | + /// Multiplies and assigns real number y to complex number z. |
| 141 | + friend complex<value_type> &operator*=(complex<value_type> &z, value_type y); |
| 142 | + /// Divides and assigns real number y to complex number z. |
| 143 | + friend complex<value_type> &operator/=(complex<value_type> &z, value_type y); |
| 144 | + |
| 145 | + /// Assigns cx.real() and cx.imag() to the real and the imaginary parts of the complex number respectively. |
| 146 | + complex<value_type> &operator=(const complex<value_type> &cx); |
| 147 | + /// Adds and assigns complex number w to complex number z. |
| 148 | + template <class X> friend complex<value_type> &operator+=(complex<value_type> &z, const complex<X> &w); |
| 149 | + /// Subtracts and assigns complex number w to complex number z. |
| 150 | + template <class X> friend complex<value_type> &operator-=(complex<value_type> &z, const complex<X> &w); |
| 151 | + /// Multiplies and assigns complex number w to complex number z. |
| 152 | + template <class X> friend complex<value_type> &operator*=(complex<value_type> &z, const complex<X> &w); |
| 153 | + /// Divides and assigns complex number w to complex number z. |
| 154 | + template <class X> friend complex<value_type> &operator/=(complex<value_type> &z, const complex<X> &w); |
| 155 | + |
| 156 | + /// Adds complex numbers z and w and returns the value. |
| 157 | + friend complex<value_type> operator+(const complex<value_type> &z, const complex<value_type> &w); |
| 158 | + /// Adds complex number z and real y and returns the value. |
| 159 | + friend complex<value_type> operator+(const complex<value_type> &z, value_type y); |
| 160 | + /// Adds real x and complex number w and returns the value. |
| 161 | + friend complex<value_type> operator+(value_type x, const complex<value_type> &w); |
| 162 | + /// Returns the value of its argument. |
| 163 | + friend complex<value_type> operator+(const complex<value_type> &); |
| 164 | + |
| 165 | + /// Subtracts complex numbers z and w and returns the value. |
| 166 | + friend complex<value_type> operator-(const complex<value_type> &z, const complex<value_type> &w); |
| 167 | + /// Subtracts complex number z and real y and returns the value. |
| 168 | + friend complex<value_type> operator-(const complex<value_type> &z, value_type y); |
| 169 | + /// Subtracts real x and complex number w and returns the value. |
| 170 | + friend complex<value_type> operator-(value_type x, const complex<value_type> &w); |
| 171 | + /// Negates the argument. |
| 172 | + friend complex<value_type> operator-(const complex<value_type> &); |
| 173 | + |
| 174 | + /// Multiplies complex numbers z and w and returns the value. |
| 175 | + friend complex<value_type> operator*(const complex<value_type> &z, const complex<value_type> &w); |
| 176 | + /// Multiplies complex number z and real y and returns the value. |
| 177 | + friend complex<value_type> operator*(const complex<value_type> &z, value_type y); |
| 178 | + /// Multiplies real x and complex number w and returns the value. |
| 179 | + friend complex<value_type> operator*(value_type x, const complex<value_type> &w); |
| 180 | + |
| 181 | + /// Divides complex numbers z and w and returns the value. |
| 182 | + friend complex<value_type> operator/(const complex<value_type> &z, const complex<value_type> &w); |
| 183 | + /// Divides complex number z and real y and returns the value. |
| 184 | + friend complex<value_type> operator/(const complex<value_type> &z, value_type y); |
| 185 | + /// Divides real x and complex number w and returns the value. |
| 186 | + friend complex<value_type> operator/(value_type x, const complex<value_type> &w); |
| 187 | + |
| 188 | + /// Compares complex numbers z and w and returns true if they are the same, otherwise false. |
| 189 | + friend constexpr bool operator==(const complex<value_type> &z, const complex<value_type> &w); |
| 190 | + /// Compares complex number z and real y and returns true if they are the same, otherwise false. |
| 191 | + friend constexpr bool operator==(const complex<value_type> &z, value_type y); |
| 192 | + /// Compares real x and complex number w and returns true if they are the same, otherwise false. |
| 193 | + friend constexpr bool operator==(value_type x, const complex<value_type> &w); |
| 194 | + |
| 195 | + /// Compares complex numbers z and w and returns true if they are different, otherwise false. |
| 196 | + friend constexpr bool operator!=(const complex<value_type> &z, const complex<value_type> &w); |
| 197 | + ///Compares complex number z and real y and returns true if they are different, otherwise false. |
| 198 | + friend constexpr bool operator!=(const complex<value_type> &z, value_type y); |
| 199 | + /// Compares real x and complex number w and returns true if they are different, otherwise false. |
| 200 | + friend constexpr bool operator!=(value_type x, const complex<value_type> &w); |
| 201 | + |
| 202 | + /// Reads a complex number from is. |
| 203 | + /// Not allowed in device code. |
| 204 | + template <class C, class T> friend std::basic_istream<C, T> &operator>>(std::basic_istream<C, T> &is, complex<value_type> &); |
| 205 | + /// Writes to os the complex number z in the form (real,imaginary). |
| 206 | + /// Not allowed in device code. |
| 207 | + template <class C, class T> friend std::basic_ostream<C, T> &operator<<(std::basic_ostream<C, T> &os, const complex<value_type> &); |
| 208 | + /// Streams the complex number z in the format "(real,imaginary)" into `sycl::stream` x and return the result. |
| 209 | + friend const sycl::stream &operator<<(const sycl::stream &x, const complex<value_type> &z); |
| 210 | + |
| 211 | +} // namespace sycl::ext::oneapi::experimental |
| 212 | +``` |
| 213 | + |
| 214 | +=== Mathematical operations |
| 215 | + |
| 216 | +This proposal adds to the `sycl::ext::oneapi::experimental` namespace, math |
| 217 | +functions accepting the complex types `complex<sycl::half>`, `complex<float>`, |
| 218 | +`complex<double>` as well as the scalar types `sycl::half`, `float` and `double` |
| 219 | +for the SYCL math functions, `abs`, `acos`, `asin`, `atan`, `acosh`, `asinh`, |
| 220 | +`atanh`, `arg`, `conj`, `cos`, `cosh`, `exp`, `log`, `log10`, `norm`, `polar`, |
| 221 | +`pow`, `proj`, `sin`, `sinh`, `sqrt`, `tan`, and `tanh`. |
| 222 | + |
| 223 | +These functions are available in both host and device code, and each math |
| 224 | +function should follow the C++ standard for handling NaN's and Inf values. |
| 225 | + |
| 226 | +Note: In the case of the `pow` function, additional overloads have been added |
| 227 | +to ensure that for their first argument `base` and second argument `exponent`: |
| 228 | + |
| 229 | +* If `base` and/or `exponent` has type `complex<double>` or `double`, |
| 230 | + then `pow(base, exponent)` has the same effect as |
| 231 | + `pow(complex<double>(base), complex<double>(exponent))`. |
| 232 | + |
| 233 | +* Otherwise, if `base` and/or `exponent` has type `complex<float>` or `float`, |
| 234 | + then `pow(base, exponent)` has the same effect as |
| 235 | + `pow(complex<float>(base), complex<float>(exponent))`. |
| 236 | + |
| 237 | +* Otherwise, if `base` and/or `exponent` has type `complex<sycl::half>` or `sycl::half`, |
| 238 | + then `pow(base, exponent)` has the same effect as |
| 239 | + `pow(complex<sycl::half>(base), complex<sycl::half>(exponent))`. |
| 240 | + |
| 241 | +```C++ |
| 242 | +namespace sycl::ext::oneapi::experimental { |
| 243 | + |
| 244 | + /// VALUES: |
| 245 | + /// Returns the real component of the complex number z. |
| 246 | + template <class T> constexpr T real(const complex<T> &); |
| 247 | + /// Returns the real component of the number y, treated as complex numbers with zero imaginary component. |
| 248 | + template <class T> constexpr T real(T); |
| 249 | + /// Returns the imaginary component of the complex number z. |
| 250 | + template <class T> constexpr T imag(const complex<T> &); |
| 251 | + /// Returns the imaginary component of the number y, treated as complex numbers with zero imaginary component. |
| 252 | + template <class T> constexpr T imag(T); |
| 253 | + |
| 254 | + /// Compute the magnitude of complex number x. |
| 255 | + template <class T> T abs(const complex<T> &); |
| 256 | + /// Compute phase angle in radians of complex number x. |
| 257 | + template <class T> T arg(const complex<T> &); |
| 258 | + /// Compute phase angle in radians of complex number x, treated as complex number with positive zero imaginary component. |
| 259 | + template <class T> T arg(T); |
| 260 | + /// Compute the squared magnitude of complex number x. |
| 261 | + template <class T> T norm(const complex<T> &); |
| 262 | + /// Compute the squared magnitude of number x, treated as complex number with positive zero imaginary component. |
| 263 | + template <class T> T norm(T); |
| 264 | + /// Compute the conjugate of complex number x. |
| 265 | + template <class T> complex<T> conj(const complex<T> &); |
| 266 | + /// Compute the conjugate of number y, treated as complex number with positive zero imaginary component. |
| 267 | + template <class T> complex<T> conj(T); |
| 268 | + /// Compute the projection of complex number x. |
| 269 | + template <class T> complex<T> proj(const complex<T> &); |
| 270 | + /// Compute the projection of number y, treated as complex number with positive zero imaginary component. |
| 271 | + template <class T> complex<T> proj(T); |
| 272 | + /// Construct a complex number from polar coordinates with mangitude rho and angle theta. |
| 273 | + template <class T> complex<T> polar(const T &rho, const T &theta = T()); |
| 274 | + |
| 275 | + /// TRANSCENDENTALS: |
| 276 | + /// Compute the natural log of complex number x. |
| 277 | + template <class T> complex<T> log(const complex<T> &); |
| 278 | + /// Compute the base-10 log of complex number x. |
| 279 | + template <class T> complex<T> log10(const complex<T> &); |
| 280 | + /// Compute the square root of complex number x. |
| 281 | + template <class T> complex<T> sqrt(const complex<T> &); |
| 282 | + /// Compute the base-e exponent of complex number x. |
| 283 | + template <class T> complex<T> exp(const complex<T> &); |
| 284 | + |
| 285 | + /// Compute complex number z raised to the power of complex number y. |
| 286 | + template <class T> complex<T> pow(const complex<T> &, const complex<T> &); |
| 287 | + /// Compute complex number z raised to the power of complex number y. |
| 288 | + template <class T, class U> complex</*Promoted*/> pow(const complex<T> &, const complex<U> &); |
| 289 | + /// Compute complex number z raised to the power of real number y. |
| 290 | + template <class T, class U> complex</*Promoted*/> pow(const complex<T> &, const U &); |
| 291 | + /// Compute real number x raised to the power of complex number y. |
| 292 | + template <class T, class U> complex</*Promoted*/> pow(const T &, const complex<U> &); |
| 293 | + |
| 294 | + /// Compute the inverse hyperbolic sine of complex number x. |
| 295 | + template <class T> complex<T> asinh(const complex<T> &); |
| 296 | + /// Compute the inverse hyperbolic cosine of complex number x. |
| 297 | + template <class T> complex<T> acosh(const complex<T> &); |
| 298 | + /// Compute the inverse hyperbolic tangent of complex number x. |
| 299 | + template <class T> complex<T> atanh(const complex<T> &); |
| 300 | + /// Compute the hyperbolic sine of complex number x. |
| 301 | + template <class T> complex<T> sinh(const complex<T> &); |
| 302 | + /// Compute the hyperbolic cosine of complex number x. |
| 303 | + template <class T> complex<T> cosh(const complex<T> &); |
| 304 | + /// Compute the hyperbolic tangent of complex number x. |
| 305 | + template <class T> complex<T> tanh(const complex<T> &); |
| 306 | + /// Compute the inverse sine of complex number x. |
| 307 | + template <class T> complex<T> asin(const complex<T> &); |
| 308 | + /// Compute the inverse cosine of complex number x. |
| 309 | + template <class T> complex<T> acos(const complex<T> &); |
| 310 | + /// Compute the inverse tangent of complex number x. |
| 311 | + template <class T> complex<T> atan(const complex<T> &); |
| 312 | + /// Compute the sine of complex number x. |
| 313 | + template <class T> complex<T> sin(const complex<T> &); |
| 314 | + /// Compute the cosine of complex number x. |
| 315 | + template <class T> complex<T> cos(const complex<T> &); |
| 316 | + // Compute the tangent of complex number x. |
| 317 | + template <class T> complex<T> tan(const complex<T> &); |
| 318 | + |
| 319 | +} // namespace sycl::ext::oneapi::experimental |
| 320 | +``` |
| 321 | + |
| 322 | +== Implementation notes |
| 323 | + |
| 324 | +The complex mathematical operations can all be defined using SYCL built-ins. |
| 325 | +Therefore, implementing complex with SYCL built-ins would allow any backend |
| 326 | +with SYCL built-ins to support `sycl::ext::oneapi::experimental::complex`. |
| 327 | +The current implementation of `std::complex` relies on `libdevice`, which |
| 328 | +requires adjusting and altering the clang driver. This additional work would not |
| 329 | +be necessary for adding complex support with this extension. |
| 330 | + |
| 331 | +== Issues |
| 332 | + |
| 333 | +The motivation for adding this extension is to allow for complex support of |
| 334 | +`marray` and `vec`. This raises the issue of if this should be represented as |
| 335 | +an array of structs or a struct of arrays. The advantage of having an array |
| 336 | +of structs is that this is the most intuitive format for the user. As the |
| 337 | +user is likely thinking about the problem as a vector of complex numbers. |
| 338 | +However, this would cause the real and imaginary vectors to be non-contiguous. |
| 339 | +Conversely, having a struct of arrays would be less intuitive but would keep |
| 340 | +the vector's memory contiguous. |
0 commit comments