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