1311118Sdim/*===---- complex - CUDA wrapper for <complex> ------------------------------=== 2311118Sdim * 3311118Sdim * Permission is hereby granted, free of charge, to any person obtaining a copy 4311118Sdim * of this software and associated documentation files (the "Software"), to deal 5311118Sdim * in the Software without restriction, including without limitation the rights 6311118Sdim * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell 7311118Sdim * copies of the Software, and to permit persons to whom the Software is 8311118Sdim * furnished to do so, subject to the following conditions: 9311118Sdim * 10311118Sdim * The above copyright notice and this permission notice shall be included in 11311118Sdim * all copies or substantial portions of the Software. 12311118Sdim * 13311118Sdim * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 14311118Sdim * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 15311118Sdim * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 16311118Sdim * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 17311118Sdim * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 18311118Sdim * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN 19311118Sdim * THE SOFTWARE. 20311118Sdim * 21311118Sdim *===-----------------------------------------------------------------------=== 22311118Sdim */ 23311118Sdim 24311118Sdim#ifndef __CLANG_CUDA_WRAPPERS_COMPLEX 25311118Sdim#define __CLANG_CUDA_WRAPPERS_COMPLEX 26311118Sdim 27311118Sdim// Wrapper around <complex> that forces its functions to be __host__ 28311118Sdim// __device__. 29311118Sdim 30311118Sdim// First, include host-only headers we think are likely to be included by 31311118Sdim// <complex>, so that the pragma below only applies to <complex> itself. 32311118Sdim#if __cplusplus >= 201103L 33311118Sdim#include <type_traits> 34311118Sdim#endif 35311118Sdim#include <stdexcept> 36311118Sdim#include <cmath> 37311118Sdim#include <sstream> 38311118Sdim 39311118Sdim// Next, include our <algorithm> wrapper, to ensure that device overloads of 40311118Sdim// std::min/max are available. 41311118Sdim#include <algorithm> 42311118Sdim 43311118Sdim#pragma clang force_cuda_host_device begin 44311118Sdim 45311118Sdim// When compiling for device, ask libstdc++ to use its own implements of 46311118Sdim// complex functions, rather than calling builtins (which resolve to library 47311118Sdim// functions that don't exist when compiling CUDA device code). 48311118Sdim// 49311118Sdim// This is a little dicey, because it causes libstdc++ to define a different 50311118Sdim// set of overloads on host and device. 51311118Sdim// 52311118Sdim// // Present only when compiling for host. 53311118Sdim// __host__ __device__ void complex<float> sin(const complex<float>& x) { 54311118Sdim// return __builtin_csinf(x); 55311118Sdim// } 56311118Sdim// 57311118Sdim// // Present when compiling for host and for device. 58311118Sdim// template <typename T> 59311118Sdim// void __host__ __device__ complex<T> sin(const complex<T>& x) { 60311118Sdim// return complex<T>(sin(x.real()) * cosh(x.imag()), 61311118Sdim// cos(x.real()), sinh(x.imag())); 62311118Sdim// } 63311118Sdim// 64311118Sdim// This is safe because when compiling for device, all function calls in 65311118Sdim// __host__ code to sin() will still resolve to *something*, even if they don't 66311118Sdim// resolve to the same function as they resolve to when compiling for host. We 67311118Sdim// don't care that they don't resolve to the right function because we won't 68311118Sdim// codegen this host code when compiling for device. 69311118Sdim 70311118Sdim#pragma push_macro("_GLIBCXX_USE_C99_COMPLEX") 71311118Sdim#pragma push_macro("_GLIBCXX_USE_C99_COMPLEX_TR1") 72311118Sdim#define _GLIBCXX_USE_C99_COMPLEX 0 73311118Sdim#define _GLIBCXX_USE_C99_COMPLEX_TR1 0 74311118Sdim 75311118Sdim#include_next <complex> 76311118Sdim 77311118Sdim#pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX_TR1") 78311118Sdim#pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX") 79311118Sdim 80311118Sdim#pragma clang force_cuda_host_device end 81311118Sdim 82311118Sdim#endif // include guard 83