From patchwork Sat Jul 20 20:42:26 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Sandra Loosemore X-Patchwork-Id: 1962839 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; dkim=pass (2048-bit key; unprotected) header.d=baylibre-com.20230601.gappssmtp.com header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256 header.s=20230601 header.b=t65t1C+A; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4WRJVT6jMSz1yZ7 for ; Sun, 21 Jul 2024 06:46:25 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 20964385DDF0 for ; Sat, 20 Jul 2024 20:46:24 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-il1-x12e.google.com (mail-il1-x12e.google.com [IPv6:2607:f8b0:4864:20::12e]) by sourceware.org (Postfix) with ESMTPS id 4B535385828B for ; Sat, 20 Jul 2024 20:42:53 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 4B535385828B Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=baylibre.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=baylibre.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 4B535385828B Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2607:f8b0:4864:20::12e ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1721508182; cv=none; b=c9sc0+zH24nKJYfZA4I7xai1/1HjBqT7BNtiBoq7MCu0oDp8bdGgYtTNKBnvbbhCh26Yg1QZiZBfQ8L/fRU1wNG+8dJEsYgKSkJjLnuOeZq/W8gktkGrDR78dEjwAgM5OLAg6prwUT7CikuaYeO9kpBLIPvtHauBLFhGijpAv+c= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1721508182; c=relaxed/simple; bh=4x36MC5ARL1QtBTzAQ1roQOW933fix5i3jOeCiKbugc=; h=DKIM-Signature:From:To:Subject:Date:Message-Id:MIME-Version; b=qSpFIjv9hLjR8tpZeR1QSFZTMbRvgCDs4f/m5KQ4itXZP+adiw+MsmvK8DoQ60QFDckGTY26X0P+4apWoUcW/sq053nSi0gQAt3Zv9X7s5YQuGfqBP9h3kIIVk00RJmxC+WzVFq2VvwbCN81xp3FzwCJ1kFcwoAv33VvbmvDJRg= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-il1-x12e.google.com with SMTP id e9e14a558f8ab-39641271f2aso11530385ab.3 for ; Sat, 20 Jul 2024 13:42:53 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1721508172; x=1722112972; darn=gcc.gnu.org; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to; bh=nk/Owtb8DyfjT/rXUN4XMpsT3PxVwEWvqdCUu2V/OGo=; b=t65t1C+AMas19wJqVRYsOfDULVzYeHHsY6CYMuxEcX12BeJtbB6hWruyc3uqpH/GQF i47aK40YWt7A3wGIAlseUspdAHFwZaleGorFp9uW7bcS5DvF2CNcMVN9QIP1tVPp2EUM ifXwV6QJBs0Uif4D0f34Te2AwNs+AnRIdiVqd9i+MUdUwVb/UkNEhHjwaZOaM+MMKF3J nvc6g5UTHUD3uJI8PJM9ua6jwpVH8JXc/HoO2wETZsSkSMgSUrkS65jmv49nPvPPLu9W lJJgBwABAad1WzC5N4vOO3dzTvAu6mox4A9IvzeqLnMXE5qR2MKbIdt4b5D15F4seKIY Vv+Q== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1721508172; x=1722112972; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:x-gm-message-state:from:to:cc :subject:date:message-id:reply-to; bh=nk/Owtb8DyfjT/rXUN4XMpsT3PxVwEWvqdCUu2V/OGo=; b=kjWjxgfXoUbpIZbl4SFE52rDwKkMZZGUNuys8FpjIJuBIfgfAQIyIDe1rRlCLgYrK2 +BeOb1/S3UjUcwfYSt9vxlDvRFXwsUwJEqk8kvj2PGQKuyQxTXWstuIFGi0SAhadisCv 6i3EYZcEcMb9QeF9SmGUlQXNvu61cGswSZPZRKMA9N5WBEH9uNUM4r9Fcae5xVE/eOiq 44wtnmY9BQPi9k0KlaFb0VgPTqn1J4AtTxra1Yx7UoWG0v4EigjZzzPEblrv9rrCaR4W bNrOBaX6EEPeKhBgpFITFGNE7g3C0fIrluWqfKAi9G0pgX3mEzIeiKGgYFZgxNRCizVt W/aQ== X-Gm-Message-State: AOJu0YzYW5oRJHHOUSjhaXWERkVVYrkAsKfnZJ0mityRduOpSlxB19gX GkoY8Zi9KDYqETzqCLLKErH9vP3t/Lfjt6WznsO7WkRwM3lYLNLI7qw5g+jE4np2O96wjl4Tjeg T3kQ= X-Google-Smtp-Source: AGHT+IG7jA/TP/lSfRkmLhid93YhNBVoX2PYnIMv7apLmfpOw97dtRYy1CUdOVrfhA2A0r9y3DWAAw== X-Received: by 2002:a92:cd8c:0:b0:399:3203:397 with SMTP id e9e14a558f8ab-3993ff97b14mr33655325ab.5.1721508170327; Sat, 20 Jul 2024 13:42:50 -0700 (PDT) Received: from pondscum.hsd1.co.comcast.net ([2601:281:d901:5620:ee50:43b7:d7db:ddf2]) by smtp.gmail.com with ESMTPSA id e9e14a558f8ab-397f5c2b74esm13133105ab.35.2024.07.20.13.42.48 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Sat, 20 Jul 2024 13:42:49 -0700 (PDT) From: Sandra Loosemore To: gcc-patches@gcc.gnu.org Cc: jakub@redhat.com, tburnus@baylibre.com Subject: [PATCH v3 07/12] OpenMP: Fortran front-end support for metadirectives. Date: Sat, 20 Jul 2024 14:42:26 -0600 Message-Id: <20240720204231.2229891-8-sloosemore@baylibre.com> X-Mailer: git-send-email 2.34.1 In-Reply-To: <20240720204231.2229891-1-sloosemore@baylibre.com> References: <20240720204231.2229891-1-sloosemore@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.1 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP 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~incoming=patchwork.ozlabs.org@gcc.gnu.org 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_variant): New. (struct gfc_st_label): Add omp_region field. (gfc_exec_op): Add EXEC_OMP_METADIRECTIVE. (struct gfc_code): Add omp_variants field. (gfc_free_omp_variants): Declare. (match_omp_directive): Declare. (is_omp_declarative_stmt): 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_variants): New. (gfc_match_omp_clauses): Remove context_selector parameter. (match_omp): Adjust call to gfc_match_omp_clauses. (gfc_match_omp_context_selector): Add metadirective_p parameter. Adjust error-checking logic and calls to gfc_match_omp_clauses. Set gfc_matching_omp_context_selector. (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 (enum gfc_compile_state): Add metadirective constructs. (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. * 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: Kwok Cheung Yeung 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 | 21 +- gcc/fortran/io.cc | 2 +- gcc/fortran/match.h | 2 + gcc/fortran/openmp.cc | 294 +++++++-- gcc/fortran/parse.cc | 578 +++++++++++------- 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 | 233 ++++--- 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, 1952 insertions(+), 365 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 b8308aeee55..b6633a913e7 100644 --- a/gcc/fortran/decl.cc +++ b/gcc/fortran/decl.cc @@ -8422,6 +8422,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; @@ -8429,6 +8430,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; } @@ -8574,6 +8597,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 80aa8ef84e7..17dda3c0fdf 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -2193,6 +2193,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; @@ -2394,6 +2395,25 @@ show_omp_node (int level, gfc_code *c) d = d->block; } } + else if (c->op == EXEC_OMP_METADIRECTIVE) + { + gfc_omp_variant *variant + = c->ext.omp_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) @@ -3537,6 +3557,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 3bdf18d6f9b..7341088af3e 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, @@ -1593,7 +1594,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, indirect:1; unsigned full:1, erroneous:1; ENUM_BITFIELD (gfc_omp_sched_kind) sched_kind:3; @@ -1714,6 +1715,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_variant +{ + struct gfc_omp_variant *next; + locus where; /* Where the metadirective clause occurred. */ + + gfc_omp_set_selector *selectors; + enum gfc_statement stmt; + struct gfc_code *code; + +} gfc_omp_variant; +#define gfc_get_omp_variant() XCNEW (gfc_omp_variant) typedef struct gfc_omp_udr { @@ -1762,6 +1774,7 @@ typedef struct gfc_st_label locus where; gfc_namespace *ns; + int omp_region; } gfc_st_label; @@ -3040,7 +3053,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_UNROLL, EXEC_OMP_TILE, + EXEC_OMP_UNROLL, EXEC_OMP_TILE, EXEC_OMP_METADIRECTIVE, EXEC_OMP_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS }; @@ -3099,6 +3112,7 @@ typedef struct gfc_code gfc_omp_clauses *omp_clauses; const char *omp_name; gfc_omp_namelist *omp_namelist; + gfc_omp_variant *omp_variants; bool omp_bool; } ext; /* Points to additional structures required by statement */ @@ -3699,6 +3713,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_variants (gfc_omp_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 *); @@ -3986,6 +4001,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); 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 c2b7d69c37c..04058458f8c 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 333f0c7fe7f..1aafbf5e511 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -114,7 +114,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) @@ -125,17 +126,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; @@ -342,6 +351,19 @@ gfc_free_omp_udr (gfc_omp_udr *omp_udr) } } +/* Free variants of an !$omp metadirective construct. */ + +void +gfc_free_omp_variants (gfc_omp_variant *variant) +{ + while (variant) + { + gfc_omp_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) @@ -1876,8 +1898,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 (); @@ -3887,9 +3908,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"); @@ -4596,7 +4615,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; @@ -5777,7 +5796,8 @@ gfc_ignore_trait_property_extension_list (void) score(score-expression) */ match -gfc_match_omp_context_selector (gfc_omp_set_selector *oss) +gfc_match_omp_context_selector (gfc_omp_set_selector *oss, + bool metadirective_p) { do { @@ -5937,14 +5957,27 @@ gfc_match_omp_context_selector (gfc_omp_set_selector *oss) || (property_kind == OMP_TRAIT_PROPERTY_DEV_NUM_EXPR && otp->expr->ts.type != BT_INTEGER) || otp->expr->rank != 0 - || otp->expr->expr_type != EXPR_CONSTANT) + || (!metadirective_p + && otp->expr->expr_type != EXPR_CONSTANT)) { - if (property_kind == OMP_TRAIT_PROPERTY_BOOL_EXPR) - gfc_error ("property must be a constant logical expression " - "at %C"); + if (metadirective_p) + { + if (property_kind == OMP_TRAIT_PROPERTY_BOOL_EXPR) + gfc_error ("property must be a " + "logical expression at %C"); + else + gfc_error ("property must be an " + "integer expression at %C"); + } else - gfc_error ("property must be a constant integer expression " - "at %C"); + { + if (property_kind == OMP_TRAIT_PROPERTY_BOOL_EXPR) + gfc_error ("property must be a constant " + "logical expression at %C"); + else + gfc_error ("property must be a constant " + "integer expression at %C"); + } return MATCH_ERROR; } /* Device number must be conforming, which includes @@ -5964,14 +5997,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) { @@ -6028,7 +6064,8 @@ 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, + bool metadirective_p) { do { @@ -6061,11 +6098,11 @@ 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) + if (gfc_match_omp_context_selector (oss, metadirective_p) != MATCH_YES) return MATCH_ERROR; m = gfc_match (" }"); @@ -6164,7 +6201,9 @@ 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, + false) + != MATCH_YES) return MATCH_ERROR; if (gfc_match (" )") != MATCH_YES) @@ -6180,6 +6219,154 @@ gfc_match_omp_declare_variant (void) } +static match +match_omp_metadirective (bool begin_p) +{ + locus old_loc = gfc_current_locus; + gfc_omp_variant *variants_head; + gfc_omp_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, true) + != 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_variant *omv = gfc_get_omp_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; + } + + *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_variant *omv = gfc_get_omp_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_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) { @@ -11213,6 +11400,19 @@ resolve_omp_do (gfc_code *code) non_generated_count); } +static void +resolve_omp_metadirective (gfc_code *code, gfc_namespace *ns) +{ + gfc_omp_variant *variant = code->ext.omp_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) @@ -11760,13 +11960,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_variant *mc + = c->ext.omp_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) @@ -11895,6 +12114,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 b28c8a94547..1102d78ef43 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); @@ -1049,6 +1053,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, @@ -1093,6 +1099,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); @@ -1176,6 +1184,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); @@ -1304,6 +1314,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 (); @@ -1421,6 +1435,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) { @@ -1945,6 +1965,44 @@ 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: \ + case ST_OMP_TILE: case ST_OMP_UNROLL + /* Block end statements. Errors associated with interchanging these are detected in gfc_match_end(). */ @@ -2582,6 +2640,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; @@ -2681,6 +2742,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; @@ -2831,6 +2895,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"; @@ -3097,6 +3164,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 (); @@ -5458,6 +5527,144 @@ 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; + case ST_OMP_TILE: + return ST_OMP_END_TILE; + case ST_OMP_UNROLL: + return ST_OMP_END_UNROLL; + 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. */ @@ -5518,94 +5725,16 @@ parse_omp_do (gfc_statement omp_st, int nested) st = next_statement (); do_end: - 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; - case ST_OMP_TILE: omp_end_st = ST_OMP_END_TILE; break; - case ST_OMP_UNROLL: omp_end_st = ST_OMP_END_UNROLL; 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) @@ -5917,80 +6046,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; @@ -6036,11 +6100,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: @@ -6213,6 +6279,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_variant *variant + = new_st.ext.omp_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, 0); + /* 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 @@ -6223,6 +6371,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) @@ -6230,6 +6379,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) @@ -6339,70 +6494,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 ST_OMP_TILE: - case ST_OMP_UNROLL: + case_omp_do: st = parse_omp_do (st, 0); if (st == ST_IMPLIED_ENDDO) return st; @@ -6416,6 +6514,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; } @@ -7181,6 +7292,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 */ @@ -7493,3 +7608,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 503029364c1..910fe339bc5 100644 --- a/gcc/fortran/resolve.cc +++ b/gcc/fortran/resolve.cc @@ -12538,6 +12538,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_variant *variant + = code->ext.omp_variants; + variant; variant = variant->next) + gfc_resolve_code (variant->code, ns); else if (code->block) { omp_workshare_save = -1; @@ -13072,6 +13077,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 0218d290782..494e82da8f5 100644 --- a/gcc/fortran/st.cc +++ b/gcc/fortran/st.cc @@ -301,6 +301,10 @@ gfc_free_statement (gfc_code *p) case EXEC_OMP_TASKYIELD: break; + case EXEC_OMP_METADIRECTIVE: + gfc_free_omp_variants (p->ext.omp_variants); + break; + default: gfc_internal_error ("gfc_free_statement(): Bad statement"); } diff --git a/gcc/fortran/symbol.cc b/gcc/fortran/symbol.cc index 2f326492d5f..769f6cf5f9c 100644 --- a/gcc/fortran/symbol.cc +++ b/gcc/fortran/symbol.cc @@ -2687,10 +2687,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; } @@ -2740,6 +2743,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; @@ -2756,10 +2760,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; @@ -2771,6 +2781,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 e6ac7f25b3b..4c1a2e65b0f 100644 --- a/gcc/fortran/trans-decl.cc +++ b/gcc/fortran/trans-decl.cc @@ -331,7 +331,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 ceeb9e366b1..e0665a803e9 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -8291,6 +8291,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: @@ -8382,6 +8384,100 @@ gfc_trans_omp_declare_simd (gfc_namespace *ns) } } +/* Translate the context selector list GFC_SELECTORS, using WHERE as the + locus for error messages. */ + +static tree +gfc_trans_omp_set_selector (gfc_omp_set_selector *gfc_selectors, locus where) +{ + 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_DEV_NUM_EXPR: + case OMP_TRAIT_PROPERTY_BOOL_EXPR: + { + tree expr = NULL_TREE; + 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) { @@ -8457,90 +8553,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_DEV_NUM_EXPR: - case OMP_TRAIT_PROPERTY_BOOL_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); 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) @@ -8641,3 +8655,54 @@ 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_variant *variant = code->ext.omp_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); + ctx = omp_check_context_selector (gfc_get_location (&variant->where), + ctx, true); + if (ctx == error_mark_node) + return error_mark_node; + + /* If the selector doesn't match, drop the whole variant. */ + if (!omp_context_selector_matches (ctx, true, true)) + { + variant = variant->next; + continue; + } + + 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; +} 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 d4c54093cbc..ea8ed548879 100644 --- a/gcc/fortran/trans.cc +++ b/gcc/fortran/trans.cc @@ -2614,6 +2614,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..3be368ef0ac --- /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 "otherwise:" 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..2707900454a --- /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 "otherwise:" 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..c017649d707 --- /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.x00\"\\\[0\\\], &\"amdgcn.x00\"\\\[0\\\], &\"gfx906.x00\"\\\[0\\\]\\)" "gimple" } } +! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(.+, &\"gpu.x00\"\\\[0\\\], &\"nvptx.x00\"\\\[0\\\], 0B\\)" "gimple" } } +! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(.+, &\"cpu.x00\"\\\[0\\\], &\"x86_64.x00\"\\\[0\\\], 0B\\)" "gimple" } } +! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(-1, &\"gpu.x00\"\\\[0\\\], &\"nvptx.x00\"\\\[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..1ebcd33a7be --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-8.f90 @@ -0,0 +1,22 @@ +! { dg-do compile } + +program test + integer :: i + integer, parameter :: N = 100 + integer :: sum = 0 + + ! The compiler should never consider a situation where both metadirectives + ! match, but that does not matter because the spec says "Replacement of + ! the metadirective with the directive variant associated with any of the + ! dynamic replacement candidates must result in a conforming OpenMP + ! program. So the second metadirective is rejected as not being + ! a valid loop-nest even if the first one does not match. + +!$omp metadirective when (implementation={vendor("ibm")}: & + !$omp& target teams distribute) + !$omp metadirective when (implementation={vendor("gnu")}: parallel do) ! { dg-error "Unexpected !.OMP METADIRECTIVE statement" } + do i = 1, N + sum = sum + i + end do +end program + 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 cdbebe215db..e5c9376796b 100644 --- a/gcc/testsuite/gfortran.dg/gomp/pure-1.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/pure-1.f90 @@ -110,3 +110,10 @@ pure integer function func_tile(n) end do func_tile = 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 35503c6a284..f6022189a68 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