Hi!
On Sat, 21 May 2016 17:59:17 +0200, I wrote:
> As discussed in <https://gcc.gnu.org/PR70945> "Offloading: compatibility
> of target and offloading toolchains", there are situations where we have
> to do more work to ensure compatibility between target and offloading
> toolchains.
The outcome of the discussion a few months ago has been that instead of
my proposed compiler patch, the "offloading port(s) of newlib [...]
should provide [a] translation layer from the host headers to the
offloading target functions". We still plan to, but have not yet gotten
to spend time on that, unfortunately.
To at least support the following use case:
> The first thing I'm working on is math functions usage in offloaded
> regions.
>
> Here is a first patch, addressing glibc's finite math optimizations: 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, obviously, is incompatible
> with offloading targets that don't use glibc, and thus don't provide
> these "__[function]_finite" entry points.
..., I have now in r241355 committed my original patch to
gomp-4_0-branch, with an (incomplete) test case added:
commit 7e178f04bf6d692a17c4be4ff050da2d23a543e7
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Oct 19 21:24:37 2016 +0000
[PR other/70945] Handle function_glibc_finite_math in offloading
gcc/
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.
libgomp/
PR other/70945
* testsuite/libgomp.oacc-c-c++-common/pr70945-1.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@241355
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 20 ++
gcc/config/darwin.c | 2 +
gcc/coretypes.h | 11 +-
gcc/doc/tm.texi | 2 +-
gcc/lto-cgraph.c | 181 +++++++++++-----
gcc/lto-streamer-out.c | 2 +-
gcc/lto-streamer.h | 6 +-
gcc/lto/lto.c | 2 +-
gcc/target.def | 2 +-
gcc/targhooks.c | 2 +-
libgomp/ChangeLog.gomp | 5 +
.../libgomp.oacc-c-c++-common/pr70945-1.c | 231 +++++++++++++++++++++
12 files changed, 408 insertions(+), 58 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 04894b8..851da81 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,23 @@
+2016-10-19 Thomas Schwinge <tho...@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 <nat...@codesourcery.com>
* tree.h (OMP_CLAUSE_TILE_ITERVAR, OMP_CLAUSE_TILE_COUNT): New.
diff --git gcc/config/darwin.c gcc/config/darwin.c
index 0055d80..92fe3e5 100644
--- gcc/config/darwin.c
+++ gcc/config/darwin.c
@@ -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;
}
diff --git gcc/coretypes.h gcc/coretypes.h
index 12067fd..e28235a 100644
--- gcc/coretypes.h
+++ gcc/coretypes.h
@@ -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
diff --git gcc/doc/tm.texi gcc/doc/tm.texi
index 745910f..3de3554 100644
--- gcc/doc/tm.texi
+++ gcc/doc/tm.texi
@@ -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
diff --git gcc/lto-cgraph.c gcc/lto-cgraph.c
index 857ce4d..bd28090 100644
--- gcc/lto-cgraph.c
+++ gcc/lto-cgraph.c
@@ -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. */
diff --git gcc/lto-streamer-out.c gcc/lto-streamer-out.c
index 6703d41..c504e7d 100644
--- gcc/lto-streamer-out.c
+++ gcc/lto-streamer-out.c
@@ -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);
diff --git gcc/lto-streamer.h gcc/lto-streamer.h
index 7374b4e..9c208ba 100644
--- gcc/lto-streamer.h
+++ gcc/lto-streamer.h
@@ -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 *,
diff --git gcc/lto/lto.c gcc/lto/lto.c
index 5c991a5..f9b1d9e 100644
--- gcc/lto/lto.c
+++ gcc/lto/lto.c
@@ -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. */
diff --git gcc/target.def gcc/target.def
index 20f2b32..bf8b7d8 100644
--- gcc/target.def
+++ gcc/target.def
@@ -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)
diff --git gcc/targhooks.c gcc/targhooks.c
index a342277..a5f4bfe 100644
--- gcc/targhooks.c
+++ gcc/targhooks.c
@@ -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)
{
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index ecd51e6..5dbcf48 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2016-10-19 Thomas Schwinge <tho...@codesourcery.com>
+
+ PR other/70945
+ * testsuite/libgomp.oacc-c-c++-common/pr70945-1.c: New file.
+
2016-10-05 Nathan Sidwell <nat...@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/tile-1.c: New.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr70945-1.c
libgomp/testsuite/libgomp.oacc-c-c++-common/pr70945-1.c
new file mode 100644
index 0000000..6ee44eb
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr70945-1.c
@@ -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;
+}
Grüße
Thomas
