https://reviews.llvm.org/D25403 for adding constexpr to the new libc++ functions.
On Sat, Oct 8, 2016 at 8:56 PM, Justin Lebar <jle...@google.com> wrote: >> > The fix is not as simple as simply changing our implementation of >> e.g. >> std::isnan to call __builtin_isnanf, because we also would want to >> fix >> ::isnanf, >> >> No, if I understand what you're saying, you specifically wouldn't. > > I understand how this is feasible on the CPU, because ::isnan is a > library function that can never be inlined. But on the GPU, these > library functions are (at the moment) always declared inline. That > seems to complicate this idea. > > Right now ::isnan(x) is going to call __nv_isnan(x), which computes > abs(x). If we pass -ffast-math, the compiler will be able to assume > that abs(x) is not nan. I guess you're saying that we would need to > special-case __nv_isnan so that -ffast-math is always off > (essentially). But, what if it gets inlined? > > It looks like libstdc++'s std::isnan calls __builtin_isnan (same for > its std::isinf), and its ::isnan is an alias for std::isnan. So > libstdc++'s isnan is going to return false with -ffast-math (or anyway > it will do the same thing as the builtin functions, which aiui is what > you're proposing libc++'s isnan *not* do). > >> This is important for use cases where, for example, even though the user >> might want fast math, they still need to check their inputs for NaNs. > > Since this isn't going to work with libstdc++, and it relies on not > doing anything that the compiler might construe as "arithmetic" on the > value, this seems pretty dicey to me. One could instead compile a > separate TU without -ffast-math and do all their validation there? > I'd have a lot more confidence in that working today, continuing to > work tomorrow, and being portable across compilers and standard > libraries. > > I don't mean to relitigate https://reviews.llvm.org/D18639, but I am > not convinced that libc++'s isnan should have a path that returns true > with -ffast-math, given that > > * libstdc++'s isnan will always return false with -ffast-math, > * it's at best complicated for us to make this work if you can inline > the body of isnan (as we can on the GPU), > * it's at best complicated for users to write "correct" C++ that > calls isnan with -ffast-math, especially if they want their code to > continue to work in the future in the face of changing compilers > (-ffast-math is not specified anywhere, so who knows what it means), > and > * there's a relatively simple workaround (use a separate TU) that > sidesteps all these problems. > > I'm not saying we should go in and change libc++'s CPU implementation > of isnan to call the builtin. I'll leave that up to people who care > about CPU code. But at least on the GPU, it still makes sense to me > to fix the problem you originally identified by making > std::/::isnan/isinf always return false/true with -ffast-math. Which > I think we should be able to do with the intrinsic upgrade I > originally suggested. > > On a separate note: Can we make __libcpp_isnan and __libcpp_isinf > constexpr? This will make them implicitly host+device functions, > solving the problem on the GPU. Otherwise I may have to reimplement > these functions in a header, and that's lame. Although I am clearly > not above that. :) > > On Sat, Oct 8, 2016 at 6:50 PM, Hal Finkel <hfin...@anl.gov> wrote: >> ----- Original Message ----- >>> From: "Justin Lebar" <jle...@google.com> >>> To: "Hal Finkel" <hfin...@anl.gov> >>> Cc: "Clang Commits" <cfe-commits@lists.llvm.org> >>> Sent: Saturday, October 8, 2016 6:16:12 PM >>> Subject: Re: r283680 - [CUDA] Support <complex> and std::min/max on the >>> device. >>> >>> Hal, >>> >>> On NVPTX, these functions eventually get resolved to function calls >>> in >>> libdevice, e.g. __nv_isinff and __nv_isnanf. >>> >>> llvm does not do a good job understanding the body of e.g. >>> __nvvm_isnanf, because it uses nvptx-specific intrinsic functions, >>> notably @llvm.nvvm.fabs.f. These are opaque to the LLVM optimizer. >>> >>> The fix is not as simple as simply changing our implementation of >>> e.g. >>> std::isnan to call __builtin_isnanf, because we also would want to >>> fix >>> ::isnanf, >> >> No, if I understand what you're saying, you specifically wouldn't. We had a >> discussion about this on the review thread(s) that led to r283051, and while >> we want to elide the checks inside the mathematical functions, we don't want >> to replace isnan itself with something that will get optimized away. We want >> to keep the ability for the user to explicitly check for NaNs, etc. even if >> we don't want those checks to appear inside of mathematical operations. This >> is important for use cases where, for example, even though the user might >> want fast math, they still need to check their inputs for NaNs. >> >> -Hal >> >>> but we can't override that implementation without some >>> major >>> surgery on the nvptx headers. >>> >>> David Majnemer and I talked about one way to fix this, namely by >>> using >>> IR intrinsic upgrades to replace the opaque nvptx intrinsics with >>> LLVM >>> intrinsics. LLVM would then be able to understand these intrinsics >>> and optimize them. We would reap benefits not just for std::isnan, >>> but also e.g. constant-folding calls like std::abs that also >>> eventually end up in libnvvm. >>> >>> I did the first half of this work, by adding lowerings for the >>> various >>> LLVM intrinsics to the NVPTX backend [1]. But David is now busy with >>> other things and hasn't been able to help with the second half, >>> namely >>> using IR upgrades to replace the nvptx target-specific intrinsics >>> with >>> generalized LLVM intrinsics. Perhaps this is something you'd be able >>> to help with? >>> >>> In any case, using builtins here without fixing std::isnan and >>> ::isnan >>> feels to me to be the wrong solution. It seems to me that we should >>> be able to rely on std::isnan and friends being fast, and if they're >>> not, we should fix that. Using builtins here would be "cheating" to >>> make our implementation faster than user code. >>> >>> I'll note, separately, that on x86, clang does not seem to >>> constant-fold std::isinf or __builtin_isinff to false with >>> -ffast-math >>> -ffinite-math-only. GCC can do it. Clang gets std::isnan. >>> https://godbolt.org/g/vZB55a >>> >>> By the way, the changes you made to libc++ unfortunately break this >>> patch with libc++, because e.g. __libcpp_isnan is not a device >>> function. I'll have to think about how to fix that -- I may send you >>> a patch. >>> >>> Regards, >>> -Justin >>> >>> [1] https://reviews.llvm.org/D24300 >>> >>> On Sat, Oct 8, 2016 at 3:36 PM, Hal Finkel <hfin...@anl.gov> wrote: >>> > Hi Justin, >>> > >>> > This is neat! >>> > >>> > I see a bunch of uses of std::isinf, etc. here. It tends to be >>> > important that, when using -ffast-math (or -ffinite-math-only) >>> > these checks get optimized away. Can you please check that they >>> > do? If not, you might mirror what I've done in r283051 for libc++, >>> > which is similar to what libstdc++ ends up doing, so that we use >>> > __builtin_isnan/isinf/isfinite. >>> > >>> > Thanks again, >>> > Hal >>> > >>> > ----- Original Message ----- >>> >> From: "Justin Lebar via cfe-commits" <cfe-commits@lists.llvm.org> >>> >> To: cfe-commits@lists.llvm.org >>> >> Sent: Saturday, October 8, 2016 5:16:13 PM >>> >> Subject: r283680 - [CUDA] Support <complex> and std::min/max on >>> >> the device. >>> >> >>> >> 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@lists.llvm.org >>> >> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits >>> >> >>> > >>> > -- >>> > Hal Finkel >>> > Lead, Compiler Technology and Programming Languages >>> > Leadership Computing Facility >>> > Argonne National Laboratory >>> >> >> -- >> Hal Finkel >> Lead, Compiler Technology and Programming Languages >> Leadership Computing Facility >> Argonne National Laboratory _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
