From patchwork Wed Oct 19 21:28:23 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 684343 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 3szlSD3GPDz9s9x for ; Thu, 20 Oct 2016 08:28:58 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=WrA72OD4; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:in-reply-to:references:date:message-id:mime-version :content-type:content-transfer-encoding; q=dns; s=default; b=O/p YZ3yW4xe424a2QGm0qG2t5byr9bEJPRisY+F385h5WnD5WNRwtPoOFQUk/xucL/E 9ZM2S/MJcf01+jcCLLqmKNNfjW9sRpdZ+YrYrGEsR9HQ4PTEDyb0x5o7lLl6GgAR 9vdQl3FqH+auswAzl89mvOQxYD/bSCvBUd2P7oEo= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:in-reply-to:references:date:message-id:mime-version :content-type:content-transfer-encoding; s=default; bh=daCXJ/8Pq pfToUi/fk/MZAa3izg=; b=WrA72OD48bibFVbxh1H8kU3wxv86RJeaH/uC9JuAy qykTre86aH7OGgUwDiD4NEQ3ElVESERTOh55DUbNaI+kd7q7KNScNC/555JrNHmi QrTonTsqShq7vzyyGFHfe9tJTIUkj+mAtySEMFUgdS0JamyKSgah9igzmUxwroIf 0Y= Received: (qmail 96070 invoked by alias); 19 Oct 2016 21:28:49 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 96060 invoked by uid 89); 19 Oct 2016 21:28:48 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.2 required=5.0 tests=AWL, BAYES_20, RCVD_IN_DNSWL_NONE, SPF_PASS, T_FILL_THIS_FORM_SHORT, URIBL_RED autolearn=ham version=3.3.2 spammy=055, diversion, Must, sk:nathan X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 19 Oct 2016 21:28:38 +0000 Received: from svr-orw-fem-04.mgc.mentorg.com ([147.34.97.41]) by relay1.mentorg.com with esmtp id 1bwyPb-0006cV-6H from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Wed, 19 Oct 2016 14:28:35 -0700 Received: from tftp-cs (147.34.91.1) by svr-orw-fem-04.mgc.mentorg.com (147.34.97.41) with Microsoft SMTP Server id 14.3.224.2; Wed, 19 Oct 2016 14:28:34 -0700 Received: by tftp-cs (Postfix, from userid 49978) id 1B638C224D; Wed, 19 Oct 2016 14:28:34 -0700 (PDT) From: Thomas Schwinge To: Subject: Re: [PR other/70945] Handle function_glibc_finite_math in offloading In-Reply-To: <87vb27tnai.fsf@hertz.schwinge.homeip.net> References: <87vb27tnai.fsf@hertz.schwinge.homeip.net> User-Agent: Notmuch/0.9-125-g4686d11 (http://notmuchmail.org) Emacs/24.5.1 (i686-pc-linux-gnu) Date: Wed, 19 Oct 2016 23:28:23 +0200 Message-ID: <877f94111k.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Sat, 21 May 2016 17:59:17 +0200, I wrote: > As discussed in "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 is known to include 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 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(-) Grüße Thomas 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 + + 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 * 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 is known to include 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 + + PR other/70945 + * testsuite/libgomp.oacc-c-c++-common/pr70945-1.c: New file. + 2016-10-05 Nathan Sidwell * 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 +# include +# include +# include +#else /* C++ */ +# include +# include +# include +# include +#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; +}