diff options
63 files changed, 1217 insertions, 134 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index d7e8702..59f447c 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,26 @@ +2025-05-29 Yuta Mukai <mukai.yuta@fujitsu.com> + + Backported from master: + 2025-05-28 Yuta Mukai <mukai.yuta@fujitsu.com> + + * config/aarch64/aarch64-cores.def (fujitsu-monaka): Update ISA + features. + +2025-05-27 Eric Botcazou <ebotcazou@adacore.com> + + * ipa-sra.cc (scan_expr_access): Also disqualify storage order + barriers from splitting. + * tree.h (storage_order_barrier_p): Also return false if the + operand of the VIEW_CONVERT_EXPR has reverse storage order. + +2025-05-27 Jonathan Wakely <jwakely@redhat.com> + + Backported from master: + 2025-05-27 Jonathan Wakely <jwakely@redhat.com> + + * doc/extend.texi (Common Variable Attributes): Fix typo in + description of nonstring. + 2025-05-25 Michael J. Eager <eager@eagercon.com> PR target/86772 diff --git a/gcc/DATESTAMP b/gcc/DATESTAMP index dbf258b..5646e6e 100644 --- a/gcc/DATESTAMP +++ b/gcc/DATESTAMP @@ -1 +1 @@ -20250526 +20250602 diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index d842028..061f1c6 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,29 @@ +2025-05-30 Sandra Loosemore <sloosemore@baylibre.com> + + Backported from master: + 2025-05-29 Sandra Loosemore <sloosemore@baylibre.com> + + * c-parser.cc (c_parser_skip_to_closing_brace): New, copied from + the equivalent function in the C++ front end. + (c_parser_skip_to_end_of_block_or_statement): Pass false to + the error flag. + (c_parser_omp_context_selector): Immediately return error_mark_node + after giving an error that the integer trait property is invalid, + similarly to C++ front end. + (c_parser_omp_context_selector_specification): Likewise handle + error return from c_parser_omp_context_selector similarly to C++. + (c_parser_omp_metadirective): Do not call + c_parser_skip_to_end_of_block_or_statement after an error. + +2025-05-30 Sandra Loosemore <sloosemore@baylibre.com> + + Backported from master: + 2025-05-29 Sandra Loosemore <sloosemore@baylibre.com> + + PR c/120180 + * c-parser.cc (c_parser_omp_metadirective): Only consume the + token if it is the expected close paren. + 2025-05-02 Jakub Jelinek <jakub@redhat.com> Backported from master: diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 9af7440..d132704 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -1418,6 +1418,51 @@ c_parser_skip_to_end_of_parameter (c_parser *parser) parser->error = false; } +/* Skip tokens until a non-nested closing curly brace is the next + token, or there are no more tokens. Return true in the first case, + false otherwise. */ + +static bool +c_parser_skip_to_closing_brace (c_parser *parser) +{ + unsigned nesting_depth = 0; + + while (true) + { + c_token *token = c_parser_peek_token (parser); + + switch (token->type) + { + case CPP_PRAGMA_EOL: + if (!parser->in_pragma) + break; + /* FALLTHRU */ + case CPP_EOF: + /* If we've run out of tokens, stop. */ + return false; + + case CPP_CLOSE_BRACE: + /* If the next token is a non-nested `}', then we have reached + the end of the current block. */ + if (nesting_depth-- == 0) + return true; + break; + + case CPP_OPEN_BRACE: + /* If it the next token is a `{', then we are entering a new + block. Consume the entire block. */ + ++nesting_depth; + break; + + default: + break; + } + + /* Consume the token. */ + c_parser_consume_token (parser); + } +} + /* Expect to be at the end of the pragma directive and consume an end of line marker. */ @@ -1582,7 +1627,7 @@ c_parser_skip_to_end_of_block_or_statement (c_parser *parser, here for secondary error recovery, after parser->error has been cleared. */ c_parser_consume_pragma (parser); - c_parser_skip_to_pragma_eol (parser); + c_parser_skip_to_pragma_eol (parser, false); parser->error = save_error; continue; @@ -28381,19 +28426,17 @@ c_parser_omp_context_selector (c_parser *parser, enum omp_tss_code set, case OMP_TRAIT_PROPERTY_DEV_NUM_EXPR: case OMP_TRAIT_PROPERTY_BOOL_EXPR: t = c_parser_expr_no_commas (parser, NULL).value; - if (t != error_mark_node) + if (t == error_mark_node) + return error_mark_node; + mark_exp_read (t); + t = c_fully_fold (t, false, NULL); + if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) { - mark_exp_read (t); - t = c_fully_fold (t, false, NULL); - if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) - error_at (token->location, - "property must be integer expression"); - else - properties = make_trait_property (NULL_TREE, t, - properties); + error_at (token->location, + "property must be integer expression"); + return error_mark_node; } - else - return error_mark_node; + properties = make_trait_property (NULL_TREE, t, properties); break; case OMP_TRAIT_PROPERTY_CLAUSE_LIST: if (sel == OMP_TRAIT_CONSTRUCT_SIMD) @@ -28495,11 +28538,14 @@ c_parser_omp_context_selector_specification (c_parser *parser, tree parms) tree selectors = c_parser_omp_context_selector (parser, set, parms); if (selectors == error_mark_node) - ret = error_mark_node; + { + c_parser_skip_to_closing_brace (parser); + ret = error_mark_node; + } else if (ret != error_mark_node) ret = make_trait_set_selector (set, selectors, ret); - braces.skip_until_found_close (parser); + braces.require_close (parser); if (c_parser_next_token_is (parser, CPP_COMMA)) c_parser_consume_token (parser); @@ -31148,7 +31194,6 @@ c_parser_omp_metadirective (c_parser *parser, bool *if_p) { error_at (match_loc, "too many %<otherwise%> or %<default%> " "clauses in %<metadirective%>"); - c_parser_skip_to_end_of_block_or_statement (parser, true); goto error; } default_seen = true; @@ -31157,14 +31202,12 @@ c_parser_omp_metadirective (c_parser *parser, bool *if_p) { error_at (match_loc, "%<otherwise%> or %<default%> clause " "must appear last in %<metadirective%>"); - c_parser_skip_to_end_of_block_or_statement (parser, true); goto error; } if (!default_p && strcmp (p, "when") != 0) { error_at (match_loc, "%qs is not valid for %qs", p, "metadirective"); - c_parser_skip_to_end_of_block_or_statement (parser, true); goto error; } @@ -31232,7 +31275,6 @@ c_parser_omp_metadirective (c_parser *parser, bool *if_p) if (i == 0) { error_at (loc, "expected directive name"); - c_parser_skip_to_end_of_block_or_statement (parser, true); goto error; } @@ -31300,7 +31342,10 @@ c_parser_omp_metadirective (c_parser *parser, bool *if_p) goto add; case CPP_CLOSE_PAREN: if (nesting_depth-- == 0) - break; + { + c_parser_consume_token (parser); + break; + } goto add; default: add: @@ -31312,8 +31357,6 @@ c_parser_omp_metadirective (c_parser *parser, bool *if_p) break; } - c_parser_consume_token (parser); - if (!skip) { c_token eol_token; @@ -31440,9 +31483,9 @@ c_parser_omp_metadirective (c_parser *parser, bool *if_p) return; error: + /* Skip the metadirective pragma. Do not skip the metadirective body. */ if (parser->in_pragma) - c_parser_skip_to_pragma_eol (parser); - c_parser_skip_to_end_of_block_or_statement (parser, true); + c_parser_skip_to_pragma_eol (parser, false); } /* Main entry point to parsing most OpenMP pragmas. */ diff --git a/gcc/config/aarch64/aarch64-cores.def b/gcc/config/aarch64/aarch64-cores.def index 1209630..24b7cd3 100644 --- a/gcc/config/aarch64/aarch64-cores.def +++ b/gcc/config/aarch64/aarch64-cores.def @@ -132,7 +132,7 @@ AARCH64_CORE("octeontx2f95mm", octeontx2f95mm, cortexa57, V8_2A, (CRYPTO, PROFI /* Fujitsu ('F') cores. */ AARCH64_CORE("a64fx", a64fx, a64fx, V8_2A, (F16, SVE), a64fx, 0x46, 0x001, -1) -AARCH64_CORE("fujitsu-monaka", fujitsu_monaka, cortexa57, V9_3A, (F16, FP8, LS64, RNG, CRYPTO, SVE2_AES, SVE2_BITPERM, SVE2_SHA3, SVE2_SM4), fujitsu_monaka, 0x46, 0x003, -1) +AARCH64_CORE("fujitsu-monaka", fujitsu_monaka, cortexa57, V9_3A, (F16, FAMINMAX, FP8FMA, FP8DOT2, FP8DOT4, LS64, LUT, RNG, CRYPTO, SVE2_AES, SVE2_BITPERM, SVE2_SHA3, SVE2_SM4), fujitsu_monaka, 0x46, 0x003, -1) /* HiSilicon ('H') cores. */ AARCH64_CORE("tsv110", tsv110, tsv110, V8_2A, (CRYPTO, F16), tsv110, 0x48, 0xd01, -1) diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 59bc179..8983abf 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,29 @@ +2025-05-30 Sandra Loosemore <sloosemore@baylibre.com> + + Backported from master: + 2025-05-29 Sandra Loosemore <sloosemore@baylibre.com> + + * parser.cc (cp_parser_omp_metadirective): Do not call + cp_parser_skip_to_end_of_block_or_statement after an error. + +2025-05-30 Sandra Loosemore <sloosemore@baylibre.com> + + Backported from master: + 2025-05-29 Sandra Loosemore <sloosemore@baylibre.com> + + PR c/120180 + * parser.cc (cp_parser_omp_metadirective): Only consume the + token if it is the expected close paren. + +2025-05-26 Tobias Burnus <tburnus@baylibre.com> + + Backported from master: + 2025-05-26 Tobias Burnus <tburnus@baylibre.com> + + PR c++/120413 + * semantics.cc (finish_omp_target_clauses_r): Handle + BIND_EXPR with empty BIND_EXPR_BLOCK. + 2025-05-23 Nathaniel Shead <nathanieloshead@gmail.com> Backported from master: diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index be86252..9c1d976 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -53390,7 +53390,6 @@ cp_parser_omp_metadirective (cp_parser *parser, cp_token *pragma_tok, { error_at (match_loc, "too many %<otherwise%> or %<default%> " "clauses in %<metadirective%>"); - cp_parser_skip_to_end_of_block_or_statement (parser, true); goto fail; } else @@ -53400,14 +53399,12 @@ cp_parser_omp_metadirective (cp_parser *parser, cp_token *pragma_tok, { error_at (match_loc, "%<otherwise%> or %<default%> clause " "must appear last in %<metadirective%>"); - cp_parser_skip_to_end_of_block_or_statement (parser, true); goto fail; } if (!default_p && strcmp (p, "when") != 0) { error_at (match_loc, "%qs is not valid for %qs", p, "metadirective"); - cp_parser_skip_to_end_of_block_or_statement (parser, true); goto fail; } @@ -53476,7 +53473,6 @@ cp_parser_omp_metadirective (cp_parser *parser, cp_token *pragma_tok, if (i == 0) { error_at (loc, "expected directive name"); - cp_parser_skip_to_end_of_block_or_statement (parser, true); goto fail; } @@ -53549,7 +53545,10 @@ cp_parser_omp_metadirective (cp_parser *parser, cp_token *pragma_tok, goto add; case CPP_CLOSE_PAREN: if (nesting_depth-- == 0) - break; + { + cp_lexer_consume_token (parser->lexer); + break; + } goto add; default: add: @@ -53561,8 +53560,6 @@ cp_parser_omp_metadirective (cp_parser *parser, cp_token *pragma_tok, break; } - cp_lexer_consume_token (parser->lexer); - if (!skip) { cp_token eol_token = {}; @@ -53694,11 +53691,8 @@ cp_parser_omp_metadirective (cp_parser *parser, cp_token *pragma_tok, return; fail: - /* Skip the metadirective pragma. */ + /* Skip the metadirective pragma. Do not skip the metadirective body. */ cp_parser_skip_to_pragma_eol (parser, pragma_tok); - - /* Skip the metadirective body. */ - cp_parser_skip_to_end_of_block_or_statement (parser, true); } diff --git a/gcc/doc/extend.texi b/gcc/doc/extend.texi index ec68c85..bad3408 100644 --- a/gcc/doc/extend.texi +++ b/gcc/doc/extend.texi @@ -7336,7 +7336,7 @@ truncate the copy without appending the terminating @code{NUL} character. Using the attribute makes it possible to suppress the warning. However, when the array is declared with the attribute the call to @code{strlen} is diagnosed because when the array doesn't contain a @code{NUL}-terminated -string the call is undefined. To copy, compare, of search non-string +string the call is undefined. To copy, compare, or search non-string character arrays use the @code{memcpy}, @code{memcmp}, @code{memchr}, and other functions that operate on arrays of bytes. In addition, calling @code{strnlen} and @code{strndup} with such arrays is safe diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index b95f7ee..2f9f5c9 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,28 @@ +2025-05-31 Jerry DeLisle <jvdelisle@gcc.gnu.org> + + Backported from master: + 2025-05-27 Jerry DeLisle <jvdelisle@gcc.gnu.org> + + PR fortran/120049 + * check.cc (gfc_check_c_associated): Use new helper functions. + Only call check_c_ptr_1 if optional c_ptr_2 tests succeed. + (check_c_ptr_1): Handle only c_ptr_1 checks. + (check_c_ptr_2): Expand checks for c_ptr_2 and handle cases + where there is no derived pointer in the gfc_expr and check + the inmod_sym_id only if it exists. Rephrase error message. + * misc.cc (gfc_typename): Handle the case for BT_VOID rather + than throw an internal error. + +2025-05-31 Thomas Koenig <tkoenig@gcc.gnu.org> + + Backported from master: + 2025-05-30 Thomas Koenig <tkoenig@gcc.gnu.org> + + PR fortran/120355 + * interface.cc (compare_parameter): If the global function has a + result clause, take typespec from there for the comparison against + the dummy argument. + 2025-05-22 Harald Anlauf <anlauf@gmx.de> Backported from master: diff --git a/gcc/fortran/check.cc b/gcc/fortran/check.cc index 0073cd0..ce9da31 100644 --- a/gcc/fortran/check.cc +++ b/gcc/fortran/check.cc @@ -5916,49 +5916,110 @@ gfc_check_c_sizeof (gfc_expr *arg) } -bool -gfc_check_c_associated (gfc_expr *c_ptr_1, gfc_expr *c_ptr_2) +/* Helper functions check_c_ptr_1 and check_c_ptr_2 + used in gfc_check_c_associated. */ + +static inline +bool check_c_ptr_1 (gfc_expr *c_ptr_1) { - if (c_ptr_1) - { - if (c_ptr_1->expr_type == EXPR_FUNCTION && c_ptr_1->ts.type == BT_VOID) - return true; + if ((c_ptr_1->ts.type == BT_VOID) + && (c_ptr_1->expr_type == EXPR_FUNCTION)) + return true; - if (c_ptr_1->ts.type != BT_DERIVED - || c_ptr_1->ts.u.derived->from_intmod != INTMOD_ISO_C_BINDING - || (c_ptr_1->ts.u.derived->intmod_sym_id != ISOCBINDING_PTR - && c_ptr_1->ts.u.derived->intmod_sym_id != ISOCBINDING_FUNPTR)) - { - gfc_error ("Argument C_PTR_1 at %L to C_ASSOCIATED shall have the " - "type TYPE(C_PTR) or TYPE(C_FUNPTR)", &c_ptr_1->where); - return false; - } - } + if (c_ptr_1->ts.type != BT_DERIVED + || c_ptr_1->ts.u.derived->from_intmod != INTMOD_ISO_C_BINDING + || (c_ptr_1->ts.u.derived->intmod_sym_id != ISOCBINDING_PTR + && c_ptr_1->ts.u.derived->intmod_sym_id != ISOCBINDING_FUNPTR)) + goto check_1_error; - if (!scalar_check (c_ptr_1, 0)) + if ((c_ptr_1->ts.type == BT_DERIVED) + && (c_ptr_1->expr_type == EXPR_STRUCTURE) + && (c_ptr_1->ts.u.derived->intmod_sym_id + == ISOCBINDING_NULL_FUNPTR)) + goto check_1_error; + + if (scalar_check (c_ptr_1, 0)) + return true; + else + /* Return since the check_1_error message may not apply here. */ return false; - if (c_ptr_2) - { - if (c_ptr_2->expr_type == EXPR_FUNCTION && c_ptr_2->ts.type == BT_VOID) - return true; +check_1_error: - if (c_ptr_2->ts.type != BT_DERIVED - || c_ptr_2->ts.u.derived->from_intmod != INTMOD_ISO_C_BINDING - || (c_ptr_1->ts.u.derived->intmod_sym_id - != c_ptr_2->ts.u.derived->intmod_sym_id)) + gfc_error ("Argument C_PTR_1 at %L to C_ASSOCIATED shall have the " + "type TYPE(C_PTR) or TYPE(C_FUNPTR)", &c_ptr_1->where); + return false; +} + +static inline +bool check_c_ptr_2 (gfc_expr *c_ptr_1, gfc_expr *c_ptr_2) +{ + switch (c_ptr_2->ts.type) + { + case BT_VOID: + if (c_ptr_2->expr_type == EXPR_FUNCTION) { - gfc_error ("Argument C_PTR_2 at %L to C_ASSOCIATED shall have the " - "same type as C_PTR_1: %s instead of %s", &c_ptr_1->where, - gfc_typename (&c_ptr_1->ts), gfc_typename (&c_ptr_2->ts)); - return false; + if ((c_ptr_1->ts.type == BT_DERIVED) + && c_ptr_1->expr_type == EXPR_STRUCTURE + && (c_ptr_1->ts.u.derived->intmod_sym_id + == ISOCBINDING_FUNPTR)) + goto check_2_error; } - } + break; + + case BT_DERIVED: + if ((c_ptr_2->expr_type == EXPR_STRUCTURE) + && (c_ptr_2->ts.u.derived->intmod_sym_id == ISOCBINDING_PTR) + && (c_ptr_1->ts.type == BT_VOID) + && (c_ptr_1->expr_type == EXPR_FUNCTION)) + return scalar_check (c_ptr_2, 1); + + if ((c_ptr_2->expr_type == EXPR_STRUCTURE) + && (c_ptr_1->ts.type == BT_VOID) + && (c_ptr_1->expr_type == EXPR_FUNCTION)) + goto check_2_error; - if (c_ptr_2 && !scalar_check (c_ptr_2, 1)) + if (c_ptr_2->ts.u.derived->from_intmod != INTMOD_ISO_C_BINDING) + goto check_2_error; + + if (c_ptr_1->ts.type == BT_DERIVED + && (c_ptr_1->ts.u.derived->intmod_sym_id + != c_ptr_2->ts.u.derived->intmod_sym_id)) + goto check_2_error; + break; + + default: + goto check_2_error; + } + + if (scalar_check (c_ptr_2, 1)) + return true; + else + /* Return since the check_2_error message may not apply here. */ return false; - return true; +check_2_error: + + gfc_error ("Argument C_PTR_2 at %L to C_ASSOCIATED shall have the " + "same type as C_PTR_1, found %s instead of %s", &c_ptr_2->where, + gfc_typename (&c_ptr_2->ts), gfc_typename (&c_ptr_1->ts)); + + return false; + } + + +bool +gfc_check_c_associated (gfc_expr *c_ptr_1, gfc_expr *c_ptr_2) +{ + if (c_ptr_2) + { + if (check_c_ptr_2 (c_ptr_1, c_ptr_2)) + return check_c_ptr_1 (c_ptr_1); + else + return false; + } + else + return check_c_ptr_1 (c_ptr_1); } diff --git a/gcc/fortran/interface.cc b/gcc/fortran/interface.cc index 753f589..b854292 100644 --- a/gcc/fortran/interface.cc +++ b/gcc/fortran/interface.cc @@ -2547,7 +2547,14 @@ compare_parameter (gfc_symbol *formal, gfc_expr *actual, } else if (formal->attr.function) { - if (!gfc_compare_types (&global_asym->ts, + gfc_typespec ts; + + if (global_asym->result) + ts = global_asym->result->ts; + else + ts = global_asym->ts; + + if (!gfc_compare_types (&ts, &formal->ts)) { gfc_error ("Type mismatch at %L passing global " diff --git a/gcc/fortran/io.cc b/gcc/fortran/io.cc index b5c9d33..7466d8f 100644 --- a/gcc/fortran/io.cc +++ b/gcc/fortran/io.cc @@ -1228,7 +1228,8 @@ between_desc: default: if (mode != MODE_FORMAT) format_locus.nextc += format_string_pos - 1; - if (!gfc_notify_std (GFC_STD_GNU, "Missing comma at %L", &format_locus)) + if (!gfc_notify_std (GFC_STD_LEGACY, + "Missing comma in FORMAT string at %L", &format_locus)) return false; /* If we do not actually return a failure, we need to unwind this before the next round. */ @@ -1290,7 +1291,8 @@ extension_optional_comma: default: if (mode != MODE_FORMAT) format_locus.nextc += format_string_pos; - if (!gfc_notify_std (GFC_STD_GNU, "Missing comma at %L", &format_locus)) + if (!gfc_notify_std (GFC_STD_LEGACY, + "Missing comma in FORMAT string at %L", &format_locus)) return false; /* If we do not actually return a failure, we need to unwind this before the next round. */ diff --git a/gcc/fortran/misc.cc b/gcc/fortran/misc.cc index 893c40f..b8bdf75 100644 --- a/gcc/fortran/misc.cc +++ b/gcc/fortran/misc.cc @@ -214,6 +214,9 @@ gfc_typename (gfc_typespec *ts, bool for_hash) case BT_UNKNOWN: strcpy (buffer, "UNKNOWN"); break; + case BT_VOID: + strcpy (buffer, "VOID"); + break; default: gfc_internal_error ("gfc_typename(): Undefined type"); } diff --git a/gcc/ipa-sra.cc b/gcc/ipa-sra.cc index 88bfae9..6e6cf89 100644 --- a/gcc/ipa-sra.cc +++ b/gcc/ipa-sra.cc @@ -1848,6 +1848,12 @@ scan_expr_access (tree expr, gimple *stmt, isra_scan_context ctx, if (!desc || !desc->split_candidate) return; + if (storage_order_barrier_p (expr)) + { + disqualify_split_candidate (desc, "Encountered a storage order barrier."); + return; + } + if (!poffset.is_constant (&offset) || !psize.is_constant (&size) || !pmax_size.is_constant (&max_size)) diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 750a1b9..5b94c72 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,75 @@ +2025-05-31 Jerry DeLisle <jvdelisle@gcc.gnu.org> + + Backported from master: + 2025-05-27 Jerry DeLisle <jvdelisle@gcc.gnu.org> + + PR fortran/120049 + * gfortran.dg/pr120049_a.f90: Update test directives. + * gfortran.dg/pr120049_b.f90: Update test directives + * gfortran.dg/pr120049_2.f90: New test. + * gfortran.dg/c_f_pointer_tests_6.f90: Adjust dg-error + directive. + Co-Authored-By: Steve Kargl <kargl@gcc.gnu.org> + +2025-05-31 Thomas Koenig <tkoenig@gcc.gnu.org> + + Backported from master: + 2025-05-30 Thomas Koenig <tkoenig@gcc.gnu.org> + + PR fortran/120355 + * gfortran.dg/interface_62.f90: New test. + +2025-05-30 Sandra Loosemore <sloosemore@baylibre.com> + + Backported from master: + 2025-05-29 Sandra Loosemore <sloosemore@baylibre.com> + + * c-c++-common/gomp/declare-variant-2.c: Adjust patterns now that + C and C++ now behave similarly. + * c-c++-common/gomp/metadirective-error-recovery.c: New. + +2025-05-30 Sandra Loosemore <sloosemore@baylibre.com> + + Backported from master: + 2025-05-29 Sandra Loosemore <sloosemore@baylibre.com> + + PR c/120180 + * c-c++-common/gomp/pr120180.c: New. + +2025-05-30 Jakub Jelinek <jakub@redhat.com> + + Backported from master: + 2025-05-30 Jakub Jelinek <jakub@redhat.com> + + PR target/120480 + * gcc.dg/pr120480.c: New test. + +2025-05-27 Eric Botcazou <ebotcazou@adacore.com> + + * gnat.dg/sso19.adb: New test. + * gnat.dg/sso19_pkg.ads, gnat.dg/sso19_pkg.adb: New helper. + +2025-05-26 Tobias Burnus <tburnus@baylibre.com> + + Backported from master: + 2025-05-26 Tobias Burnus <tburnus@baylibre.com> + + PR middle-end/118694 + * c-c++-common/gomp/attrs-metadirective-3.c: Change to never + expect 'omp metadirective' in the dump. If !offload_nvptx, check + that no 'teams' shows up in the dump; for offload_nvptx, expect + OMP_NEXT_VARIANT and an error about directive between 'target' + and 'teams'. + * c-c++-common/gomp/metadirective-3.c: Likewise. + +2025-05-26 Tobias Burnus <tburnus@baylibre.com> + + Backported from master: + 2025-05-26 Tobias Burnus <tburnus@baylibre.com> + + PR c++/120413 + * g++.dg/gomp/target-4.C: New test. + 2025-05-23 Nathaniel Shead <nathanieloshead@gmail.com> Backported from master: diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c b/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c index 7711dbc..f8f5143 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c @@ -47,10 +47,9 @@ void f23 (void); #pragma omp declare variant (f1) match(construct={teams,parallel,master,for}) /* { dg-warning "unknown selector 'master' for context selector set 'construct'" } */ void f24 (void); #pragma omp declare variant (f1) match(construct={parallel(1 /* { dg-error "selector 'parallel' does not accept any properties" } */ -void f25 (void); /* { dg-error "expected '\\\}' before end of line" "" { target c++ } .-1 } */ - /* { dg-error "expected '\\\}' before '\\(' token" "" { target c } .-2 } */ +void f25 (void); /* { dg-error "expected '\\\}' before end of line" "" { target *-*-* } .-1 } */ #pragma omp declare variant (f1) match(construct={parallel(1)}) /* { dg-error "selector 'parallel' does not accept any properties" } */ -void f26 (void); /* { dg-error "expected '\\\}' before '\\(' token" "" { target c } .-1 } */ +void f26 (void); #pragma omp declare variant (f0) match(construct={simd(12)}) /* { dg-error "expected \[^\n\r]* clause before" } */ void f27 (void); /* { dg-error "'\\)' before numeric constant" "" { target c++ } .-1 } */ #pragma omp declare variant (f1) match(construct={parallel},construct={for}) /* { dg-error "selector set 'construct' specified more than once" } */ @@ -96,13 +95,13 @@ void f46 (void); #pragma omp declare variant (f1) match(implementation={vendor("foobar")}) /* { dg-warning "unknown property '.foobar.' of 'vendor' selector" } */ void f47 (void); #pragma omp declare variant (f1) match(implementation={unified_address(yes)}) /* { dg-error "selector 'unified_address' does not accept any properties" } */ -void f48 (void); /* { dg-error "expected '\\\}' before '\\(' token" "" { target c } .-1 } */ +void f48 (void); #pragma omp declare variant (f1) match(implementation={unified_shared_memory(no)}) /* { dg-error "selector 'unified_shared_memory' does not accept any properties" } */ -void f49 (void); /* { dg-error "expected '\\\}' before '\\(' token" "" { target c } .-1 } */ +void f49 (void); #pragma omp declare variant (f1) match(implementation={dynamic_allocators(42)}) /* { dg-error "selector 'dynamic_allocators' does not accept any properties" } */ -void f50 (void); /* { dg-error "expected '\\\}' before '\\(' token" "" { target c } .-1 } */ +void f50 (void); #pragma omp declare variant (f1) match(implementation={reverse_offload()}) /* { dg-error "selector 'reverse_offload' does not accept any properties" } */ -void f51 (void); /* { dg-error "expected '\\\}' before '\\(' token" "" { target c } .-1 } */ +void f51 (void); #pragma omp declare variant (f1) match(implementation={atomic_default_mem_order}) /* { dg-error "expected '\\(' before '\\\}' token" } */ void f52 (void); #pragma omp declare variant (f1) match(implementation={atomic_default_mem_order(acquire)}) diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-error-recovery.c b/gcc/testsuite/c-c++-common/gomp/metadirective-error-recovery.c new file mode 100644 index 0000000..3242281 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-error-recovery.c @@ -0,0 +1,20 @@ +/* { dg-do compile } */ + +/* This test used to ICE in C and only diagnose the first error in C++. */ + +struct s { + int a, b; +}; + +void f (int aa, int bb) +{ + struct s s1, s2; + s1.a = aa; + s1.b = bb; + s2.a = aa + 1; + s2.b = bb + 1; + + /* A struct is not a valid argument for the condition selector. */ + #pragma omp metadirective when(user={condition(s1)} : nothing) otherwise(nothing) /* { dg-error "property must be integer expression" } */ + #pragma omp metadirective when(user={condition(s2)} : nothing) otherwise(nothing) /* { dg-error "property must be integer expression" } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/pr120180.c b/gcc/testsuite/c-c++-common/gomp/pr120180.c new file mode 100644 index 0000000..cb5a0d5 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/pr120180.c @@ -0,0 +1,22 @@ +/* { dg-do compile } */ + +/* This test used to ICE after erroring on the metadirective in the + loop nest. */ + +int main() +{ + int blksize = 15000; + double *qq; + int i, k, nq; + + #pragma omp metadirective when(user={condition(0)}: target teams distribute parallel for collapse(2) map(qq[:0]) private(i)) \ + when(user={condition(0)}: target teams distribute parallel for map(qq[:0]) private(i)) \ + when(user={condition(1)}: target teams loop collapse(2) map(qq[:0]) private(i)) + for(k=0; k<blksize; k++) + { +#pragma omp metadirective when(user={condition(0)}: simd) default() // { dg-error "intervening code must not contain OpenMP directives" } + for (i=0; i<nq; i++) + qq[k*nq + i] = 0.0; + } + return 0; +} diff --git a/gcc/testsuite/gcc.dg/pr120480.c b/gcc/testsuite/gcc.dg/pr120480.c new file mode 100644 index 0000000..cf7b47a --- /dev/null +++ b/gcc/testsuite/gcc.dg/pr120480.c @@ -0,0 +1,11 @@ +/* PR target/120480 */ +/* { dg-do compile } */ +/* { dg-options "-O0" } */ + +struct S { int a, b, c; } s; + +void +foo (void) +{ + struct S t = s; +} diff --git a/gcc/testsuite/gfortran.dg/c_f_pointer_tests_6.f90 b/gcc/testsuite/gfortran.dg/c_f_pointer_tests_6.f90 index 23ca88b..bc2206d 100644 --- a/gcc/testsuite/gfortran.dg/c_f_pointer_tests_6.f90 +++ b/gcc/testsuite/gfortran.dg/c_f_pointer_tests_6.f90 @@ -38,6 +38,6 @@ contains type(my_c_ptr_0) :: my_ptr2 type(c_funptr) :: myfun print *,c_associated(my_ptr,my_ptr2) - print *,c_associated(my_ptr,myfun) ! { dg-error "Argument C_PTR_2 at .1. to C_ASSOCIATED shall have the same type as C_PTR_1: TYPE.c_ptr. instead of TYPE.c_funptr." } + print *,c_associated(my_ptr,myfun) ! { dg-error "Argument C_PTR_2 at .1. to C_ASSOCIATED shall have the same type as C_PTR_1, found TYPE.c_funptr. instead of TYPE.c_ptr." } end subroutine end diff --git a/gcc/testsuite/gfortran.dg/comma_format_extension_1.f b/gcc/testsuite/gfortran.dg/comma_format_extension_1.f index a3a5a98..c4b43f0 100644 --- a/gcc/testsuite/gfortran.dg/comma_format_extension_1.f +++ b/gcc/testsuite/gfortran.dg/comma_format_extension_1.f @@ -1,5 +1,5 @@ ! { dg-do compile } -! { dg-options "" } +! { dg-options "-std=legacy" } ! test that the extension for a missing comma is accepted subroutine mysub diff --git a/gcc/testsuite/gfortran.dg/comma_format_extension_3.f b/gcc/testsuite/gfortran.dg/comma_format_extension_3.f index 0b00224..9d974d6 100644 --- a/gcc/testsuite/gfortran.dg/comma_format_extension_3.f +++ b/gcc/testsuite/gfortran.dg/comma_format_extension_3.f @@ -3,7 +3,7 @@ ! did do the correct thing at runtime. ! Note the missing , before i1 in the format. ! { dg-do run } -! { dg-options "" } +! { dg-options "-std=legacy" } character*12 c write (c,100) 0, 1 diff --git a/gcc/testsuite/gfortran.dg/continuation_13.f90 b/gcc/testsuite/gfortran.dg/continuation_13.f90 index 9799b59e..475c896 100644 --- a/gcc/testsuite/gfortran.dg/continuation_13.f90 +++ b/gcc/testsuite/gfortran.dg/continuation_13.f90 @@ -1,5 +1,5 @@ ! { dg-do run } -! { dg-options "-std=gnu" } +! { dg-options "-std=legacy" } ! PR64506 character(25) :: astring diff --git a/gcc/testsuite/gfortran.dg/interface_62.f90 b/gcc/testsuite/gfortran.dg/interface_62.f90 new file mode 100644 index 0000000..19d4325 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/interface_62.f90 @@ -0,0 +1,39 @@ +! { dg-do compile } +! PR fortran/120355 - this was rejected because the typespec from +! the RESULT clause was not picked up. +! Test case jsberg@bnl.gov. + +program p + implicit none + integer :: i,j + interface + function s(x) result(y) + implicit none + integer, intent(in) :: x + integer :: y + end function s + end interface + i = 0 + call t(s,i,j) +contains + subroutine t(f,x,y) + implicit none + integer, intent(in) :: x + integer, intent(out) :: y + interface + function f(x) result(y) + implicit none + integer, intent(in) :: x + integer :: y + end function f + end interface + y = f(x) + end subroutine t +end program p + +function s(x) result(y) + implicit none + integer, intent(in) :: x + integer :: y + y = 1 - x +end function s diff --git a/gcc/testsuite/gfortran.dg/pr119856.f90 b/gcc/testsuite/gfortran.dg/pr119856.f90 new file mode 100644 index 0000000..60ada0a --- /dev/null +++ b/gcc/testsuite/gfortran.dg/pr119856.f90 @@ -0,0 +1,15 @@ +! { dg-do run } +! PR119856, the error should occur in both write statements. +program badfmt + implicit none + + character(10):: fmt = "(AI5)" ! Not a PARAMETER so not examined + ! at compile time + integer :: ioerr + ioerr = 0 + write (*, fmt, iostat=ioerr) 'value =', 42 + if (ioerr /= 5006) stop 10 +! + write (*, fmt, iostat=ioerr) 'value =', 43 + if (ioerr /= 5006) stop 13 +end program badfmt diff --git a/gcc/testsuite/gfortran.dg/pr120049_2.f90 b/gcc/testsuite/gfortran.dg/pr120049_2.f90 new file mode 100644 index 0000000..1f91e06 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/pr120049_2.f90 @@ -0,0 +1,62 @@ +! Compiled with pr120049_b.f90 +! { dg-options -O0 } +! { dg-do compile } +! { dg-compile-aux-modules "pr120049_b.f90" } +! +! Test the fix for PR120049 +program tests_gtk_sup + use gtk_sup + implicit none + + type mytype + integer :: myint + end type mytype + type(mytype) :: ijkl = mytype(42) + logical :: truth + real :: var1 + type(c_ptr), target :: val + type(c_funptr), target :: fptr + character(15) :: stringy + complex :: certainly + truth = .true. + var1 = 86. + stringy = "what the hay!" + certainly = (3.14,-4.13) + if (c_associated(val, c_loc(val))) then + stop 1 + endif + if (c_associated(c_loc(val), val)) then + stop 2 + endif + print *, c_associated(fptr, C_NULL_FUNPTR) + print *, c_associated(c_loc(val), C_NULL_PTR) + print *, c_associated(C_NULL_PTR, c_loc(val)) + print *, c_associated(c_loc(val), 42) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(c_loc(val), .42) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(c_loc(val), truth) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(c_loc(val), .false.) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(c_loc(val), var1) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(c_loc(val), stringy) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(c_loc(val), certainly) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(42) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(.42) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(truth) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(.false.) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(var1) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(stringy) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(certainly) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(.42) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(val, testit(val)) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(testit(val), val) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(testit(val)) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(c_loc(val), C_NULL_FUNPTR) ! { dg-error "C_ASSOCIATED shall have the" } + print *, c_associated(C_NULL_FUNPTR, c_loc(val)) ! { dg-error "C_ASSOCIATED shall have the" } +contains + + function testit (avalue) result(res) + type(c_ptr) :: avalue + type(mytype) :: res + res%myint = 42 + end function + +end program tests_gtk_sup diff --git a/gcc/testsuite/gfortran.dg/pr120049_a.f90 b/gcc/testsuite/gfortran.dg/pr120049_a.f90 index c404a4d..7095314 100644 --- a/gcc/testsuite/gfortran.dg/pr120049_a.f90 +++ b/gcc/testsuite/gfortran.dg/pr120049_a.f90 @@ -1,5 +1,8 @@ -! { dg-do preprocess } -! { dg-additional-options "-cpp" } +! Compiled with pr120049_b.f90 +! { dg-options -O0 } +! { dg-do run } +! { dg-compile-aux-modules "pr120049_b.f90" } +! { dg-additional-sources pr120049_b.f90 } ! ! Test the fix for PR86248 program tests_gtk_sup diff --git a/gcc/testsuite/gfortran.dg/pr120049_b.f90 b/gcc/testsuite/gfortran.dg/pr120049_b.f90 index 127db98..28a2783 100644 --- a/gcc/testsuite/gfortran.dg/pr120049_b.f90 +++ b/gcc/testsuite/gfortran.dg/pr120049_b.f90 @@ -1,5 +1,3 @@ -! { dg-do run } -! { dg-additional-sources pr120049_a.f90 } ! ! Module for pr120049.f90 ! diff --git a/gcc/testsuite/gnat.dg/sso19.adb b/gcc/testsuite/gnat.dg/sso19.adb new file mode 100644 index 0000000..497d9874 --- /dev/null +++ b/gcc/testsuite/gnat.dg/sso19.adb @@ -0,0 +1,13 @@ +-- { dg-do run } +-- { dg-options "-O2" } + +with SSO19_Pkg; use SSO19_Pkg; + +procedure SSO19 is + R : constant Rec := (D => (I => 8, F => 4.6095713E-41)); + +begin + if not Is_Valid (R) then + raise Program_Error; + end if; +end; diff --git a/gcc/testsuite/gnat.dg/sso19_pkg.adb b/gcc/testsuite/gnat.dg/sso19_pkg.adb new file mode 100644 index 0000000..cbcb2f9 --- /dev/null +++ b/gcc/testsuite/gnat.dg/sso19_pkg.adb @@ -0,0 +1,13 @@ +package body SSO19_Pkg is + + function Is_Valid_Private (Item : Data) return Boolean is + begin + return Item.I'Valid and Item.F'Valid; + end Is_Valid_Private; + + function Is_Valid (Item : Rec) return Boolean is + begin + return Is_Valid_Private (Item.D); + end Is_Valid; + +end SSO19_Pkg; diff --git a/gcc/testsuite/gnat.dg/sso19_pkg.ads b/gcc/testsuite/gnat.dg/sso19_pkg.ads new file mode 100644 index 0000000..cad5368 --- /dev/null +++ b/gcc/testsuite/gnat.dg/sso19_pkg.ads @@ -0,0 +1,24 @@ +with System; + +package SSO19_Pkg is + + subtype Small_Int is Short_Integer range -1000 .. 1000; + + type Data is record + I : Small_Int; + F : Float; + end record; + for Data use record + I at 0 range 0 .. 15; + F at 4 range 0 .. 31; + end record; + for Data'Bit_Order use System.High_Order_First; + for Data'Scalar_Storage_Order use System.High_Order_First; + + type Rec is record + D : Data; + end record; + + function Is_Valid (Item : Rec) return Boolean; + +end SSO19_Pkg; @@ -5572,7 +5572,7 @@ storage_order_barrier_p (const_tree t) && TYPE_REVERSE_STORAGE_ORDER (TREE_TYPE (op))) return true; - return false; + return reverse_storage_order_for_component_p (op); } /* Given a DECL or TYPE, return the scope in which it was declared, or diff --git a/libgcc/config/avr/libf7/ChangeLog b/libgcc/config/avr/libf7/ChangeLog index d7637ee..0e6afa0 100644 --- a/libgcc/config/avr/libf7/ChangeLog +++ b/libgcc/config/avr/libf7/ChangeLog @@ -1,3 +1,24 @@ +2025-05-27 Georg-Johann Lay <avr@gjlay.de> + + Backported from master: + 2025-05-27 Georg-Johann Lay <avr@gjlay.de> + + PR target/120442 + * libf7-common.mk (LIBF_C_PARTS, m_ddd): Add fdim. + * libf7.h (f7_fdim): New proto. + * libf7.c (f7_fdim): New function. + * f7renames.sh (f7_fdim): Add rename. + * f7-wraps.h: Rebuild + * f7-renames.h: Rebuild + +2025-05-27 Georg-Johann Lay <avr@gjlay.de> + + Backported from master: + 2025-05-27 Georg-Johann Lay <avr@gjlay.de> + + PR target/120441 + * libf7.c (f7_exp): Limit aa->expo to 10 (not to 9). + 2025-04-25 Release Manager * GCC 15.1.0 released. diff --git a/libgcc/config/avr/libf7/f7-renames.h b/libgcc/config/avr/libf7/f7-renames.h index bbe571a..bce2dd3 100644 --- a/libgcc/config/avr/libf7/f7-renames.h +++ b/libgcc/config/avr/libf7/f7-renames.h @@ -97,6 +97,7 @@ #define f7_acos __f7_acos #define f7_atan __f7_atan #define f7_atan2 __f7_atan2 +#define f7_fdim __f7_fdim #define f7_mul_noround __f7_mul_noround #define f7_sqrt16_round __f7_sqrt16_round #define f7_sqrt16_floor __f7_sqrt16_floor diff --git a/libgcc/config/avr/libf7/f7-wraps.h b/libgcc/config/avr/libf7/f7-wraps.h index a455b7d..409492e 100644 --- a/libgcc/config/avr/libf7/f7-wraps.h +++ b/libgcc/config/avr/libf7/f7-wraps.h @@ -239,7 +239,7 @@ _ENDF __extendsfdf2 ;; Functions that usually live in libm: Depending on [long] double layout, ;; define <name> and <name>l as weak alias(es) of __<name> for <name> in: -;; pow fmin fmax fmod hypot atan2 +;; pow fmin fmax fmod hypot atan2 fdim ;; double __pow (double, double) #ifdef F7MOD_D_pow_ @@ -313,6 +313,18 @@ _DEFUN __atan2 _ENDF __atan2 #endif /* F7MOD_D_atan2_ */ +;; double __fdim (double, double) +#ifdef F7MOD_D_fdim_ +_DEFUN __fdim + DALIAS fdim + LALIAS fdiml + .global F7_NAME(fdim) + ldi ZH, hi8(gs(F7_NAME(fdim))) + ldi ZL, lo8(gs(F7_NAME(fdim))) + F7jmp call_ddd +_ENDF __fdim +#endif /* F7MOD_D_fdim_ */ + ;; Functions that usually live in libm: Depending on [long] double layout, ;; define <name> and <name>l as weak alias(es) of __<name> for <name> in: ;; ldexp frexp diff --git a/libgcc/config/avr/libf7/f7renames.sh b/libgcc/config/avr/libf7/f7renames.sh index 7ef251e..4ced423 100755 --- a/libgcc/config/avr/libf7/f7renames.sh +++ b/libgcc/config/avr/libf7/f7renames.sh @@ -35,9 +35,9 @@ EOF c) if [ x${PRE} != xf7_ ]; then - echo " " + echo "" echo "/* Renames for libf7.c, libf7.h. */" - echo " " + echo "" for x in $*; do echo "#define f7_$x ${PRE}$x" done @@ -46,9 +46,9 @@ EOF cst) if [ x${PRE} != xf7_ ]; then - echo " " + echo "" echo "/* Renames for libf7.c, libf7.h. */" - echo " " + echo "" for x in $*; do echo "#define f7_const_${x} ${PRE}const_${x}" echo "#define f7_const_${x}_P ${PRE}const_${x}_P" @@ -58,9 +58,9 @@ EOF asm) if [ x${PRE} != xf7_ ]; then - echo " " + echo "" echo "/* Renames for libf7-asm.sx, f7-wraps.h. */" - echo " " + echo "" for x in $*; do echo "#define f7_${x}_asm ${PRE}${x}_asm" done diff --git a/libgcc/config/avr/libf7/libf7-common.mk b/libgcc/config/avr/libf7/libf7-common.mk index 5d41107..644be2c 100644 --- a/libgcc/config/avr/libf7/libf7-common.mk +++ b/libgcc/config/avr/libf7/libf7-common.mk @@ -8,7 +8,7 @@ F7_C_PARTS += set_float get_float get_double set_double set_pdouble F7_C_PARTS += fabs neg fmin fmax minmax truncx trunc floor ceil round lround F7_C_PARTS += horner logx log log10 log2 exp pow10 pow powi F7_C_PARTS += sin cos tan cotan sincos sinh cosh tanh sinhcosh -F7_C_PARTS += asinacos asin acos atan atan2 +F7_C_PARTS += asinacos asin acos atan atan2 fdim F7_C_PARTS += abscmp_msb_ge cmp cmp_abs cmp_unordered F7_C_PARTS += const_1 const_1_2 const_1_3 @@ -34,7 +34,7 @@ g_xdd_cmp += le lt ge gt ne eq unord g_dx += floatunsidf floatsidf extendsfdf2 g_xd += fixdfsi fixdfdi fixunsdfdi fixunsdfsi truncdfsf2 -m_ddd += pow fmin fmax fmod hypot atan2 +m_ddd += pow fmin fmax fmod hypot atan2 fdim m_ddx += ldexp frexp m_dd += sqrt cbrt exp exp10 pow10 log log10 log2 sin cos tan cotan asin acos atan m_dd += ceil floor trunc round sinh cosh tanh @@ -59,7 +59,7 @@ F7F += lrint ldexp frexp exp logx log log10 log2 F7F += minmax fmax fmin floor ceil round lround trunc truncx F7F += horner pow10 exp10 pow powi F7F += sin cos tan cotan sincos sinh cosh tanh sinhcosh -F7F += asinacos asin acos atan atan2 +F7F += asinacos asin acos atan atan2 fdim F7F += mul_noround sqrt16_round sqrt16_floor F7F += clr_mant_lsbs abscmp_msb_ge lshrdi3 ashldi3 F7F += assert diff --git a/libgcc/config/avr/libf7/libf7.c b/libgcc/config/avr/libf7/libf7.c index a64554c..78c218a 100644 --- a/libgcc/config/avr/libf7/libf7.c +++ b/libgcc/config/avr/libf7/libf7.c @@ -928,6 +928,21 @@ void f7_sub (f7_t *cc, const f7_t *aa, const f7_t *bb) #endif // F7MOD_sub_ +#ifdef F7MOD_fdim_ +F7_WEAK +void f7_fdim (f7_t *cc, const f7_t *aa, const f7_t *bb) +{ + int8_t cmp = f7_cmp_unordered (aa, bb, true /*with_sign*/); + if (cmp == INT8_MIN) + return f7_set_nan (cc); + if (cmp < 0) + return f7_clr (cc); + + f7_sub (cc, aa, bb); +} +#endif // F7MOD_fdim_ + + #ifdef F7MOD_addsub_ static void return_with_sign (f7_t *cc, const f7_t *aa, int8_t c_sign) { @@ -1649,10 +1664,10 @@ void f7_exp (f7_t *cc, const f7_t *aa) return f7_set_nan (cc); /* The maximal exponent of 2 for a double is 1023, hence we may limit - to |A| < 1023 * ln2 ~ 709. We limit to 1024 ~ 1.99 * 2^9 */ + to |A| < 1023 * ln2 ~ 709. We limit to 1024 = 2^10 */ if (f7_class_inf (a_class) - || (f7_class_nonzero (a_class) && aa->expo >= 9)) + || (f7_class_nonzero (a_class) && aa->expo >= 10)) { if (f7_class_sign (a_class)) return f7_clr (cc); diff --git a/libgcc/config/avr/libf7/libf7.h b/libgcc/config/avr/libf7/libf7.h index 8aa91c7..786e141 100644 --- a/libgcc/config/avr/libf7/libf7.h +++ b/libgcc/config/avr/libf7/libf7.h @@ -612,6 +612,7 @@ extern void f7_cos (f7_t*, const f7_t*); extern void f7_tan (f7_t*, const f7_t*); extern void f7_atan (f7_t*, const f7_t*); extern void f7_atan2 (f7_t*, const f7_t*, const f7_t*); +extern void f7_fdim (f7_t*, const f7_t*, const f7_t*); extern void f7_asin (f7_t*, const f7_t*); extern void f7_acos (f7_t*, const f7_t*); extern void f7_tanh (f7_t*, const f7_t*); diff --git a/libgfortran/io/format.c b/libgfortran/io/format.c index eef1d34..87e21a9 100644 --- a/libgfortran/io/format.c +++ b/libgfortran/io/format.c @@ -1235,9 +1235,9 @@ parse_format_list (st_parameter_dt *dtp, bool *seen_dd) default: /* Assume a missing comma with -std=legacy, GNU extension. */ - if (compile_options.warn_std == 0) - goto format_item_1; - format_error (dtp, tail, comma_missing); + if (compile_options.warn_std != 0) + fmt->error = comma_missing; + goto format_item_1; } /* Optional comma is a weird between state where we've just finished @@ -1252,7 +1252,7 @@ parse_format_list (st_parameter_dt *dtp, bool *seen_dd) case FMT_RPAREN: goto finished; - default: /* Assume that we have another format item */ + default: /* Assume that we have another format item */ fmt->saved_token = t; break; } @@ -1419,7 +1419,7 @@ parse_format (st_parameter_dt *dtp) else fmt->error = "Missing initial left parenthesis in format"; - if (format_cache_ok) + if (format_cache_ok && !fmt->error) save_parsed_format (dtp); else dtp->u.p.format_not_saved = 1; diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 6cef1dd..c9035c3 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,21 @@ +2025-05-28 Tobias Burnus <tburnus@baylibre.com> + + Backported from master: + 2025-05-28 Tobias Burnus <tburnus@baylibre.com> + + PR middle-end/118694 + * testsuite/libgomp.fortran/metadirective-1.f90: xfail when + compiling (also) for nvptx offloading as an error is then expected. + +2025-05-26 Tobias Burnus <tburnus@baylibre.com> + + Backported from master: + 2025-05-23 Tobias Burnus <tburnus@baylibre.com> + + PR middle-end/118694 + * testsuite/libgomp.c-c++-common/metadirective-1.c: xfail when + compiling (also) for nvptx offloading as an error is then expected. + 2025-05-19 Tobias Burnus <tburnus@baylibre.com> Backported from master: diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index afc8184..5ecab6d 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -122,6 +122,42 @@ * testsuite/libgomp.c-c++-common/target-abi-struct-1.c: New. * testsuite/libgomp.c-c++-common/target-abi-struct-1-O0.c: Adjust. +2025-05-30 Tobias Burnus <tburnus@baylibre.com> + + Backported from master: + 2025-05-29 Tobias Burnus <tburnus@baylibre.com> + + PR libgomp/93226 + * libgomp-plugin.h (GOMP_OFFLOAD_openacc_async_dev2dev): New + prototype. + * libgomp.h (struct acc_dispatch_t): Add dev2dev_func. + (gomp_copy_dev2dev): New prototype. + * libgomp.map (OACC_2.6.1): New; add acc_memcpy_device{,_async}. + * libgomp.texi (acc_memcpy_device): New. + * oacc-mem.c (memcpy_tofrom_device): Change to take from/to + device boolean; use memcpy not memmove; add early return if + size == 0 or same device + same ptr. + (acc_memcpy_to_device, acc_memcpy_to_device_async, + acc_memcpy_from_device, acc_memcpy_from_device_async): Update. + (acc_memcpy_device, acc_memcpy_device_async): New. + * openacc.f90 (acc_memcpy_device, acc_memcpy_device_async): + Add interface. + * openacc_lib.h (acc_memcpy_device, acc_memcpy_device_async): + Likewise. + * openacc.h (acc_memcpy_device, acc_memcpy_device_async): Add + prototype. + * plugin/plugin-gcn.c (GOMP_OFFLOAD_openacc_async_host2dev): + Update comment. + (GOMP_OFFLOAD_openacc_async_dev2host): Update call. + (GOMP_OFFLOAD_openacc_async_dev2dev): New. + * plugin/plugin-nvptx.c (cuda_memcpy_dev_sanity_check): New. + (GOMP_OFFLOAD_dev2dev): Call it. + (GOMP_OFFLOAD_openacc_async_dev2dev): New. + * target.c (gomp_copy_dev2dev): New. + (gomp_load_plugin_for_device): Load dev2dev and async_dev2dev. + * testsuite/libgomp.oacc-c-c++-common/acc_memcpy_device-1.c: New test. + * testsuite/libgomp.oacc-fortran/acc_memcpy_device-1.f90: New test. + 2025-05-22 Thomas Schwinge <tschwinge@baylibre.com> Backported from master: @@ -1014,4 +1050,4 @@ * testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h: Support - header for new tests.
\ No newline at end of file + header for new tests. diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h index 479264b..3c7741b 100644 --- a/libgomp/libgomp-plugin.h +++ b/libgomp/libgomp-plugin.h @@ -202,6 +202,8 @@ extern bool GOMP_OFFLOAD_openacc_async_dev2host (int, void *, const void *, size struct goacc_asyncqueue *); extern bool GOMP_OFFLOAD_openacc_async_host2dev (int, void *, const void *, size_t, struct goacc_asyncqueue *); +extern bool GOMP_OFFLOAD_openacc_async_dev2dev (int, void *, const void *, size_t, + struct goacc_asyncqueue *); extern void *GOMP_OFFLOAD_openacc_cuda_get_current_device (void); extern void *GOMP_OFFLOAD_openacc_cuda_get_current_context (void); extern void *GOMP_OFFLOAD_openacc_cuda_get_stream (struct goacc_asyncqueue *); diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 04f3c6d..571ac62c 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1378,6 +1378,7 @@ typedef struct acc_dispatch_t __typeof (GOMP_OFFLOAD_openacc_async_exec) *exec_func; __typeof (GOMP_OFFLOAD_openacc_async_dev2host) *dev2host_func; __typeof (GOMP_OFFLOAD_openacc_async_host2dev) *host2dev_func; + __typeof (GOMP_OFFLOAD_openacc_async_dev2dev) *dev2dev_func; } async; __typeof (GOMP_OFFLOAD_openacc_get_property) *get_property_func; @@ -1485,6 +1486,9 @@ extern void gomp_copy_host2dev (struct gomp_device_descr *, extern void gomp_copy_dev2host (struct gomp_device_descr *, struct goacc_asyncqueue *, void *, const void *, size_t); +extern void gomp_copy_dev2dev (struct gomp_device_descr *, + struct goacc_asyncqueue *, void *, const void *, + size_t); extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t); extern bool gomp_attach_pointer (struct gomp_device_descr *, struct goacc_asyncqueue *, splay_tree, diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index 6e2cdbf..bc2de6b 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -610,6 +610,12 @@ OACC_2.6 { acc_get_property_string_h_; } OACC_2.5.1; +OACC_2.6.1 { + global: + acc_memcpy_device; + acc_memcpy_device_async; +} OACC_2.6; + GOACC_2.0 { global: GOACC_data_end; diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 658df0e..e1b70b0 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -4780,6 +4780,7 @@ acceleration device. present on device. * acc_memcpy_to_device:: Copy host memory to device memory. * acc_memcpy_from_device:: Copy device memory to host memory. +* acc_memcpy_device:: Copy memory within a device. * acc_attach:: Let device pointer point to device-pointer target. * acc_detach:: Let device pointer point to host-pointer target. @@ -5854,6 +5855,44 @@ This function copies device memory specified by device address of +@node acc_memcpy_device +@section @code{acc_memcpy_device} -- Copy memory within a device. +@table @asis +@item @emph{Description} +This function copies device memory from one memory location to another +on the current device. It copies @var{bytes} bytes of data from the device +address, specified by @var{data_dev_src}, to the device address +@var{data_dev_dest}. The @code{_async} version performs the transfer +asnychronously using the queue associated with @var{async_arg}. + +@item @emph{C/C++}: +@multitable @columnfractions .20 .80 +@item @emph{Prototype}: @tab @code{void acc_memcpy_device(d_void* data_dev_dest,} +@item @tab @code{d_void* data_dev_src, size_t bytes);} +@item @emph{Prototype}: @tab @code{void acc_memcpy_device_async(d_void* data_dev_dest,} +@item @tab @code{d_void* data_dev_src, size_t bytes, int async_arg);} +@end multitable + +@item @emph{Fortran}: +@multitable @columnfractions .20 .80 +@item @emph{Interface}: @tab @code{subroutine acc_memcpy_device(data_dev_dest, &} +@item @tab @code{data_dev_src, bytes)} +@item @emph{Interface}: @tab @code{subroutine acc_memcpy_device_async(data_dev_dest, &} +@item @tab @code{data_dev_src, bytes, async_arg)} +@item @tab @code{type(c_ptr), value :: data_dev_dest} +@item @tab @code{type(c_ptr), value :: data_dev_src} +@item @tab @code{integer(c_size_t), value :: bytes} +@item @tab @code{integer(acc_handle_kind), value :: async_arg} +@end multitable + +@item @emph{Reference}: +@uref{https://www.openacc.org, OpenACC specification v2.6}, section +3.2.33. @uref{https://www.openacc.org, OpenACC specification v3.3}, section +3.2.28. +@end table + + + @node acc_attach @section @code{acc_attach} -- Let device pointer point to device-pointer target. @table @asis diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index a00ea16..e40b41b 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -171,21 +171,22 @@ acc_free (void *d) } static void -memcpy_tofrom_device (bool from, void *d, void *h, size_t s, int async, - const char *libfnname) +memcpy_tofrom_device (bool dev_to, bool dev_from, void *dst, void *src, + size_t s, int async, const char *libfnname) { /* No need to call lazy open here, as the device pointer must have been obtained from a routine that did that. */ struct goacc_thread *thr = goacc_thread (); assert (thr && thr->dev); + if (s == 0) + return; if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) { - if (from) - memmove (h, d, s); - else - memmove (d, h, s); + if (src == dst) + return; + memcpy (dst, src, s); return; } @@ -199,10 +200,15 @@ memcpy_tofrom_device (bool from, void *d, void *h, size_t s, int async, } goacc_aq aq = get_goacc_asyncqueue (async); - if (from) - gomp_copy_dev2host (thr->dev, aq, h, d, s); + if (dev_to && dev_from) + { + if (dst != src) + gomp_copy_dev2dev (thr->dev, aq, dst, src, s); + } + else if (dev_from) + gomp_copy_dev2host (thr->dev, aq, dst, src, s); else - gomp_copy_host2dev (thr->dev, aq, d, h, s, false, /* TODO: cbuf? */ NULL); + gomp_copy_host2dev (thr->dev, aq, dst, src, s, false, /* TODO: cbuf? */ NULL); if (profiling_p) { @@ -214,25 +220,37 @@ memcpy_tofrom_device (bool from, void *d, void *h, size_t s, int async, void acc_memcpy_to_device (void *d, void *h, size_t s) { - memcpy_tofrom_device (false, d, h, s, acc_async_sync, __FUNCTION__); + memcpy_tofrom_device (true, false, d, h, s, acc_async_sync, __FUNCTION__); } void acc_memcpy_to_device_async (void *d, void *h, size_t s, int async) { - memcpy_tofrom_device (false, d, h, s, async, __FUNCTION__); + memcpy_tofrom_device (true, false, d, h, s, async, __FUNCTION__); } void acc_memcpy_from_device (void *h, void *d, size_t s) { - memcpy_tofrom_device (true, d, h, s, acc_async_sync, __FUNCTION__); + memcpy_tofrom_device (false, true, h, d, s, acc_async_sync, __FUNCTION__); } void acc_memcpy_from_device_async (void *h, void *d, size_t s, int async) { - memcpy_tofrom_device (true, d, h, s, async, __FUNCTION__); + memcpy_tofrom_device (false, true, h, d, s, async, __FUNCTION__); +} + +void +acc_memcpy_device (void *dst, void *src, size_t s) +{ + memcpy_tofrom_device (true, true, dst, src, s, acc_async_sync, __FUNCTION__); +} + +void +acc_memcpy_device_async (void *dst, void *src, size_t s, int async) +{ + memcpy_tofrom_device (true, true, dst, src, s, async, __FUNCTION__); } /* Return the device pointer that corresponds to host data H. Or NULL diff --git a/libgomp/openacc.f90 b/libgomp/openacc.f90 index a3d7bcb..55894df 100644 --- a/libgomp/openacc.f90 +++ b/libgomp/openacc.f90 @@ -797,6 +797,7 @@ module openacc public :: acc_copyout_finalize, acc_delete_finalize public :: acc_memcpy_to_device, acc_memcpy_to_device_async public :: acc_memcpy_from_device, acc_memcpy_from_device_async + public :: acc_memcpy_device, acc_memcpy_device_async integer, parameter :: openacc_version = 201811 @@ -1046,6 +1047,27 @@ module openacc end subroutine end interface + interface + subroutine acc_memcpy_device (data_dev_dest, data_dev_src, bytes) bind(C) + use iso_c_binding, only: c_ptr, c_size_t + type(c_ptr), value :: data_dev_dest + type(c_ptr), value :: data_dev_src + integer(c_size_t), value :: bytes + end subroutine + end interface + + interface + subroutine acc_memcpy_device_async (data_dev_dest, data_dev_src, & + bytes, async_arg) bind(C) + use iso_c_binding, only: c_ptr, c_size_t + import :: acc_handle_kind + type(c_ptr), value :: data_dev_dest + type(c_ptr), value :: data_dev_src + integer(c_size_t), value :: bytes + integer(acc_handle_kind), value :: async_arg + end subroutine + end interface + interface acc_copyin_async procedure :: acc_copyin_async_32_h procedure :: acc_copyin_async_64_h diff --git a/libgomp/openacc.h b/libgomp/openacc.h index a520bbe..3085b00 100644 --- a/libgomp/openacc.h +++ b/libgomp/openacc.h @@ -123,6 +123,7 @@ void *acc_hostptr (void *) __GOACC_NOTHROW; int acc_is_present (void *, size_t) __GOACC_NOTHROW; void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW; void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW; +void acc_memcpy_device (void *, void *, size_t) __GOACC_NOTHROW; void acc_attach (void **) __GOACC_NOTHROW; void acc_attach_async (void **, int) __GOACC_NOTHROW; void acc_detach (void **) __GOACC_NOTHROW; @@ -136,7 +137,7 @@ void acc_delete_finalize_async (void *, size_t, int) __GOACC_NOTHROW; void acc_detach_finalize (void **) __GOACC_NOTHROW; void acc_detach_finalize_async (void **, int) __GOACC_NOTHROW; -/* Async functions, specified in OpenACC 2.5. */ +/* Async functions, specified in OpenACC 2.5, acc_memcpy_device in 2.6. */ void acc_copyin_async (void *, size_t, int) __GOACC_NOTHROW; void acc_create_async (void *, size_t, int) __GOACC_NOTHROW; void acc_copyout_async (void *, size_t, int) __GOACC_NOTHROW; @@ -145,6 +146,7 @@ void acc_update_device_async (void *, size_t, int) __GOACC_NOTHROW; void acc_update_self_async (void *, size_t, int) __GOACC_NOTHROW; void acc_memcpy_to_device_async (void *, void *, size_t, int) __GOACC_NOTHROW; void acc_memcpy_from_device_async (void *, void *, size_t, int) __GOACC_NOTHROW; +void acc_memcpy_device_async (void *, void *, size_t, int) __GOACC_NOTHROW; /* CUDA-specific routines. */ void *acc_get_current_cuda_device (void) __GOACC_NOTHROW; diff --git a/libgomp/openacc_lib.h b/libgomp/openacc_lib.h index d830574..e0e7788 100644 --- a/libgomp/openacc_lib.h +++ b/libgomp/openacc_lib.h @@ -528,6 +528,30 @@ end subroutine end interface + interface + subroutine acc_memcpy_device(data_dev_dest, data_dev_src, & + & bytes) bind(C) + use iso_c_binding, only: c_ptr, c_size_t + type(c_ptr), value :: data_dev_dest + type(c_ptr), value :: data_dev_src + integer(c_size_t), value :: bytes + end subroutine + end interface + + interface + subroutine acc_memcpy_device_async(data_dev_dest, & + & data_dev_src, bytes, & + & async_arg) bind(C) + use iso_c_binding, only: c_ptr, c_size_t + import :: acc_handle_kind + type(c_ptr), value :: data_dev_dest + type(c_ptr), value :: data_dev_src + integer(c_size_t), value :: bytes + integer(acc_handle_kind), value :: async_arg + end subroutine + end interface + + interface acc_copyin_async subroutine acc_copyin_async_32_h (a, len, async) use iso_c_binding, only: c_int32_t diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c index b39a94b..f823b27 100644 --- a/libgomp/plugin/plugin-gcn.c +++ b/libgomp/plugin/plugin-gcn.c @@ -5081,7 +5081,8 @@ GOMP_OFFLOAD_openacc_async_queue_callback (struct goacc_asyncqueue *aq, queue_push_callback (aq, fn, data); } -/* Queue up an asynchronous data copy from host to DEVICE. */ +/* Queue up an asynchronous data copy from host to DEVICE. + (Also handles dev2host and dev2dev.) */ bool GOMP_OFFLOAD_openacc_async_host2dev (int device, void *dst, const void *src, @@ -5099,10 +5100,16 @@ bool GOMP_OFFLOAD_openacc_async_dev2host (int device, void *dst, const void *src, size_t n, struct goacc_asyncqueue *aq) { - struct agent_info *agent = get_agent_info (device); - assert (agent == aq->agent); - queue_push_copy (aq, dst, src, n); - return true; + return GOMP_OFFLOAD_openacc_async_host2dev (device, dst, src, n, aq); +} + +/* Queue up an asynchronous data copy from DEVICE to DEVICE. */ + +bool +GOMP_OFFLOAD_openacc_async_dev2dev (int device, void *dst, const void *src, + size_t n, struct goacc_asyncqueue *aq) +{ + return GOMP_OFFLOAD_openacc_async_host2dev (device, dst, src, n, aq); } union goacc_property_value diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c index a6c8198..712c8b7 100644 --- a/libgomp/plugin/plugin-nvptx.c +++ b/libgomp/plugin/plugin-nvptx.c @@ -2060,6 +2060,34 @@ GOMP_OFFLOAD_openacc_async_queue_callback (struct goacc_asyncqueue *aq, } static bool +cuda_memcpy_dev_sanity_check (const void *d1, const void *d2, size_t s) +{ + CUdeviceptr pb1, pb2; + size_t ps1, ps2; + if (!s) + return true; + if (!d1 || !d2) + { + GOMP_PLUGIN_error ("invalid device address"); + return false; + } + CUDA_CALL (cuMemGetAddressRange, &pb1, &ps1, (CUdeviceptr) d1); + CUDA_CALL (cuMemGetAddressRange, &pb2, &ps2, (CUdeviceptr) d2); + if (!pb1 || !pb2) + { + GOMP_PLUGIN_error ("invalid device address"); + return false; + } + if ((void *)(d1 + s) > (void *)(pb1 + ps1) + || (void *)(d2 + s) > (void *)(pb2 + ps2)) + { + GOMP_PLUGIN_error ("invalid size"); + return false; + } + return true; +} + +static bool cuda_memcpy_sanity_check (const void *h, const void *d, size_t s) { CUdeviceptr pb; @@ -2118,6 +2146,9 @@ GOMP_OFFLOAD_dev2host (int ord, void *dst, const void *src, size_t n) bool GOMP_OFFLOAD_dev2dev (int ord, void *dst, const void *src, size_t n) { + if (!nvptx_attach_host_thread_to_device (ord) + || !cuda_memcpy_dev_sanity_check (dst, src, n)) + return false; CUDA_CALL (cuMemcpyDtoDAsync, (CUdeviceptr) dst, (CUdeviceptr) src, n, NULL); return true; } @@ -2329,6 +2360,18 @@ GOMP_OFFLOAD_openacc_async_dev2host (int ord, void *dst, const void *src, return true; } +bool +GOMP_OFFLOAD_openacc_async_dev2dev (int ord, void *dst, const void *src, + size_t n, struct goacc_asyncqueue *aq) +{ + if (!nvptx_attach_host_thread_to_device (ord) + || !cuda_memcpy_dev_sanity_check (dst, src, n)) + return false; + CUDA_CALL (cuMemcpyDtoDAsync, (CUdeviceptr) dst, (CUdeviceptr) src, n, + aq->cuda_stream); + return true; +} + union goacc_property_value GOMP_OFFLOAD_openacc_get_property (int n, enum goacc_property prop) { diff --git a/libgomp/target.c b/libgomp/target.c index 01434f8..4ad803a 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -461,6 +461,19 @@ gomp_copy_dev2host (struct gomp_device_descr *devicep, gomp_device_copy (devicep, devicep->dev2host_func, "host", h, "dev", d, sz); } +attribute_hidden void +gomp_copy_dev2dev (struct gomp_device_descr *devicep, + struct goacc_asyncqueue *aq, + void *dst, const void *src, size_t sz) +{ + if (__builtin_expect (aq != NULL, 0)) + goacc_device_copy_async (devicep, devicep->openacc.async.dev2dev_func, + "dev", dst, "dev", src, NULL, sz, aq); + else + gomp_device_copy (devicep, devicep->dev2dev_func, "dev", dst, + "dev", src, sz); +} + static void gomp_free_device_memory (struct gomp_device_descr *devicep, void *devptr) { @@ -6312,6 +6325,7 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device, || !DLSYM_OPT (openacc.async.exec, openacc_async_exec) || !DLSYM_OPT (openacc.async.dev2host, openacc_async_dev2host) || !DLSYM_OPT (openacc.async.host2dev, openacc_async_host2dev) + || !DLSYM_OPT (openacc.async.dev2dev, openacc_async_dev2dev) || !DLSYM_OPT (openacc.get_property, openacc_get_property)) { /* Require all the OpenACC handlers if we have diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 index 7b3e09f..d6f4d5b 100644 --- a/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 +++ b/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 @@ -1,4 +1,5 @@ -! { dg-do run } +! { dg-do run { target { ! offload_target_nvptx } } } +! { dg-do compile { target offload_target_nvptx } } program test implicit none @@ -33,6 +34,10 @@ program test contains subroutine f (x, y, z) integer :: x(N), y(N), z(N) + ! The following fails as on the host the target side cannot be + ! resolved - and the 'teams' or not status affects how 'target' + ! is called. -> See PR118694, esp. comment 9. + ! Note also the dg-do compile above for offload_target_nvptx !$omp target map (to: x, y) map(from: z) block @@ -43,6 +48,7 @@ contains z(i) = x(i) * y(i) enddo end block + ! { dg-bogus "'target' construct with nested 'teams' construct contains directives outside of the 'teams' construct" "PR118694" { xfail offload_target_nvptx } .-9 } */ end subroutine subroutine g (x, y, z) integer :: x(N), y(N), z(N) @@ -56,6 +62,7 @@ contains z(i) = x(i) * y(i) enddo end block + ! { dg-bogus "'target' construct with nested 'teams' construct contains directives outside of the 'teams' construct" "PR118694" { xfail offload_target_nvptx } .-9 } */ !$omp end target end subroutine end program diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_memcpy_device-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_memcpy_device-1.c new file mode 100644 index 0000000..eda651d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_memcpy_device-1.c @@ -0,0 +1,96 @@ +/* { dg-prune-output "using .vector_length \\(32\\)" } */ + +/* PR libgomp/93226 */ + +#include <stdlib.h> +#include <stdint.h> +#include <string.h> +#include <openacc.h> + +enum { N = 1024 }; + +static int D[N]; +#pragma acc declare device_resident(D) + +#pragma acc routine +intptr_t init_d() +{ + for (int i = 0; i < N; i++) + D[i] = 27*i; + return (intptr_t) &D[0]; +} + +int +main () +{ + int *a, *b, *e; + void *d_a, *d_b, *d_c, *d_d, *d_e, *d_f; + intptr_t intptr; + bool fail = false; + + a = (int *) malloc (N*sizeof (int)); + b = (int *) malloc (N*sizeof (int)); + e = (int *) malloc (N*sizeof (int)); + d_c = acc_malloc (N*sizeof (int)); + d_f = acc_malloc (N*sizeof (int)); + + memset (e, 0xff, N*sizeof (int)); + d_e = acc_copyin (e, N*sizeof (int)); + + #pragma acc serial copyout(intptr) + intptr = init_d (); + d_d = (void*) intptr; + acc_memcpy_device (d_c, d_d, N*sizeof (int)); + + #pragma acc serial copy(fail) deviceptr(d_c) firstprivate(intptr) + { + int *cc = (int *) d_c; + int *dd = (int *) intptr; + for (int i = 0; i < N; i++) + if (dd[i] != 27*i || cc[i] != 27*i) + { + fail = true; + __builtin_abort (); + } + } + if (fail) __builtin_abort (); + + for (int i = 0; i < N; i++) + a[i] = 11*i; + for (int i = 0; i < N; i++) + b[i] = 31*i; + + d_a = acc_copyin (a, N*sizeof (int)); + acc_copyin_async (b, N*sizeof (int), acc_async_noval); + + #pragma acc parallel deviceptr(d_c) async + { + int *cc = (int *) d_c; + #pragma acc loop + for (int i = 0; i < N; i++) + cc[i] = -17*i; + } + + acc_memcpy_device_async (d_d, d_a, N*sizeof (int), acc_async_noval); + acc_memcpy_device_async (d_f, d_c, N*sizeof (int), acc_async_noval); + acc_wait (acc_async_noval); + d_b = acc_deviceptr (b); + acc_memcpy_device_async (d_e, d_b, N*sizeof (int), acc_async_noval); + acc_wait (acc_async_noval); + + #pragma acc serial deviceptr(d_d, d_e, d_f) copy(fail) + { + int *dd = (int *) d_d; + int *ee = (int *) d_e; + int *ff = (int *) d_f; + for (int i = 0; i < N; i++) + if (dd[i] != 11*i + || ee[i] != 31*i + || ff[i] != -17*i) + { + fail = true; + __builtin_abort (); + } + } + if (fail) __builtin_abort (); +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_memcpy_device-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/acc_memcpy_device-1.f90 new file mode 100644 index 0000000..8f3a8f0 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_memcpy_device-1.f90 @@ -0,0 +1,113 @@ +! { dg-prune-output "using .vector_length \\(32\\)" } + +! PR libgomp/93226 */ + +module m + use iso_c_binding + use openacc + implicit none (external, type) + + integer, parameter :: N = 1024 + + integer :: D(N) + !$acc declare device_resident(D) + +contains + + integer(c_intptr_t) function init_d() + !$acc routine + integer :: i + do i = 1, N + D(i) = 27*i + end do + init_d = loc(D) + end +end module + +program main + use m + implicit none (external, type) + + integer, allocatable, target :: a(:), b(:), e(:) + type(c_ptr) :: d_a, d_b, d_c, d_d, d_e, d_f + integer(c_intptr_t) intptr + integer :: i + logical fail + + fail = .false. + + allocate(a(N), b(N), e(N)) + d_c = acc_malloc (N*c_sizeof (i)) + d_f = acc_malloc (N*c_sizeof (i)) + + e = huge(e) + call acc_copyin (e, N*c_sizeof (i)); + d_e = acc_deviceptr (e); + + !$acc serial copyout(intptr) + intptr = init_d () + !$acc end serial + d_d = transfer(intptr, d_d) + call acc_memcpy_device (d_c, d_d, N*c_sizeof (i)) + + !$acc serial copy(fail) copy(a) deviceptr(d_c, d_d) firstprivate(intptr) + block + integer, pointer :: cc(:), dd(:) + call c_f_pointer (d_c, cc, [N]) + call c_f_pointer (d_d, dd, [N]) + a = cc + do i = 1, N + if (dd(i) /= 27*i .or. cc(i) /= 27*i) then + fail = .true. + stop 1 + end if + end do + end block + !$acc end serial + if (fail) error stop 1 + + do i = 1, N + a(i) = 11*i + b(i) = 31*i + end do + + call acc_copyin (a, N*c_sizeof (i)) + d_a = acc_deviceptr (a) + call acc_copyin_async (b, N*c_sizeof (i), acc_async_noval) + + !$acc parallel deviceptr(d_c) private(i) async + block + integer, pointer :: cc(:) + call c_f_pointer (d_c, cc, [N]) + !$acc loop + do i = 1, N + cc(i) = -17*i + end do + end block + !$acc end parallel + + call acc_memcpy_device_async (d_d, d_a, N*c_sizeof (i), acc_async_noval) + call acc_memcpy_device_async (d_f, d_c, N*c_sizeof (i), acc_async_noval) + call acc_wait (acc_async_noval) + d_b = acc_deviceptr (b) + call acc_memcpy_device_async (d_e, d_b, N*c_sizeof (i), acc_async_noval) + call acc_wait (acc_async_noval) + + !$acc serial deviceptr(d_d, d_e, d_f) private(i) copy(fail) + block + integer, pointer :: dd(:), ee(:), ff(:) + call c_f_pointer (d_d, dd, [N]) + call c_f_pointer (d_e, ee, [N]) + call c_f_pointer (d_f, ff, [N]) + do i = 1, N + if (dd(i) /= 11*i & + .or. ee(i) /= 31*i & + .or. ff(i) /= -17*i) then + fail = .true. + stop 2 + end if + end do + end block + !$acc end serial + if (fail) error stop 2 +end diff --git a/libstdc++-v3/ChangeLog b/libstdc++-v3/ChangeLog index d7cd8f4..4dfe958 100644 --- a/libstdc++-v3/ChangeLog +++ b/libstdc++-v3/ChangeLog @@ -1,3 +1,47 @@ +2025-06-01 Patrick Palka <ppalka@redhat.com> + + Backported from master: + 2025-05-29 Patrick Palka <ppalka@redhat.com> + + * include/std/flat_map (_Flat_map_impl::operator==): Compare + keys and values separately. + +2025-06-01 Patrick Palka <ppalka@redhat.com> + + Backported from master: + 2025-05-29 Patrick Palka <ppalka@redhat.com> + Jonathan Wakely <jwakely@redhat.com> + + PR libstdc++/120465 + * include/std/flat_map (_Flat_map_impl::_M_erase_if): Use a + projection with ranges::remove_if to pass a pair instead of + a tuple to the predicate. + * testsuite/23_containers/flat_map/1.cc (test07): Strengthen + to expect the argument passed to the predicate is a pair. + * testsuite/23_containers/flat_multimap/1.cc (test07): Likewise. + +2025-05-27 Jonathan Wakely <jwakely@redhat.com> + + Backported from master: + 2025-05-22 Jonathan Wakely <jwakely@redhat.com> + + * testsuite/util/pstl/test_utils.h (ForwardIterator::operator++): + Fix return type. + (BidirectionalIterator::operator++): Likewise. + (BidirectionalIterator::operator--): Likewise. + +2025-05-27 Jonathan Wakely <jwakely@redhat.com> + + Backported from master: + 2025-05-22 Jonathan Wakely <jwakely@redhat.com> + + PR libstdc++/120367 + * include/bits/stl_vector.h (_M_range_initialize): Initialize + _M_impl._M_finish. + * testsuite/23_containers/vector/cons/from_range.cc: Check with + a type that throws on construction. + exceptions during construction. + 2025-05-20 Jonathan Wakely <jwakely@redhat.com> Backported from master: diff --git a/libstdc++-v3/include/bits/stl_vector.h b/libstdc++-v3/include/bits/stl_vector.h index aff9d5d..4861edb 100644 --- a/libstdc++-v3/include/bits/stl_vector.h +++ b/libstdc++-v3/include/bits/stl_vector.h @@ -1981,8 +1981,9 @@ _GLIBCXX_BEGIN_NAMESPACE_CONTAINER _M_range_initialize_n(_Iterator __first, _Sentinel __last, size_type __n) { - pointer __start = this->_M_impl._M_start = + pointer __start = this->_M_allocate(_S_check_init_len(__n, _M_get_Tp_allocator())); + this->_M_impl._M_start = this->_M_impl._M_finish = __start; this->_M_impl._M_end_of_storage = __start + __n; this->_M_impl._M_finish = std::__uninitialized_copy_a(_GLIBCXX_MOVE(__first), __last, diff --git a/libstdc++-v3/include/std/flat_map b/libstdc++-v3/include/std/flat_map index 6593988..4bd4963 100644 --- a/libstdc++-v3/include/std/flat_map +++ b/libstdc++-v3/include/std/flat_map @@ -873,7 +873,10 @@ _GLIBCXX_BEGIN_NAMESPACE_VERSION [[nodiscard]] friend bool operator==(const _Derived& __x, const _Derived& __y) - { return std::equal(__x.begin(), __x.end(), __y.begin(), __y.end()); } + { + return __x._M_cont.keys == __y._M_cont.keys + && __x._M_cont.values == __y._M_cont.values; + } template<typename _Up = value_type> [[nodiscard]] @@ -895,7 +898,10 @@ _GLIBCXX_BEGIN_NAMESPACE_VERSION { auto __guard = _M_make_clear_guard(); auto __zv = views::zip(_M_cont.keys, _M_cont.values); - auto __sr = ranges::remove_if(__zv, __pred); + auto __sr = ranges::remove_if(__zv, __pred, + [](const auto& __e) { + return const_reference(__e); + }); auto __erased = __sr.size(); erase(end() - __erased, end()); __guard._M_disable(); diff --git a/libstdc++-v3/testsuite/23_containers/flat_map/1.cc b/libstdc++-v3/testsuite/23_containers/flat_map/1.cc index 38fd449..9d99796 100644 --- a/libstdc++-v3/testsuite/23_containers/flat_map/1.cc +++ b/libstdc++-v3/testsuite/23_containers/flat_map/1.cc @@ -245,8 +245,9 @@ void test07() { // PR libstdc++/119427 - std::erase_if(std::flat_foo) does not work + // PR libstdc++/120465 - erase_if for flat_map calls predicate with incorrect type std::flat_map<int, int> m = {std::pair{1, 2}, {3, 4}, {5, 6}}; - auto n = std::erase_if(m, [](auto x) { auto [k,v] = x; return k == 1 || v == 6; }); + auto n = std::erase_if(m, [](auto x) { return x.first == 1 || x.second == 6; }); VERIFY( n == 2 ); VERIFY( std::ranges::equal(m, (std::pair<int,int>[]){{3,4}}) ); } diff --git a/libstdc++-v3/testsuite/23_containers/flat_multimap/1.cc b/libstdc++-v3/testsuite/23_containers/flat_multimap/1.cc index 79fbc1a..0010955 100644 --- a/libstdc++-v3/testsuite/23_containers/flat_multimap/1.cc +++ b/libstdc++-v3/testsuite/23_containers/flat_multimap/1.cc @@ -223,8 +223,9 @@ void test07() { // PR libstdc++/119427 - std::erase_if(std::flat_foo) does not work + // PR libstdc++/120465 - erase_if for flat_map calls predicate with incorrect type std::flat_multimap<int, int> m = {std::pair{1, 2}, {3, 4}, {3, 3}, {5, 6}, {6, 6}}; - auto n = std::erase_if(m, [](auto x) { auto [k,v] = x; return k == 1 || v == 6; }); + auto n = std::erase_if(m, [](auto x) { return x.first == 1 || x.second == 6; }); VERIFY( n == 3 ); VERIFY( std::ranges::equal(m, (std::pair<int,int>[]){{3,4},{3,3}}) ); } diff --git a/libstdc++-v3/testsuite/23_containers/vector/cons/from_range.cc b/libstdc++-v3/testsuite/23_containers/vector/cons/from_range.cc index 7a62645..3784b9c 100644 --- a/libstdc++-v3/testsuite/23_containers/vector/cons/from_range.cc +++ b/libstdc++-v3/testsuite/23_containers/vector/cons/from_range.cc @@ -106,8 +106,30 @@ test_constexpr() return true; } +void +test_pr120367() +{ +#ifdef __cpp_exceptions + struct X + { + X(int) { throw 1; } // Cannot successfully construct an X. + ~X() { VERIFY(false); } // So should never need to destroy one. + }; + + try + { + int i[1]{}; + std::vector<X> v(std::from_range, i); + } + catch (int) + { + } +#endif +} + int main() { test_ranges(); static_assert( test_constexpr() ); + test_pr120367(); } diff --git a/libstdc++-v3/testsuite/util/pstl/test_utils.h b/libstdc++-v3/testsuite/util/pstl/test_utils.h index 55b5100..9c61a714 100644 --- a/libstdc++-v3/testsuite/util/pstl/test_utils.h +++ b/libstdc++-v3/testsuite/util/pstl/test_utils.h @@ -154,7 +154,7 @@ class ForwardIterator explicit ForwardIterator(Iterator i) : my_iterator(i) {} reference operator*() const { return *my_iterator; } Iterator operator->() const { return my_iterator; } - ForwardIterator + ForwardIterator& operator++() { ++my_iterator; @@ -194,13 +194,13 @@ class BidirectionalIterator : public ForwardIterator<Iterator, IteratorTag> explicit BidirectionalIterator(Iterator i) : base_type(i) {} BidirectionalIterator(const base_type& i) : base_type(i.iterator()) {} - BidirectionalIterator + BidirectionalIterator& operator++() { ++base_type::my_iterator; return *this; } - BidirectionalIterator + BidirectionalIterator& operator--() { --base_type::my_iterator; |