[SYCL][COMPAT] Added Complex muladd to syclcompat (#12969)

Adds the documentation, functionality and tests for complex
multiplication addition to SYCLcompat.

Author: Alcpz

sycl/doc/syclcompat/README.md

@@ -1229,7 +1229,7 @@ as a vector of elements, and returning `0` for vector components for which
 `vectorized_sum_abs_diff` calculates the absolute difference for two values
 without modulo overflow for vector types.
 
-The functions `cmul`,`cdiv`,`cabs`, and `conj` define complex math operations
+The functions `cmul`,`cdiv`,`cabs`, `cmul_add`, and `conj` define complex math operations
 which accept `sycl::vec<T,2>` arguments representing complex values.
 
 ```cpp
@@ -1259,6 +1259,16 @@ sycl::vec<T, 2> cdiv(sycl::vec<T, 2> x, sycl::vec<T, 2> y);
 
 template <typename T> T cabs(sycl::vec<T, 2> x);
 
+template <typename ValueT>
+inline sycl::vec<ValueT, 2> cmul_add(const sycl::vec<ValueT, 2> a,
+                                     const sycl::vec<ValueT, 2> b,
+                                     const sycl::vec<ValueT, 2> c);
+
+template <typename ValueT>
+inline sycl::marray<ValueT, 2> cmul_add(const sycl::marray<ValueT, 2> a,
+                                        const sycl::marray<ValueT, 2> b,
+                                        const sycl::marray<ValueT, 2> c);
+
 template <typename T> sycl::vec<T, 2> conj(sycl::vec<T, 2> x);
 
 template <typename ValueT> inline ValueT reverse_bits(ValueT a);

sycl/include/syclcompat/math.hpp

@@ -251,6 +251,32 @@ template <typename T> sycl::vec<T, 2> conj(sycl::vec<T, 2> x) {
   return sycl::vec<T, 2>(t.real(), t.imag());
 }
 
+/// Performs complex number multiply addition.
+/// \param [in] a The first value
+/// \param [in] b The second value
+/// \param [in] c The third value
+/// \returns the operation result
+template <typename ValueT>
+inline sycl::vec<ValueT, 2> cmul_add(const sycl::vec<ValueT, 2> a,
+                                     const sycl::vec<ValueT, 2> b,
+                                     const sycl::vec<ValueT, 2> c) {
+  sycl::ext::oneapi::experimental::complex<ValueT> t(a[0], a[1]);
+  sycl::ext::oneapi::experimental::complex<ValueT> u(b[0], b[1]);
+  sycl::ext::oneapi::experimental::complex<ValueT> v(c[0], c[1]);
+  t = t * u + v;
+  return sycl::vec<ValueT, 2>{t.real(), t.imag()};
+}
+template <typename ValueT>
+inline sycl::marray<ValueT, 2> cmul_add(const sycl::marray<ValueT, 2> a,
+                                        const sycl::marray<ValueT, 2> b,
+                                        const sycl::marray<ValueT, 2> c) {
+  sycl::ext::oneapi::experimental::complex<ValueT> t(a[0], a[1]);
+  sycl::ext::oneapi::experimental::complex<ValueT> u(b[0], b[1]);
+  sycl::ext::oneapi::experimental::complex<ValueT> v(c[0], c[1]);
+  t = t * u + v;
+  return sycl::marray<ValueT, 2>{t.real(), t.imag()};
+}
+
 /// A sycl::abs wrapper functors.
 struct abs {
   template <typename ValueT> auto operator()(const ValueT x) const {

sycl/test-e2e/syclcompat/math/math_complex.cpp

@@ -174,6 +174,46 @@ void kernel_mul(int *result) {
   *result = r;
 }
 
+void kernel_mul_add(int *result) {
+  sycl::double2 d1, d2, d3;
+  sycl::float2 f1, f2, f3;
+  sycl::marray<double, 2> m_d1, m_d2, m_d3;
+  sycl::marray<float, 2> m_f1, m_f2, m_f3;
+
+  d1 = sycl::double2(5.4, -6.3);
+  d2 = sycl::double2(-7.2, 8.1);
+  d3 = sycl::double2(1.0, -1.0);
+
+  f1 = sycl::float2(1.8, -2.7);
+  f2 = sycl::float2(-3.6, 4.5);
+  f3 = sycl::float2(1.0, -1.0);
+
+  bool r = true;
+  float expect[4] = {13.150000, 88.100000, 6.670001, 16.820000};
+
+  auto a1 = syclcompat::cmul_add(d1, d2, d3);
+  r = r && check(a1, expect);
+
+  auto a2 = syclcompat::cmul_add(f1, f2, f3);
+  r = r && check(a2, expect + 2);
+
+  m_d1 = sycl::marray<double, 2>(5.4, -6.3);
+  m_d2 = sycl::marray<double, 2>(-7.2, 8.1);
+  m_d3 = sycl::marray<double, 2>(1.0, -1.0);
+
+  m_f1 = sycl::marray<float, 2>(1.8, -2.7);
+  m_f2 = sycl::marray<float, 2>(-3.6, 4.5);
+  m_f3 = sycl::marray<float, 2>(1.0, -1.0);
+
+  auto a3 = syclcompat::cmul_add(d1, d2, d3);
+  r = r && check(a3, expect);
+
+  auto a4 = syclcompat::cmul_add(f1, f2, f3);
+  r = r && check(a4, expect + 2);
+
+  *result = r;
+}
+
 void test_abs() {
   std::cout << __PRETTY_FUNCTION__ << std::endl;
   ComplexLauncher<kernel_abs>().launch();
@@ -191,11 +231,17 @@ void test_conj() {
   ComplexLauncher<kernel_conj>().launch();
 }
 
+void test_mul_add() {
+  std::cout << __PRETTY_FUNCTION__ << std::endl;
+  ComplexLauncher<kernel_mul_add>().launch();
+}
+
 int main() {
   test_abs();
   test_mul();
   test_div();
   test_conj();
+  test_mul_add();
 
   return 0;
 }