Skip to content

Commit e8de527

Browse files
author
Thomas Grützmacher
authored
Merge improvements to isfinite
Improves and simplifies the implementation of the hand-written `isfinite` function. Now, it produces the same PTX code as the actual function (at least in a small example). Related PR: #465
2 parents 46a7381 + 67abdb9 commit e8de527

File tree

5 files changed

+140
-149
lines changed

5 files changed

+140
-149
lines changed

‎common/base/math.hpp.inc

Lines changed: 48 additions & 100 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,16 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
3030
OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
3131
******************************<GINKGO LICENSE>*******************************/
3232

33+
// We need this struct, because otherwise we would call a __host__ function in a
34+
// __device__ function (even though it is constexpr)
35+
template <typename T>
36+
struct device_numeric_limits {
37+
static constexpr auto inf = std::numeric_limits<T>::infinity();
38+
static constexpr auto max = std::numeric_limits<T>::max();
39+
static constexpr auto min = std::numeric_limits<T>::min();
40+
};
41+
42+
3343
namespace detail {
3444

3545

@@ -50,6 +60,38 @@ struct truncate_type_impl<thrust::complex<T>> {
5060
};
5161

5262

63+
/**
64+
* Checks if a given value is finite, meaning it is neither +/- infinity
65+
* nor NaN.
66+
*
67+
* @internal Should only be used if the provided one (from CUDA or HIP) can
68+
* not be used.
69+
* Designed to mirror the math function of CUDA (PTX code was
70+
* identical in the testcase).
71+
*
72+
* @tparam T type of the value to check
73+
*
74+
* @param value value to check
75+
*
76+
* returns `true` if the given value is finite, meaning it is neither
77+
* +/- infinity nor NaN.
78+
*/
79+
template <typename T>
80+
GKO_INLINE __device__ xstd::enable_if_t<!is_complex_s<T>::value, bool>
81+
custom_isfinite(T value)
82+
{
83+
constexpr T infinity{device_numeric_limits<T>::inf};
84+
return abs(value) < infinity;
85+
}
86+
87+
template <typename T>
88+
GKO_INLINE __device__ xstd::enable_if_t<is_complex_s<T>::value, bool>
89+
custom_isfinite(T value)
90+
{
91+
return custom_isfinite(value.real()) && custom_isfinite(value.imag());
92+
}
93+
94+
5395
} // namespace detail
5496

5597

@@ -66,103 +108,19 @@ struct truncate_type_impl<thrust::complex<T>> {
66108
(defined(__clang__) || defined(__ICC) || defined(__ICL))))
67109

68110

69-
namespace detail {
70-
71-
72-
/**
73-
* This structure can be used to get the exponent mask of a given floating
74-
* point type. Uses specialization to implement different types.
75-
*/
76-
template <typename T>
77-
struct mask_creator {};
78-
79-
template <>
80-
struct mask_creator<float> {
81-
using int_type = int32;
82-
static constexpr int_type number_exponent_bits = 8;
83-
static constexpr int_type number_significand_bits = 23;
84-
// integer representation of a floating point number, where all exponent
85-
// bits are set
86-
static constexpr int_type exponent_mask =
87-
((int_type{1} << number_exponent_bits) - 1) << number_significand_bits;
88-
static __device__ int_type reinterpret_int(const float &value)
89-
{
90-
return __float_as_int(value);
91-
}
92-
};
93-
94-
template <>
95-
struct mask_creator<double> {
96-
using int_type = int64;
97-
static constexpr int_type number_exponent_bits = 11;
98-
static constexpr int_type number_significand_bits = 52;
99-
// integer representation of a floating point number, where all exponent
100-
// bits are set
101-
static constexpr int_type exponent_mask =
102-
((int_type{1} << number_exponent_bits) - 1) << number_significand_bits;
103-
static __device__ int_type reinterpret_int(const double &value)
104-
{
105-
return __double_as_longlong(value);
106-
}
107-
};
108-
109-
110-
} // namespace detail
111-
112-
113-
/**
114-
* Checks if a given value is finite, meaning it is neither +/- infinity
115-
* nor NaN.
116-
*
117-
* @internal It checks if all exponent bits are set. If all are set, the
118-
* number either represents NaN or +/- infinity, meaning it is a
119-
* non-finite number.
120-
*
121-
* @param value value to check
122-
*
123-
* returns `true` if the given value is finite, meaning it is neither
124-
* +/- infinity nor NaN.
125-
*/
126-
#define GKO_DEFINE_ISFINITE_FOR_TYPE(_type) \
127-
GKO_INLINE __device__ bool isfinite(const _type &value) \
128-
{ \
129-
constexpr auto mask = detail::mask_creator<_type>::exponent_mask; \
130-
const auto re_int = \
131-
detail::mask_creator<_type>::reinterpret_int(value); \
132-
return (re_int & mask) != mask; \
111+
#define GKO_DEFINE_ISFINITE_FOR_TYPE(_type) \
112+
GKO_INLINE __device__ bool isfinite(const _type &value) \
113+
{ \
114+
return detail::custom_isfinite(value); \
133115
}
134116

135117
GKO_DEFINE_ISFINITE_FOR_TYPE(float)
136118
GKO_DEFINE_ISFINITE_FOR_TYPE(double)
119+
GKO_DEFINE_ISFINITE_FOR_TYPE(thrust::complex<float>)
120+
GKO_DEFINE_ISFINITE_FOR_TYPE(thrust::complex<double>)
137121
#undef GKO_DEFINE_ISFINITE_FOR_TYPE
138122

139123

140-
/**
141-
* Checks if all components of a complex value are finite, meaning they are
142-
* neither +/- infinity nor NaN.
143-
*
144-
* @internal required for the clang compiler. This function will be used rather
145-
* than the `isfinite` function in the public `math.hpp` because
146-
* there is no template parameter, so it is prefered during lookup.
147-
*
148-
* @tparam T complex type of the value to check
149-
*
150-
* @param value complex value to check
151-
*
152-
* returns `true` if both components of the given value are finite, meaning
153-
* they are neither +/- infinity nor NaN.
154-
*/
155-
#define GKO_DEFINE_ISFINITE_FOR_COMPLEX_TYPE(_type) \
156-
GKO_INLINE __device__ bool isfinite(const _type &value) \
157-
{ \
158-
return isfinite(value.real()) && isfinite(value.imag()); \
159-
}
160-
161-
GKO_DEFINE_ISFINITE_FOR_COMPLEX_TYPE(thrust::complex<float>)
162-
GKO_DEFINE_ISFINITE_FOR_COMPLEX_TYPE(thrust::complex<double>)
163-
#undef GKO_DEFINE_ISFINITE_FOR_COMPLEX_TYPE
164-
165-
166124
// For all other compiler in combination with CUDA or HIP, just use the provided
167125
// `isfinite` function
168126
#elif defined(__CUDA_ARCH__) || __HIP_DEVICE_COMPILE__
@@ -173,13 +131,3 @@ using ::isfinite;
173131

174132

175133
#endif // defined(__CUDA_ARCH__) || __HIP_DEVICE_COMPILE__
176-
177-
178-
// We need this struct, because otherwise we would call a __host__ function in a
179-
// __device__ function (even though it is constexpr)
180-
template <typename T>
181-
struct device_numeric_limits {
182-
static constexpr auto inf = std::numeric_limits<T>::infinity();
183-
static constexpr auto max = std::numeric_limits<T>::max();
184-
static constexpr auto min = std::numeric_limits<T>::min();
185-
};

‎common/factorization/par_ilu_kernels.hpp.inc

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -317,12 +317,12 @@ __global__ __launch_bounds__(default_block_size) void compute_l_u_factors(
317317
sum += last_operation; // undo the last operation
318318
if (row > col) {
319319
auto to_write = sum / u_values[u_row_ptrs[col + 1] - 1];
320-
if (::gko::isfinite(to_write)) {
320+
if (gko::isfinite(to_write)) {
321321
l_values[l_idx - 1] = to_write;
322322
}
323323
} else {
324324
auto to_write = sum;
325-
if (::gko::isfinite(to_write)) {
325+
if (gko::isfinite(to_write)) {
326326
u_values[u_idx - 1] = to_write;
327327
}
328328
}

‎cuda/test/base/math.cu

Lines changed: 44 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -50,47 +50,68 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
5050

5151

5252
namespace {
53+
namespace kernel {
5354

5455

55-
template <typename T>
56-
__global__ void test_real_isfinite(bool *result)
56+
template <typename T, typename FuncType>
57+
__device__ bool test_real_isfinite_function(FuncType isfin)
5758
{
58-
constexpr T inf = INFINITY;
59+
constexpr T inf = gko::device_numeric_limits<T>::inf;
60+
constexpr T quiet_nan = NAN;
5961
bool test_true{};
6062
bool test_false{};
6163

62-
test_true =
63-
gko::isfinite(T{0}) && gko::isfinite(-T{0}) && gko::isfinite(T{1});
64-
test_false = gko::isfinite(inf) || gko::isfinite(-inf) ||
65-
gko::isfinite(NAN) || gko::isfinite(inf - inf) ||
66-
gko::isfinite(inf / inf) || gko::isfinite(inf * T{2}) ||
67-
gko::isfinite(T{1} / T{0}) || gko::isfinite(T{0} / T{0});
68-
*result = test_true && !test_false;
64+
test_true = isfin(T{0}) && isfin(-T{0}) && isfin(T{1});
65+
test_false = isfin(inf) || isfin(-inf) || isfin(quiet_nan) ||
66+
isfin(inf - inf) || isfin(inf / inf) || isfin(inf * T{2}) ||
67+
isfin(T{1} / T{0}) || isfin(T{0} / T{0});
68+
return test_true && !test_false;
6969
}
7070

7171

72-
template <typename ComplexType>
73-
__global__ void test_complex_isfinite(bool *result)
72+
template <typename ComplexType, typename FuncType>
73+
__device__ bool test_complex_isfinite_function(FuncType isfin)
7474
{
7575
static_assert(gko::is_complex_s<ComplexType>::value,
7676
"Template type must be a complex type.");
7777
using T = gko::remove_complex<ComplexType>;
7878
using c_type = gko::kernels::cuda::cuda_type<ComplexType>;
79-
constexpr T inf = INFINITY;
79+
constexpr T inf = gko::device_numeric_limits<T>::inf;
8080
constexpr T quiet_nan = NAN;
8181
bool test_true{};
8282
bool test_false{};
8383

84-
test_true = gko::isfinite(c_type{T{0}, T{0}}) &&
85-
gko::isfinite(c_type{-T{0}, -T{0}}) &&
86-
gko::isfinite(c_type{T{1}, T{0}}) &&
87-
gko::isfinite(c_type{T{0}, T{1}});
88-
test_false =
89-
gko::isfinite(c_type{inf, T{0}}) || gko::isfinite(c_type{-inf, T{0}}) ||
90-
gko::isfinite(c_type{quiet_nan, T{0}}) ||
91-
gko::isfinite(c_type{T{0}, inf}) || gko::isfinite(c_type{T{0}, -inf}) ||
92-
gko::isfinite(c_type{T{0}, quiet_nan});
93-
*result = test_true && !test_false;
84+
test_true = isfin(c_type{T{0}, T{0}}) && isfin(c_type{-T{0}, -T{0}}) &&
85+
isfin(c_type{T{1}, T{0}}) && isfin(c_type{T{0}, T{1}});
86+
test_false = isfin(c_type{inf, T{0}}) || isfin(c_type{-inf, T{0}}) ||
87+
isfin(c_type{quiet_nan, T{0}}) || isfin(c_type{T{0}, inf}) ||
88+
isfin(c_type{T{0}, -inf}) || isfin(c_type{T{0}, quiet_nan});
89+
return test_true && !test_false;
90+
}
91+
92+
93+
} // namespace kernel
94+
95+
96+
template <typename T>
97+
__global__ void test_real_isfinite(bool *result)
98+
{
99+
bool gko_isfinite = kernel::test_real_isfinite_function<T>(
100+
[](T val) { return gko::isfinite(val); });
101+
bool custom_isfinite = kernel::test_real_isfinite_function<T>(
102+
[](T val) { return gko::detail::custom_isfinite(val); });
103+
*result = gko_isfinite && custom_isfinite;
104+
}
105+
106+
107+
template <typename ComplexType>
108+
__global__ void test_complex_isfinite(bool *result)
109+
{
110+
bool gko_isfinite = kernel::test_complex_isfinite_function<ComplexType>(
111+
[](ComplexType val) { return gko::isfinite(val); });
112+
bool custom_isfinite = kernel::test_complex_isfinite_function<ComplexType>(
113+
[](ComplexType val) { return gko::detail::custom_isfinite(val); });
114+
*result = gko_isfinite && custom_isfinite;
94115
}
95116

96117

‎hip/test/base/math.hip.cpp

Lines changed: 44 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -51,47 +51,68 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
5151

5252

5353
namespace {
54+
namespace kernel {
5455

5556

56-
template <typename T>
57-
__global__ void test_real_isfinite(bool *result)
57+
template <typename T, typename FuncType>
58+
__device__ bool test_real_isfinite_function(FuncType isfin)
5859
{
59-
constexpr T inf = INFINITY;
60+
constexpr T inf = gko::device_numeric_limits<T>::inf;
61+
constexpr T quiet_nan = NAN;
6062
bool test_true{};
6163
bool test_false{};
6264

63-
test_true =
64-
gko::isfinite(T{0}) && gko::isfinite(-T{0}) && gko::isfinite(T{1});
65-
test_false = gko::isfinite(inf) || gko::isfinite(-inf) ||
66-
gko::isfinite(NAN) || gko::isfinite(inf - inf) ||
67-
gko::isfinite(inf / inf) || gko::isfinite(inf * T{2}) ||
68-
gko::isfinite(T{1} / T{0}) || gko::isfinite(T{0} / T{0});
69-
*result = test_true && !test_false;
65+
test_true = isfin(T{0}) && isfin(-T{0}) && isfin(T{1});
66+
test_false = isfin(inf) || isfin(-inf) || isfin(quiet_nan) ||
67+
isfin(inf - inf) || isfin(inf / inf) || isfin(inf * T{2}) ||
68+
isfin(T{1} / T{0}) || isfin(T{0} / T{0});
69+
return test_true && !test_false;
7070
}
7171

7272

73-
template <typename ComplexType>
74-
__global__ void test_complex_isfinite(bool *result)
73+
template <typename ComplexType, typename FuncType>
74+
__device__ bool test_complex_isfinite_function(FuncType isfin)
7575
{
7676
static_assert(gko::is_complex_s<ComplexType>::value,
7777
"Template type must be a complex type.");
7878
using T = gko::remove_complex<ComplexType>;
7979
using c_type = gko::kernels::hip::hip_type<ComplexType>;
80-
constexpr T inf = INFINITY;
80+
constexpr T inf = gko::device_numeric_limits<T>::inf;
8181
constexpr T quiet_nan = NAN;
8282
bool test_true{};
8383
bool test_false{};
8484

85-
test_true = gko::isfinite(c_type{T{0}, T{0}}) &&
86-
gko::isfinite(c_type{-T{0}, -T{0}}) &&
87-
gko::isfinite(c_type{T{1}, T{0}}) &&
88-
gko::isfinite(c_type{T{0}, T{1}});
89-
test_false =
90-
gko::isfinite(c_type{inf, T{0}}) || gko::isfinite(c_type{-inf, T{0}}) ||
91-
gko::isfinite(c_type{quiet_nan, T{0}}) ||
92-
gko::isfinite(c_type{T{0}, inf}) || gko::isfinite(c_type{T{0}, -inf}) ||
93-
gko::isfinite(c_type{T{0}, quiet_nan});
94-
*result = test_true && !test_false;
85+
test_true = isfin(c_type{T{0}, T{0}}) && isfin(c_type{-T{0}, -T{0}}) &&
86+
isfin(c_type{T{1}, T{0}}) && isfin(c_type{T{0}, T{1}});
87+
test_false = isfin(c_type{inf, T{0}}) || isfin(c_type{-inf, T{0}}) ||
88+
isfin(c_type{quiet_nan, T{0}}) || isfin(c_type{T{0}, inf}) ||
89+
isfin(c_type{T{0}, -inf}) || isfin(c_type{T{0}, quiet_nan});
90+
return test_true && !test_false;
91+
}
92+
93+
94+
} // namespace kernel
95+
96+
97+
template <typename T>
98+
__global__ void test_real_isfinite(bool *result)
99+
{
100+
bool gko_isfinite = kernel::test_real_isfinite_function<T>(
101+
[](T val) { return gko::isfinite(val); });
102+
bool custom_isfinite = kernel::test_real_isfinite_function<T>(
103+
[](T val) { return gko::detail::custom_isfinite(val); });
104+
*result = gko_isfinite && custom_isfinite;
105+
}
106+
107+
108+
template <typename ComplexType>
109+
__global__ void test_complex_isfinite(bool *result)
110+
{
111+
bool gko_isfinite = kernel::test_complex_isfinite_function<ComplexType>(
112+
[](ComplexType val) { return gko::isfinite(val); });
113+
bool custom_isfinite = kernel::test_complex_isfinite_function<ComplexType>(
114+
[](ComplexType val) { return gko::detail::custom_isfinite(val); });
115+
*result = gko_isfinite && custom_isfinite;
95116
}
96117

97118

‎include/ginkgo/core/base/math.hpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -730,7 +730,8 @@ template <typename T>
730730
GKO_INLINE GKO_ATTRIBUTES xstd::enable_if_t<!is_complex_s<T>::value, bool>
731731
isfinite(const T &value)
732732
{
733-
return std::isfinite(value);
733+
using std::isfinite;
734+
return isfinite(value);
734735
}
735736

736737
#endif // defined(__CUDA_ARCH__)

0 commit comments

Comments
 (0)