These were reverted in r283753 and r283747.
The first patch added a header to the root 'Headers' install directory,
instead of into 'Headers/cuda_wrappers'. This was fixed in the second
patch, but by then the damage was done: The bad header stayed in the
'Headers' directory, continuing to break the build.
We reverted both patches in an attempt to fix things, but that still
didn't get rid of the header, so the Windows boostrap build remained
broken.
It's probably worth fixing up our cmake logic to remove things from the
install dirs, but in the meantime, re-land these patches, since we
believe they no longer have this bug.
llvm-svn: 283907
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;
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
xtestintrin.h
)
+set(cuda_wrapper_files
+ cuda_wrappers/algorithm
+ cuda_wrappers/complex
+)
+
set(output_dir ${LLVM_LIBRARY_OUTPUT_INTDIR}/clang/${CLANG_VERSION}/include)
# Generate arm_neon.h
SOURCE ${CLANG_SOURCE_DIR}/include/clang/Basic/arm_neon.td)
set(out_files)
-foreach( f ${files} )
+foreach( f ${files} ${cuda_wrapper_files} )
set( src ${CMAKE_CURRENT_SOURCE_DIR}/${f} )
set( dst ${output_dir}/${f} )
add_custom_command(OUTPUT ${dst}
PERMISSIONS OWNER_READ OWNER_WRITE GROUP_READ WORLD_READ
DESTINATION lib${LLVM_LIBDIR_SUFFIX}/clang/${CLANG_VERSION}/include)
+install(
+ FILES ${cuda_wrapper_files}
+ COMPONENT clang-headers
+ PERMISSIONS OWNER_READ OWNER_WRITE GROUP_READ WORLD_READ
+ DESTINATION lib${LLVM_LIBDIR_SUFFIX}/clang/${CLANG_VERSION}/include/cuda_wrappers)
+
if (NOT CMAKE_CONFIGURATION_TYPES) # don't add this for IDE's.
add_custom_target(install-clang-headers
DEPENDS clang-headers
--- /dev/null
+/*===-- __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
#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
--- /dev/null
+/*===---- 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
--- /dev/null
+/*===---- 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
projects / platform / upstream / llvm.git / commitdiff