From patchwork Sat Jan 6 18:52:55 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Sandra Loosemore X-Patchwork-Id: 83462 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 32150385782D for ; Sat, 6 Jan 2024 18:55:35 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 23F843858C60 for ; Sat, 6 Jan 2024 18:54:41 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 23F843858C60 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 23F843858C60 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=68.232.129.153 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1704567291; cv=none; b=mFM0nczANesmXXDNzgncWI456ZuGe7MVIFm831vrd2vQCRmqqOYQHD2cg1QBqUwifrSGz0tDre8WW3uOjs+phZPdbhQtLlMIiAsNHFh7+bkNBfKmXGp8BLrReKBF0Gsi/VI/dBqddaqH8lks611M5hOizOWRPXbPCE5chqq1MDY= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1704567291; c=relaxed/simple; bh=rl7uDk+7kCz8flTCdtZdlQuMCFOPC+k5uN7BJRzf3lM=; h=From:To:Subject:Date:Message-ID:MIME-Version; b=AYjaP5dKgHQs4C30PXIaTCGk3XNaAg3gr3MR6AE5C/WBhV3Rlw/WeKaLxInWCvQv4y1JJMrf7pJL7JRaH3M6DPOIB5iXCjo7Dm3veuDr5bKA2CduK7GxOGn+dkefEn6RkfU63gZWTD53VMiMvtEH7RcKjAl5AH9/wwevnOpmlE0= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: A0ufJUGjQOqi70e/GHZ7eg== X-CSE-MsgGUID: jSlgr5CnS4KiKioU/GBgsg== X-IronPort-AV: E=Sophos;i="6.04,337,1695715200"; d="scan'208";a="30827357" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 06 Jan 2024 10:54:40 -0800 IronPort-SDR: Wc3caJnwMRanPCdrHUrHtCH4N/T9OiNH40t53A4d5E01hRox0/C5unhKWFJm0yisdFKp07SF/W 62iKlFjKcRXNmI8Z/awsnFCFUaEHo6zlepi3Zx2QFg7gtd5B4kF/EKc6bOGIykx6V3el6v+Fh2 TYqTHeBC9V7Q8L473x5qLA2TYtIqiQROt76RJlcl60v6fAtWfr95BfxzghHJBVLdpFA0hsRt0d jFDCNI2L1DukSfMt+5uVfZLBsMX/FrLIGAWglsCY60jUyxA/u157zys1Ng4B1Ng5pEwluwowUf U0o= From: Sandra Loosemore To: CC: , , Subject: [PATCH 7/8] OpenMP: Fortran front-end support for metadirectives. Date: Sat, 6 Jan 2024 11:52:55 -0700 Message-ID: <20240106185257.126445-8-sandra@codesourcery.com> X-Mailer: git-send-email 2.34.1 In-Reply-To: <20240106185257.126445-1-sandra@codesourcery.com> References: <20240106185257.126445-1-sandra@codesourcery.com> MIME-Version: 1.0 X-ClientProxiedBy: svr-orw-mbx-10.mgc.mentorg.com (147.34.90.210) To svr-orw-mbx-13.mgc.mentorg.com (147.34.90.213) X-Spam-Status: No, score=-10.1 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org From: Kwok Cheung Yeung This patch adds support for metadirectives to the Fortran front end. gcc/fortran/ChangeLog * decl.cc (gfc_match_end): Handle metadirectives. * dump-parse-tree.cc (show_omp_node): Likewise. (show_code_node): Likewise. * gfortran.h (enum gfc_statement): Add ST_OMP_METADIRECTIVE. (struct gfc_omp_clauses): Rename target_first_st_is_teams field to target_first_st_is_teams_or_meta. (struct gfc_omp_metadirective_variant): New. (struct gfc_st_label): Add omp_region field. (gfc_exec_op): Add EXEC_OMP_METADIRECTIVE. (struct gfc_code): Add omp_metadirective_variants field. (gfc_free_omp_metadirective_variants): Declare. (match_omp_directive): Declare. (is_omp_declarative_stmt): Declare. (gfc_skip_omp_metadirective_variant): Declare. * io.cc (format_asterisk): Add initializer for new omp_region field. * match.h (gfc_match_omp_begin_metadirective): Declare. (gfc_match_omp_metadirective): Declare. * openmp.cc (gfc_match_omp_eos): Special case for matching an OpenMP context selector. (gfc_free_omp_metadirective_variants): New. (gfc_match_omp_clauses): Remove context_selector parameter, use global state variable gfc_matching_omp_context_selector instead. (match_omp): Adjust call to gfc_match_omp_clauses, set gfc_matching_omp_context_selector instead. (gfc_match_omp_context_selector): Likewise. (gfc_match_omp_context_selector_specification): Generalize to take a set selector list pointer as parameter, instead of a declare variant pointer. (gfc_match_omp_declare_variant): Adjust call to match above change. (match_omp_metadirective): New. (gfc_match_omp_begin_metadirective): New. (gfc_match_omp_metadirective): New. (resolve_omp_metadirective): New. (resolve_omp_target): Handle metadirectives. (gfc_resolve_omp_directive): Handle metadirectives. * parse.cc (gfc_matching_omp_context_selector): New. (gfc_in_metadirective_body): New. (gfc_omp_region_count): New. (decode_omp_directive): Handle "begin metadirective", "end metadirective", and "metadirective". (match_omp_directive): New. (case_omp_structured_block): New define. (case_omp_do): New define. (gfc_ascii_statement): Handle ST_OMP_BEGIN_METADIRECTIVE, ST_OMP_END_METADIRECTIVE, and ST_OMP_METADIRECTIVE. (accept_statement): Handle ST_OMP_BEGIN_METADIRECTIVE and ST_OMP_METADIRECTIVE. (gfc_omp_end_stmt): New. (parse_omp_do): Use gfc_omp_end_stmt. Special-case "omp end metadirective" to end the current construct. (parse_omp_structured_block): Likewise. Adjust setting of target_first_st_is_teams_or_meta flag. (parse_omp_metadirective_body): New. (parse_executable): Handle metadirectives. Use case_omp_structured_block and case_omp_do here. (gfc_parse_file): Initialize gfc_omp_region_count, gfc_in_metadirective_body, and gfc_matching_omp_context_selector. (is_omp_declarative_stmt): New. * parse.h (gfc_omp_end_stmt): Declare. (gfc_matching_omp_context_selector): Declare. (gfc_in_metadirective_body): Declare. (gfc_omp_region_count): Declare. * resolve.cc (gfc_resolve_code): Handle EXEC_OMP_METADIRECTIVE. * st.cc (gfc_free_statement): Handle EXEC_OMP_METADIRECTIVE. * symbol.cc (compare_st_labels): Compare omp_region, not just the value. (gfc_get_st_label): Likewise. Initialize the omp_region field when creating a new label. * trans-decl.cc (gfc_get_label_decl): Encode the omp_region in the label name. * trans-openmp.cc (gfc_trans_omp_directive): Handle EXEC_OMP_METADIRECTIVE. (gfc_trans_omp_set_selector): New, split from... (gfc_trans_omp_declare_variant): ...here. (gfc_trans_omp_metadirective): New. (gfc_skip_omp_metadirective_variant): New. * trans-stmt.h (gfc_trans_omp_metadirective): Declare. * trans.cc (trans_code): Handle EXEC_OMP_METADIRECTIVE. gcc/testsuite/ChangeLog * gfortran.dg/gomp/metadirective-1.f90: New. * gfortran.dg/gomp/metadirective-10.f90: New. * gfortran.dg/gomp/metadirective-11.f90: New. * gfortran.dg/gomp/metadirective-2.f90: New. * gfortran.dg/gomp/metadirective-3.f90: New. * gfortran.dg/gomp/metadirective-4.f90: New. * gfortran.dg/gomp/metadirective-5.f90: New. * gfortran.dg/gomp/metadirective-6.f90: New. * gfortran.dg/gomp/metadirective-7.f90: New. * gfortran.dg/gomp/metadirective-8.f90: New. * gfortran.dg/gomp/metadirective-9.f90: New. * gfortran.dg/gomp/metadirective-construct.f90: New. * gfortran.dg/gomp/metadirective-no-score.f90: New. * gfortran.dg/gomp/pure-1.f90: Add metadirective test. * gfortran.dg/gomp/pure-2.f90: Remove metadirective test. libgomp/ChangeLog * testsuite/libgomp.fortran/metadirective-1.f90: New. * testsuite/libgomp.fortran/metadirective-2.f90: New. * testsuite/libgomp.fortran/metadirective-3.f90: New. * testsuite/libgomp.fortran/metadirective-4.f90: New. * testsuite/libgomp.fortran/metadirective-5.f90: New. * testsuite/libgomp.fortran/metadirective-6.f90: New. Co-Authored-By: Sandra Loosemore Co-Authored-By: Tobias Burnus Co-Authored-By: Paul-Antoine Arras --- gcc/fortran/decl.cc | 29 + gcc/fortran/dump-parse-tree.cc | 21 + gcc/fortran/gfortran.h | 24 +- gcc/fortran/io.cc | 2 +- gcc/fortran/match.h | 2 + gcc/fortran/openmp.cc | 265 +++++++- gcc/fortran/parse.cc | 571 +++++++++++------- gcc/fortran/parse.h | 8 +- gcc/fortran/resolve.cc | 6 + gcc/fortran/st.cc | 4 + gcc/fortran/symbol.cc | 25 +- gcc/fortran/trans-decl.cc | 5 +- gcc/fortran/trans-openmp.cc | 239 +++++--- gcc/fortran/trans-stmt.h | 1 + gcc/fortran/trans.cc | 1 + .../gfortran.dg/gomp/metadirective-1.f90 | 55 ++ .../gfortran.dg/gomp/metadirective-10.f90 | 40 ++ .../gfortran.dg/gomp/metadirective-11.f90 | 33 + .../gfortran.dg/gomp/metadirective-2.f90 | 62 ++ .../gfortran.dg/gomp/metadirective-3.f90 | 34 ++ .../gfortran.dg/gomp/metadirective-4.f90 | 39 ++ .../gfortran.dg/gomp/metadirective-5.f90 | 30 + .../gfortran.dg/gomp/metadirective-6.f90 | 31 + .../gfortran.dg/gomp/metadirective-7.f90 | 36 ++ .../gfortran.dg/gomp/metadirective-8.f90 | 22 + .../gfortran.dg/gomp/metadirective-9.f90 | 30 + .../gomp/metadirective-construct.f90 | 260 ++++++++ .../gomp/metadirective-no-score.f90 | 122 ++++ gcc/testsuite/gfortran.dg/gomp/pure-1.f90 | 7 + gcc/testsuite/gfortran.dg/gomp/pure-2.f90 | 8 - .../libgomp.fortran/metadirective-1.f90 | 61 ++ .../libgomp.fortran/metadirective-2.f90 | 40 ++ .../libgomp.fortran/metadirective-3.f90 | 29 + .../libgomp.fortran/metadirective-4.f90 | 46 ++ .../libgomp.fortran/metadirective-5.f90 | 44 ++ .../libgomp.fortran/metadirective-6.f90 | 58 ++ 36 files changed, 1937 insertions(+), 353 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-10.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-11.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-7.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-8.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-9.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-construct.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-no-score.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/metadirective-1.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/metadirective-2.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/metadirective-3.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/metadirective-4.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/metadirective-5.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/metadirective-6.f90 diff --git a/gcc/fortran/decl.cc b/gcc/fortran/decl.cc index 503ecb8d9b5..0ab2e710e78 100644 --- a/gcc/fortran/decl.cc +++ b/gcc/fortran/decl.cc @@ -8404,6 +8404,7 @@ gfc_match_end (gfc_statement *st) case COMP_CONTAINS: case COMP_DERIVED_CONTAINS: + case COMP_OMP_BEGIN_METADIRECTIVE: state = gfc_state_stack->previous->state; block_name = gfc_state_stack->previous->sym == NULL ? NULL : gfc_state_stack->previous->sym->name; @@ -8411,6 +8412,28 @@ gfc_match_end (gfc_statement *st) && gfc_state_stack->previous->sym->abr_modproc_decl; break; + case COMP_OMP_METADIRECTIVE: + { + /* Metadirectives can be nested, so we need to drill down to the + first state that is not COMP_OMP_METADIRECTIVE. */ + gfc_state_data *state_data = gfc_state_stack; + + do + { + state_data = state_data->previous; + state = state_data->state; + block_name = (state_data->sym == NULL + ? NULL : state_data->sym->name); + abbreviated_modproc_decl = (state_data->sym + && state_data->sym->abr_modproc_decl); + } + while (state == COMP_OMP_METADIRECTIVE); + + if (block_name && startswith (block_name, "block@")) + block_name = NULL; + } + break; + default: break; } @@ -8556,6 +8579,12 @@ gfc_match_end (gfc_statement *st) gfc_free_enum_history (); break; + case COMP_OMP_BEGIN_METADIRECTIVE: + *st = ST_OMP_END_METADIRECTIVE; + target = " metadirective"; + eos_ok = 0; + break; + default: gfc_error ("Unexpected END statement at %C"); goto cleanup; diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 1563b810b98..204c6b60a1f 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -2162,6 +2162,7 @@ show_omp_node (int level, gfc_code *c) case EXEC_OMP_MASTER: name = "MASTER"; break; case EXEC_OMP_MASTER_TASKLOOP: name = "MASTER TASKLOOP"; break; case EXEC_OMP_MASTER_TASKLOOP_SIMD: name = "MASTER TASKLOOP SIMD"; break; + case EXEC_OMP_METADIRECTIVE: name = "METADIRECTIVE"; break; case EXEC_OMP_ORDERED: name = "ORDERED"; break; case EXEC_OMP_DEPOBJ: name = "DEPOBJ"; break; case EXEC_OMP_PARALLEL: name = "PARALLEL"; break; @@ -2359,6 +2360,25 @@ show_omp_node (int level, gfc_code *c) d = d->block; } } + else if (c->op == EXEC_OMP_METADIRECTIVE) + { + gfc_omp_metadirective_variant *variant + = c->ext.omp_metadirective_variants; + + while (variant) + { + code_indent (level + 1, 0); + if (variant->selectors) + fputs ("WHEN ()\n", dumpfile); + else + fputs ("DEFAULT ()\n", dumpfile); + /* TODO: Print selector. */ + show_code (level + 2, variant->code); + if (variant->next) + fputs ("\n", dumpfile); + variant = variant->next; + } + } else show_code (level + 1, c->block->next); if (c->op == EXEC_OMP_ATOMIC) @@ -3493,6 +3513,7 @@ show_code_node (int level, gfc_code *c) case EXEC_OMP_MASTER: case EXEC_OMP_MASTER_TASKLOOP: case EXEC_OMP_MASTER_TASKLOOP_SIMD: + case EXEC_OMP_METADIRECTIVE: case EXEC_OMP_ORDERED: case EXEC_OMP_PARALLEL: case EXEC_OMP_PARALLEL_DO: diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 82f388c05f8..7adc98fccf1 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -317,6 +317,7 @@ enum gfc_statement ST_OMP_END_PARALLEL_MASKED_TASKLOOP_SIMD, ST_OMP_MASKED_TASKLOOP, ST_OMP_END_MASKED_TASKLOOP, ST_OMP_MASKED_TASKLOOP_SIMD, ST_OMP_END_MASKED_TASKLOOP_SIMD, ST_OMP_SCOPE, ST_OMP_END_SCOPE, + ST_OMP_METADIRECTIVE, ST_OMP_BEGIN_METADIRECTIVE, ST_OMP_END_METADIRECTIVE, ST_OMP_ERROR, ST_OMP_ASSUME, ST_OMP_END_ASSUME, ST_OMP_ASSUMES, ST_OMP_ALLOCATE, ST_OMP_ALLOCATE_EXEC, ST_OMP_ALLOCATORS, ST_OMP_END_ALLOCATORS, @@ -1583,7 +1584,7 @@ typedef struct gfc_omp_clauses unsigned order_unconstrained:1, order_reproducible:1, capture:1; unsigned grainsize_strict:1, num_tasks_strict:1, compare:1, weak:1; unsigned non_rectangular:1, order_concurrent:1; - unsigned contains_teams_construct:1, target_first_st_is_teams:1; + unsigned contains_teams_construct:1, target_first_st_is_teams_or_meta:1; unsigned contained_in_target_construct:1; ENUM_BITFIELD (gfc_omp_sched_kind) sched_kind:3; ENUM_BITFIELD (gfc_omp_device_type) device_type:2; @@ -1703,6 +1704,17 @@ typedef struct gfc_omp_declare_variant gfc_omp_declare_variant; #define gfc_get_omp_declare_variant() XCNEW (gfc_omp_declare_variant) +typedef struct gfc_omp_metadirective_variant +{ + struct gfc_omp_metadirective_variant *next; + locus where; /* Where the metadirective clause occurred. */ + + gfc_omp_set_selector *selectors; + enum gfc_statement stmt; + struct gfc_code *code; + +} gfc_omp_metadirective_variant; +#define gfc_get_omp_metadirective_variant() XCNEW (gfc_omp_metadirective_variant) typedef struct gfc_omp_udr { @@ -1751,6 +1763,7 @@ typedef struct gfc_st_label locus where; gfc_namespace *ns; + int omp_region; } gfc_st_label; @@ -3009,6 +3022,7 @@ enum gfc_exec_op EXEC_OMP_TARGET_TEAMS_LOOP, EXEC_OMP_MASKED, EXEC_OMP_PARALLEL_MASKED, EXEC_OMP_PARALLEL_MASKED_TASKLOOP, EXEC_OMP_PARALLEL_MASKED_TASKLOOP_SIMD, EXEC_OMP_MASKED_TASKLOOP, EXEC_OMP_MASKED_TASKLOOP_SIMD, EXEC_OMP_SCOPE, + EXEC_OMP_METADIRECTIVE, EXEC_OMP_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS }; @@ -3067,6 +3081,7 @@ typedef struct gfc_code gfc_omp_clauses *omp_clauses; const char *omp_name; gfc_omp_namelist *omp_namelist; + gfc_omp_metadirective_variant *omp_metadirective_variants; bool omp_bool; } ext; /* Points to additional structures required by statement */ @@ -3661,6 +3676,7 @@ void gfc_free_omp_declare_variant_list (gfc_omp_declare_variant *list); void gfc_free_omp_declare_simd (gfc_omp_declare_simd *); void gfc_free_omp_declare_simd_list (gfc_omp_declare_simd *); void gfc_free_omp_udr (gfc_omp_udr *); +void gfc_free_omp_metadirective_variants (gfc_omp_metadirective_variant *); gfc_omp_udr *gfc_omp_udr_find (gfc_symtree *, gfc_typespec *); void gfc_resolve_omp_allocate (gfc_namespace *, gfc_omp_namelist *); void gfc_resolve_omp_assumptions (gfc_omp_assumptions *); @@ -3942,6 +3958,8 @@ void debug (gfc_expr *); bool gfc_parse_file (void); void gfc_global_used (gfc_gsymbol *, locus *); gfc_namespace* gfc_build_block_ns (gfc_namespace *); +gfc_statement match_omp_directive (void); +bool is_omp_declarative_stmt (gfc_statement); /* dependency.cc */ int gfc_dep_compare_functions (gfc_expr *, gfc_expr *, bool); @@ -4038,4 +4056,8 @@ bool gfc_is_reallocatable_lhs (gfc_expr *); void finish_oacc_declare (gfc_namespace *, gfc_symbol *, bool); void gfc_adjust_builtins (void); +/* trans-openmp.c */ + +bool gfc_skip_omp_metadirective_variant (gfc_omp_metadirective_variant *); + #endif /* GCC_GFORTRAN_H */ diff --git a/gcc/fortran/io.cc b/gcc/fortran/io.cc index 6fd69f7c9a8..9dfb4e1ef25 100644 --- a/gcc/fortran/io.cc +++ b/gcc/fortran/io.cc @@ -29,7 +29,7 @@ along with GCC; see the file COPYING3. If not see gfc_st_label format_asterisk = {0, NULL, NULL, -1, ST_LABEL_FORMAT, ST_LABEL_FORMAT, NULL, - 0, {NULL, NULL}, NULL}; + 0, {NULL, NULL}, NULL, 0}; typedef struct { diff --git a/gcc/fortran/match.h b/gcc/fortran/match.h index b09921357fd..b0f5528ee72 100644 --- a/gcc/fortran/match.h +++ b/gcc/fortran/match.h @@ -155,6 +155,7 @@ match gfc_match_omp_assume (void); match gfc_match_omp_assumes (void); match gfc_match_omp_atomic (void); match gfc_match_omp_barrier (void); +match gfc_match_omp_begin_metadirective (void); match gfc_match_omp_cancel (void); match gfc_match_omp_cancellation_point (void); match gfc_match_omp_critical (void); @@ -178,6 +179,7 @@ match gfc_match_omp_masked_taskloop_simd (void); match gfc_match_omp_master (void); match gfc_match_omp_master_taskloop (void); match gfc_match_omp_master_taskloop_simd (void); +match gfc_match_omp_metadirective (void); match gfc_match_omp_nothing (void); match gfc_match_omp_ordered (void); match gfc_match_omp_ordered_depend (void); diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 0af80d54fad..b602d9aa34d 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -113,7 +113,8 @@ static const struct gfc_omp_directive gfc_omp_directives[] = { /* Match an end of OpenMP directive. End of OpenMP directive is optional - whitespace, followed by '\n' or comment '!'. */ + whitespace, followed by '\n' or comment '!'. In the special case where a + context selector is being matched, match against ')' instead. */ static match gfc_match_omp_eos (void) @@ -124,17 +125,25 @@ gfc_match_omp_eos (void) old_loc = gfc_current_locus; gfc_gobble_whitespace (); - c = gfc_next_ascii_char (); - switch (c) + if (gfc_matching_omp_context_selector) { - case '!': - do - c = gfc_next_ascii_char (); - while (c != '\n'); - /* Fall through */ + if (gfc_peek_ascii_char () == ')') + return MATCH_YES; + } + else + { + c = gfc_next_ascii_char (); + switch (c) + { + case '!': + do + c = gfc_next_ascii_char (); + while (c != '\n'); + /* Fall through */ - case '\n': - return MATCH_YES; + case '\n': + return MATCH_YES; + } } gfc_current_locus = old_loc; @@ -339,6 +348,19 @@ gfc_free_omp_udr (gfc_omp_udr *omp_udr) } } +/* Free variants of an !$omp metadirective construct. */ + +void +gfc_free_omp_metadirective_variants (gfc_omp_metadirective_variant *variant) +{ + while (variant) + { + gfc_omp_metadirective_variant *next_variant = variant->next; + gfc_free_omp_set_selector_list (variant->selectors); + free (variant); + variant = next_variant; + } +} static gfc_omp_udr * gfc_find_omp_udr (gfc_namespace *ns, const char *name, gfc_typespec *ts) @@ -1866,8 +1888,7 @@ gfc_match_dupl_atomic (bool not_dupl, const char *name) static match gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, bool first = true, bool needs_space = true, - bool openacc = false, bool context_selector = false, - bool openmp_target = false) + bool openacc = false, bool openmp_target = false) { bool error = false; gfc_omp_clauses *c = gfc_get_omp_clauses (); @@ -3789,9 +3810,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, } end: - if (error - || (context_selector && gfc_peek_ascii_char () != ')') - || (!context_selector && gfc_match_omp_eos () != MATCH_YES)) + if (error || gfc_match_omp_eos () != MATCH_YES) { if (!gfc_error_flag_test ()) gfc_error ("Failed to match clause at %C"); @@ -4479,7 +4498,7 @@ static match match_omp (gfc_exec_op op, const omp_mask mask) { gfc_omp_clauses *c; - if (gfc_match_omp_clauses (&c, mask, true, true, false, false, + if (gfc_match_omp_clauses (&c, mask, true, true, false, op == EXEC_OMP_TARGET) != MATCH_YES) return MATCH_ERROR; new_st.op = op; @@ -5810,14 +5829,17 @@ gfc_match_omp_context_selector (gfc_omp_set_selector *oss) { if (os->code == OMP_TRAIT_CONSTRUCT_SIMD) { + gfc_matching_omp_context_selector = true; if (gfc_match_omp_clauses (&otp->clauses, OMP_DECLARE_SIMD_CLAUSES, - true, false, false, true) + true, false, false) != MATCH_YES) { + gfc_matching_omp_context_selector = false; gfc_error ("expected simd clause at %C"); return MATCH_ERROR; } + gfc_matching_omp_context_selector = false; } else if (os->code == OMP_TRAIT_IMPLEMENTATION_REQUIRES) { @@ -5874,7 +5896,7 @@ gfc_match_omp_context_selector (gfc_omp_set_selector *oss) user */ match -gfc_match_omp_context_selector_specification (gfc_omp_declare_variant *odv) +gfc_match_omp_context_selector_specification (gfc_omp_set_selector **oss_head) { do { @@ -5907,9 +5929,9 @@ gfc_match_omp_context_selector_specification (gfc_omp_declare_variant *odv) } gfc_omp_set_selector *oss = gfc_get_omp_set_selector (); - oss->next = odv->set_selectors; + oss->next = *oss_head; oss->code = set; - odv->set_selectors = oss; + *oss_head = oss; if (gfc_match_omp_context_selector (oss) != MATCH_YES) return MATCH_ERROR; @@ -6010,7 +6032,8 @@ gfc_match_omp_declare_variant (void) return MATCH_ERROR; } - if (gfc_match_omp_context_selector_specification (odv) != MATCH_YES) + if (gfc_match_omp_context_selector_specification (&odv->set_selectors) + != MATCH_YES) return MATCH_ERROR; if (gfc_match (" )") != MATCH_YES) @@ -6026,6 +6049,157 @@ gfc_match_omp_declare_variant (void) } +static match +match_omp_metadirective (bool begin_p) +{ + locus old_loc = gfc_current_locus; + gfc_omp_metadirective_variant *variants_head; + gfc_omp_metadirective_variant **next_variant = &variants_head; + bool default_seen = false; + + /* Parse the context selectors. */ + for (;;) + { + bool default_p = false; + gfc_omp_set_selector *selectors = NULL; + locus variant_locus = gfc_current_locus; + + if (gfc_match (" default ( ") == MATCH_YES) + default_p = true; + else if (gfc_match (" otherwise ( ") == MATCH_YES) + default_p = true; + else if (gfc_match_eos () == MATCH_YES) + break; + else if (gfc_match (" when ( ") != MATCH_YES) + { + gfc_error ("expected %, %, or % at %C"); + gfc_current_locus = old_loc; + return MATCH_ERROR; + } + + if (default_p && default_seen) + { + gfc_error ("too many % or % clauses " + "in % at %C"); + gfc_current_locus = old_loc; + return MATCH_ERROR; + } + + if (!default_p) + { + if (gfc_match_omp_context_selector_specification (&selectors) + != MATCH_YES) + return MATCH_ERROR; + + if (gfc_match (" : ") != MATCH_YES) + { + gfc_error ("expected %<:%> at %C"); + gfc_current_locus = old_loc; + return MATCH_ERROR; + } + + gfc_commit_symbols (); + } + + gfc_matching_omp_context_selector = true; + gfc_statement directive = match_omp_directive (); + gfc_matching_omp_context_selector = false; + + if (is_omp_declarative_stmt (directive)) + sorry ("declarative directive variants are not supported"); + + if (gfc_error_flag_test ()) + { + gfc_current_locus = old_loc; + return MATCH_ERROR; + } + + if (gfc_match (" )") != MATCH_YES) + { + gfc_error ("Expected %<)%> at %C"); + gfc_current_locus = old_loc; + return MATCH_ERROR; + } + + gfc_commit_symbols (); + + if (begin_p && directive != ST_NONE + && gfc_omp_end_stmt (directive) == ST_NONE) + { + gfc_error ("variant directive used in OMP BEGIN METADIRECTIVE " + "at %C must have a corresponding end directive"); + gfc_current_locus = old_loc; + return MATCH_ERROR; + } + + if (default_p) + default_seen = true; + + gfc_omp_metadirective_variant *omv = gfc_get_omp_metadirective_variant (); + omv->selectors = selectors; + omv->stmt = directive; + omv->where = variant_locus; + + if (directive == ST_NONE) + { + /* The directive was a 'nothing' directive. */ + omv->code = gfc_get_code (EXEC_CONTINUE); + omv->code->ext.omp_clauses = NULL; + } + else + { + omv->code = gfc_get_code (new_st.op); + omv->code->ext.omp_clauses = new_st.ext.omp_clauses; + /* Prevent the OpenMP clauses from being freed via NEW_ST. */ + new_st.ext.omp_clauses = NULL; + } + + if (!gfc_skip_omp_metadirective_variant (omv)) + { + *next_variant = omv; + next_variant = &omv->next; + } + } + + if (gfc_match_omp_eos () != MATCH_YES) + { + gfc_error ("Unexpected junk after OMP METADIRECTIVE at %C"); + gfc_current_locus = old_loc; + return MATCH_ERROR; + } + + /* Add a 'default (nothing)' clause if no default is explicitly given. */ + if (!default_seen) + { + gfc_omp_metadirective_variant *omv = gfc_get_omp_metadirective_variant (); + omv->stmt = ST_NONE; + omv->code = gfc_get_code (EXEC_CONTINUE); + omv->code->ext.omp_clauses = NULL; + omv->where = old_loc; + omv->selectors = NULL; + + *next_variant = omv; + next_variant = &omv->next; + } + + new_st.op = EXEC_OMP_METADIRECTIVE; + new_st.ext.omp_metadirective_variants = variants_head; + + return MATCH_YES; +} + +match +gfc_match_omp_begin_metadirective (void) +{ + return match_omp_metadirective (true); +} + +match +gfc_match_omp_metadirective (void) +{ + return match_omp_metadirective (false); +} + match gfc_match_omp_threadprivate (void) { @@ -10919,6 +11093,19 @@ resolve_omp_do (gfc_code *code) restructure_intervening_code (&(code->block->next), code, count); } +static void +resolve_omp_metadirective (gfc_code *code, gfc_namespace *ns) +{ + gfc_omp_metadirective_variant *variant = code->ext.omp_metadirective_variants; + + while (variant) + { + gfc_code *variant_code = variant->code; + gfc_resolve_code (variant_code, ns); + variant = variant->next; + } +} + static gfc_statement omp_code_to_statement (gfc_code *code) @@ -11462,13 +11649,32 @@ resolve_omp_target (gfc_code *code) gfc_code *c = code->block->next; if (c->op == EXEC_BLOCK) c = c->ext.block.ns->code; - if (code->ext.omp_clauses->target_first_st_is_teams - && ((GFC_IS_TEAMS_CONSTRUCT (c->op) && c->next == NULL) - || (c->op == EXEC_BLOCK - && c->next - && GFC_IS_TEAMS_CONSTRUCT (c->next->op) - && c->next->next == NULL))) - return; + if (code->ext.omp_clauses->target_first_st_is_teams_or_meta) + { + if (c->op == EXEC_OMP_METADIRECTIVE) + { + struct gfc_omp_metadirective_variant *mc + = c->ext.omp_metadirective_variants; + /* All mc->(next...->)code should be identical with regards + to the diagnostic below. */ + do + { + if (mc->stmt != ST_NONE + && GFC_IS_TEAMS_CONSTRUCT (mc->code->op)) + { + if (c->next == NULL && mc->code->next == NULL) + return; + c = mc->code; + break; + } + mc = mc->next; + } + while (mc); + } + else if (GFC_IS_TEAMS_CONSTRUCT (c->op) && c->next == NULL) + return; + } + while (c && !GFC_IS_TEAMS_CONSTRUCT (c->op)) c = c->next; if (c) @@ -11595,6 +11801,9 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns) code->ext.omp_clauses->if_present = false; resolve_omp_clauses (code, code->ext.omp_clauses, ns); break; + case EXEC_OMP_METADIRECTIVE: + resolve_omp_metadirective (code, ns); + break; default: break; } diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc index d8b38cfb5ac..f23045636c9 100644 --- a/gcc/fortran/parse.cc +++ b/gcc/fortran/parse.cc @@ -48,6 +48,10 @@ gfc_state_data *gfc_state_stack; static bool last_was_use_stmt = false; bool in_exec_part; +bool gfc_matching_omp_context_selector; +bool gfc_in_metadirective_body; +int gfc_omp_region_count; + /* TODO: Re-order functions to kill these forward decls. */ static void check_statement_label (gfc_statement); static void undo_new_statement (void); @@ -1041,6 +1045,8 @@ decode_omp_directive (void) break; case 'b': matcho ("barrier", gfc_match_omp_barrier, ST_OMP_BARRIER); + matcho ("begin metadirective", gfc_match_omp_begin_metadirective, + ST_OMP_BEGIN_METADIRECTIVE); break; case 'c': matcho ("cancellation% point", gfc_match_omp_cancellation_point, @@ -1085,6 +1091,8 @@ decode_omp_directive (void) matcho ("end master taskloop", gfc_match_omp_eos_error, ST_OMP_END_MASTER_TASKLOOP); matcho ("end master", gfc_match_omp_eos_error, ST_OMP_END_MASTER); + matcho ("end metadirective", gfc_match_omp_eos_error, + ST_OMP_END_METADIRECTIVE); matchs ("end ordered", gfc_match_omp_eos_error, ST_OMP_END_ORDERED); matchs ("end parallel do simd", gfc_match_omp_eos_error, ST_OMP_END_PARALLEL_DO_SIMD); @@ -1168,6 +1176,8 @@ decode_omp_directive (void) matcho ("master taskloop", gfc_match_omp_master_taskloop, ST_OMP_MASTER_TASKLOOP); matcho ("master", gfc_match_omp_master, ST_OMP_MASTER); + matcho ("metadirective", gfc_match_omp_metadirective, + ST_OMP_METADIRECTIVE); break; case 'n': matcho ("nothing", gfc_match_omp_nothing, ST_NONE); @@ -1296,6 +1306,10 @@ decode_omp_directive (void) gfc_error_now ("Unclassifiable OpenMP directive at %C"); } + /* If parsing a metadirective, let the caller deal with the cleanup. */ + if (gfc_matching_omp_context_selector) + return ST_NONE; + reject_statement (); gfc_error_recovery (); @@ -1413,6 +1427,12 @@ decode_omp_directive (void) return ST_GET_FCN_CHARACTERISTICS; } +gfc_statement +match_omp_directive (void) +{ + return decode_omp_directive (); +} + static gfc_statement decode_gcc_attribute (void) { @@ -1935,6 +1955,43 @@ next_statement (void) case ST_OMP_DECLARE_VARIANT: case ST_OMP_ALLOCATE: case ST_OMP_ASSUMES: \ case ST_OMP_REQUIRES: case ST_OACC_ROUTINE: case ST_OACC_DECLARE +/* OpenMP statements that are followed by a structured block. */ + +#define case_omp_structured_block case ST_OMP_ASSUME: case ST_OMP_PARALLEL: \ + case ST_OMP_PARALLEL_MASKED: case ST_OMP_PARALLEL_MASTER: \ + case ST_OMP_PARALLEL_SECTIONS: case ST_OMP_ORDERED: \ + case ST_OMP_CRITICAL: case ST_OMP_MASKED: case ST_OMP_MASTER: \ + case ST_OMP_SCOPE: case ST_OMP_SECTIONS: case ST_OMP_SINGLE: \ + case ST_OMP_TARGET: case ST_OMP_TARGET_DATA: case ST_OMP_TARGET_PARALLEL: \ + case ST_OMP_TARGET_TEAMS: case ST_OMP_TEAMS: case ST_OMP_TASK: \ + case ST_OMP_TASKGROUP: \ + case ST_OMP_WORKSHARE: case ST_OMP_PARALLEL_WORKSHARE + +/* OpenMP statements that are followed by a do loop. */ + +#define case_omp_do case ST_OMP_DISTRIBUTE: \ + case ST_OMP_DISTRIBUTE_PARALLEL_DO: \ + case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: case ST_OMP_DISTRIBUTE_SIMD: \ + case ST_OMP_DO: case ST_OMP_DO_SIMD: case ST_OMP_LOOP: \ + case ST_OMP_PARALLEL_DO: case ST_OMP_PARALLEL_DO_SIMD: \ + case ST_OMP_PARALLEL_LOOP: case ST_OMP_PARALLEL_MASKED_TASKLOOP: \ + case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD: \ + case ST_OMP_PARALLEL_MASTER_TASKLOOP: \ + case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD: \ + case ST_OMP_MASKED_TASKLOOP: case ST_OMP_MASKED_TASKLOOP_SIMD: \ + case ST_OMP_MASTER_TASKLOOP: case ST_OMP_MASTER_TASKLOOP_SIMD: \ + case ST_OMP_SIMD: \ + case ST_OMP_TARGET_PARALLEL_DO: case ST_OMP_TARGET_PARALLEL_DO_SIMD: \ + case ST_OMP_TARGET_PARALLEL_LOOP: case ST_OMP_TARGET_SIMD: \ + case ST_OMP_TARGET_TEAMS_DISTRIBUTE: \ + case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: \ + case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: \ + case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: case ST_OMP_TARGET_TEAMS_LOOP: \ + case ST_OMP_TASKLOOP: case ST_OMP_TASKLOOP_SIMD: \ + case ST_OMP_TEAMS_DISTRIBUTE: case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: \ + case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: \ + case ST_OMP_TEAMS_DISTRIBUTE_SIMD: case ST_OMP_TEAMS_LOOP + /* Block end statements. Errors associated with interchanging these are detected in gfc_match_end(). */ @@ -2572,6 +2629,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel) case ST_OMP_BARRIER: p = "!$OMP BARRIER"; break; + case ST_OMP_BEGIN_METADIRECTIVE: + p = "!$OMP BEGIN METADIRECTIVE"; + break; case ST_OMP_CANCEL: p = "!$OMP CANCEL"; break; @@ -2671,6 +2731,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel) case ST_OMP_END_MASTER_TASKLOOP_SIMD: p = "!$OMP END MASTER TASKLOOP SIMD"; break; + case ST_OMP_END_METADIRECTIVE: + p = "!$OMP END METADIRECTIVE"; + break; case ST_OMP_END_ORDERED: p = "!$OMP END ORDERED"; break; @@ -2815,6 +2878,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel) case ST_OMP_MASTER_TASKLOOP_SIMD: p = "!$OMP MASTER TASKLOOP SIMD"; break; + case ST_OMP_METADIRECTIVE: + p = "!$OMP METADIRECTIVE"; + break; case ST_OMP_ORDERED: case ST_OMP_ORDERED_DEPEND: p = "!$OMP ORDERED"; @@ -3075,6 +3141,8 @@ accept_statement (gfc_statement st) break; case ST_ENTRY: + case ST_OMP_METADIRECTIVE: + case ST_OMP_BEGIN_METADIRECTIVE: case_executable: case_exec_markers: add_statement (); @@ -5386,6 +5454,140 @@ loop: accept_statement (st); } +/* Get the corresponding ending statement type for the OpenMP directive + OMP_ST. If it does not have one, return ST_NONE. */ + +gfc_statement +gfc_omp_end_stmt (gfc_statement omp_st, + bool omp_do_p, bool omp_structured_p) +{ + if (omp_do_p) + { + switch (omp_st) + { + case ST_OMP_DISTRIBUTE: return ST_OMP_END_DISTRIBUTE; + case ST_OMP_DISTRIBUTE_PARALLEL_DO: + return ST_OMP_END_DISTRIBUTE_PARALLEL_DO; + case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: + return ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD; + case ST_OMP_DISTRIBUTE_SIMD: + return ST_OMP_END_DISTRIBUTE_SIMD; + case ST_OMP_DO: return ST_OMP_END_DO; + case ST_OMP_DO_SIMD: return ST_OMP_END_DO_SIMD; + case ST_OMP_LOOP: return ST_OMP_END_LOOP; + case ST_OMP_PARALLEL_DO: return ST_OMP_END_PARALLEL_DO; + case ST_OMP_PARALLEL_DO_SIMD: + return ST_OMP_END_PARALLEL_DO_SIMD; + case ST_OMP_PARALLEL_LOOP: + return ST_OMP_END_PARALLEL_LOOP; + case ST_OMP_SIMD: return ST_OMP_END_SIMD; + case ST_OMP_TARGET_PARALLEL_DO: + return ST_OMP_END_TARGET_PARALLEL_DO; + case ST_OMP_TARGET_PARALLEL_DO_SIMD: + return ST_OMP_END_TARGET_PARALLEL_DO_SIMD; + case ST_OMP_TARGET_PARALLEL_LOOP: + return ST_OMP_END_TARGET_PARALLEL_LOOP; + case ST_OMP_TARGET_SIMD: return ST_OMP_END_TARGET_SIMD; + case ST_OMP_TARGET_TEAMS_DISTRIBUTE: + return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE; + case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: + return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO; + case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: + return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD; + case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: + return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD; + case ST_OMP_TARGET_TEAMS_LOOP: + return ST_OMP_END_TARGET_TEAMS_LOOP; + case ST_OMP_TASKLOOP: return ST_OMP_END_TASKLOOP; + case ST_OMP_TASKLOOP_SIMD: return ST_OMP_END_TASKLOOP_SIMD; + case ST_OMP_MASKED_TASKLOOP: return ST_OMP_END_MASKED_TASKLOOP; + case ST_OMP_MASKED_TASKLOOP_SIMD: + return ST_OMP_END_MASKED_TASKLOOP_SIMD; + case ST_OMP_MASTER_TASKLOOP: return ST_OMP_END_MASTER_TASKLOOP; + case ST_OMP_MASTER_TASKLOOP_SIMD: + return ST_OMP_END_MASTER_TASKLOOP_SIMD; + case ST_OMP_PARALLEL_MASKED_TASKLOOP: + return ST_OMP_END_PARALLEL_MASKED_TASKLOOP; + case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD: + return ST_OMP_END_PARALLEL_MASKED_TASKLOOP_SIMD; + case ST_OMP_PARALLEL_MASTER_TASKLOOP: + return ST_OMP_END_PARALLEL_MASTER_TASKLOOP; + case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD: + return ST_OMP_END_PARALLEL_MASTER_TASKLOOP_SIMD; + case ST_OMP_TEAMS_DISTRIBUTE: + return ST_OMP_END_TEAMS_DISTRIBUTE; + case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: + return ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO; + case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: + return ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD; + case ST_OMP_TEAMS_DISTRIBUTE_SIMD: + return ST_OMP_END_TEAMS_DISTRIBUTE_SIMD; + case ST_OMP_TEAMS_LOOP: + return ST_OMP_END_TEAMS_LOOP; + default: + break; + } + } + + if (omp_structured_p) + { + switch (omp_st) + { + case ST_OMP_ASSUME: + return ST_OMP_END_ASSUME; + case ST_OMP_PARALLEL: + return ST_OMP_END_PARALLEL; + case ST_OMP_PARALLEL_MASKED: + return ST_OMP_END_PARALLEL_MASKED; + case ST_OMP_PARALLEL_MASTER: + return ST_OMP_END_PARALLEL_MASTER; + case ST_OMP_PARALLEL_SECTIONS: + return ST_OMP_END_PARALLEL_SECTIONS; + case ST_OMP_SCOPE: + return ST_OMP_END_SCOPE; + case ST_OMP_SECTIONS: + return ST_OMP_END_SECTIONS; + case ST_OMP_ORDERED: + return ST_OMP_END_ORDERED; + case ST_OMP_CRITICAL: + return ST_OMP_END_CRITICAL; + case ST_OMP_MASKED: + return ST_OMP_END_MASKED; + case ST_OMP_MASTER: + return ST_OMP_END_MASTER; + case ST_OMP_SINGLE: + return ST_OMP_END_SINGLE; + case ST_OMP_TARGET: + return ST_OMP_END_TARGET; + case ST_OMP_TARGET_DATA: + return ST_OMP_END_TARGET_DATA; + case ST_OMP_TARGET_PARALLEL: + return ST_OMP_END_TARGET_PARALLEL; + case ST_OMP_TARGET_TEAMS: + return ST_OMP_END_TARGET_TEAMS; + case ST_OMP_TASK: + return ST_OMP_END_TASK; + case ST_OMP_TASKGROUP: + return ST_OMP_END_TASKGROUP; + case ST_OMP_TEAMS: + return ST_OMP_END_TEAMS; + case ST_OMP_TEAMS_DISTRIBUTE: + return ST_OMP_END_TEAMS_DISTRIBUTE; + case ST_OMP_DISTRIBUTE: + return ST_OMP_END_DISTRIBUTE; + case ST_OMP_WORKSHARE: + return ST_OMP_END_WORKSHARE; + case ST_OMP_PARALLEL_WORKSHARE: + return ST_OMP_END_PARALLEL_WORKSHARE; + case ST_OMP_BEGIN_METADIRECTIVE: + return ST_OMP_END_METADIRECTIVE; + default: + break; + } + } + + return ST_NONE; +} /* Parse the statements of OpenMP do/parallel do. */ @@ -5436,94 +5638,16 @@ parse_omp_do (gfc_statement omp_st) pop_state (); st = next_statement (); - gfc_statement omp_end_st = ST_OMP_END_DO; - switch (omp_st) - { - case ST_OMP_DISTRIBUTE: omp_end_st = ST_OMP_END_DISTRIBUTE; break; - case ST_OMP_DISTRIBUTE_PARALLEL_DO: - omp_end_st = ST_OMP_END_DISTRIBUTE_PARALLEL_DO; - break; - case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: - omp_end_st = ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD; - break; - case ST_OMP_DISTRIBUTE_SIMD: - omp_end_st = ST_OMP_END_DISTRIBUTE_SIMD; - break; - case ST_OMP_DO: omp_end_st = ST_OMP_END_DO; break; - case ST_OMP_DO_SIMD: omp_end_st = ST_OMP_END_DO_SIMD; break; - case ST_OMP_LOOP: omp_end_st = ST_OMP_END_LOOP; break; - case ST_OMP_PARALLEL_DO: omp_end_st = ST_OMP_END_PARALLEL_DO; break; - case ST_OMP_PARALLEL_DO_SIMD: - omp_end_st = ST_OMP_END_PARALLEL_DO_SIMD; - break; - case ST_OMP_PARALLEL_LOOP: - omp_end_st = ST_OMP_END_PARALLEL_LOOP; - break; - case ST_OMP_SIMD: omp_end_st = ST_OMP_END_SIMD; break; - case ST_OMP_TARGET_PARALLEL_DO: - omp_end_st = ST_OMP_END_TARGET_PARALLEL_DO; - break; - case ST_OMP_TARGET_PARALLEL_DO_SIMD: - omp_end_st = ST_OMP_END_TARGET_PARALLEL_DO_SIMD; - break; - case ST_OMP_TARGET_PARALLEL_LOOP: - omp_end_st = ST_OMP_END_TARGET_PARALLEL_LOOP; - break; - case ST_OMP_TARGET_SIMD: omp_end_st = ST_OMP_END_TARGET_SIMD; break; - case ST_OMP_TARGET_TEAMS_DISTRIBUTE: - omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE; - break; - case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: - omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO; - break; - case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: - omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD; - break; - case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: - omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD; - break; - case ST_OMP_TARGET_TEAMS_LOOP: - omp_end_st = ST_OMP_END_TARGET_TEAMS_LOOP; - break; - case ST_OMP_TASKLOOP: omp_end_st = ST_OMP_END_TASKLOOP; break; - case ST_OMP_TASKLOOP_SIMD: omp_end_st = ST_OMP_END_TASKLOOP_SIMD; break; - case ST_OMP_MASKED_TASKLOOP: omp_end_st = ST_OMP_END_MASKED_TASKLOOP; break; - case ST_OMP_MASKED_TASKLOOP_SIMD: - omp_end_st = ST_OMP_END_MASKED_TASKLOOP_SIMD; - break; - case ST_OMP_MASTER_TASKLOOP: omp_end_st = ST_OMP_END_MASTER_TASKLOOP; break; - case ST_OMP_MASTER_TASKLOOP_SIMD: - omp_end_st = ST_OMP_END_MASTER_TASKLOOP_SIMD; - break; - case ST_OMP_PARALLEL_MASKED_TASKLOOP: - omp_end_st = ST_OMP_END_PARALLEL_MASKED_TASKLOOP; - break; - case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD: - omp_end_st = ST_OMP_END_PARALLEL_MASKED_TASKLOOP_SIMD; - break; - case ST_OMP_PARALLEL_MASTER_TASKLOOP: - omp_end_st = ST_OMP_END_PARALLEL_MASTER_TASKLOOP; - break; - case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD: - omp_end_st = ST_OMP_END_PARALLEL_MASTER_TASKLOOP_SIMD; - break; - case ST_OMP_TEAMS_DISTRIBUTE: - omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE; - break; - case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: - omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO; - break; - case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: - omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD; - break; - case ST_OMP_TEAMS_DISTRIBUTE_SIMD: - omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_SIMD; - break; - case ST_OMP_TEAMS_LOOP: - omp_end_st = ST_OMP_END_TEAMS_LOOP; - break; - default: gcc_unreachable (); - } + gfc_statement omp_end_st = gfc_omp_end_stmt (omp_st, true, false); + if (omp_st == ST_NONE) + gcc_unreachable (); + + /* If handling a metadirective variant, treat 'omp end metadirective' + as the expected end statement for the current construct. */ + if (st == ST_OMP_END_METADIRECTIVE + && gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE) + st = omp_end_st; + if (st == omp_end_st) { if (new_st.op == EXEC_OMP_END_NOWAIT) @@ -5835,80 +5959,15 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only) np->op = cp->op; np->block = NULL; - switch (omp_st) - { - case ST_OMP_ASSUME: - omp_end_st = ST_OMP_END_ASSUME; - break; - case ST_OMP_PARALLEL: - omp_end_st = ST_OMP_END_PARALLEL; - break; - case ST_OMP_PARALLEL_MASKED: - omp_end_st = ST_OMP_END_PARALLEL_MASKED; - break; - case ST_OMP_PARALLEL_MASTER: - omp_end_st = ST_OMP_END_PARALLEL_MASTER; - break; - case ST_OMP_PARALLEL_SECTIONS: - omp_end_st = ST_OMP_END_PARALLEL_SECTIONS; - break; - case ST_OMP_SCOPE: - omp_end_st = ST_OMP_END_SCOPE; - break; - case ST_OMP_SECTIONS: - omp_end_st = ST_OMP_END_SECTIONS; - break; - case ST_OMP_ORDERED: - omp_end_st = ST_OMP_END_ORDERED; - break; - case ST_OMP_CRITICAL: - omp_end_st = ST_OMP_END_CRITICAL; - break; - case ST_OMP_MASKED: - omp_end_st = ST_OMP_END_MASKED; - break; - case ST_OMP_MASTER: - omp_end_st = ST_OMP_END_MASTER; - break; - case ST_OMP_SINGLE: - omp_end_st = ST_OMP_END_SINGLE; - break; - case ST_OMP_TARGET: - omp_end_st = ST_OMP_END_TARGET; - break; - case ST_OMP_TARGET_DATA: - omp_end_st = ST_OMP_END_TARGET_DATA; - break; - case ST_OMP_TARGET_PARALLEL: - omp_end_st = ST_OMP_END_TARGET_PARALLEL; - break; - case ST_OMP_TARGET_TEAMS: - omp_end_st = ST_OMP_END_TARGET_TEAMS; - break; - case ST_OMP_TASK: - omp_end_st = ST_OMP_END_TASK; - break; - case ST_OMP_TASKGROUP: - omp_end_st = ST_OMP_END_TASKGROUP; - break; - case ST_OMP_TEAMS: - omp_end_st = ST_OMP_END_TEAMS; - break; - case ST_OMP_TEAMS_DISTRIBUTE: - omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE; - break; - case ST_OMP_DISTRIBUTE: - omp_end_st = ST_OMP_END_DISTRIBUTE; - break; - case ST_OMP_WORKSHARE: - omp_end_st = ST_OMP_END_WORKSHARE; - break; - case ST_OMP_PARALLEL_WORKSHARE: - omp_end_st = ST_OMP_END_PARALLEL_WORKSHARE; - break; - default: - gcc_unreachable (); - } + omp_end_st = gfc_omp_end_stmt (omp_st, false, true); + if (omp_end_st == ST_NONE) + gcc_unreachable (); + + /* If handling a metadirective variant, treat 'omp end metadirective' + as the expected end statement for the current construct. */ + if (gfc_state_stack->previous != NULL + && gfc_state_stack->previous->state == COMP_OMP_BEGIN_METADIRECTIVE) + omp_end_st = ST_OMP_END_METADIRECTIVE; bool block_construct = false; gfc_namespace *my_ns = NULL; @@ -5954,11 +6013,13 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only) case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: case ST_OMP_TEAMS_LOOP: + case ST_OMP_METADIRECTIVE: + case ST_OMP_BEGIN_METADIRECTIVE: { gfc_state_data *stk = gfc_state_stack->previous; if (stk->state == COMP_OMP_STRICTLY_STRUCTURED_BLOCK) stk = stk->previous; - stk->tail->ext.omp_clauses->target_first_st_is_teams = true; + stk->tail->ext.omp_clauses->target_first_st_is_teams_or_meta = true; break; } default: @@ -6131,6 +6192,88 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only) return st; } +static gfc_statement +parse_omp_metadirective_body (gfc_statement omp_st) +{ + gfc_omp_metadirective_variant *variant + = new_st.ext.omp_metadirective_variants; + locus body_locus = gfc_current_locus; + + accept_statement (omp_st); + + gfc_statement next_st = ST_NONE; + + while (variant) + { + gfc_current_locus = body_locus; + gfc_state_data s; + bool workshare_p + = (variant->stmt == ST_OMP_WORKSHARE + || variant->stmt == ST_OMP_PARALLEL_WORKSHARE); + enum gfc_compile_state new_state + = (omp_st == ST_OMP_METADIRECTIVE + ? COMP_OMP_METADIRECTIVE : COMP_OMP_BEGIN_METADIRECTIVE); + + new_st = *variant->code; + push_state (&s, new_state, NULL); + + gfc_statement st; + bool old_in_metadirective_body = gfc_in_metadirective_body; + gfc_in_metadirective_body = true; + + gfc_omp_region_count++; + switch (variant->stmt) + { + case_omp_structured_block: + st = parse_omp_structured_block (variant->stmt, workshare_p); + break; + case_omp_do: + st = parse_omp_do (variant->stmt); + /* TODO: Does st == ST_IMPLIED_ENDDO need special handling? */ + break; + default: + accept_statement (variant->stmt); + st = parse_executable (next_statement ()); + break; + } + + if (gfc_state_stack->state == COMP_OMP_METADIRECTIVE + && startswith (gfc_ascii_statement (st), "!$OMP END ")) + { + for (gfc_state_data *p = gfc_state_stack; p; p = p->previous) + if (p->state == COMP_OMP_STRUCTURED_BLOCK + || p->state == COMP_OMP_BEGIN_METADIRECTIVE) + goto finish; + gfc_error ( + "Unexpected %s statement in an OMP METADIRECTIVE block at %C", + gfc_ascii_statement (st)); + reject_statement (); + st = next_statement (); + } + finish: + + gfc_in_metadirective_body = old_in_metadirective_body; + + if (gfc_state_stack->head) + *variant->code = *gfc_state_stack->head; + pop_state (); + + gfc_commit_symbols (); + gfc_warning_check (); + if (variant->next) + gfc_clear_new_st (); + + /* Sanity-check that each variant finishes parsing at the same place. */ + if (next_st == ST_NONE) + next_st = st; + else + gcc_assert (st == next_st); + + variant = variant->next; + } + + return next_st; +} /* Accept a series of executable statements. We return the first statement that doesn't fit to the caller. Any block statements are @@ -6141,6 +6284,7 @@ static gfc_statement parse_executable (gfc_statement st) { int close_flag; + bool one_stmt_p = false; in_exec_part = true; if (st == ST_NONE) @@ -6148,6 +6292,12 @@ parse_executable (gfc_statement st) for (;;) { + /* Only parse one statement for the form of metadirective without + an explicit begin..end. */ + if (gfc_state_stack->state == COMP_OMP_METADIRECTIVE && one_stmt_p) + return st; + one_stmt_p = true; + close_flag = check_do_closure (); if (close_flag) switch (st) @@ -6257,68 +6407,13 @@ parse_executable (gfc_statement st) st = parse_openmp_allocate_block (st); continue; - case ST_OMP_ASSUME: - case ST_OMP_PARALLEL: - case ST_OMP_PARALLEL_MASKED: - case ST_OMP_PARALLEL_MASTER: - case ST_OMP_PARALLEL_SECTIONS: - case ST_OMP_ORDERED: - case ST_OMP_CRITICAL: - case ST_OMP_MASKED: - case ST_OMP_MASTER: - case ST_OMP_SCOPE: - case ST_OMP_SECTIONS: - case ST_OMP_SINGLE: - case ST_OMP_TARGET: - case ST_OMP_TARGET_DATA: - case ST_OMP_TARGET_PARALLEL: - case ST_OMP_TARGET_TEAMS: - case ST_OMP_TEAMS: - case ST_OMP_TASK: - case ST_OMP_TASKGROUP: - st = parse_omp_structured_block (st, false); + case_omp_structured_block: + st = parse_omp_structured_block (st, + st == ST_OMP_WORKSHARE + || st == ST_OMP_PARALLEL_WORKSHARE); continue; - case ST_OMP_WORKSHARE: - case ST_OMP_PARALLEL_WORKSHARE: - st = parse_omp_structured_block (st, true); - continue; - - case ST_OMP_DISTRIBUTE: - case ST_OMP_DISTRIBUTE_PARALLEL_DO: - case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: - case ST_OMP_DISTRIBUTE_SIMD: - case ST_OMP_DO: - case ST_OMP_DO_SIMD: - case ST_OMP_LOOP: - case ST_OMP_PARALLEL_DO: - case ST_OMP_PARALLEL_DO_SIMD: - case ST_OMP_PARALLEL_LOOP: - case ST_OMP_PARALLEL_MASKED_TASKLOOP: - case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD: - case ST_OMP_PARALLEL_MASTER_TASKLOOP: - case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD: - case ST_OMP_MASKED_TASKLOOP: - case ST_OMP_MASKED_TASKLOOP_SIMD: - case ST_OMP_MASTER_TASKLOOP: - case ST_OMP_MASTER_TASKLOOP_SIMD: - case ST_OMP_SIMD: - case ST_OMP_TARGET_PARALLEL_DO: - case ST_OMP_TARGET_PARALLEL_DO_SIMD: - case ST_OMP_TARGET_PARALLEL_LOOP: - case ST_OMP_TARGET_SIMD: - case ST_OMP_TARGET_TEAMS_DISTRIBUTE: - case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: - case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: - case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: - case ST_OMP_TARGET_TEAMS_LOOP: - case ST_OMP_TASKLOOP: - case ST_OMP_TASKLOOP_SIMD: - case ST_OMP_TEAMS_DISTRIBUTE: - case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: - case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: - case ST_OMP_TEAMS_DISTRIBUTE_SIMD: - case ST_OMP_TEAMS_LOOP: + case_omp_do: st = parse_omp_do (st); if (st == ST_IMPLIED_ENDDO) return st; @@ -6332,6 +6427,19 @@ parse_executable (gfc_statement st) st = parse_omp_oacc_atomic (true); continue; + case ST_OMP_METADIRECTIVE: + case ST_OMP_BEGIN_METADIRECTIVE: + st = parse_omp_metadirective_body (st); + continue; + + case ST_OMP_END_METADIRECTIVE: + if (gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE) + { + st = next_statement (); + return st; + } + /* FALLTHRU */ + default: return st; } @@ -7097,6 +7205,10 @@ gfc_parse_file (void) gfc_statement_label = NULL; + gfc_omp_region_count = 0; + gfc_in_metadirective_body = false; + gfc_matching_omp_context_selector = false; + if (setjmp (eof_buf)) return false; /* Come here on unexpected EOF */ @@ -7409,3 +7521,16 @@ is_oacc (gfc_state_data *sd) return false; } } + +/* Return true if ST is a declarative OpenMP statement. */ +bool +is_omp_declarative_stmt (gfc_statement st) +{ + switch (st) + { + case_omp_decl: + return true; + default: + return false; + } +} diff --git a/gcc/fortran/parse.h b/gcc/fortran/parse.h index ce19d4deb07..2c5b3adec3b 100644 --- a/gcc/fortran/parse.h +++ b/gcc/fortran/parse.h @@ -31,7 +31,8 @@ enum gfc_compile_state COMP_STRUCTURE, COMP_UNION, COMP_MAP, COMP_DO, COMP_SELECT, COMP_FORALL, COMP_WHERE, COMP_CONTAINS, COMP_ENUM, COMP_SELECT_TYPE, COMP_SELECT_RANK, COMP_OMP_STRUCTURED_BLOCK, COMP_CRITICAL, - COMP_DO_CONCURRENT, COMP_OMP_STRICTLY_STRUCTURED_BLOCK + COMP_DO_CONCURRENT, COMP_OMP_STRICTLY_STRUCTURED_BLOCK, + COMP_OMP_METADIRECTIVE, COMP_OMP_BEGIN_METADIRECTIVE }; /* Stack element for the current compilation state. These structures @@ -67,10 +68,15 @@ bool gfc_check_do_variable (gfc_symtree *); bool gfc_find_state (gfc_compile_state); gfc_state_data *gfc_enclosing_unit (gfc_compile_state *); const char *gfc_ascii_statement (gfc_statement, bool strip_sentinel = false) ; +gfc_statement gfc_omp_end_stmt (gfc_statement, bool = true, bool = true); match gfc_match_enum (void); match gfc_match_enumerator_def (void); void gfc_free_enum_history (void); extern bool gfc_matching_function; +extern bool gfc_matching_omp_context_selector; +extern bool gfc_in_metadirective_body; +extern int gfc_omp_region_count; + match gfc_match_prefix (gfc_typespec *); bool is_oacc (gfc_state_data *); #endif /* GFC_PARSE_H */ diff --git a/gcc/fortran/resolve.cc b/gcc/fortran/resolve.cc index 6148ed94fe1..5858b3a88dc 100644 --- a/gcc/fortran/resolve.cc +++ b/gcc/fortran/resolve.cc @@ -12261,6 +12261,11 @@ gfc_resolve_code (gfc_code *code, gfc_namespace *ns) gfc_resolve_forall (code, ns, forall_save); forall_flag = 2; } + else if (code->op == EXEC_OMP_METADIRECTIVE) + for (gfc_omp_metadirective_variant *variant + = code->ext.omp_metadirective_variants; + variant; variant = variant->next) + gfc_resolve_code (variant->code, ns); else if (code->block) { omp_workshare_save = -1; @@ -12793,6 +12798,7 @@ start: case EXEC_OMP_MASKED: case EXEC_OMP_MASKED_TASKLOOP: case EXEC_OMP_MASKED_TASKLOOP_SIMD: + case EXEC_OMP_METADIRECTIVE: case EXEC_OMP_ORDERED: case EXEC_OMP_SCAN: case EXEC_OMP_SCOPE: diff --git a/gcc/fortran/st.cc b/gcc/fortran/st.cc index 6a605ad91d4..9d6cc974083 100644 --- a/gcc/fortran/st.cc +++ b/gcc/fortran/st.cc @@ -299,6 +299,10 @@ gfc_free_statement (gfc_code *p) case EXEC_OMP_TASKYIELD: break; + case EXEC_OMP_METADIRECTIVE: + gfc_free_omp_metadirective_variants (p->ext.omp_metadirective_variants); + break; + default: gfc_internal_error ("gfc_free_statement(): Bad statement"); } diff --git a/gcc/fortran/symbol.cc b/gcc/fortran/symbol.cc index fddf68f8398..cd1de989eab 100644 --- a/gcc/fortran/symbol.cc +++ b/gcc/fortran/symbol.cc @@ -2630,10 +2630,13 @@ free_components (gfc_component *p) static int compare_st_labels (void *a1, void *b1) { - int a = ((gfc_st_label *) a1)->value; - int b = ((gfc_st_label *) b1)->value; + gfc_st_label *a = (gfc_st_label *) a1; + gfc_st_label *b = (gfc_st_label *) b1; - return (b - a); + if (a->omp_region == b->omp_region) + return b->value - a->value; + else + return b->omp_region - a->omp_region; } @@ -2683,6 +2686,7 @@ gfc_get_st_label (int labelno) { gfc_st_label *lp; gfc_namespace *ns; + int omp_region = gfc_in_metadirective_body ? gfc_omp_region_count : 0; if (gfc_current_state () == COMP_DERIVED) ns = gfc_current_block ()->f2k_derived; @@ -2699,10 +2703,16 @@ gfc_get_st_label (int labelno) lp = ns->st_labels; while (lp) { - if (lp->value == labelno) - return lp; - - if (lp->value < labelno) + if (lp->omp_region == omp_region) + { + if (lp->value == labelno) + return lp; + if (lp->value < labelno) + lp = lp->left; + else + lp = lp->right; + } + else if (lp->omp_region < omp_region) lp = lp->left; else lp = lp->right; @@ -2714,6 +2724,7 @@ gfc_get_st_label (int labelno) lp->defined = ST_LABEL_UNKNOWN; lp->referenced = ST_LABEL_UNKNOWN; lp->ns = ns; + lp->omp_region = omp_region; gfc_insert_bbt (&ns->st_labels, lp, compare_st_labels); diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc index de162f6cc75..a17f6582d81 100644 --- a/gcc/fortran/trans-decl.cc +++ b/gcc/fortran/trans-decl.cc @@ -327,7 +327,10 @@ gfc_get_label_decl (gfc_st_label * lp) gcc_assert (lp != NULL && lp->value <= MAX_LABEL_VALUE); /* Build a mangled name for the label. */ - sprintf (label_name, "__label_%.6d", lp->value); + if (lp->omp_region) + sprintf (label_name, "__label_%d_%.6d", lp->omp_region, lp->value); + else + sprintf (label_name, "__label_%.6d", lp->value); /* Build the LABEL_DECL node. */ label_decl = gfc_build_label_decl (get_identifier (label_name)); diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 439ae5f9f55..2b72862d62b 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -8228,6 +8228,8 @@ gfc_trans_omp_directive (gfc_code *code) case EXEC_OMP_MASTER_TASKLOOP: case EXEC_OMP_MASTER_TASKLOOP_SIMD: return gfc_trans_omp_master_masked_taskloop (code, code->op); + case EXEC_OMP_METADIRECTIVE: + return gfc_trans_omp_metadirective (code); case EXEC_OMP_ORDERED: return gfc_trans_omp_ordered (code); case EXEC_OMP_PARALLEL: @@ -8319,6 +8321,104 @@ gfc_trans_omp_declare_simd (gfc_namespace *ns) } } +/* Translate the context selector list GFC_SELECTORS, using WHERE as the + locus for error messages. CONV_EXPR_P is true for normal translation + and false for early conversion during parsing. */ + +static tree +gfc_trans_omp_set_selector (gfc_omp_set_selector *gfc_selectors, locus where, + bool conv_expr_p) +{ + tree set_selectors = NULL_TREE; + gfc_omp_set_selector *oss; + + for (oss = gfc_selectors; oss; oss = oss->next) + { + tree selectors = NULL_TREE; + gfc_omp_selector *os; + enum omp_tss_code set = oss->code; + gcc_assert (set != OMP_TRAIT_SET_INVALID); + + for (os = oss->trait_selectors; os; os = os->next) + { + tree scoreval = NULL_TREE; + tree properties = NULL_TREE; + gfc_omp_trait_property *otp; + enum omp_ts_code sel = os->code; + + /* Per the spec, "Implementations can ignore specified + selectors that are not those described in this section"; + however, we must record such selectors because they + cause match failures. */ + if (sel == OMP_TRAIT_INVALID) + { + selectors = make_trait_selector (sel, NULL_TREE, NULL_TREE, + selectors); + continue; + } + + for (otp = os->properties; otp; otp = otp->next) + { + switch (otp->property_kind) + { + case OMP_TRAIT_PROPERTY_EXPR: + { + tree expr = NULL_TREE; + if (conv_expr_p) + { + gfc_se se; + gfc_init_se (&se, NULL); + gfc_conv_expr (&se, otp->expr); + expr = se.expr; + } + properties = make_trait_property (NULL_TREE, expr, + properties); + } + break; + case OMP_TRAIT_PROPERTY_ID: + properties + = make_trait_property (get_identifier (otp->name), + NULL_TREE, properties); + break; + case OMP_TRAIT_PROPERTY_NAME_LIST: + { + tree prop = OMP_TP_NAMELIST_NODE; + tree value = NULL_TREE; + if (otp->is_name) + value = get_identifier (otp->name); + else + value = gfc_conv_constant_to_tree (otp->expr); + + properties = make_trait_property (prop, value, + properties); + } + break; + case OMP_TRAIT_PROPERTY_CLAUSE_LIST: + properties = gfc_trans_omp_clauses (NULL, otp->clauses, + where, true); + break; + default: + gcc_unreachable (); + } + } + + if (os->score) + { + gfc_se se; + gfc_init_se (&se, NULL); + gfc_conv_expr (&se, os->score); + scoreval = se.expr; + } + + selectors = make_trait_selector (sel, scoreval, + properties, selectors); + } + set_selectors = make_trait_set_selector (set, selectors, set_selectors); + } + return set_selectors; +} + + void gfc_trans_omp_declare_variant (gfc_namespace *ns) { @@ -8394,89 +8494,8 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns) && strcmp (odv->base_proc_symtree->name, ns->proc_name->name))) continue; - tree set_selectors = NULL_TREE; - gfc_omp_set_selector *oss; - - for (oss = odv->set_selectors; oss; oss = oss->next) - { - tree selectors = NULL_TREE; - gfc_omp_selector *os; - enum omp_tss_code set = oss->code; - gcc_assert (set != OMP_TRAIT_SET_INVALID); - - for (os = oss->trait_selectors; os; os = os->next) - { - tree scoreval = NULL_TREE; - tree properties = NULL_TREE; - gfc_omp_trait_property *otp; - enum omp_ts_code sel = os->code; - - /* Per the spec, "Implementations can ignore specified - selectors that are not those described in this section"; - however, we must record such selectors because they - cause match failures. */ - if (sel == OMP_TRAIT_INVALID) - { - selectors = make_trait_selector (sel, NULL_TREE, NULL_TREE, - selectors); - continue; - } - - for (otp = os->properties; otp; otp = otp->next) - { - switch (otp->property_kind) - { - case OMP_TRAIT_PROPERTY_EXPR: - { - gfc_se se; - gfc_init_se (&se, NULL); - gfc_conv_expr (&se, otp->expr); - properties = make_trait_property (NULL_TREE, se.expr, - properties); - } - break; - case OMP_TRAIT_PROPERTY_ID: - properties - = make_trait_property (get_identifier (otp->name), - NULL_TREE, properties); - break; - case OMP_TRAIT_PROPERTY_NAME_LIST: - { - tree prop = OMP_TP_NAMELIST_NODE; - tree value = NULL_TREE; - if (otp->is_name) - value = get_identifier (otp->name); - else - value = gfc_conv_constant_to_tree (otp->expr); - - properties = make_trait_property (prop, value, - properties); - } - break; - case OMP_TRAIT_PROPERTY_CLAUSE_LIST: - properties = gfc_trans_omp_clauses (NULL, otp->clauses, - odv->where, true); - break; - default: - gcc_unreachable (); - } - } - - if (os->score) - { - gfc_se se; - gfc_init_se (&se, NULL); - gfc_conv_expr (&se, os->score); - scoreval = se.expr; - } - - selectors = make_trait_selector (sel, scoreval, - properties, selectors); - } - set_selectors = make_trait_set_selector (set, selectors, - set_selectors); - } - + tree set_selectors = gfc_trans_omp_set_selector (odv->set_selectors, + odv->where, true); const char *variant_proc_name = odv->variant_proc_symtree->name; gfc_symbol *variant_proc_sym = odv->variant_proc_symtree->n.sym; if (variant_proc_sym == NULL || variant_proc_sym->attr.implicit_type) @@ -8577,3 +8596,57 @@ gfc_omp_call_is_alloc (tree ptr) } return build_call_expr_loc (input_location, fn, 1, ptr); } + +tree +gfc_trans_omp_metadirective (gfc_code *code) +{ + gfc_omp_metadirective_variant *variant = code->ext.omp_metadirective_variants; + + tree metadirective_tree = make_node (OMP_METADIRECTIVE); + SET_EXPR_LOCATION (metadirective_tree, gfc_get_location (&code->loc)); + TREE_TYPE (metadirective_tree) = void_type_node; + OMP_METADIRECTIVE_VARIANTS (metadirective_tree) = NULL_TREE; + + tree tree_body = NULL_TREE; + + while (variant) + { + tree ctx = gfc_trans_omp_set_selector (variant->selectors, + variant->where, true); + ctx = omp_check_context_selector (gfc_get_location (&variant->where), + ctx, true); + if (ctx == error_mark_node) + return error_mark_node; + + gfc_code *next_code = variant->code->next; + if (next_code && tree_body == NULL_TREE) + tree_body = gfc_trans_code (next_code); + + if (next_code) + variant->code->next = NULL; + tree directive = gfc_trans_code (variant->code); + if (next_code) + variant->code->next = next_code; + + tree body = next_code ? tree_body : NULL_TREE; + tree omp_variant = make_omp_metadirective_variant (ctx, directive, body); + OMP_METADIRECTIVE_VARIANTS (metadirective_tree) + = chainon (OMP_METADIRECTIVE_VARIANTS (metadirective_tree), + omp_variant); + variant = variant->next; + } + + /* TODO: Resolve the metadirective here if possible. */ + + return metadirective_tree; +} + +/* Hook to check during parsing whether a metadirective is both resolvable + and non-matching; if so, it can be eliminated immediately. */ +bool +gfc_skip_omp_metadirective_variant (gfc_omp_metadirective_variant *variant) +{ + tree selector = gfc_trans_omp_set_selector (variant->selectors, + variant->where, false); + return omp_context_selector_matches (selector, true, true) == 0; +} diff --git a/gcc/fortran/trans-stmt.h b/gcc/fortran/trans-stmt.h index 0f0f99931ca..d19e161cf11 100644 --- a/gcc/fortran/trans-stmt.h +++ b/gcc/fortran/trans-stmt.h @@ -71,6 +71,7 @@ tree gfc_trans_deallocate (gfc_code *); tree gfc_trans_omp_directive (gfc_code *); void gfc_trans_omp_declare_simd (gfc_namespace *); void gfc_trans_omp_declare_variant (gfc_namespace *); +tree gfc_trans_omp_metadirective (gfc_code *code); tree gfc_trans_oacc_directive (gfc_code *); tree gfc_trans_oacc_declare (gfc_namespace *); diff --git a/gcc/fortran/trans.cc b/gcc/fortran/trans.cc index 1a03a4a91d4..7e3089dfc8b 100644 --- a/gcc/fortran/trans.cc +++ b/gcc/fortran/trans.cc @@ -2611,6 +2611,7 @@ trans_code (gfc_code * code, tree cond) case EXEC_OMP_MASTER: case EXEC_OMP_MASTER_TASKLOOP: case EXEC_OMP_MASTER_TASKLOOP_SIMD: + case EXEC_OMP_METADIRECTIVE: case EXEC_OMP_ORDERED: case EXEC_OMP_PARALLEL: case EXEC_OMP_PARALLEL_DO: diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90 new file mode 100644 index 00000000000..c5b3946341d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90 @@ -0,0 +1,55 @@ +! { dg-do compile } + +program main + integer, parameter :: N = 10 + integer, dimension(N) :: a + integer, dimension(N) :: b + integer, dimension(N) :: c + integer :: i + + do i = 1, N + a(i) = i * 2 + b(i) = i * 3 + end do + + !$omp metadirective & + !$omp& default (teams loop) & + !$omp& default (parallel loop) ! { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective' at .1." } + do i = 1, N + c(i) = a(i) * b(i) + end do + + !$omp metadirective & + !$omp& otherwise (teams loop) & + !$omp& default (parallel loop) ! { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective' at .1." } + do i = 1, N + c(i) = a(i) * b(i) + end do + + !$omp metadirective & + !$omp& otherwise (teams loop) & + !$omp& otherwise (parallel loop) ! { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective' at .1." } + do i = 1, N + c(i) = a(i) * b(i) + end do + + !$omp metadirective default (xyz) ! { dg-error "Unclassifiable OpenMP directive at .1." } + do i = 1, N + c(i) = a(i) * b(i) + end do + + !$omp metadirective & + !$omp& default (teams loop) & ! { dg-error "expected 'when', 'otherwise', or 'default' at .1." } + !$omp& where (device={arch("nvptx")}: parallel loop) + do i = 1, N + c(i) = a(i) * b(i) + end do + + !$omp begin metadirective & + !$omp& when (device={arch("nvptx")}: parallel do) & + !$omp& default (barrier) ! { dg-error "variant directive used in OMP BEGIN METADIRECTIVE at .1. must have a corresponding end directive" } + do i = 1, N + c(i) = a(i) * b(i) + end do + !$omp end metadirective ! { dg-error "Unexpected !.OMP END METADIRECTIVE statement at .1." } +end program diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-10.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-10.f90 new file mode 100644 index 00000000000..5dad5d29eb6 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-10.f90 @@ -0,0 +1,40 @@ +! { dg-do compile } + +program metadirectives + implicit none + logical :: UseDevice + + !$OMP metadirective & + !$OMP when ( user = { condition ( UseDevice ) } & + !$OMP : parallel ) & + !$OMP default ( parallel ) + block + call bar() + end block + + !$OMP metadirective & + !$OMP when ( user = { condition ( UseDevice ) } & + !$OMP : parallel ) & + !$OMP default ( parallel ) + call bar() + !$omp end parallel ! Accepted, because all cases have 'parallel' + + !$OMP begin metadirective & + !$OMP when ( user = { condition ( UseDevice ) } & + !$OMP : nothing ) & + !$OMP default ( parallel ) + call bar() + block + call foo() + end block + !$OMP end metadirective + + !$OMP begin metadirective & + !$OMP when ( user = { condition ( UseDevice ) } & + !$OMP : parallel ) & + !$OMP default ( parallel ) + call bar() + !$omp end parallel ! { dg-error "Unexpected !.OMP END PARALLEL statement at .1." } +end program ! { dg-error "Unexpected END statement at .1." } + +! { dg-error "Unexpected end of file" "" { target *-*-* } 0 } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-11.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-11.f90 new file mode 100644 index 00000000000..e7de70e6259 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-11.f90 @@ -0,0 +1,33 @@ +! { dg-do compile } +! { dg-ice "Statements following a block in a metadirective" } +! PR fortran/107067 + +program metadirectives + implicit none + logical :: UseDevice + + !$OMP begin metadirective & + !$OMP when ( user = { condition ( UseDevice ) } & + !$OMP : nothing ) & + !$OMP default ( parallel ) + block + call foo() + end block + call bar() ! FIXME/XFAIL ICE in parse_omp_metadirective_body() + !$omp end metadirective + + + !$OMP begin metadirective & + !$OMP when ( user = { condition ( UseDevice ) } & + !$OMP : nothing ) & + !$OMP default ( parallel ) + block + call bar() + end block + block ! FIXME/XFAIL ICE in parse_omp_metadirective_body() + call foo() + end block + !$omp end metadirective +end program + + diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90 new file mode 100644 index 00000000000..cdd5e85068e --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90 @@ -0,0 +1,62 @@ +! { dg-do compile } + +program main + integer, parameter :: N = 100 + integer :: x = 0 + integer :: y = 0 + integer :: i + + ! Test implicit default directive + !$omp metadirective & + !$omp& when (device={arch("nvptx")}: barrier) + x = 1 + + ! Test implicit default directive combined with a directive that takes a + ! do loop. + !$omp metadirective & + !$omp& when (device={arch("nvptx")}: parallel do) + do i = 1, N + x = x + i + end do + + ! Test with multiple standalone directives. + !$omp metadirective & + !$omp& when (device={arch("nvptx")}: barrier) & + !$omp& default (flush) + x = 1 + + ! Test combining a standalone directive with one that takes a do loop. + !$omp metadirective & + !$omp& when (device={arch("nvptx")}: parallel do) & + !$omp& default (barrier) + do i = 1, N + x = x + i + end do + + ! Test combining a directive that takes a do loop with one that takes + ! a statement body. + !$omp begin metadirective & + !$omp& when (device={arch("nvptx")}: parallel do) & + !$omp& default (parallel) + do i = 1, N + x = x + i + end do + !$omp end metadirective + + ! Test labels in the body. + !$omp begin metadirective & + !$omp& when (device={arch("nvptx")}: parallel do) & + !$omp& when (device={arch("gcn")}: parallel) + do i = 1, N + x = x + i + if (x .gt. N/2) goto 10 +10 x = x + 1 + goto 20 + x = x + 2 +20 continue + end do + !$omp end metadirective + + ! Test empty metadirective. + !$omp metadirective +end program diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90 new file mode 100644 index 00000000000..eca389a7842 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90 @@ -0,0 +1,34 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original" } +! { dg-additional-options "-fdump-tree-gimple" } +! { dg-additional-options "-fdump-tree-optimized" } + +module test + integer, parameter :: N = 100 +contains + subroutine f (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: v1, v2) map(from: v3) + !$omp metadirective & + !$omp& when(device={arch("nvptx")}: teams loop) & + !$omp& default(parallel loop) + do i = 1, N + z(i) = x(i) * y(i) + enddo + !$omp end target + end subroutine +end module + +! The metadirective should be resolved after Gimplification. + +! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } +! { dg-final { scan-tree-dump-times "when \\(device =.*arch.*nvptx.*\\):" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } +! { dg-final { scan-tree-dump-times "default:" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp loop" 2 "original" } } + +! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "gimple" } } + +! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "optimized" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 new file mode 100644 index 00000000000..2fdc7f3c515 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 @@ -0,0 +1,39 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original" } +! { dg-additional-options "-fdump-tree-gimple" } + +program test + implicit none + integer, parameter :: N = 100 + real :: a(N) + + !$omp target map(from: a) + call f (a, 3.14159) + !$omp end target + + ! TODO: This does not execute a version of f with the default clause + ! active as might be expected. + call f (a, 2.71828) ! { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } +contains + subroutine f (a, x) + integer :: i + real :: a(N), x + !$omp declare target + + !$omp metadirective & + !$omp& when (construct={target}: distribute parallel do ) & + !$omp& default(parallel do simd) + do i = 1, N + a(i) = x * i + end do + end subroutine +end program + +! The metadirective should be resolved during Gimplification. + +! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } +! { dg-final { scan-tree-dump-times "when \\(construct = .*target.*\\):" 1 "original" } } +! { dg-final { scan-tree-dump-times "default:" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } + +! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90 new file mode 100644 index 00000000000..03970393eb4 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90 @@ -0,0 +1,30 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module test + integer, parameter :: N = 100 +contains + subroutine f (a, flag) + integer :: a(N) + logical :: flag + integer :: i + + !$omp metadirective & + !$omp& when (user={condition(flag)}: & + !$omp& target teams distribute parallel do map(from: a(1:N))) & + !$omp& default(parallel do) + do i = 1, N + a(i) = i + end do + end subroutine +end module + +! The metadirective should be resolved at parse time, but is currently +! resolved during Gimplification + +! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp target" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp distribute" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp for" 2 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90 new file mode 100644 index 00000000000..9b6c371296f --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90 @@ -0,0 +1,31 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module test + integer, parameter :: N = 100 +contains + subroutine f (a, run_parallel, run_guided) + integer :: a(N) + logical :: run_parallel, run_guided + integer :: i + + !$omp begin metadirective when(user={condition(run_parallel)}: parallel) + !$omp metadirective & + !$omp& when(construct={parallel}, user={condition(run_guided)}: & + !$omp& do schedule(guided)) & + !$omp& when(construct={parallel}: do schedule(static)) + do i = 1, N + a(i) = i + end do + !$omp end metadirective + end subroutine +end module + +! The outer metadirective should be resolved at parse time, but is +! currently resolved during Gimplification. + +! The inner metadirective should be resolved during Gimplificiation. + +! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp for" 2 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-7.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-7.f90 new file mode 100644 index 00000000000..870ea192fbc --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-7.f90 @@ -0,0 +1,36 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +program main + integer, parameter :: N = 256 +contains + subroutine f (a, num) + integer :: a(N) + integer :: num + integer :: i + + !$omp metadirective & + !$omp& when (target_device={device_num(num), kind("gpu"), arch("nvptx")}: & + !$omp& target parallel do map(tofrom: a(1:N))) & + !$omp& when (target_device={device_num(num), kind("gpu"), & + !$omp& arch("amdgcn"), isa("gfx906")}: & + !$omp& target parallel do) & + !$omp& when (target_device={device_num(num), kind("cpu"), arch("x86_64")}: & + !$omp& parallel do) + do i = 1, N + a(i) = a(i) + i + end do + + !$omp metadirective & + !$omp& when (target_device={kind("gpu"), arch("nvptx")}: & + !$omp& target parallel do map(tofrom: a(1:N))) + do i = 1, N + a(i) = a(i) + i + end do + end subroutine +end program + +! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(.+, &\"gpu\"\\\[0\\\], &\"amdgcn\"\\\[0\\\], &\"gfx906\"\\\[0\\\]\\)" "gimple" } } +! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(.+, &\"gpu\"\\\[0\\\], &\"nvptx\"\\\[0\\\], 0B\\)" "gimple" } } +! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(.+, &\"cpu\"\\\[0\\\], &\"x86_64\"\\\[0\\\], 0B\\)" "gimple" } } +! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(-1, &\"gpu\"\\\[0\\\], &\"nvptx\"\\\[0\\\], 0B\\)" "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-8.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-8.f90 new file mode 100644 index 00000000000..3f46801e044 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-8.f90 @@ -0,0 +1,22 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original" } + +program test + integer :: i + integer, parameter :: N = 100 + integer :: sum = 0 + + ! The compiler should never consider a situation where both metadirectives + ! match. If it does, then the nested metadirective would be an error + ! as it is not a loop-nest as per the OpenMP specification. + + !$omp metadirective when (implementation={vendor("ibm")}: & + !$omp& target teams distribute) + !$omp metadirective when (implementation={vendor("gnu")}: parallel do) + do i = 1, N + sum = sum + i + end do +end program + +! { dg-final { scan-tree-dump-not "when \\(implementation = .*vendor.*ibm.*\\):" "original" } } +! { dg-final { scan-tree-dump-times "when \\(implementation = .*vendor.*gnu.*\\):" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-9.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-9.f90 new file mode 100644 index 00000000000..e6ab3fc0a65 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-9.f90 @@ -0,0 +1,30 @@ +! { dg-do compile } + +program OpenMP_Metadirective_WrongEnd_Test + implicit none + + integer :: & + iaVS, iV, jV, kV + integer, dimension ( 3 ) :: & + lV, uV + logical :: & + UseDevice + + !$OMP metadirective & + !$OMP when ( user = { condition ( UseDevice ) } & + !$OMP : target teams distribute parallel do simd collapse ( 3 ) & + !$OMP private ( iaVS ) ) & + !$OMP default ( parallel do simd collapse ( 3 ) private ( iaVS ) ) + do kV = lV ( 3 ), uV ( 3 ) + do jV = lV ( 2 ), uV ( 2 ) + do iV = lV ( 1 ), uV ( 1 ) + + + end do + end do + end do + !$OMP end target teams distribute parallel do simd ! { dg-error "Unexpected !.OMP END TARGET TEAMS DISTRIBUTE PARALLEL DO SIMD statement in an OMP METADIRECTIVE block at .1." } + + +end program + diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-construct.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-construct.f90 new file mode 100644 index 00000000000..ec1f0ee3d9d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-construct.f90 @@ -0,0 +1,260 @@ +! { dg-do compile } +! { dg-additional-options "-foffload=disable -fdump-tree-original -fdump-tree-gimple" } + +program main +implicit none + +integer, parameter :: N = 10 +double precision, parameter :: S = 2.0 +double precision :: a(N) + +call init (N, a) +call f1 (N, a, S) +call check (N, a, S) + +call init (N, a) +call f2 (N, a, S) +call check (N, a, S) + +call init (N, a) +call f3 (N, a, S) +call check (N, a, S) + +call init (N, a) +call f4 (N, a, S) +call check (N, a, S) + +call init (N, a) +call f5 (N, a, S) +call check (N, a, S) + +call init (N, a) +call f6 (N, a, S) +call check (N, a, S) + +call init (N, a) +call f7 (N, a, S) +call check (N, a, S) + +call init (N, a) +call f8 (N, a, S) +call check (N, a, S) + +call init (N, a) +call f9 (N, a, S) +call check (N, a, S) + +contains + +subroutine init (n, a) + implicit none + integer :: n + double precision :: a(n) + integer :: i + do i = 1, n + a(i) = i + end do +end subroutine + +subroutine check (n, a, s) + implicit none + integer :: n + double precision :: a(n) + double precision :: s + integer :: i + do i = 1, n + if (a(i) /= i * s) error stop + end do +end subroutine + +! Check various combinations for enforcing correct ordering of +! construct matches. +subroutine f1 (n, a, s) + implicit none + integer :: n + double precision :: a(n) + double precision :: s + integer :: i +!$omp target teams +!$omp parallel +!$omp metadirective & +!$omp & when (construct={target} & +!$omp & : do) & +!$omp & default (error at(execution) message("f1 match failed")) + do i = 1, n + a(i) = a(i) * s + end do +!$omp end parallel +!$omp end target teams +end subroutine + +subroutine f2 (n, a, s) + implicit none + integer :: n + double precision :: a(n) + double precision :: s + integer :: i +!$omp target teams +!$omp parallel +!$omp metadirective & +!$omp & when (construct={teams, parallel} & +!$omp & : do) & +!$omp & default (error at(execution) message("f2 match failed")) + do i = 1, n + a(i) = a(i) * s + end do +!$omp end parallel +!$omp end target teams +end subroutine + +subroutine f3 (n, a, s) + implicit none + integer :: n + double precision :: a(n) + double precision :: s + integer :: i +!$omp target teams +!$omp parallel +!$omp metadirective & +!$omp & when (construct={target, teams, parallel} & +!$omp & : do) & +!$omp & default (error at(execution) message("f3 match failed")) + do i = 1, n + a(i) = a(i) * s + end do +!$omp end parallel +!$omp end target teams +end subroutine + +subroutine f4 (n, a, s) + implicit none + integer :: n + double precision :: a(n) + double precision :: s + integer :: i +!$omp target teams +!$omp parallel +!$omp metadirective & +!$omp & when (construct={target, parallel} & +!$omp & : do) & +!$omp & default (error at(execution) message("f4 match failed")) + do i = 1, n + a(i) = a(i) * s + end do +!$omp end parallel +!$omp end target teams +end subroutine + +subroutine f5 (n, a, s) + implicit none + integer :: n + double precision :: a(n) + double precision :: s + integer :: i +!$omp target teams +!$omp parallel +!$omp metadirective & +!$omp & when (construct={target, teams} & +!$omp & : do) & +!$omp & default (error at(execution) message("f5 match failed")) + do i = 1, n + a(i) = a(i) * s + end do +!$omp end parallel +!$omp end target teams +end subroutine + +! Next batch is for things where the construct doesn't match the context. +subroutine f6 (n, a, s) + implicit none + integer :: n + double precision :: a(n) + double precision :: s + integer :: i +!$omp target +!$omp teams +!$omp metadirective & +!$omp & when (construct={parallel} & +!$omp & : error at(execution) message("f6 match failed")) & +!$omp & default (parallel do) + do i = 1, n + a(i) = a(i) * s + end do +!$omp end teams +!$omp end target +end subroutine + +subroutine f7 (n, a, s) + implicit none + integer :: n + double precision :: a(n) + double precision :: s + integer :: i +!$omp target +!$omp teams +!$omp metadirective & +!$omp & when (construct={target, parallel} & +!$omp & : error at(execution) message("f7 match failed")) & +!$omp & default (parallel do) + do i = 1, n + a(i) = a(i) * s + end do +!$omp end teams +!$omp end target +end subroutine + +subroutine f8 (n, a, s) + implicit none + integer :: n + double precision :: a(n) + double precision :: s + integer :: i +!$omp target +!$omp teams +!$omp metadirective & +!$omp & when (construct={parallel, target} & +!$omp & : error at(execution) message("f8 match failed")) & +!$omp & default (parallel do) + do i = 1, n + a(i) = a(i) * s + end do +!$omp end teams +!$omp end target +end subroutine + +! Next test choosing the best alternative when there are multiple +! matches. +subroutine f9 (n, a, s) + implicit none + integer :: n + double precision :: a(n) + double precision :: s + integer :: i +!$omp target teams +!$omp parallel +!$omp metadirective & +!$omp & when (construct={teams, parallel} & +!$omp & : error at(execution) message("f9 match incorrect 1")) & +!$omp & when (construct={target, teams, parallel} & +!$omp & : do) & +!$omp & when (construct={target, teams} & +!$omp & : error at(execution) message("f9 match incorrect 2")) & +!$omp & default (error at(execution) message("f9 match failed")) + do i = 1, n + a(i) = a(i) * s + end do +!$omp end parallel +!$omp end target teams +end subroutine + +end program + +! Note there are no tests for the matching the extended simd clause +! syntax, which is only useful for "declare variant". + + +! After parsing, there should be a runtime error call for each of the +! failure cases, but they should all be optimized away during OMP +! lowering. +! { dg-final { scan-tree-dump-times "__builtin_GOMP_error" 11 "original" } } +! { dg-final { scan-tree-dump-not "__builtin_GOMP_error" "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-no-score.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-no-score.f90 new file mode 100644 index 00000000000..968ce609b10 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-no-score.f90 @@ -0,0 +1,122 @@ +! { dg-do compile { target x86_64-*-* } } +! { dg-additional-options "-foffload=disable" } + +! This test is expected to fail with compile-time errors: +! "A trait-score cannot be specified in traits from the construct, +! device or target_device trait-selector-sets." + + +subroutine f1 (n, a, s) + implicit none + integer :: n + double precision :: a(*) + double precision :: s + integer :: i +!$omp metadirective & +!$omp& when (device={kind (score(5) : host)} & +!$omp& : parallel do) + ! { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-2 } + do i = 1, n + a(i) = a(i) * s; + end do +end subroutine + +subroutine f2 (n, a, s) + implicit none + integer :: n + double precision :: a(*) + double precision :: s + integer :: i +!$omp metadirective & +!$omp& when (device={kind (host), arch (score(6) : x86_64), isa (avx512f)} & +!$omp& : parallel do) + ! { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-2 } + do i = 1, n + a(i) = a(i) * s; + end do +end subroutine + +subroutine f3 (n, a, s) + implicit none + integer :: n + double precision :: a(*) + double precision :: s + integer :: i +!$omp metadirective & +!$omp& when (device={kind (host), arch (score(6) : x86_64), & +!$omp& isa (score(7): avx512f)} & +!$omp& : parallel do) + ! { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-3 } + do i = 1, n + a(i) = a(i) * s; + end do +end subroutine + +subroutine f4 (n, a, s) + implicit none + integer :: n + double precision :: a(*) + double precision :: s + integer :: i + integer, parameter :: omp_initial_device = -1 +!$omp metadirective & +!$omp& when (target_device={device_num (score(42) : omp_initial_device), & +!$omp& kind (host)} & +!$omp& : parallel do) + ! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 } + do i = 1, n + a(i) = a(i) * s; + end do +end subroutine + +subroutine f5 (n, a, s) + implicit none + integer :: n + double precision :: a(*) + double precision :: s + integer :: i + integer, parameter :: omp_initial_device = -1 +!$omp metadirective & +!$omp& when (target_device={device_num(omp_initial_device), & +!$omp& kind (score(5) : host)} & +!$omp& : parallel do) + ! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-2 } + do i = 1, n + a(i) = a(i) * s; + end do +end subroutine + +subroutine f6 (n, a, s) + implicit none + integer :: n + double precision :: a(*) + double precision :: s + integer :: i + integer, parameter :: omp_initial_device = -1 +!$omp metadirective & +!$omp& when (target_device={device_num(omp_initial_device), kind (host), & +!$omp& arch (score(6) : x86_64), isa (avx512f)} & +!$omp& : parallel do) + ! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-2 } + do i = 1, n + a(i) = a(i) * s; + end do +end subroutine + +subroutine f7 (n, a, s) + implicit none + integer :: n + double precision :: a(*) + double precision :: s + integer :: i + integer, parameter :: omp_initial_device = -1 +!$omp metadirective & +!$omp& when (target_device={device_num(omp_initial_device), kind (host), & +!$omp& arch (score(6) : x86_64), & +!$omp& isa (score(7): avx512f)} & +!$omp& : parallel do) + ! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 } + do i = 1, n + a(i) = a(i) * s; + end do +end subroutine diff --git a/gcc/testsuite/gfortran.dg/gomp/pure-1.f90 b/gcc/testsuite/gfortran.dg/gomp/pure-1.f90 index 598e455d2e9..15681ac604a 100644 --- a/gcc/testsuite/gfortran.dg/gomp/pure-1.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/pure-1.f90 @@ -86,3 +86,10 @@ pure integer function func_simd(n) end do func_simd = r end + +!pure logical function func_metadirective() +logical function func_metadirective() + implicit none + !$omp metadirective + func_metadirective = .false. +end diff --git a/gcc/testsuite/gfortran.dg/gomp/pure-2.f90 b/gcc/testsuite/gfortran.dg/gomp/pure-2.f90 index 1e3cf8c9416..25d6069b236 100644 --- a/gcc/testsuite/gfortran.dg/gomp/pure-2.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/pure-2.f90 @@ -26,14 +26,6 @@ logical function func_interchange(n) end do end - -!pure logical function func_metadirective() -logical function func_metadirective() - implicit none - !$omp metadirective ! { dg-error "Unclassifiable OpenMP directive" } - func_metadirective = .false. -end - !pure logical function func_reverse(n) logical function func_reverse(n) implicit none diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 new file mode 100644 index 00000000000..7b3e09f7c2a --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 @@ -0,0 +1,61 @@ +! { dg-do run } + +program test + implicit none + + integer, parameter :: N = 100 + integer :: x(N), y(N), z(N) + integer :: i + + do i = 1, N + x(i) = i; + y(i) = -i; + end do + + call f (x, y, z) + + do i = 1, N + if (z(i) .ne. x(i) * y(i)) stop 1 + end do + + ! ----- + do i = 1, N + x(i) = i; + y(i) = -i; + end do + + call g (x, y, z) + + do i = 1, N + if (z(i) .ne. x(i) * y(i)) stop 1 + end do + +contains + subroutine f (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: x, y) map(from: z) + block + !$omp metadirective & + !$omp& when(device={arch("nvptx")}: teams loop) & + !$omp& default(parallel loop) + do i = 1, N + z(i) = x(i) * y(i) + enddo + end block + end subroutine + subroutine g (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: x, y) map(from: z) + block + !$omp metadirective & + !$omp& when(device={arch("nvptx")}: teams loop) & + !$omp& default(parallel loop) + do i = 1, N + z(i) = x(i) * y(i) + enddo + end block + !$omp end target + end subroutine +end program diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 new file mode 100644 index 00000000000..d83474cf2db --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 @@ -0,0 +1,40 @@ +! { dg-do run } + +program test + implicit none + integer, parameter :: N = 100 + real, parameter :: PI_CONST = 3.14159 + real, parameter :: E_CONST = 2.71828 + real, parameter :: EPSILON = 0.001 + integer :: i + real :: a(N) + + !$omp target map(from: a) + call f (a, PI_CONST) + !$omp end target + + do i = 1, N + if (abs (a(i) - (PI_CONST * i)) .gt. EPSILON) stop 1 + end do + + ! TODO: This does not execute a version of f with the default clause + ! active as might be expected. + call f (a, E_CONST) ! { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } + + do i = 1, N + if (abs (a(i) - (E_CONST * i)) .gt. EPSILON) stop 2 + end do +contains + subroutine f (a, x) + integer :: i + real :: a(N), x + !$omp declare target + + !$omp metadirective & + !$omp& when (construct={target}: distribute parallel do ) & + !$omp& default(parallel do simd) + do i = 1, N + a(i) = x * i + end do + end subroutine +end program diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-3.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-3.f90 new file mode 100644 index 00000000000..693c40bca5a --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/metadirective-3.f90 @@ -0,0 +1,29 @@ +! { dg-do run } + +program test + implicit none + + integer, parameter :: N = 100 + integer :: a(N) + integer :: res + + if (f (a, .false.)) stop 1 + if (.not. f (a, .true.)) stop 2 +contains + logical function f (a, flag) + integer :: a(N) + logical :: flag + logical :: res = .false. + integer :: i + f = .false. + !$omp metadirective & + !$omp& when (user={condition(.not. flag)}: & + !$omp& target teams distribute parallel do & + !$omp& map(from: a(1:N)) private(res)) & + !$omp& default(parallel do) + do i = 1, N + a(i) = i + f = .true. + end do + end function +end program diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-4.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-4.f90 new file mode 100644 index 00000000000..04fdf61489c --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/metadirective-4.f90 @@ -0,0 +1,46 @@ +! { dg-do run } + +program test + use omp_lib + + implicit none + integer, parameter :: N = 100 + integer :: a(N) + logical :: is_parallel, is_static + + ! is_static is always set if run_parallel is false. + call f (a, .false., .false., is_parallel, is_static) + if (is_parallel .or. .not. is_static) stop 1 + + call f (a, .false., .true., is_parallel, is_static) + if (is_parallel .or. .not. is_static) stop 2 + + call f (a, .true., .false., is_parallel, is_static) + if (.not. is_parallel .or. is_static) stop 3 + + call f (a, .true., .true., is_parallel, is_static) + if (.not. is_parallel .or. .not. is_static) stop 4 +contains + subroutine f (a, run_parallel, run_static, is_parallel, is_static) + integer :: a(N) + logical, intent(in) :: run_parallel, run_static + logical, intent(out) :: is_parallel, is_static + integer :: i + + is_parallel = .false. + is_static = .false. + + !$omp begin metadirective when(user={condition(run_parallel)}: parallel) + if (omp_in_parallel ()) is_parallel = .true. + + !$omp metadirective & + !$omp& when(construct={parallel}, user={condition(.not. run_static)}: & + !$omp& do schedule(guided) private(is_static)) & + !$omp& when(construct={parallel}: do schedule(static)) + do i = 1, N + a(i) = i + is_static = .true. + end do + !$omp end metadirective + end subroutine +end program diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-5.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-5.f90 new file mode 100644 index 00000000000..3992286dc08 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/metadirective-5.f90 @@ -0,0 +1,44 @@ +! { dg-do run } + +program main + use omp_lib + + implicit none + + integer, parameter :: N = 100 + integer :: a(N) + integer :: on_device_count = 0 + integer :: i + + do i = 1, N + a(i) = i + end do + + do i = 0, omp_get_num_devices () + on_device_count = on_device_count + f (a, i) + end do + + if (on_device_count .ne. omp_get_num_devices ()) stop 1 + + do i = 1, N + if (a(i) .ne. 2 * i) stop 2; + end do +contains + integer function f (a, num) + integer, intent(inout) :: a(N) + integer, intent(in) :: num + integer :: on_device + integer :: i + + on_device = 0 + !$omp metadirective & + !$omp& when (target_device={device_num(num), kind("gpu")}: & + !$omp& target parallel do map(to: a(1:N)), map(from: on_device)) & + !$omp& default (parallel do private(on_device)) + do i = 1, N + a(i) = a(i) + i + on_device = 1 + end do + f = on_device; + end function +end program diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-6.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-6.f90 new file mode 100644 index 00000000000..436fdbade2f --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/metadirective-6.f90 @@ -0,0 +1,58 @@ +! { dg-do compile } + +program test + implicit none + + integer, parameter :: N = 100 + integer :: x(N), y(N), z(N) + integer :: i + +contains + subroutine f (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: x, y) map(from: z) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp metadirective & + !$omp& when(device={arch("nvptx")}: teams loop) & + !$omp& default(parallel loop) ! { dg-error "\\(1\\)" } + ! FIXME: The line above should be the same error as above but some fails here with -fno-diagnostics-show-caret + ! Seems as if some gcc/testsuite/ fix is missing for libgomp/testsuite + do i = 1, N + z(i) = x(i) * y(i) + enddo + z(N) = z(N) + 1 ! <<< invalid + end block + end subroutine + + subroutine f2 (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: x, y) map(from: z) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + integer :: i ! << invalid + !$omp metadirective & + !$omp& when(device={arch("nvptx")}: teams loop) & + !$omp& default(parallel loop) + do i = 1, N + z(i) = x(i) * y(i) + enddo + end block + end subroutine + subroutine g (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: x, y) map(from: z) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp metadirective & ! <<<< invalid + !$omp& when(device={arch("nvptx")}: flush) & + !$omp& default(nothing) + !$omp teams loop + do i = 1, N + z(i) = x(i) * y(i) + enddo + end block + !$omp end target + end subroutine + +end program