[test-suite] [CUDA] Delete nexttoward and nextafter CUDA tests, because these functions have no implementations.

[jlebar]

Summary

These functions never actually worked -- if you called them, you'd get a
ptxas error when llvm tried to make a libcall to e.g. nexttowardf.

This test only compiled because the calls were DCE'ed. Sadly my
laziness in not checking the return value of these functions resulted in
this long-standing bug.

On the upside, we check the return values of all other math functions,
so this bug is limited to nexttoward and nextafter.

[jlebar]

Abandoning because it looks like we *can* implement nexttoward in CUDA, see D39703. Will send a different patch.

External/CUDA/cmath.cu

 __device__ Ambiguous llrint(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous llround(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous log1p(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous log2(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous logb(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous lrint(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous lround(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous nearbyint(Ambiguous){ return Ambiguous(); }
-__device__ Ambiguous nextafter(Ambiguous, Ambiguous){ return Ambiguous(); }
-__device__ Ambiguous nexttoward(Ambiguous, Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous remainder(Ambiguous, Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous remquo(Ambiguous, Ambiguous, int*){ return Ambiguous(); }
 __device__ Ambiguous rint(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous round(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous scalbln(Ambiguous, Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous scalbn(Ambiguous, Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous tgamma(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous trunc(Ambiguous){ return Ambiguous(); }
__device__ void test_nearbyint()
     static_assert((std::is_same<decltype(std::nearbyint((double)0)), double>::value), "");
     static_assert((std::is_same<decltype(std::nearbyintf(0)), float>::value), "");
     static_assert((std::is_same<decltype(nearbyint(Ambiguous())), Ambiguous>::value), "");
     assert(std::nearbyint(1) == 1);
     assert(std::nearbyint(1.) == 1);
     assert(std::nearbyint(1.f) == 1);
 }
 
-__device__ void test_nextafter()
-{
-    static_assert((std::is_same<decltype(std::nextafter((float)0, (float)0)), float>::value), "");
-    static_assert((std::is_same<decltype(std::nextafter((bool)0, (float)0)), double>::value), "");
-    static_assert((std::is_same<decltype(std::nextafter((unsigned short)0, (double)0)), double>::value), "");
-    static_assert((std::is_same<decltype(std::nextafter((float)0, (unsigned int)0)), double>::value), "");
-    static_assert((std::is_same<decltype(std::nextafter((double)0, (long)0)), double>::value), "");
-    static_assert((std::is_same<decltype(std::nextafter((int)0, (long long)0)), double>::value), "");
-    static_assert((std::is_same<decltype(std::nextafter((int)0, (unsigned long long)0)), double>::value), "");
-    static_assert((std::is_same<decltype(std::nextafter((double)0, (double)0)), double>::value), "");
-    static_assert((std::is_same<decltype(std::nextafter((float)0, (double)0)), double>::value), "");
-    static_assert((std::is_same<decltype(std::nextafterf(0,0)), float>::value), "");
-    static_assert((std::is_same<decltype(std::nextafter((int)0, (int)0)), double>::value), "");
-    static_assert((std::is_same<decltype(nextafter(Ambiguous(), Ambiguous())), Ambiguous>::value), "");
-
-    // Invoke all our overloads, even if we can't be bothered to check the
-    // results.
-    std::nextafter(0, 1);
-    std::nextafter(0, 1.);
-    std::nextafter(0, 1.f);
-
-    std::nextafter(0., 1);
-    std::nextafter(0., 1.);
-    std::nextafter(0., 1.f);
-
-    std::nextafter(0.f, 1);
-    std::nextafter(0.f, 1.);
-    std::nextafter(0.f, 1.f);
-}
-
-__device__ void test_nexttoward()
-{
-    static_assert((std::is_same<decltype(nexttoward(Ambiguous(), Ambiguous())), Ambiguous>::value), "");
-
-    // Invoke all our overloads, even if we can't be bothered to check the
-    // results.
-    std::nexttoward(0, 1);
-    std::nexttoward(0, 1.);
-    std::nexttoward(0, 1.f);
-
-    std::nexttoward(0., 1);
-    std::nexttoward(0., 1.);
-    std::nexttoward(0., 1.f);
-
-    std::nexttoward(0.f, 1);
-    std::nexttoward(0.f, 1.);
-    std::nexttoward(0.f, 1.f);
-}
-
 __device__ void test_remainder()
 {
     static_assert((std::is_same<decltype(std::remainder((float)0, (float)0)), float>::value), "");
     static_assert((std::is_same<decltype(std::remainder((bool)0, (float)0)), double>::value), "");
     static_assert((std::is_same<decltype(std::remainder((unsigned short)0, (double)0)), double>::value), "");
     static_assert((std::is_same<decltype(std::remainder((float)0, (unsigned int)0)), double>::value), "");
     static_assert((std::is_same<decltype(std::remainder((double)0, (long)0)), double>::value), "");
     static_assert((std::is_same<decltype(std::remainder((int)0, (long long)0)), double>::value), "");
__global__ void tests()
     test_llround();
     test_log1p();
     test_log2();
     test_logb();
     test_lrint();
     test_lround();
     test_nan();
     test_nearbyint();
-    test_nextafter();
-    test_nexttoward();
     test_remainder();
     test_remquo();
     test_rint();
     test_round();
     test_scalbln();
     test_scalbn();
     test_tgamma();
     test_trunc();

External/CUDA/math_h.cu

 __device__ Ambiguous llrint(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous llround(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous log1p(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous log2(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous logb(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous lrint(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous lround(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous nearbyint(Ambiguous){ return Ambiguous(); }
-__device__ Ambiguous nextafter(Ambiguous, Ambiguous){ return Ambiguous(); }
-__device__ Ambiguous nexttoward(Ambiguous, Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous remainder(Ambiguous, Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous remquo(Ambiguous, Ambiguous, int*){ return Ambiguous(); }
 __device__ Ambiguous rint(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous round(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous scalbln(Ambiguous, Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous scalbn(Ambiguous, Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous tgamma(Ambiguous){ return Ambiguous(); }
 __device__ Ambiguous trunc(Ambiguous){ return Ambiguous(); }
__device__ void test_nearbyint()
     static_assert((std::is_same<decltype(nearbyint((double)0)), double>::value), "");
     static_assert((std::is_same<decltype(nearbyintf(0)), float>::value), "");
     static_assert((std::is_same<decltype(nearbyint(Ambiguous())), Ambiguous>::value), "");
     assert(nearbyint(1) == 1);
     assert(nearbyint(1.) == 1);
     assert(nearbyint(1.f) == 1);
 }
 
-__device__ void test_nextafter()
-{
-    static_assert((std::is_same<decltype(nextafter((float)0, (float)0)), float>::value), "");
-    static_assert((std::is_same<decltype(nextafter((bool)0, (float)0)), double>::value), "");
-    static_assert((std::is_same<decltype(nextafter((unsigned short)0, (double)0)), double>::value), "");
-    static_assert((std::is_same<decltype(nextafter((float)0, (unsigned int)0)), double>::value), "");
-    static_assert((std::is_same<decltype(nextafter((double)0, (long)0)), double>::value), "");
-    static_assert((std::is_same<decltype(nextafter((int)0, (long long)0)), double>::value), "");
-    static_assert((std::is_same<decltype(nextafter((int)0, (unsigned long long)0)), double>::value), "");
-    static_assert((std::is_same<decltype(nextafter((double)0, (double)0)), double>::value), "");
-    static_assert((std::is_same<decltype(nextafter((float)0, (double)0)), double>::value), "");
-    static_assert((std::is_same<decltype(nextafterf(0,0)), float>::value), "");
-    static_assert((std::is_same<decltype(nextafter((int)0, (int)0)), double>::value), "");
-    static_assert((std::is_same<decltype(nextafter(Ambiguous(), Ambiguous())), Ambiguous>::value), "");
-    //assert(nextafter(0,1) == hexfloat<double>(0x1, 0, -1074));
-
-    // Invoke all our overloads, even if we can't be bothered to check the
-    // results.
-    nextafter(0, 1);
-    nextafter(0, 1.);
-    nextafter(0, 1.f);
-
-    nextafter(0., 1);
-    nextafter(0., 1.);
-    nextafter(0., 1.f);
-
-    nextafter(0.f, 1);
-    nextafter(0.f, 1.);
-    nextafter(0.f, 1.f);
-}
-
-__device__ void test_nexttoward()
-{
-    static_assert((std::is_same<decltype(nexttoward(Ambiguous(), Ambiguous())), Ambiguous>::value), "");
-    //assert(nexttoward(0, 1) == hexfloat<double>(0x1, 0, -1074));
-
-    // Invoke all our overloads, even if we can't be bothered to check the
-    // results.
-    nexttoward(0, 1);
-    nexttoward(0, 1.);
-    nexttoward(0, 1.f);
-
-    nexttoward(0., 1);
-    nexttoward(0., 1.);
-    nexttoward(0., 1.f);
-
-    nexttoward(0.f, 1);
-    nexttoward(0.f, 1.);
-    nexttoward(0.f, 1.f);
-}
-
 __device__ void test_remainder()
 {
     static_assert((std::is_same<decltype(remainder((float)0, (float)0)), float>::value), "");
     static_assert((std::is_same<decltype(remainder((bool)0, (float)0)), double>::value), "");
     static_assert((std::is_same<decltype(remainder((unsigned short)0, (double)0)), double>::value), "");
     static_assert((std::is_same<decltype(remainder((float)0, (unsigned int)0)), double>::value), "");
     static_assert((std::is_same<decltype(remainder((double)0, (long)0)), double>::value), "");
     static_assert((std::is_same<decltype(remainder((int)0, (long long)0)), double>::value), "");
__global__ void tests()
     test_llround();
     test_log1p();
     test_log2();
     test_logb();
     test_lrint();
     test_lround();
     test_nan();
     test_nearbyint();
-    test_nextafter();
-    test_nexttoward();
     test_remainder();
     test_remquo();
     test_rint();
     test_round();
     test_scalbln();
     test_scalbn();
     test_tgamma();
     test_trunc();