1/*===-- __clang_cuda_complex_builtins - CUDA impls of runtime complex fns ---=== 2 * 3 * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4 * See https://llvm.org/LICENSE.txt for license information. 5 * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6 * 7 *===-----------------------------------------------------------------------=== 8 */ 9 10#ifndef __CLANG_CUDA_COMPLEX_BUILTINS 11#define __CLANG_CUDA_COMPLEX_BUILTINS 12 13// This header defines __muldc3, __mulsc3, __divdc3, and __divsc3. These are 14// libgcc functions that clang assumes are available when compiling c99 complex 15// operations. (These implementations come from libc++, and have been modified 16// to work with CUDA.) 17 18extern "C" inline __device__ double _Complex __muldc3(double __a, double __b, 19 double __c, double __d) { 20 double __ac = __a * __c; 21 double __bd = __b * __d; 22 double __ad = __a * __d; 23 double __bc = __b * __c; 24 double _Complex z; 25 __real__(z) = __ac - __bd; 26 __imag__(z) = __ad + __bc; 27 if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) { 28 int __recalc = 0; 29 if (std::isinf(__a) || std::isinf(__b)) { 30 __a = std::copysign(std::isinf(__a) ? 1 : 0, __a); 31 __b = std::copysign(std::isinf(__b) ? 1 : 0, __b); 32 if (std::isnan(__c)) 33 __c = std::copysign(0, __c); 34 if (std::isnan(__d)) 35 __d = std::copysign(0, __d); 36 __recalc = 1; 37 } 38 if (std::isinf(__c) || std::isinf(__d)) { 39 __c = std::copysign(std::isinf(__c) ? 1 : 0, __c); 40 __d = std::copysign(std::isinf(__d) ? 1 : 0, __d); 41 if (std::isnan(__a)) 42 __a = std::copysign(0, __a); 43 if (std::isnan(__b)) 44 __b = std::copysign(0, __b); 45 __recalc = 1; 46 } 47 if (!__recalc && (std::isinf(__ac) || std::isinf(__bd) || 48 std::isinf(__ad) || std::isinf(__bc))) { 49 if (std::isnan(__a)) 50 __a = std::copysign(0, __a); 51 if (std::isnan(__b)) 52 __b = std::copysign(0, __b); 53 if (std::isnan(__c)) 54 __c = std::copysign(0, __c); 55 if (std::isnan(__d)) 56 __d = std::copysign(0, __d); 57 __recalc = 1; 58 } 59 if (__recalc) { 60 // Can't use std::numeric_limits<double>::infinity() -- that doesn't have 61 // a device overload (and isn't constexpr before C++11, naturally). 62 __real__(z) = __builtin_huge_valf() * (__a * __c - __b * __d); 63 __imag__(z) = __builtin_huge_valf() * (__a * __d + __b * __c); 64 } 65 } 66 return z; 67} 68 69extern "C" inline __device__ float _Complex __mulsc3(float __a, float __b, 70 float __c, float __d) { 71 float __ac = __a * __c; 72 float __bd = __b * __d; 73 float __ad = __a * __d; 74 float __bc = __b * __c; 75 float _Complex z; 76 __real__(z) = __ac - __bd; 77 __imag__(z) = __ad + __bc; 78 if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) { 79 int __recalc = 0; 80 if (std::isinf(__a) || std::isinf(__b)) { 81 __a = std::copysign(std::isinf(__a) ? 1 : 0, __a); 82 __b = std::copysign(std::isinf(__b) ? 1 : 0, __b); 83 if (std::isnan(__c)) 84 __c = std::copysign(0, __c); 85 if (std::isnan(__d)) 86 __d = std::copysign(0, __d); 87 __recalc = 1; 88 } 89 if (std::isinf(__c) || std::isinf(__d)) { 90 __c = std::copysign(std::isinf(__c) ? 1 : 0, __c); 91 __d = std::copysign(std::isinf(__d) ? 1 : 0, __d); 92 if (std::isnan(__a)) 93 __a = std::copysign(0, __a); 94 if (std::isnan(__b)) 95 __b = std::copysign(0, __b); 96 __recalc = 1; 97 } 98 if (!__recalc && (std::isinf(__ac) || std::isinf(__bd) || 99 std::isinf(__ad) || std::isinf(__bc))) { 100 if (std::isnan(__a)) 101 __a = std::copysign(0, __a); 102 if (std::isnan(__b)) 103 __b = std::copysign(0, __b); 104 if (std::isnan(__c)) 105 __c = std::copysign(0, __c); 106 if (std::isnan(__d)) 107 __d = std::copysign(0, __d); 108 __recalc = 1; 109 } 110 if (__recalc) { 111 __real__(z) = __builtin_huge_valf() * (__a * __c - __b * __d); 112 __imag__(z) = __builtin_huge_valf() * (__a * __d + __b * __c); 113 } 114 } 115 return z; 116} 117 118extern "C" inline __device__ double _Complex __divdc3(double __a, double __b, 119 double __c, double __d) { 120 int __ilogbw = 0; 121 // Can't use std::max, because that's defined in <algorithm>, and we don't 122 // want to pull that in for every compile. The CUDA headers define 123 // ::max(float, float) and ::max(double, double), which is sufficient for us. 124 double __logbw = std::logb(max(std::abs(__c), std::abs(__d))); 125 if (std::isfinite(__logbw)) { 126 __ilogbw = (int)__logbw; 127 __c = std::scalbn(__c, -__ilogbw); 128 __d = std::scalbn(__d, -__ilogbw); 129 } 130 double __denom = __c * __c + __d * __d; 131 double _Complex z; 132 __real__(z) = std::scalbn((__a * __c + __b * __d) / __denom, -__ilogbw); 133 __imag__(z) = std::scalbn((__b * __c - __a * __d) / __denom, -__ilogbw); 134 if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) { 135 if ((__denom == 0.0) && (!std::isnan(__a) || !std::isnan(__b))) { 136 __real__(z) = std::copysign(__builtin_huge_valf(), __c) * __a; 137 __imag__(z) = std::copysign(__builtin_huge_valf(), __c) * __b; 138 } else if ((std::isinf(__a) || std::isinf(__b)) && std::isfinite(__c) && 139 std::isfinite(__d)) { 140 __a = std::copysign(std::isinf(__a) ? 1.0 : 0.0, __a); 141 __b = std::copysign(std::isinf(__b) ? 1.0 : 0.0, __b); 142 __real__(z) = __builtin_huge_valf() * (__a * __c + __b * __d); 143 __imag__(z) = __builtin_huge_valf() * (__b * __c - __a * __d); 144 } else if (std::isinf(__logbw) && __logbw > 0.0 && std::isfinite(__a) && 145 std::isfinite(__b)) { 146 __c = std::copysign(std::isinf(__c) ? 1.0 : 0.0, __c); 147 __d = std::copysign(std::isinf(__d) ? 1.0 : 0.0, __d); 148 __real__(z) = 0.0 * (__a * __c + __b * __d); 149 __imag__(z) = 0.0 * (__b * __c - __a * __d); 150 } 151 } 152 return z; 153} 154 155extern "C" inline __device__ float _Complex __divsc3(float __a, float __b, 156 float __c, float __d) { 157 int __ilogbw = 0; 158 float __logbw = std::logb(max(std::abs(__c), std::abs(__d))); 159 if (std::isfinite(__logbw)) { 160 __ilogbw = (int)__logbw; 161 __c = std::scalbn(__c, -__ilogbw); 162 __d = std::scalbn(__d, -__ilogbw); 163 } 164 float __denom = __c * __c + __d * __d; 165 float _Complex z; 166 __real__(z) = std::scalbn((__a * __c + __b * __d) / __denom, -__ilogbw); 167 __imag__(z) = std::scalbn((__b * __c - __a * __d) / __denom, -__ilogbw); 168 if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) { 169 if ((__denom == 0) && (!std::isnan(__a) || !std::isnan(__b))) { 170 __real__(z) = std::copysign(__builtin_huge_valf(), __c) * __a; 171 __imag__(z) = std::copysign(__builtin_huge_valf(), __c) * __b; 172 } else if ((std::isinf(__a) || std::isinf(__b)) && std::isfinite(__c) && 173 std::isfinite(__d)) { 174 __a = std::copysign(std::isinf(__a) ? 1 : 0, __a); 175 __b = std::copysign(std::isinf(__b) ? 1 : 0, __b); 176 __real__(z) = __builtin_huge_valf() * (__a * __c + __b * __d); 177 __imag__(z) = __builtin_huge_valf() * (__b * __c - __a * __d); 178 } else if (std::isinf(__logbw) && __logbw > 0 && std::isfinite(__a) && 179 std::isfinite(__b)) { 180 __c = std::copysign(std::isinf(__c) ? 1 : 0, __c); 181 __d = std::copysign(std::isinf(__d) ? 1 : 0, __d); 182 __real__(z) = 0 * (__a * __c + __b * __d); 183 __imag__(z) = 0 * (__b * __c - __a * __d); 184 } 185 } 186 return z; 187} 188 189#endif // __CLANG_CUDA_COMPLEX_BUILTINS 190