1 /*===---- complex - CUDA wrapper for <complex> ------------------------------=== 2 * 3 * Permission is hereby granted, free of charge, to any person obtaining a copy 4 * of this software and associated documentation files (the "Software"), to deal 5 * in the Software without restriction, including without limitation the rights 6 * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell 7 * copies of the Software, and to permit persons to whom the Software is 8 * furnished to do so, subject to the following conditions: 9 * 10 * The above copyright notice and this permission notice shall be included in 11 * all copies or substantial portions of the Software. 12 * 13 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 14 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 15 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 16 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 17 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 18 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN 19 * THE SOFTWARE. 20 * 21 *===-----------------------------------------------------------------------=== 22 */ 23 24 #ifndef __CLANG_CUDA_WRAPPERS_COMPLEX 25 #define __CLANG_CUDA_WRAPPERS_COMPLEX 26 27 // Wrapper around <complex> that forces its functions to be __host__ 28 // __device__. 29 30 // First, include host-only headers we think are likely to be included by 31 // <complex>, so that the pragma below only applies to <complex> itself. 32 #if __cplusplus >= 201103L 33 #include <type_traits> 34 #endif 35 #include <stdexcept> 36 #include <cmath> 37 #include <sstream> 38 39 // Next, include our <algorithm> wrapper, to ensure that device overloads of 40 // std::min/max are available. 41 #include <algorithm> 42 43 #pragma clang force_cuda_host_device begin 44 45 // When compiling for device, ask libstdc++ to use its own implements of 46 // complex functions, rather than calling builtins (which resolve to library 47 // functions that don't exist when compiling CUDA device code). 48 // 49 // This is a little dicey, because it causes libstdc++ to define a different 50 // set of overloads on host and device. 51 // 52 // // Present only when compiling for host. 53 // __host__ __device__ void complex<float> sin(const complex<float>& x) { 54 // return __builtin_csinf(x); 55 // } 56 // 57 // // Present when compiling for host and for device. 58 // template <typename T> 59 // void __host__ __device__ complex<T> sin(const complex<T>& x) { 60 // return complex<T>(sin(x.real()) * cosh(x.imag()), 61 // cos(x.real()), sinh(x.imag())); 62 // } 63 // 64 // This is safe because when compiling for device, all function calls in 65 // __host__ code to sin() will still resolve to *something*, even if they don't 66 // resolve to the same function as they resolve to when compiling for host. We 67 // don't care that they don't resolve to the right function because we won't 68 // codegen this host code when compiling for device. 69 70 #pragma push_macro("_GLIBCXX_USE_C99_COMPLEX") 71 #pragma push_macro("_GLIBCXX_USE_C99_COMPLEX_TR1") 72 #define _GLIBCXX_USE_C99_COMPLEX 0 73 #define _GLIBCXX_USE_C99_COMPLEX_TR1 0 74 75 #include_next <complex> 76 77 #pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX_TR1") 78 #pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX") 79 80 #pragma clang force_cuda_host_device end 81 82 #endif // include guard 83
