@@ -1,3 +1,23 @@
+2016-10-19 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR other/70945
+ * targhooks.c (default_libc_has_function): Update comment.
+ * target.def (libc_has_function): Likewise.
+ * doc/tm.texi: Regenerate.
+ * coretypes.h (enum function_class): Add
+ function_glibc_finite_math.
+ * config/darwin.c (darwin_libc_has_function): Handle it.
+ * lto-streamer.h (enum lto_section_type): Rename
+ LTO_section_offload_table to LTO_section_offload_data. Adjust all
+ users.
+ * lto-cgraph.c (void output_offload_data): New function, split out
+ of output_offload_tables. Adjust all users. Stream the target's
+ function_glibc_finite_math property.
+ (input_offload_data): New function, split out of
+ input_offload_tables. Adjust all users. Handle mismatch between
+ the target's and the offloading target's
+ function_glibc_finite_math property.
+
2016-10-05 Nathan Sidwell <nathan@codesourcery.com>
* tree.h (OMP_CLAUSE_TILE_ITERVAR, OMP_CLAUSE_TILE_COUNT): New.
@@ -3401,6 +3401,8 @@ darwin_libc_has_function (enum function_class fn_class)
|| fn_class == function_c99_misc)
return (TARGET_64BIT
|| strverscmp (darwin_macosx_version_min, "10.3") >= 0);
+ if (fn_class == function_glibc_finite_math)
+ return false;
return true;
}
@@ -281,14 +281,21 @@ union _dont_use_tree_here_;
#endif
-/* Classes of functions that compiler needs to check
+/* Properties, such as classes of functions that the compiler can check
whether they are present at the runtime or not. */
enum function_class {
function_c94,
function_c99_misc,
function_c99_math_complex,
function_sincos,
- function_c11_misc
+ function_c11_misc,
+ /* If -ffinite-math-only (as implied by -ffast-math, or -Ofast) is in effect,
+ glibc's <math.h> is known to include <bits/math-finite.h> for "special
+ entry points to use when the compiler got told to only expect finite
+ results". This divertes the math functions' assembler names from
+ "[function]" to "__[function]_finite". This property indicates whether
+ such diversion may occur, not whether it actually has. */
+ function_glibc_finite_math
};
/* Enumerate visibility settings. This is deliberately ordered from most
@@ -5308,7 +5308,7 @@ macro, a reasonable default is used.
@end defmac
@deftypefn {Target Hook} bool TARGET_LIBC_HAS_FUNCTION (enum function_class @var{fn_class})
-This hook determines whether a function from a class of functions
+This hook determines properties, such as whether a class of functions
@var{fn_class} is present at the runtime.
@end deftypefn
@@ -38,6 +38,9 @@ along with GCC; see the file COPYING3. If not see
#include "ipa-utils.h"
#include "omp-low.h"
#include "ipa-chkp.h"
+#include "target.h"
+#include "output.h"
+#include "builtins.h"
/* True when asm nodes has been output. */
bool asm_nodes_output = false;
@@ -1091,21 +1094,37 @@ read_string (struct lto_input_block *ib)
return str;
}
+/* Output offload data. */
+
+static void output_offload_tables (struct lto_simple_output_block *);
+
+void output_offload_data (void)
+{
+ /* Return early if there is no offload data. */
+ if (vec_safe_is_empty (offload_funcs) && vec_safe_is_empty (offload_vars))
+ return;
+
+ struct lto_simple_output_block *ob
+ = lto_create_simple_output_block (LTO_section_offload_data);
+
+ /* Stream the target's function_glibc_finite_math property. */
+ bool g_f_m = targetm.libc_has_function (function_glibc_finite_math);
+ streamer_write_hwi_stream (ob->main_stream, g_f_m);
+
+ output_offload_tables (ob);
+
+ lto_destroy_simple_output_block (ob);
+}
+
/* Output function/variable tables that will allow libgomp to look up offload
target code.
OFFLOAD_FUNCS is filled in expand_omp_target, OFFLOAD_VARS is filled in
varpool_node::get_create. In WHOPR (partitioned) mode during the WPA stage
both OFFLOAD_FUNCS and OFFLOAD_VARS are filled by input_offload_tables. */
-void
-output_offload_tables (void)
+static void
+output_offload_tables (struct lto_simple_output_block *ob)
{
- if (vec_safe_is_empty (offload_funcs) && vec_safe_is_empty (offload_vars))
- return;
-
- struct lto_simple_output_block *ob
- = lto_create_simple_output_block (LTO_section_offload_table);
-
for (unsigned i = 0; i < vec_safe_length (offload_funcs); i++)
{
streamer_write_enum (ob->main_stream, LTO_symtab_tags,
@@ -1123,7 +1142,6 @@ output_offload_tables (void)
}
streamer_write_uhwi_stream (ob->main_stream, 0);
- lto_destroy_simple_output_block (ob);
/* In WHOPR mode during the WPA stage the joint offload tables need to be
streamed to one partition only. That's why we free offload_funcs and
@@ -1885,65 +1903,132 @@ input_symtab (void)
}
}
-/* Input function/variable tables that will allow libgomp to look up offload
- target code, and store them into OFFLOAD_FUNCS and OFFLOAD_VARS. */
+/* Input offload data. */
+
+static void input_offload_tables (struct lto_input_block *,
+ struct lto_file_decl_data *, bool);
void
-input_offload_tables (bool do_force_output)
+input_offload_data (bool do_force_output)
{
struct lto_file_decl_data **file_data_vec = lto_get_file_decl_data ();
struct lto_file_decl_data *file_data;
unsigned int j = 0;
+ bool g_f_m_target = false;
while ((file_data = file_data_vec[j++]))
{
const char *data;
size_t len;
struct lto_input_block *ib
- = lto_create_simple_input_block (file_data, LTO_section_offload_table,
+ = lto_create_simple_input_block (file_data, LTO_section_offload_data,
&data, &len);
if (!ib)
continue;
- enum LTO_symtab_tags tag
- = streamer_read_enum (ib, LTO_symtab_tags, LTO_symtab_last_tag);
- while (tag)
- {
- if (tag == LTO_symtab_unavail_node)
- {
- int decl_index = streamer_read_uhwi (ib);
- tree fn_decl
- = lto_file_decl_data_get_fn_decl (file_data, decl_index);
- vec_safe_push (offload_funcs, fn_decl);
+ /* Merge the target's function_glibc_finite_math property. */
+ g_f_m_target |= streamer_read_hwi (ib);
- /* Prevent IPA from removing fn_decl as unreachable, since there
- may be no refs from the parent function to child_fn in offload
- LTO mode. */
- if (do_force_output)
- cgraph_node::get (fn_decl)->mark_force_output ();
- }
- else if (tag == LTO_symtab_variable)
- {
- int decl_index = streamer_read_uhwi (ib);
- tree var_decl
- = lto_file_decl_data_get_var_decl (file_data, decl_index);
- vec_safe_push (offload_vars, var_decl);
+ input_offload_tables (ib, file_data, do_force_output);
- /* Prevent IPA from removing var_decl as unused, since there
- may be no refs to var_decl in offload LTO mode. */
- if (do_force_output)
- varpool_node::get (var_decl)->force_output = 1;
- }
- else
- fatal_error (input_location,
- "invalid offload table in %s", file_data->file_name);
-
- tag = streamer_read_enum (ib, LTO_symtab_tags, LTO_symtab_last_tag);
- }
-
- lto_destroy_simple_input_block (file_data, LTO_section_offload_table,
+ lto_destroy_simple_input_block (file_data, LTO_section_offload_data,
ib, data, len);
}
+
+ /* Take action if the target has the function_glibc_finite_math property set,
+ and that doesn't match the current (that is, offloading target's). */
+ bool g_f_m = targetm.libc_has_function (function_glibc_finite_math);
+ if (g_f_m_target && !g_f_m)
+ {
+ struct cgraph_node *node;
+ FOR_EACH_FUNCTION (node)
+ {
+ /* This only applies to references to external math functions. */
+ if (!DECL_EXTERNAL (node->decl))
+ continue;
+ /* All the relevant math functions are registered as GCC builtins. */
+ if (!DECL_BUILT_IN (node->decl)
+ || (mathfn_built_in (TREE_TYPE (TREE_TYPE (node->decl)),
+ DECL_FUNCTION_CODE (node->decl))
+ == NULL_TREE))
+ continue;
+ /* Check whether the assembler name for "[function]" has been set to
+ "__[function]_finite". */
+ if (!DECL_ASSEMBLER_NAME_SET_P (node->decl))
+ continue;
+ const char *asm_name
+ = IDENTIFIER_POINTER (DECL_ASSEMBLER_NAME (node->decl));
+ if (*asm_name++ != '*')
+ continue;
+ size_t ulp_len = strlen (user_label_prefix);
+ if (ulp_len == 0)
+ ;
+ else if (strncmp (asm_name, user_label_prefix, ulp_len) == 0)
+ asm_name += ulp_len;
+ else
+ continue;
+ if (*asm_name++ != '_')
+ continue;
+ if (*asm_name++ != '_')
+ continue;
+ const char *name = IDENTIFIER_POINTER (DECL_NAME (node->decl));
+ size_t name_len = strlen (name);
+ if (strncmp (asm_name, name, name_len) == 0)
+ asm_name += name_len;
+ else
+ continue;
+ if (strcmp (asm_name, "_finite") != 0)
+ continue;
+ /* ..., and if yes, reset it. */
+ symtab->change_decl_assembler_name (node->decl,
+ DECL_NAME (node->decl));
+ }
+ }
+}
+
+/* Input function/variable tables that will allow libgomp to look up offload
+ target code, and store them into OFFLOAD_FUNCS and OFFLOAD_VARS. */
+
+static void
+input_offload_tables (struct lto_input_block *ib,
+ struct lto_file_decl_data *file_data,
+ bool do_force_output)
+{
+ enum LTO_symtab_tags tag
+ = streamer_read_enum (ib, LTO_symtab_tags, LTO_symtab_last_tag);
+ while (tag)
+ {
+ if (tag == LTO_symtab_unavail_node)
+ {
+ int decl_index = streamer_read_uhwi (ib);
+ tree fn_decl
+ = lto_file_decl_data_get_fn_decl (file_data, decl_index);
+ vec_safe_push (offload_funcs, fn_decl);
+
+ /* Prevent IPA from removing fn_decl as unreachable, since there
+ may be no refs from the parent function to child_fn in offload
+ LTO mode. */
+ if (do_force_output)
+ cgraph_node::get (fn_decl)->mark_force_output ();
+ }
+ else if (tag == LTO_symtab_variable)
+ {
+ int decl_index = streamer_read_uhwi (ib);
+ tree var_decl
+ = lto_file_decl_data_get_var_decl (file_data, decl_index);
+ vec_safe_push (offload_vars, var_decl);
+
+ /* Prevent IPA from removing var_decl as unused, since there
+ may be no refs to var_decl in offload LTO mode. */
+ if (do_force_output)
+ varpool_node::get (var_decl)->force_output = 1;
+ }
+ else
+ fatal_error (input_location,
+ "invalid offload table in %s", file_data->file_name);
+
+ tag = streamer_read_enum (ib, LTO_symtab_tags, LTO_symtab_last_tag);
+ }
}
/* True when we need optimization summary for NODE. */
@@ -2381,7 +2381,7 @@ lto_output (void)
statements using the statement UIDs. */
output_symtab ();
- output_offload_tables ();
+ output_offload_data ();
#if CHECKING_P
lto_bitmap_free (output);
@@ -242,7 +242,7 @@ enum lto_section_type
LTO_section_inline_summary,
LTO_section_ipcp_transform,
LTO_section_ipa_icf,
- LTO_section_offload_table,
+ LTO_section_offload_data,
LTO_section_mode_table,
LTO_section_ipa_hsa,
LTO_N_SECTION_TYPES /* Must be last. */
@@ -914,8 +914,8 @@ bool lto_symtab_encoder_encode_initializer_p (lto_symtab_encoder_t,
varpool_node *);
void output_symtab (void);
void input_symtab (void);
-void output_offload_tables (void);
-void input_offload_tables (bool);
+void output_offload_data (void);
+void input_offload_data (bool);
bool referenced_from_other_partition_p (struct ipa_ref_list *,
lto_symtab_encoder_t);
bool reachable_from_other_partition_p (struct cgraph_node *,
@@ -2856,7 +2856,7 @@ read_cgraph_and_symbols (unsigned nfiles, const char **fnames)
/* Read the symtab. */
input_symtab ();
- input_offload_tables (!flag_ltrans);
+ input_offload_data (!flag_ltrans);
/* Store resolutions into the symbol table. */
@@ -2533,7 +2533,7 @@ set via @code{__attribute__}.",
DEFHOOK
(libc_has_function,
- "This hook determines whether a function from a class of functions\n\
+ "This hook determines properties, such as whether a class of functions\n\
@var{fn_class} is present at the runtime.",
bool, (enum function_class fn_class),
default_libc_has_function)
@@ -1389,7 +1389,7 @@ default_have_conditional_execution (void)
}
/* By default we assume that c99 functions are present at the runtime,
- but sincos is not. */
+ but others are not. */
bool
default_libc_has_function (enum function_class fn_class)
{
@@ -1,3 +1,8 @@
+2016-10-19 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR other/70945
+ * testsuite/libgomp.oacc-c-c++-common/pr70945-1.c: New file.
+
2016-10-05 Nathan Sidwell <nathan@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/tile-1.c: New.
new file mode 100644
@@ -0,0 +1,231 @@
+/* Verify that target-side header-foo done in glibc for finite-only math
+ functions rewriting doesn't cause offloading-side confusion with newlib,
+ such as "unresolved symbol __atanh_finite", etc. */
+
+/* { dg-additional-options "-lm -foffload=-lm -ffast-math" } */
+
+#ifndef __cplusplus /* C */
+# include <stdlib.h>
+# include <math.h>
+# include <float.h>
+# include <complex.h>
+#else /* C++ */
+# include <cstdlib>
+# include <cmath>
+# include <cfloat>
+# include <complex>
+#endif
+
+/* Assign "var <= val", and make sure the compiler can't tell. */
+#define LOAD(var, val) \
+ do { \
+ (var) = (val); \
+ asm volatile ("" : : "g" (&(var)) : "memory"); \
+ } while (0)
+
+/* Floating point, you know... Let's keep it simple. */
+#define EPSILON 0.001f
+/* These evaluate macro arguments more than once. */
+#define EQUALSf(a, b) \
+ (((a) > (b)) ? (((a) - (b)) < (EPSILON)) : (((b) - (a)) < (EPSILON)))
+#define EQUALS(a, b) \
+ EQUALSf((a), (b))
+
+
+int main(int argc, char *argv[])
+{
+#pragma acc parallel
+ {
+ int i, i_;
+ long l, l_;
+ long long ll, ll_;
+ float f, f_, f__;
+ double d, d_, d__;
+ const char *s;
+#ifndef __cplusplus /* C */
+ div_t div_i;
+ ldiv_t div_l;
+ lldiv_t div_ll;
+#else /* C++ */
+ std::div_t div_i;
+ std::ldiv_t div_l;
+ std::lldiv_t div_ll;
+#endif
+
+ LOAD(i, -10); i = abs (i); if (i != 10) __builtin_abort();
+ LOAD(l, -9); l = abs (l); if (l != 9) __builtin_abort();
+ LOAD(ll, -8); ll = abs (ll); if (ll != 8) __builtin_abort();
+ LOAD(f, -7); f = fabsf (f); if (!EQUALSf(f, 7)) __builtin_abort();
+ LOAD(d, 6); d = fabs (d); if (!EQUALS(d, 6)) __builtin_abort();
+
+ LOAD(i, -10); LOAD(i_, -3); div_i = div (i, i_); if (div_i.quot != 3 && div_i.rem != -1) __builtin_abort();
+ LOAD(l, -11); LOAD(l_, -4); div_l = ldiv (l, l_); if (div_l.quot != 2 && div_l.rem != -3) __builtin_abort();
+ LOAD(ll, -12); LOAD(ll_, -5); div_ll = lldiv (ll, ll_); if (div_ll.quot != 2 && div_ll.rem != -2) __builtin_abort();
+
+ LOAD(f, -7); LOAD(f_, -2.5); f = fmodf (f, f_); if (!EQUALSf(f, -2)) __builtin_abort();
+ LOAD(d, -8); LOAD(d_, -2.6); d = fmod (d, d_); if (!EQUALS(d, -0.2)) __builtin_abort();
+
+ LOAD(f, -8); LOAD(f_, -2.5); f = remainderf (f, f_); if (!EQUALSf(f, -0.5)) __builtin_abort();
+ LOAD(d, -7); LOAD(d_, -2.6); d = remainder (d, d_); if (!EQUALS(d, 0.8)) __builtin_abort();
+
+ LOAD(f, -8); LOAD(f_, -2.5); f = remquof (f, f_, &i); if (!EQUALSf(f, -0.5) || i < 0) __builtin_abort();
+ LOAD(d, -7); LOAD(d_, -2.6); d = remquo (d, d_, &i); if (!EQUALS(d, 0.8) || i < 0) __builtin_abort();
+
+ LOAD(f, -8); LOAD(f_, -2.3); LOAD(f__, 2.6); f = fmaf (f, f_, f__); if (!EQUALSf(f, 21)) __builtin_abort();
+ LOAD(d, -7); LOAD(d_, -2.6); LOAD(d__, 1.8); d = fma (d, d_, d__); if (!EQUALS(d, 20)) __builtin_abort();
+
+ LOAD(f, -3); LOAD(f_, -2.5); f = fmaxf (f, f_); if (!EQUALSf(f, -2.5)) __builtin_abort();
+ LOAD(d, -4); LOAD(d_, 2.6); d = fmax (d, d_); if (!EQUALS(d, 2.6)) __builtin_abort();
+
+ LOAD(f, 3); LOAD(f_, -2.5); f = fminf (f, f_); if (!EQUALSf(f, -2.5)) __builtin_abort();
+ LOAD(d, -4); LOAD(d_, 2.6); d = fmin (d, d_); if (!EQUALS(d, -4)) __builtin_abort();
+
+ LOAD(f, 3); LOAD(f_, -2.5); f = fdimf (f, f_); if (!EQUALSf(f, 5.5)) __builtin_abort();
+ LOAD(d, -4); LOAD(d_, 2.6); d = fdim (d, d_); if (!EQUALS(d, 0)) __builtin_abort();
+
+ LOAD(f, 3.3); f = expf (f); if (!EQUALSf(f, 27.1126)) __builtin_abort();
+ LOAD(d, -0.24); d = exp (d); if (!EQUALS(d, 0.7866)) __builtin_abort();
+
+ LOAD(f, 3.3); f = exp2f (f); if (!EQUALSf(f, 9.8492)) __builtin_abort();
+ LOAD(d, -0.24); d = exp2 (d); if (!EQUALS(d, 0.8467)) __builtin_abort();
+
+ LOAD(f, 3.3); f = expm1f (f); if (!EQUALSf(f, 26.1126)) __builtin_abort();
+ LOAD(d, -0.24); d = expm1 (d); if (!EQUALS(d, -0.2134)) __builtin_abort();
+
+ LOAD(f, 10.3); f = logf (f); if (!EQUALSf(f, 2.3321)) __builtin_abort();
+ LOAD(d, 0.55); d = log (d); if (!EQUALS(d, -0.5978)) __builtin_abort();
+
+ LOAD(f, 1); f = log2f (f); if (!EQUALSf(f, 0)) __builtin_abort();
+ LOAD(d, 32768); d = log2 (d); if (!EQUALS(d, 15)) __builtin_abort();
+
+ LOAD(f, 100); f = log10f (f); if (!EQUALSf(f, 2)) __builtin_abort();
+ LOAD(d, 0.3162); d = log10 (d); if (!EQUALS(d, -0.5000)) __builtin_abort();
+
+ LOAD(f, 4); f = log1pf (f); if (!EQUALSf(f, 1.6094)) __builtin_abort();
+ LOAD(d, -0); d = log1p (d); if (!EQUALS(d, 0)) __builtin_abort();
+
+ LOAD(f, 4); i = ilogbf (f); if (i != 2) __builtin_abort();
+ LOAD(d, 987.55); i = ilogb (d); if (i != 9) __builtin_abort();
+
+ LOAD(f, 987.55); f = logbf (f); if (!EQUALSf(f, 9)) __builtin_abort();
+ LOAD(d, 4); d = logb (d); if (!EQUALS(d, 2)) __builtin_abort();
+
+ LOAD(f, 987.55); f = sqrtf (f); if (!EQUALSf(f, 31.4253)) __builtin_abort();
+ LOAD(d, 4); d = sqrt (d); if (!EQUALS(d, 2)) __builtin_abort();
+
+ LOAD(f, 31034.0387); f = cbrtf (f); if (!EQUALSf(f, 31.4253)) __builtin_abort();
+ LOAD(d, 8); d = cbrt (d); if (!EQUALS(d, 2)) __builtin_abort();
+
+ LOAD(f, -8); LOAD(f_, -2.5); f = hypotf (f, f_); if (!EQUALSf(f, 8.3815)) __builtin_abort();
+ LOAD(d, -7); LOAD(d_, -2.6); d = hypot (d, d_); if (!EQUALS(d, 7.4673)) __builtin_abort();
+
+ LOAD(f, 8); LOAD(f_, -2.5); f = powf (f, f_); if (!EQUALSf(f, 0.0055)) __builtin_abort();
+ LOAD(d, 7); LOAD(d_, -2.6); d = pow (d, d_); if (!EQUALS(d, 0.0063)) __builtin_abort();
+
+ LOAD(f, 8); f = sinf (f); if (!EQUALSf(f, 0.9894)) __builtin_abort();
+ LOAD(d, 7); d = sin (d); if (!EQUALS(d, 0.6570)) __builtin_abort();
+
+ LOAD(f, 8); f = cosf (f); if (!EQUALSf(f, -0.1455)) __builtin_abort();
+ LOAD(d, 7); d = cos (d); if (!EQUALS(d, 0.7539)) __builtin_abort();
+
+ LOAD(f, 8); f = tanf (f); if (!EQUALSf(f, -6.7997)) __builtin_abort();
+ LOAD(d, 7); d = tan (d); if (!EQUALS(d, 0.8714)) __builtin_abort();
+
+ LOAD(f, 0.8); f = asinf (f); if (!EQUALSf(f, 0.9273)) __builtin_abort();
+ LOAD(d, 0.7); d = asin (d); if (!EQUALS(d, 0.7754)) __builtin_abort();
+
+ LOAD(f, 0.8); f = acosf (f); if (!EQUALSf(f, 0.6435)) __builtin_abort();
+ LOAD(d, 0.7); d = acos (d); if (!EQUALS(d, 0.7954)) __builtin_abort();
+
+ LOAD(f, 0.8); f = atanf (f); if (!EQUALSf(f, 0.6747)) __builtin_abort();
+ LOAD(d, 0.7); d = atan (d); if (!EQUALS(d, 0.6107)) __builtin_abort();
+
+ LOAD(f, 0.8); LOAD(f_, -0.7); f = atan2f (f, f_); if (!EQUALSf(f, 2.2896)) __builtin_abort();
+ LOAD(d, -0.7); LOAD(d_, 0.8); d = atan2 (d, d_); if (!EQUALS(d, -0.7188)) __builtin_abort();
+
+ LOAD(f, 0.8); f = sinhf (f); if (!EQUALSf(f, 0.8881)) __builtin_abort();
+ LOAD(d, 0.7); d = sinh (d); if (!EQUALS(d, 0.7585)) __builtin_abort();
+
+ LOAD(f, 0.8); f = coshf (f); if (!EQUALSf(f, 1.3374)) __builtin_abort();
+ LOAD(d, 0.7); d = cosh (d); if (!EQUALS(d, 1.2551)) __builtin_abort();
+
+ LOAD(f, 0.8); f = tanhf (f); if (!EQUALSf(f, 0.6640)) __builtin_abort();
+ LOAD(d, 0.7); d = tanh (d); if (!EQUALS(d, 0.6044)) __builtin_abort();
+
+ LOAD(f, 0.8); f = asinhf (f); if (!EQUALSf(f, 0.7327)) __builtin_abort();
+ LOAD(d, 0.7); d = asinh (d); if (!EQUALS(d, 0.6527)) __builtin_abort();
+
+ LOAD(f, 1.8); f = acoshf (f); if (!EQUALSf(f, 1.1929)) __builtin_abort();
+ LOAD(d, 1.7); d = acosh (d); if (!EQUALS(d, 1.1232)) __builtin_abort();
+
+ LOAD(f, 0.8); f = atanhf (f); if (!EQUALSf(f, 1.0986)) __builtin_abort();
+ LOAD(d, 0.7); d = atanh (d); if (!EQUALS(d, 0.8673)) __builtin_abort();
+
+ LOAD(f, 0.8); f = erff (f); if (!EQUALSf(f, 0.7421)) __builtin_abort();
+ LOAD(d, 0.7); d = erf (d); if (!EQUALS(d, 0.6778)) __builtin_abort();
+
+ LOAD(f, 0.8); f = erfcf (f); if (!EQUALSf(f, 1 - 0.7421)) __builtin_abort();
+ LOAD(d, 0.7); d = erfc (d); if (!EQUALS(d, 1 - 0.6778)) __builtin_abort();
+
+#if 0
+ /* TODO: incompatible inline function. */
+ LOAD(f, 0.8); f = lgammaf (f); if (!EQUALSf(f, TODO)) __builtin_abort();
+ LOAD(d, 0.7); d = lgamma (d); if (!EQUALS(d, TODO)) __builtin_abort();
+#endif
+
+#if 0
+ /* TODO: incompatible inline function. */
+ LOAD(f, 0.8); f = tgammaf (f); if (!EQUALSf(f, TODO)) __builtin_abort();
+ LOAD(d, 0.7); d = tgamma (d); if (!EQUALS(d, TODO)) __builtin_abort();
+#endif
+
+ LOAD(f, -0.8); f = ceilf (f); if (!EQUALSf(f, -0)) __builtin_abort();
+ LOAD(d, 0.7); d = ceil (d); if (!EQUALS(d, 1)) __builtin_abort();
+
+ LOAD(f, -0.8); f = floorf (f); if (!EQUALSf(f, -1)) __builtin_abort();
+ LOAD(d, 0.7); d = floor (d); if (!EQUALS(d, 0)) __builtin_abort();
+
+ LOAD(f, -0.8); f = truncf (f); if (!EQUALSf(f, -0)) __builtin_abort();
+ LOAD(d, 0.7); d = trunc (d); if (!EQUALS(d, 0)) __builtin_abort();
+
+ LOAD(f, -0.8); f = roundf (f); if (!EQUALSf(f, -1)) __builtin_abort();
+ LOAD(d, 0.7); d = round (d); if (!EQUALS(d, 1)) __builtin_abort();
+ LOAD(f, -0.8); l = lroundf (f); if (l != -1) __builtin_abort();
+ LOAD(d, 0.7); l = lround (d); if (l != 1) __builtin_abort();
+ LOAD(f, -0.8); ll = llroundf (f); if (ll != -1) __builtin_abort();
+ LOAD(d, 0.7); ll = llround (d); if (ll != 1) __builtin_abort();
+
+#if 0
+ /* TODO: current rounding mode. */
+
+ LOAD(f, -0.8); f = nearbyintf (f); if (!EQUALSf(f, TODO)) __builtin_abort();
+ LOAD(d, 0.7); d = nearbyint (d); if (!EQUALS(d, TODO)) __builtin_abort();
+
+ LOAD(f, -0.8); f = rintf (f); if (!EQUALSf(f, TODO)) __builtin_abort();
+ LOAD(d, 0.7); d = rint (d); if (!EQUALS(d, TODO)) __builtin_abort();
+ LOAD(f, -0.8); l = lrintf (f); if (l != TODO) __builtin_abort();
+ LOAD(d, 0.7); l = lrint (d); if (l != TODO) __builtin_abort();
+ LOAD(f, -0.8); ll = llrintf (f); if (ll != TODO) __builtin_abort();
+ LOAD(d, 0.7); ll = llrint (d); if (ll != TODO) __builtin_abort();
+#endif
+
+ LOAD(f, -8.88); f = frexpf (f, &i); if (!EQUALSf(f, -0.5550) || i != 4) __builtin_abort();
+ LOAD(d, -7.77); d = frexp (d, &i); if (!EQUALS(d, -0.9712) || i != 3) __builtin_abort();
+
+ LOAD(f, -8.88); LOAD(i, 5); f = ldexpf (f, i); if (!EQUALSf(f, -284.16)) __builtin_abort();
+ LOAD(d, -7.77); LOAD(i, 6); d = ldexp (d, i); if (!EQUALS(d, -497.28)) __builtin_abort();
+
+ LOAD(f, -8.88); f = modff (f, &f_); if (!EQUALSf(f, -0.88) || !EQUALSf(f_, -8)) __builtin_abort();
+ LOAD(d, -7.77); d = modf (d, &d_); if (!EQUALS(d, -0.77) || !EQUALS(d_, -7)) __builtin_abort();
+
+#if FLT_RADIX != 2
+# error
+#endif
+ LOAD(f, -8.88); LOAD(i, 5); f = scalbnf (f, i); if (!EQUALSf(f, -284.16)) __builtin_abort();
+ LOAD(d, -7.77); LOAD(i, 6); d = scalbn (d, i); if (!EQUALS(d, -497.28)) __builtin_abort();
+ LOAD(f, -8.88); LOAD(l, 5); f = scalblnf (f, l); if (!EQUALSf(f, -284.16)) __builtin_abort();
+ LOAD(d, -7.77); LOAD(l, 6); d = scalbln (d, l); if (!EQUALS(d, -497.28)) __builtin_abort();
+ }
+
+ return 0;
+}