r283680 - [CUDA] Support <complex> and std::min/max on the device.
Reid Kleckner via cfe-commits
cfe-commits at lists.llvm.org
Tue Oct 11 09:31:40 PDT 2016
Can we figure out why having an incremental stage1 is problematic? I want
to keep it that way to speed up the bot.
It sounds like the problem is that we do not delete the stage1 install
directory before installing to it, or our install step uses a glob that may
include stale files from previous builds.
On Mon, Oct 10, 2016 at 8:56 AM, Justin Lebar <jlebar at google.com> wrote:
> As discussed over IM, builds are still failing after reverting this.
>
> The issue seems to be that the stage 1 build in bootstrap compiles is
> not a clean build. Therefore the errant <algorithm> and <cmake>
> headers are never removed from the install directory.
>
> I believe I fixed the cmake in r283683 (sent a few hours after this
> one), but because the stage 1 builds are not clean, they were still
> broken after that change.
>
> Easiest fix at this point would be to find someone with ssh access to
> the buildbots to rm -rf the stage 1 directory.
>
> On Mon, Oct 10, 2016 at 7:19 AM, Nico Weber <thakis at chromium.org> wrote:
> > This broke bootstrap builds, I reverted it for now in r283747.
> >
> > On Sat, Oct 8, 2016 at 6:16 PM, Justin Lebar via cfe-commits
> > <cfe-commits at lists.llvm.org> wrote:
> >>
> >> Author: jlebar
> >> Date: Sat Oct 8 17:16:12 2016
> >> New Revision: 283680
> >>
> >> URL: http://llvm.org/viewvc/llvm-project?rev=283680&view=rev
> >> Log:
> >> [CUDA] Support <complex> and std::min/max on the device.
> >>
> >> Summary:
> >> We do this by wrapping <complex> and <algorithm>.
> >>
> >> Tests are in the test-suite.
> >>
> >> Reviewers: tra
> >>
> >> Subscribers: jhen, beanz, cfe-commits, mgorny
> >>
> >> Differential Revision: https://reviews.llvm.org/D24979
> >>
> >> Added:
> >> cfe/trunk/lib/Headers/__clang_cuda_complex_builtins.h
> >> cfe/trunk/lib/Headers/cuda_wrappers/
> >> cfe/trunk/lib/Headers/cuda_wrappers/algorithm
> >> cfe/trunk/lib/Headers/cuda_wrappers/complex
> >> Modified:
> >> cfe/trunk/lib/Driver/ToolChains.cpp
> >> cfe/trunk/lib/Headers/CMakeLists.txt
> >> cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
> >>
> >> Modified: cfe/trunk/lib/Driver/ToolChains.cpp
> >> URL:
> >> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Driver/
> ToolChains.cpp?rev=283680&r1=283679&r2=283680&view=diff
> >>
> >> ============================================================
> ==================
> >> --- cfe/trunk/lib/Driver/ToolChains.cpp (original)
> >> +++ cfe/trunk/lib/Driver/ToolChains.cpp Sat Oct 8 17:16:12 2016
> >> @@ -4694,6 +4694,15 @@ void Linux::AddClangCXXStdlibIncludeArgs
> >>
> >> void Linux::AddCudaIncludeArgs(const ArgList &DriverArgs,
> >> ArgStringList &CC1Args) const {
> >> + if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
> >> + // Add cuda_wrappers/* to our system include path. This lets us
> wrap
> >> + // standard library headers.
> >> + SmallString<128> P(getDriver().ResourceDir);
> >> + llvm::sys::path::append(P, "include");
> >> + llvm::sys::path::append(P, "cuda_wrappers");
> >> + addSystemInclude(DriverArgs, CC1Args, P);
> >> + }
> >> +
> >> if (DriverArgs.hasArg(options::OPT_nocudainc))
> >> return;
> >>
> >>
> >> Modified: cfe/trunk/lib/Headers/CMakeLists.txt
> >> URL:
> >> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/
> CMakeLists.txt?rev=283680&r1=283679&r2=283680&view=diff
> >>
> >> ============================================================
> ==================
> >> --- cfe/trunk/lib/Headers/CMakeLists.txt (original)
> >> +++ cfe/trunk/lib/Headers/CMakeLists.txt Sat Oct 8 17:16:12 2016
> >> @@ -24,10 +24,13 @@ set(files
> >> bmiintrin.h
> >> __clang_cuda_builtin_vars.h
> >> __clang_cuda_cmath.h
> >> + __clang_cuda_complex_builtins.h
> >> __clang_cuda_intrinsics.h
> >> __clang_cuda_math_forward_declares.h
> >> __clang_cuda_runtime_wrapper.h
> >> cpuid.h
> >> + cuda_wrappers/algorithm
> >> + cuda_wrappers/complex
> >> clflushoptintrin.h
> >> emmintrin.h
> >> f16cintrin.h
> >>
> >> Added: cfe/trunk/lib/Headers/__clang_cuda_complex_builtins.h
> >> URL:
> >> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/
> __clang_cuda_complex_builtins.h?rev=283680&view=auto
> >>
> >> ============================================================
> ==================
> >> --- cfe/trunk/lib/Headers/__clang_cuda_complex_builtins.h (added)
> >> +++ cfe/trunk/lib/Headers/__clang_cuda_complex_builtins.h Sat Oct 8
> >> 17:16:12 2016
> >> @@ -0,0 +1,203 @@
> >> +/*===-- __clang_cuda_complex_builtins - CUDA impls of runtime complex
> fns
> >> ---===
> >> + *
> >> + * Permission is hereby granted, free of charge, to any person
> obtaining
> >> a copy
> >> + * of this software and associated documentation files (the
> "Software"),
> >> to deal
> >> + * in the Software without restriction, including without limitation
> the
> >> rights
> >> + * to use, copy, modify, merge, publish, distribute, sublicense, and/or
> >> sell
> >> + * copies of the Software, and to permit persons to whom the Software
> is
> >> + * furnished to do so, subject to the following conditions:
> >> + *
> >> + * The above copyright notice and this permission notice shall be
> >> included in
> >> + * all copies or substantial portions of the Software.
> >> + *
> >> + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
> >> EXPRESS OR
> >> + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
> >> MERCHANTABILITY,
> >> + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT
> >> SHALL THE
> >> + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
> OTHER
> >> + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
> >> ARISING FROM,
> >> + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
> DEALINGS
> >> IN
> >> + * THE SOFTWARE.
> >> + *
> >> +
> >> *===--------------------------------------------------------
> ---------------===
> >> + */
> >> +
> >> +#ifndef __CLANG_CUDA_COMPLEX_BUILTINS
> >> +#define __CLANG_CUDA_COMPLEX_BUILTINS
> >> +
> >> +// This header defines __muldc3, __mulsc3, __divdc3, and __divsc3.
> These
> >> are
> >> +// libgcc functions that clang assumes are available when compiling c99
> >> complex
> >> +// operations. (These implementations come from libc++, and have been
> >> modified
> >> +// to work with CUDA.)
> >> +
> >> +extern "C" inline __device__ double _Complex __muldc3(double __a,
> double
> >> __b,
> >> + double __c,
> double
> >> __d) {
> >> + double __ac = __a * __c;
> >> + double __bd = __b * __d;
> >> + double __ad = __a * __d;
> >> + double __bc = __b * __c;
> >> + double _Complex z;
> >> + __real__(z) = __ac - __bd;
> >> + __imag__(z) = __ad + __bc;
> >> + if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) {
> >> + int __recalc = 0;
> >> + if (std::isinf(__a) || std::isinf(__b)) {
> >> + __a = std::copysign(std::isinf(__a) ? 1 : 0, __a);
> >> + __b = std::copysign(std::isinf(__b) ? 1 : 0, __b);
> >> + if (std::isnan(__c))
> >> + __c = std::copysign(0, __c);
> >> + if (std::isnan(__d))
> >> + __d = std::copysign(0, __d);
> >> + __recalc = 1;
> >> + }
> >> + if (std::isinf(__c) || std::isinf(__d)) {
> >> + __c = std::copysign(std::isinf(__c) ? 1 : 0, __c);
> >> + __d = std::copysign(std::isinf(__d) ? 1 : 0, __d);
> >> + if (std::isnan(__a))
> >> + __a = std::copysign(0, __a);
> >> + if (std::isnan(__b))
> >> + __b = std::copysign(0, __b);
> >> + __recalc = 1;
> >> + }
> >> + if (!__recalc && (std::isinf(__ac) || std::isinf(__bd) ||
> >> + std::isinf(__ad) || std::isinf(__bc))) {
> >> + if (std::isnan(__a))
> >> + __a = std::copysign(0, __a);
> >> + if (std::isnan(__b))
> >> + __b = std::copysign(0, __b);
> >> + if (std::isnan(__c))
> >> + __c = std::copysign(0, __c);
> >> + if (std::isnan(__d))
> >> + __d = std::copysign(0, __d);
> >> + __recalc = 1;
> >> + }
> >> + if (__recalc) {
> >> + // Can't use std::numeric_limits<double>::infinity() -- that
> >> doesn't have
> >> + // a device overload (and isn't constexpr before C++11,
> naturally).
> >> + __real__(z) = __builtin_huge_valf() * (__a * __c - __b * __d);
> >> + __imag__(z) = __builtin_huge_valf() * (__a * __d + __b * __c);
> >> + }
> >> + }
> >> + return z;
> >> +}
> >> +
> >> +extern "C" inline __device__ float _Complex __mulsc3(float __a, float
> >> __b,
> >> + float __c, float
> >> __d) {
> >> + float __ac = __a * __c;
> >> + float __bd = __b * __d;
> >> + float __ad = __a * __d;
> >> + float __bc = __b * __c;
> >> + float _Complex z;
> >> + __real__(z) = __ac - __bd;
> >> + __imag__(z) = __ad + __bc;
> >> + if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) {
> >> + int __recalc = 0;
> >> + if (std::isinf(__a) || std::isinf(__b)) {
> >> + __a = std::copysign(std::isinf(__a) ? 1 : 0, __a);
> >> + __b = std::copysign(std::isinf(__b) ? 1 : 0, __b);
> >> + if (std::isnan(__c))
> >> + __c = std::copysign(0, __c);
> >> + if (std::isnan(__d))
> >> + __d = std::copysign(0, __d);
> >> + __recalc = 1;
> >> + }
> >> + if (std::isinf(__c) || std::isinf(__d)) {
> >> + __c = std::copysign(std::isinf(__c) ? 1 : 0, __c);
> >> + __d = std::copysign(std::isinf(__d) ? 1 : 0, __d);
> >> + if (std::isnan(__a))
> >> + __a = std::copysign(0, __a);
> >> + if (std::isnan(__b))
> >> + __b = std::copysign(0, __b);
> >> + __recalc = 1;
> >> + }
> >> + if (!__recalc && (std::isinf(__ac) || std::isinf(__bd) ||
> >> + std::isinf(__ad) || std::isinf(__bc))) {
> >> + if (std::isnan(__a))
> >> + __a = std::copysign(0, __a);
> >> + if (std::isnan(__b))
> >> + __b = std::copysign(0, __b);
> >> + if (std::isnan(__c))
> >> + __c = std::copysign(0, __c);
> >> + if (std::isnan(__d))
> >> + __d = std::copysign(0, __d);
> >> + __recalc = 1;
> >> + }
> >> + if (__recalc) {
> >> + __real__(z) = __builtin_huge_valf() * (__a * __c - __b * __d);
> >> + __imag__(z) = __builtin_huge_valf() * (__a * __d + __b * __c);
> >> + }
> >> + }
> >> + return z;
> >> +}
> >> +
> >> +extern "C" inline __device__ double _Complex __divdc3(double __a,
> double
> >> __b,
> >> + double __c,
> double
> >> __d) {
> >> + int __ilogbw = 0;
> >> + // Can't use std::max, because that's defined in <algorithm>, and we
> >> don't
> >> + // want to pull that in for every compile. The CUDA headers define
> >> + // ::max(float, float) and ::max(double, double), which is sufficient
> >> for us.
> >> + double __logbw = std::logb(max(std::abs(__c), std::abs(__d)));
> >> + if (std::isfinite(__logbw)) {
> >> + __ilogbw = (int)__logbw;
> >> + __c = std::scalbn(__c, -__ilogbw);
> >> + __d = std::scalbn(__d, -__ilogbw);
> >> + }
> >> + double __denom = __c * __c + __d * __d;
> >> + double _Complex z;
> >> + __real__(z) = std::scalbn((__a * __c + __b * __d) / __denom,
> >> -__ilogbw);
> >> + __imag__(z) = std::scalbn((__b * __c - __a * __d) / __denom,
> >> -__ilogbw);
> >> + if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) {
> >> + if ((__denom == 0.0) && (!std::isnan(__a) || !std::isnan(__b))) {
> >> + __real__(z) = std::copysign(__builtin_huge_valf(), __c) * __a;
> >> + __imag__(z) = std::copysign(__builtin_huge_valf(), __c) * __b;
> >> + } else if ((std::isinf(__a) || std::isinf(__b)) &&
> std::isfinite(__c)
> >> &&
> >> + std::isfinite(__d)) {
> >> + __a = std::copysign(std::isinf(__a) ? 1.0 : 0.0, __a);
> >> + __b = std::copysign(std::isinf(__b) ? 1.0 : 0.0, __b);
> >> + __real__(z) = __builtin_huge_valf() * (__a * __c + __b * __d);
> >> + __imag__(z) = __builtin_huge_valf() * (__b * __c - __a * __d);
> >> + } else if (std::isinf(__logbw) && __logbw > 0.0 &&
> std::isfinite(__a)
> >> &&
> >> + std::isfinite(__b)) {
> >> + __c = std::copysign(std::isinf(__c) ? 1.0 : 0.0, __c);
> >> + __d = std::copysign(std::isinf(__d) ? 1.0 : 0.0, __d);
> >> + __real__(z) = 0.0 * (__a * __c + __b * __d);
> >> + __imag__(z) = 0.0 * (__b * __c - __a * __d);
> >> + }
> >> + }
> >> + return z;
> >> +}
> >> +
> >> +extern "C" inline __device__ float _Complex __divsc3(float __a, float
> >> __b,
> >> + float __c, float
> >> __d) {
> >> + int __ilogbw = 0;
> >> + float __logbw = std::logb(max(std::abs(__c), std::abs(__d)));
> >> + if (std::isfinite(__logbw)) {
> >> + __ilogbw = (int)__logbw;
> >> + __c = std::scalbn(__c, -__ilogbw);
> >> + __d = std::scalbn(__d, -__ilogbw);
> >> + }
> >> + float __denom = __c * __c + __d * __d;
> >> + float _Complex z;
> >> + __real__(z) = std::scalbn((__a * __c + __b * __d) / __denom,
> >> -__ilogbw);
> >> + __imag__(z) = std::scalbn((__b * __c - __a * __d) / __denom,
> >> -__ilogbw);
> >> + if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) {
> >> + if ((__denom == 0) && (!std::isnan(__a) || !std::isnan(__b))) {
> >> + __real__(z) = std::copysign(__builtin_huge_valf(), __c) * __a;
> >> + __imag__(z) = std::copysign(__builtin_huge_valf(), __c) * __b;
> >> + } else if ((std::isinf(__a) || std::isinf(__b)) &&
> std::isfinite(__c)
> >> &&
> >> + std::isfinite(__d)) {
> >> + __a = std::copysign(std::isinf(__a) ? 1 : 0, __a);
> >> + __b = std::copysign(std::isinf(__b) ? 1 : 0, __b);
> >> + __real__(z) = __builtin_huge_valf() * (__a * __c + __b * __d);
> >> + __imag__(z) = __builtin_huge_valf() * (__b * __c - __a * __d);
> >> + } else if (std::isinf(__logbw) && __logbw > 0 && std::isfinite(__a)
> >> &&
> >> + std::isfinite(__b)) {
> >> + __c = std::copysign(std::isinf(__c) ? 1 : 0, __c);
> >> + __d = std::copysign(std::isinf(__d) ? 1 : 0, __d);
> >> + __real__(z) = 0 * (__a * __c + __b * __d);
> >> + __imag__(z) = 0 * (__b * __c - __a * __d);
> >> + }
> >> + }
> >> + return z;
> >> +}
> >> +
> >> +#endif // __CLANG_CUDA_COMPLEX_BUILTINS
> >>
> >> Modified: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
> >> URL:
> >> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/
> __clang_cuda_runtime_wrapper.h?rev=283680&r1=283679&r2=283680&view=diff
> >>
> >> ============================================================
> ==================
> >> --- cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h (original)
> >> +++ cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h Sat Oct 8
> >> 17:16:12 2016
> >> @@ -312,6 +312,7 @@ __device__ inline __cuda_builtin_gridDim
> >>
> >> #include <__clang_cuda_cmath.h>
> >> #include <__clang_cuda_intrinsics.h>
> >> +#include <__clang_cuda_complex_builtins.h>
> >>
> >> // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in
> >> host
> >> // mode, giving them their "proper" types of dim3 and uint3. This is
> >>
> >> Added: cfe/trunk/lib/Headers/cuda_wrappers/algorithm
> >> URL:
> >> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/
> cuda_wrappers/algorithm?rev=283680&view=auto
> >>
> >> ============================================================
> ==================
> >> --- cfe/trunk/lib/Headers/cuda_wrappers/algorithm (added)
> >> +++ cfe/trunk/lib/Headers/cuda_wrappers/algorithm Sat Oct 8 17:16:12
> 2016
> >> @@ -0,0 +1,96 @@
> >> +/*===---- complex - CUDA wrapper for <algorithm>
> >> ----------------------------===
> >> + *
> >> + * Permission is hereby granted, free of charge, to any person
> obtaining
> >> a copy
> >> + * of this software and associated documentation files (the
> "Software"),
> >> to deal
> >> + * in the Software without restriction, including without limitation
> the
> >> rights
> >> + * to use, copy, modify, merge, publish, distribute, sublicense, and/or
> >> sell
> >> + * copies of the Software, and to permit persons to whom the Software
> is
> >> + * furnished to do so, subject to the following conditions:
> >> + *
> >> + * The above copyright notice and this permission notice shall be
> >> included in
> >> + * all copies or substantial portions of the Software.
> >> + *
> >> + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
> >> EXPRESS OR
> >> + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
> >> MERCHANTABILITY,
> >> + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT
> >> SHALL THE
> >> + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
> OTHER
> >> + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
> >> ARISING FROM,
> >> + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
> DEALINGS
> >> IN
> >> + * THE SOFTWARE.
> >> + *
> >> +
> >> *===--------------------------------------------------------
> ---------------===
> >> + */
> >> +
> >> +#ifndef __CLANG_CUDA_WRAPPERS_ALGORITHM
> >> +#define __CLANG_CUDA_WRAPPERS_ALGORITHM
> >> +
> >> +// This header defines __device__ overloads of std::min/max, but only
> if
> >> we're
> >> +// <= C++11. In C++14, these functions are constexpr, and so are
> >> implicitly
> >> +// __host__ __device__.
> >> +//
> >> +// We don't support the initializer_list overloads because
> >> +// initializer_list::begin() and end() are not __host__ __device__
> >> functions.
> >> +//
> >> +// When compiling in C++14 mode, we could force std::min/max to have
> >> different
> >> +// implementations for host and device, by declaring the device
> overloads
> >> +// before the constexpr overloads appear. We choose not to do this
> >> because
> >> +
> >> +// a) why write our own implementation when we can use one from the
> >> standard
> >> +// library? and
> >> +// b) libstdc++ is evil and declares min/max inside a header that is
> >> included
> >> +// *before* we include <algorithm>. So we'd have to
> unconditionally
> >> +// declare our __device__ overloads of min/max, but that would
> >> pollute
> >> +// things for people who choose not to include <algorithm>.
> >> +
> >> +#include_next <algorithm>
> >> +
> >> +#if __cplusplus <= 201103L
> >> +
> >> +// We need to define these overloads in exactly the namespace our
> >> standard
> >> +// library uses (including the right inline namespace), otherwise they
> >> won't be
> >> +// picked up by other functions in the standard library (e.g. functions
> >> in
> >> +// <complex>). Thus the ugliness below.
> >> +#ifdef _LIBCPP_BEGIN_NAMESPACE_STD
> >> +_LIBCPP_BEGIN_NAMESPACE_STD
> >> +#else
> >> +namespace std {
> >> +#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
> >> +_GLIBCXX_BEGIN_NAMESPACE_VERSION
> >> +#endif
> >> +#endif
> >> +
> >> +template <class __T, class __Cmp>
> >> +inline __device__ const __T &
> >> +max(const __T &__a, const __T &__b, __Cmp __cmp) {
> >> + return __cmp(__a, __b) ? __b : __a;
> >> +}
> >> +
> >> +template <class __T>
> >> +inline __device__ const __T &
> >> +max(const __T &__a, const __T &__b) {
> >> + return __a < __b ? __b : __a;
> >> +}
> >> +
> >> +template <class __T, class __Cmp>
> >> +inline __device__ const __T &
> >> +min(const __T &__a, const __T &__b, __Cmp __cmp) {
> >> + return __cmp(__b, __a) ? __b : __a;
> >> +}
> >> +
> >> +template <class __T>
> >> +inline __device__ const __T &
> >> +min(const __T &__a, const __T &__b) {
> >> + return __a < __b ? __b : __a;
> >> +}
> >> +
> >> +#ifdef _LIBCPP_END_NAMESPACE_STD
> >> +_LIBCPP_END_NAMESPACE_STD
> >> +#else
> >> +#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
> >> +_GLIBCXX_END_NAMESPACE_VERSION
> >> +#endif
> >> +} // namespace std
> >> +#endif
> >> +
> >> +#endif // __cplusplus <= 201103L
> >> +#endif // __CLANG_CUDA_WRAPPERS_ALGORITHM
> >>
> >> Added: cfe/trunk/lib/Headers/cuda_wrappers/complex
> >> URL:
> >> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/
> cuda_wrappers/complex?rev=283680&view=auto
> >>
> >> ============================================================
> ==================
> >> --- cfe/trunk/lib/Headers/cuda_wrappers/complex (added)
> >> +++ cfe/trunk/lib/Headers/cuda_wrappers/complex Sat Oct 8 17:16:12
> 2016
> >> @@ -0,0 +1,79 @@
> >> +/*===---- complex - CUDA wrapper for <complex>
> >> ------------------------------===
> >> + *
> >> + * Permission is hereby granted, free of charge, to any person
> obtaining
> >> a copy
> >> + * of this software and associated documentation files (the
> "Software"),
> >> to deal
> >> + * in the Software without restriction, including without limitation
> the
> >> rights
> >> + * to use, copy, modify, merge, publish, distribute, sublicense, and/or
> >> sell
> >> + * copies of the Software, and to permit persons to whom the Software
> is
> >> + * furnished to do so, subject to the following conditions:
> >> + *
> >> + * The above copyright notice and this permission notice shall be
> >> included in
> >> + * all copies or substantial portions of the Software.
> >> + *
> >> + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
> >> EXPRESS OR
> >> + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
> >> MERCHANTABILITY,
> >> + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT
> >> SHALL THE
> >> + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
> OTHER
> >> + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
> >> ARISING FROM,
> >> + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
> DEALINGS
> >> IN
> >> + * THE SOFTWARE.
> >> + *
> >> +
> >> *===--------------------------------------------------------
> ---------------===
> >> + */
> >> +
> >> +#pragma once
> >> +
> >> +// Wrapper around <complex> that forces its functions to be __host__
> >> +// __device__.
> >> +
> >> +// First, include host-only headers we think are likely to be included
> by
> >> +// <complex>, so that the pragma below only applies to <complex>
> itself.
> >> +#if __cplusplus >= 201103L
> >> +#include <type_traits>
> >> +#endif
> >> +#include <stdexcept>
> >> +#include <cmath>
> >> +#include <sstream>
> >> +
> >> +// Next, include our <algorithm> wrapper, to ensure that device
> overloads
> >> of
> >> +// std::min/max are available.
> >> +#include <algorithm>
> >> +
> >> +#pragma clang force_cuda_host_device begin
> >> +
> >> +// When compiling for device, ask libstdc++ to use its own implements
> of
> >> +// complex functions, rather than calling builtins (which resolve to
> >> library
> >> +// functions that don't exist when compiling CUDA device code).
> >> +//
> >> +// This is a little dicey, because it causes libstdc++ to define a
> >> different
> >> +// set of overloads on host and device.
> >> +//
> >> +// // Present only when compiling for host.
> >> +// __host__ __device__ void complex<float> sin(const complex<float>&
> x)
> >> {
> >> +// return __builtin_csinf(x);
> >> +// }
> >> +//
> >> +// // Present when compiling for host and for device.
> >> +// template <typename T>
> >> +// void __host__ __device__ complex<T> sin(const complex<T>& x) {
> >> +// return complex<T>(sin(x.real()) * cosh(x.imag()),
> >> +// cos(x.real()), sinh(x.imag()));
> >> +// }
> >> +//
> >> +// This is safe because when compiling for device, all function calls
> in
> >> +// __host__ code to sin() will still resolve to *something*, even if
> they
> >> don't
> >> +// resolve to the same function as they resolve to when compiling for
> >> host. We
> >> +// don't care that they don't resolve to the right function because we
> >> won't
> >> +// codegen this host code when compiling for device.
> >> +
> >> +#pragma push_macro("_GLIBCXX_USE_C99_COMPLEX")
> >> +#pragma push_macro("_GLIBCXX_USE_C99_COMPLEX_TR1")
> >> +#define _GLIBCXX_USE_C99_COMPLEX 0
> >> +#define _GLIBCXX_USE_C99_COMPLEX_TR1 0
> >> +
> >> +#include_next <complex>
> >> +
> >> +#pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX_TR1")
> >> +#pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX")
> >> +
> >> +#pragma clang force_cuda_host_device end
> >>
> >>
> >> _______________________________________________
> >> cfe-commits mailing list
> >> cfe-commits at lists.llvm.org
> >> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
> >
> >
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20161011/3fc7e997/attachment-0001.html>
More information about the cfe-commits
mailing list