This revision was automatically updated to reflect the committed changes.
Closed by commit rL287288: [CUDA] Wrapper header changes necessary to support
MacOS. (authored by jlebar).
Changed prior to commit:
https://reviews.llvm.org/D26780?vs=78298&id=78438#toc
Repository:
rL LLVM
https://reviews.llvm.org/D26780
Files:
cfe/trunk/lib/Headers/__clang_cuda_cmath.h
cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
Index: cfe/trunk/lib/Headers/__clang_cuda_cmath.h
===================================================================
--- cfe/trunk/lib/Headers/__clang_cuda_cmath.h
+++ cfe/trunk/lib/Headers/__clang_cuda_cmath.h
@@ -75,7 +75,10 @@
__DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }
__DEVICE__ bool isinf(double __x) { return ::__isinf(__x); }
__DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); }
-__DEVICE__ bool isfinite(double __x) { return ::__finite(__x); }
+// For inscrutable reasons, __finite(), the double-precision version of
+// __finitef, does not exist when compiling for MacOS. __isfinited is
available
+// everywhere and is just as good.
+__DEVICE__ bool isfinite(double __x) { return ::__isfinited(__x); }
__DEVICE__ bool isgreater(float __x, float __y) {
return __builtin_isgreater(__x, __y);
}
@@ -141,7 +144,7 @@
return ::powi(__base, __iexp);
}
__DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); }
-__DEVICE__ bool signbit(double __x) { return ::__signbit(__x); }
+__DEVICE__ bool signbit(double __x) { return ::__signbitd(__x); }
__DEVICE__ float sin(float __x) { return ::sinf(__x); }
__DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
__DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); }
Index: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
===================================================================
--- cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
+++ cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
@@ -121,6 +121,15 @@
#undef __cxa_vec_delete3
#undef __cxa_pure_virtual
+// math_functions.hpp expects this host function be defined on MacOS, but it
+// ends up not being there because of the games we play here. Just define it
+// ourselves; it's simple enough.
+#ifdef __APPLE__
+inline __host__ double __signbitd(double x) {
+ return std::signbit(x);
+}
+#endif
+
// We need decls for functions in CUDA's libdevice with __device__
// attribute only. Alas they come either as __host__ __device__ or
// with no attributes at all. To work around that, define __CUDA_RTC__
Index: cfe/trunk/lib/Headers/__clang_cuda_cmath.h
===================================================================
--- cfe/trunk/lib/Headers/__clang_cuda_cmath.h
+++ cfe/trunk/lib/Headers/__clang_cuda_cmath.h
@@ -75,7 +75,10 @@
__DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }
__DEVICE__ bool isinf(double __x) { return ::__isinf(__x); }
__DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); }
-__DEVICE__ bool isfinite(double __x) { return ::__finite(__x); }
+// For inscrutable reasons, __finite(), the double-precision version of
+// __finitef, does not exist when compiling for MacOS. __isfinited is available
+// everywhere and is just as good.
+__DEVICE__ bool isfinite(double __x) { return ::__isfinited(__x); }
__DEVICE__ bool isgreater(float __x, float __y) {
return __builtin_isgreater(__x, __y);
}
@@ -141,7 +144,7 @@
return ::powi(__base, __iexp);
}
__DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); }
-__DEVICE__ bool signbit(double __x) { return ::__signbit(__x); }
+__DEVICE__ bool signbit(double __x) { return ::__signbitd(__x); }
__DEVICE__ float sin(float __x) { return ::sinf(__x); }
__DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
__DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); }
Index: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
===================================================================
--- cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
+++ cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
@@ -121,6 +121,15 @@
#undef __cxa_vec_delete3
#undef __cxa_pure_virtual
+// math_functions.hpp expects this host function be defined on MacOS, but it
+// ends up not being there because of the games we play here. Just define it
+// ourselves; it's simple enough.
+#ifdef __APPLE__
+inline __host__ double __signbitd(double x) {
+ return std::signbit(x);
+}
+#endif
+
// We need decls for functions in CUDA's libdevice with __device__
// attribute only. Alas they come either as __host__ __device__ or
// with no attributes at all. To work around that, define __CUDA_RTC__
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
