[CUDA] Remove implementations of nexttoward and nextafter.

[jlebar]

Summary

__builtin_nexttoward &co lower to a libcall, e.g. nexttowardf(), that
CUDA does not have.

Rather than try to implement it, we simply remove these calls -- nvcc
doesn't support them either.

[tra]

Libdevice does provide implementation for __nv_nextafterf() and __nv_nextafter() and it has corresponding wrappers in math_functions.h[pp].

Perhaps we should keep nextafter around.

[jlebar]

Libdevice does provide implementation for nv_nextafterf() and nv_nextafter() and it has corresponding wrappers in math_functions.h[pp].

Perhaps we should keep nextafter around.

Oh, that changes things, since nexttoward is basically the same thing as nextafter. Thanks.

Let's scrap this and try again.

clang/lib/Headers/__clang_cuda_cmath.h

__DEVICE__ bool isunordered(double __x, double __y) {
   return __builtin_isunordered(__x, __y);
 }
 __DEVICE__ float ldexp(float __arg, int __exp) {
   return ::ldexpf(__arg, __exp);
 }
 __DEVICE__ float log(float __x) { return ::logf(__x); }
 __DEVICE__ float log10(float __x) { return ::log10f(__x); }
 __DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); }
-__DEVICE__ float nexttoward(float __from, double __to) {
-  return __builtin_nexttowardf(__from, __to);
-}
-__DEVICE__ double nexttoward(double __from, double __to) {
-  return __builtin_nexttoward(__from, __to);
-}
-__DEVICE__ float nexttowardf(float __from, double __to) {
-  return __builtin_nexttowardf(__from, __to);
-}
 __DEVICE__ float pow(float __base, float __exp) {
   return ::powf(__base, __exp);
 }
 __DEVICE__ float pow(float __base, int __iexp) {
   return ::powif(__base, __iexp);
 }
 __DEVICE__ double pow(double __base, int __iexp) {
   return ::powi(__base, __iexp);
 }
 __DEVICE__ bool signbit(float __x) { return ::__signbitf(__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); }
 __DEVICE__ float tan(float __x) { return ::tanf(__x); }
 __DEVICE__ float tanh(float __x) { return ::tanhf(__x); }
 
+// Notably missing above are nexttoward and nextafter.  We omit them because
+// libdevice doesn't provide an implementation, and we don't want to be in the
+// business of implementing tricky libm functions in this header.
+
 // Now we've defined everything we promised we'd define in
 // __clang_cuda_math_forward_declares.h.  We need to do two additional things to
 // fix up our math functions.
 //
 // 1) Define __device__ overloads for e.g. sin(int).  The CUDA headers define
 //    only sin(float) and sin(double), which means that e.g. sin(0) is
 //    ambiguous.
 //

 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, log1p)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, log2)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, logb)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(long long, llrint)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(long long, llround)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(long, lrint)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(long, lround)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, nearbyint);
-__CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, nextafter);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, pow);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, remainder);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, rint);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, round);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(bool, signbit)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, sin)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, sinh)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, sqrt)

 
 template <typename __T>
 __DEVICE__ typename __clang_cuda_enable_if<std::numeric_limits<__T>::is_integer,
                                            double>::type
 ldexp(__T __x, int __exp) {
   return std::ldexp((double)__x, __exp);
 }
 
-template <typename __T>
-__DEVICE__ typename __clang_cuda_enable_if<std::numeric_limits<__T>::is_integer,
-                                           double>::type
-nexttoward(__T __from, double __to) {
-  return std::nexttoward((double)__from, __to);
-}
-
 template <typename __T1, typename __T2>
 __DEVICE__ typename __clang_cuda_enable_if<
     std::numeric_limits<__T1>::is_specialized &&
         std::numeric_limits<__T2>::is_specialized,
     double>::type
 remquo(__T1 __x, __T2 __y, int *__quo) {
   return std::remquo((double)__x, (double)__y, __quo);
 }

 using ::log;
 using ::log10;
 using ::log1p;
 using ::log2;
 using ::logb;
 using ::lrint;
 using ::lround;
 using ::nearbyint;
-using ::nextafter;
-using ::nexttoward;
 using ::pow;
 using ::remainder;
 using ::remquo;
 using ::rint;
 using ::round;
 using ::scalbln;
 using ::scalbn;
 using ::signbit;

 using ::log1pf;
 using ::log2f;
 using ::logbf;
 using ::logf;
 using ::lrintf;
 using ::lroundf;
 using ::modff;
 using ::nearbyintf;
-using ::nextafterf;
-using ::nexttowardf;
-using ::nexttowardf;
 using ::powf;
 using ::remainderf;
 using ::remquof;
 using ::rintf;
 using ::roundf;
 using ::scalblnf;
 using ::scalbnf;
 using ::sinf;

clang/lib/Headers/__clang_cuda_math_forward_declares.h

 __DEVICE__ long lround(float);
 __DEVICE__ long long llround(float); // No llround(double).
 __DEVICE__ double modf(double, double *);
 __DEVICE__ float modf(float, float *);
 __DEVICE__ double nan(const char *);
 __DEVICE__ float nanf(const char *);
 __DEVICE__ double nearbyint(double);
 __DEVICE__ float nearbyint(float);
-__DEVICE__ double nextafter(double, double);
-__DEVICE__ float nextafter(float, float);
-__DEVICE__ double nexttoward(double, double);
-__DEVICE__ float nexttoward(float, double);
-__DEVICE__ float nexttowardf(float, double);
 __DEVICE__ double pow(double, double);
 __DEVICE__ double pow(double, int);
 __DEVICE__ float pow(float, float);
 __DEVICE__ float pow(float, int);
 __DEVICE__ double remainder(double, double);
 __DEVICE__ float remainder(float, float);
 __DEVICE__ double remquo(double, double, int *);
 __DEVICE__ float remquo(float, float, int *);

 __DEVICE__ float tan(float);
 __DEVICE__ double tanh(double);
 __DEVICE__ float tanh(float);
 __DEVICE__ double tgamma(double);
 __DEVICE__ float tgamma(float);
 __DEVICE__ double trunc(double);
 __DEVICE__ float trunc(float);
 
+// Notably missing above are nexttoward and nextafter, which we don't define on
+// the device side, because libdevice doesn't give us an implementation, and we
+// don't want to be in the business of writing one ourselves.
+
 // 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 {

 using ::logb;
 using ::lrint;
 using ::lround;
 using ::llround;
 using ::modf;
 using ::nan;
 using ::nanf;
 using ::nearbyint;
-using ::nextafter;
-using ::nexttoward;
 using ::pow;
 using ::remainder;
 using ::remquo;
 using ::rint;
 using ::round;
 using ::scalbln;
 using ::scalbn;
 using ::signbit;