91 lines · plain
1/*===---- complex - CUDA wrapper for <complex> ------------------------------===2 *3 * Permission is hereby granted, free of charge, to any person obtaining a copy4 * of this software and associated documentation files (the "Software"), to deal5 * in the Software without restriction, including without limitation the rights6 * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell7 * copies of the Software, and to permit persons to whom the Software is8 * furnished to do so, subject to the following conditions:9 *10 * The above copyright notice and this permission notice shall be included in11 * all copies or substantial portions of the Software.12 *13 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR14 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,15 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE16 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER17 * 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 IN19 * THE SOFTWARE.20 *21 *===-----------------------------------------------------------------------===22 */23 24#ifndef __CLANG_CUDA_WRAPPERS_COMPLEX25#define __CLANG_CUDA_WRAPPERS_COMPLEX26 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 by31// <complex>, so that the pragma below only applies to <complex> itself.32#if __cplusplus >= 201103L33#include <type_traits>34#endif35#include <stdexcept>36#include <cmath>37#include <sstream>38 39// Next, include our <algorithm> wrapper, to ensure that device overloads of40// std::min/max are available.41#include <algorithm>42 43#pragma clang force_cuda_host_device begin44 45// When compiling for device, ask libstdc++ to use its own implements of46// complex functions, rather than calling builtins (which resolve to library47// functions that don't exist when compiling CUDA device code).48//49// This is a little dicey, because it causes libstdc++ to define a different50// 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 in65// __host__ code to sin() will still resolve to *something*, even if they don't66// resolve to the same function as they resolve to when compiling for host. We67// don't care that they don't resolve to the right function because we won't68// 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 073#define _GLIBCXX_USE_C99_COMPLEX_TR1 074 75// Work around a compatibility issue with libstdc++ 11.1.076// https://bugs.llvm.org/show_bug.cgi?id=5038377#pragma push_macro("__failed_assertion")78#if _GLIBCXX_RELEASE == 1179#define __failed_assertion __cuda_failed_assertion80#endif81 82#include_next <complex>83 84#pragma pop_macro("__failed_assertion")85#pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX_TR1")86#pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX")87 88#pragma clang force_cuda_host_device end89 90#endif // include guard91