|
| 1 | +#include <cmath> |
| 2 | +#include <iomanip> |
| 3 | + |
| 4 | +#define SYCL_EXT_ONEAPI_COMPLEX |
| 5 | +#include <sycl/ext/oneapi/experimental/sycl_complex.hpp> |
| 6 | +#include <sycl/sycl.hpp> |
| 7 | + |
| 8 | +using namespace sycl::ext::oneapi; |
| 9 | + |
| 10 | +#define SYCL_CPLX_TOL_ULP 5 |
| 11 | + |
| 12 | +#define PI 3.14159265358979323846 |
| 13 | + |
| 14 | +constexpr double INFINITYd(std::numeric_limits<double>::infinity()); |
| 15 | +constexpr double NANd(std::numeric_limits<double>::quiet_NaN()); |
| 16 | + |
| 17 | +template <typename T = double> struct cmplx { |
| 18 | + cmplx(T real, T imag) : re(real), im(imag) {} |
| 19 | + |
| 20 | + template <typename X> cmplx(cmplx<X> c) { |
| 21 | + re = c.re; |
| 22 | + im = c.im; |
| 23 | + } |
| 24 | + |
| 25 | + T re; |
| 26 | + T im; |
| 27 | +}; |
| 28 | + |
| 29 | +// Helpers for displaying results |
| 30 | + |
| 31 | +template <typename T> const char *get_typename() { return "Unknown type"; } |
| 32 | +template <> const char *get_typename<double>() { return "double"; } |
| 33 | +template <> const char *get_typename<float>() { return "float"; } |
| 34 | +template <> const char *get_typename<sycl::half>() { return "sycl::half"; } |
| 35 | + |
| 36 | +// Helper to test each complex specilization |
| 37 | +template <template <typename> typename action, typename... argsT> |
| 38 | +bool test_valid_types(sycl::queue &Q, argsT... args) { |
| 39 | + bool test_passes = true; |
| 40 | + |
| 41 | + if (Q.get_device().has(sycl::aspect::fp64)) { |
| 42 | + action<double> test; |
| 43 | + test_passes &= test(Q, args...); |
| 44 | + } |
| 45 | + |
| 46 | + { |
| 47 | + action<float> test; |
| 48 | + test_passes &= test(Q, args...); |
| 49 | + } |
| 50 | + |
| 51 | + if (Q.get_device().has(sycl::aspect::fp16)) { |
| 52 | + action<sycl::half> test; |
| 53 | + test_passes &= test(Q, args...); |
| 54 | + } |
| 55 | + |
| 56 | + return test_passes; |
| 57 | +} |
| 58 | + |
| 59 | +// Overload for host only tests |
| 60 | +template <template <typename> typename action, typename... argsT> |
| 61 | +bool test_valid_types(argsT... args) { |
| 62 | + bool test_passes = true; |
| 63 | + |
| 64 | + { |
| 65 | + action<double> test; |
| 66 | + test_passes &= test(args...); |
| 67 | + } |
| 68 | + |
| 69 | + { |
| 70 | + action<float> test; |
| 71 | + test_passes &= test(args...); |
| 72 | + } |
| 73 | + |
| 74 | + { |
| 75 | + action<sycl::half> test; |
| 76 | + test_passes &= test(args...); |
| 77 | + } |
| 78 | + |
| 79 | + return test_passes; |
| 80 | +} |
| 81 | + |
| 82 | +// Helpers for comparison |
| 83 | + |
| 84 | +template <typename T> bool almost_equal_scalar(T x, T y, int ulp) { |
| 85 | + if (std::isnan(x) && std::isnan(y)) |
| 86 | + return true; |
| 87 | + |
| 88 | + if (std::isinf(x) && std::isinf(y)) |
| 89 | + return true; |
| 90 | + |
| 91 | + return std::abs(x - y) <= |
| 92 | + std::numeric_limits<T>::epsilon() * std::abs(x + y) * ulp || |
| 93 | + std::abs(x - y) < std::numeric_limits<T>::min(); |
| 94 | +} |
| 95 | + |
| 96 | +template <typename T> T complex_magnitude(std::complex<T> a) { |
| 97 | + return std::sqrt(a.real() * a.real() + a.imag() * a.imag()); |
| 98 | +} |
| 99 | + |
| 100 | +template <typename T> T complex_magnitude(experimental::complex<T> a) { |
| 101 | + return std::sqrt(a.real() * a.real() + a.imag() * a.imag()); |
| 102 | +} |
| 103 | + |
| 104 | +template <typename T, typename U> inline bool is_nan_or_inf(T x, U y) { |
| 105 | + return (std::isnan(x.real()) && std::isnan(y.real())) || |
| 106 | + (std::isnan(x.imag()) && std::isnan(y.imag())) || |
| 107 | + (std::isinf(x.real()) && std::isinf(y.real())) || |
| 108 | + (std::isinf(x.imag()) && std::isinf(y.imag())); |
| 109 | +} |
| 110 | + |
| 111 | +template <typename T> |
| 112 | +bool almost_equal_cplx(experimental::complex<T> x, experimental::complex<T> y, |
| 113 | + int ulp) { |
| 114 | + auto diff = complex_magnitude(x - y); |
| 115 | + return diff <= std::numeric_limits<T>::epsilon() * std::abs(x + y) * ulp || |
| 116 | + diff < std::numeric_limits<T>::min() || is_nan_or_inf(x, y); |
| 117 | +} |
| 118 | + |
| 119 | +template <typename T> |
| 120 | +bool almost_equal_cplx(experimental::complex<T> x, std::complex<T> y, int ulp) { |
| 121 | + auto stdx = std::complex{x.real(), x.imag()}; |
| 122 | + auto diff = complex_magnitude(stdx - y); |
| 123 | + return diff <= std::numeric_limits<T>::epsilon() * std::abs(stdx + y) * ulp || |
| 124 | + diff < std::numeric_limits<T>::min() || is_nan_or_inf(x, y); |
| 125 | +} |
| 126 | + |
| 127 | +template <typename T> |
| 128 | +bool almost_equal_cplx(std::complex<T> x, experimental::complex<T> y, int ulp) { |
| 129 | + auto stdy = std::complex{y.real(), y.imag()}; |
| 130 | + auto diff = complex_magnitude(x - stdy); |
| 131 | + return diff <= std::numeric_limits<T>::epsilon() * std::abs(x + stdy) * ulp || |
| 132 | + diff < std::numeric_limits<T>::min() || is_nan_or_inf(x, y); |
| 133 | +} |
| 134 | + |
| 135 | +template <typename T> |
| 136 | +bool almost_equal_cplx(std::complex<T> x, std::complex<T> y, int ulp) { |
| 137 | + auto diff = complex_magnitude(x - y); |
| 138 | + return diff <= std::numeric_limits<T>::epsilon() * std::abs(x + y) * ulp || |
| 139 | + diff < std::numeric_limits<T>::min() || is_nan_or_inf(x, y); |
| 140 | +} |
| 141 | + |
| 142 | +// Helpers for testing half |
| 143 | + |
| 144 | +std::complex<float> sycl_half_to_float(experimental::complex<sycl::half> c) { |
| 145 | + auto c_sycl_float = static_cast<experimental::complex<float>>(c); |
| 146 | + return static_cast<std::complex<float>>(c_sycl_float); |
| 147 | +} |
| 148 | + |
| 149 | +std::complex<sycl::half> sycl_float_to_half(std::complex<float> c) { |
| 150 | + auto c_sycl_half = static_cast<experimental::complex<sycl::half>>(c); |
| 151 | + return static_cast<std::complex<sycl::half>>(c_sycl_half); |
| 152 | +} |
| 153 | + |
| 154 | +std::complex<float> trunc_float(std::complex<float> c) { |
| 155 | + auto c_sycl_half = static_cast<experimental::complex<sycl::half>>(c); |
| 156 | + return sycl_half_to_float(c_sycl_half); |
| 157 | +} |
| 158 | + |
| 159 | +// Helper for initializing std::complex values for tests only needed because |
| 160 | +// sycl::half cases are emulated with float for std::complex class |
| 161 | + |
| 162 | +template <typename T_in> auto constexpr init_std_complex(T_in re, T_in im) { |
| 163 | + return std::complex<T_in>(re, im); |
| 164 | +} |
| 165 | + |
| 166 | +template <> auto constexpr init_std_complex(sycl::half re, sycl::half im) { |
| 167 | + return trunc_float(std::complex<float>(re, im)); |
| 168 | +} |
| 169 | + |
| 170 | +template <typename T_in> auto constexpr init_deci(T_in re) { return re; } |
| 171 | + |
| 172 | +template <> auto constexpr init_deci(sycl::half re) { |
| 173 | + return static_cast<float>(re); |
| 174 | +} |
| 175 | + |
| 176 | +// Helper for comparing SyclCPLX and standard c++ results |
| 177 | + |
| 178 | +template <typename T> |
| 179 | +bool check_results(experimental::complex<T> output, std::complex<T> reference, |
| 180 | + bool is_device) { |
| 181 | + if (!almost_equal_cplx(output, reference, SYCL_CPLX_TOL_ULP)) { |
| 182 | + std::cerr << std::setprecision(std::numeric_limits<T>::max_digits10) |
| 183 | + << "Test failed with complex_type: " << get_typename<T>() |
| 184 | + << " Computed on " << (is_device ? "device" : "host") |
| 185 | + << " Output: " << output << " Reference: " << reference |
| 186 | + << std::endl; |
| 187 | + return false; |
| 188 | + } |
| 189 | + return true; |
| 190 | +} |
| 191 | + |
| 192 | +template <typename T> |
| 193 | +bool check_results(T output, T reference, bool is_device) { |
| 194 | + if (!almost_equal_scalar(output, reference, SYCL_CPLX_TOL_ULP)) { |
| 195 | + std::cerr << std::setprecision(std::numeric_limits<T>::max_digits10) |
| 196 | + << "Test failed with complex_type: " << get_typename<T>() |
| 197 | + << " Computed on " << (is_device ? "device" : "host") |
| 198 | + << " Output: " << output << " Reference: " << reference |
| 199 | + << std::endl; |
| 200 | + return false; |
| 201 | + } |
| 202 | + return true; |
| 203 | +} |
0 commit comments