complex 3.3 KB

12345678910111213141516171819202122232425262728293031323334353637383940414243444546474849505152535455565758596061626364656667686970717273747576777879808182
  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. #ifndef __CLANG_CUDA_WRAPPERS_COMPLEX
  24. #define __CLANG_CUDA_WRAPPERS_COMPLEX
  25. // Wrapper around <complex> that forces its functions to be __host__
  26. // __device__.
  27. // First, include host-only headers we think are likely to be included by
  28. // <complex>, so that the pragma below only applies to <complex> itself.
  29. #if __cplusplus >= 201103L
  30. #include <type_traits>
  31. #endif
  32. #include <stdexcept>
  33. #include <cmath>
  34. #include <sstream>
  35. // Next, include our <algorithm> wrapper, to ensure that device overloads of
  36. // std::min/max are available.
  37. #include <algorithm>
  38. #pragma clang force_cuda_host_device begin
  39. // When compiling for device, ask libstdc++ to use its own implements of
  40. // complex functions, rather than calling builtins (which resolve to library
  41. // functions that don't exist when compiling CUDA device code).
  42. //
  43. // This is a little dicey, because it causes libstdc++ to define a different
  44. // set of overloads on host and device.
  45. //
  46. // // Present only when compiling for host.
  47. // __host__ __device__ void complex<float> sin(const complex<float>& x) {
  48. // return __builtin_csinf(x);
  49. // }
  50. //
  51. // // Present when compiling for host and for device.
  52. // template <typename T>
  53. // void __host__ __device__ complex<T> sin(const complex<T>& x) {
  54. // return complex<T>(sin(x.real()) * cosh(x.imag()),
  55. // cos(x.real()), sinh(x.imag()));
  56. // }
  57. //
  58. // This is safe because when compiling for device, all function calls in
  59. // __host__ code to sin() will still resolve to *something*, even if they don't
  60. // resolve to the same function as they resolve to when compiling for host. We
  61. // don't care that they don't resolve to the right function because we won't
  62. // codegen this host code when compiling for device.
  63. #pragma push_macro("_GLIBCXX_USE_C99_COMPLEX")
  64. #pragma push_macro("_GLIBCXX_USE_C99_COMPLEX_TR1")
  65. #define _GLIBCXX_USE_C99_COMPLEX 0
  66. #define _GLIBCXX_USE_C99_COMPLEX_TR1 0
  67. #include_next <complex>
  68. #pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX_TR1")
  69. #pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX")
  70. #pragma clang force_cuda_host_device end
  71. #endif // include guard