jlebar created this revision.
Herald added a subscriber: sanjoy.
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.
https://reviews.llvm.org/D39702
Files:
External/CUDA/cmath.cu
External/CUDA/math_h.cu
Index: External/CUDA/math_h.cu
===================================================================
--- External/CUDA/math_h.cu
+++ External/CUDA/math_h.cu
@@ -86,8 +86,6 @@
__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(); }
@@ -1347,57 +1345,6 @@
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), "");
@@ -1647,8 +1594,6 @@
test_lround();
test_nan();
test_nearbyint();
- test_nextafter();
- test_nexttoward();
test_remainder();
test_remquo();
test_rint();
Index: External/CUDA/cmath.cu
===================================================================
--- External/CUDA/cmath.cu
+++ External/CUDA/cmath.cu
@@ -88,8 +88,6 @@
__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(); }
@@ -1372,55 +1370,6 @@
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), "");
@@ -1670,8 +1619,6 @@
test_lround();
test_nan();
test_nearbyint();
- test_nextafter();
- test_nexttoward();
test_remainder();
test_remquo();
test_rint();
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
