From patchwork Mon May 27 11:54:33 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 1939908 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=aQr8s/Xg; 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 4VnvJj5zKfz20KL for ; Mon, 27 May 2024 21:57:09 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 11A793884509 for ; Mon, 27 May 2024 11:57:08 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lj1-x22c.google.com (mail-lj1-x22c.google.com [IPv6:2a00:1450:4864:20::22c]) by sourceware.org (Postfix) with ESMTPS id 9295F386C580 for ; Mon, 27 May 2024 11:55:31 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 9295F386C580 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 9295F386C580 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::22c ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810934; cv=none; b=IyWC4NlRoJVhTE/0qR72lW/QhKmqQIhchBgU/EG3ZnwRTxQuNcdwvIwmmbxXoWhbK+/a/FG140zdWuzxT7ZVQ3PwZexZy8teb9gWHTm8Fnf8rHgsTmm0MSRfXriUZbbxiCl8F+cCqrkK0UAp5e94e9HFf2nx+pEaCZ42JeBosaY= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810934; c=relaxed/simple; bh=Q/yB5kdfCurd/Hp0Kut48dU5+upTLwZf8VNyUauIeiw=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=ljnzAV2pmKLZBLZaobBFg+8d9r7h0SnpeSnysCYhl48rRxdzxmFVrTSVuBaNsMlsvo2gtGZ8Z+MtE1UizVMpyyHjnPQL+6PN7xqTWSEtXivPEsXs8CW3KbtxV2t3+PMLM+TRrxcaaXSEOZXmYQc2+XtD9f0oFR3SiOGsqxVIyqA= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lj1-x22c.google.com with SMTP id 38308e7fff4ca-2e73359b900so77786511fa.2 for ; Mon, 27 May 2024 04:55:31 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1716810930; x=1717415730; 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=o8f1LecRbm8L/xewCy9nBx8zEScVJTGvFvFOUtsrm0o=; b=aQr8s/XgIdFktja1Nw9zPkTx1D5ibhuJyxyElQ2lgRT5k1OQnkKfZRoY8Hsc5xzhLb Hrq7p+H5eifWPy+6BaKbAuQ0oQceXBTvtxI/BvpMLjgHgX+7MjsDt9EKUSrUiEZCozTk WYtEY0Open8e0Fnd+lC/L9kAaqjfHUPMkF5xN97+wEShakJ7M9stc+bLpumeBS614+WU gZ3tzL4DhNy2CzgOqh9i9nUIEfxwFY05UIsq6GC7D0HSYpNT7qOZxCumHtVm0IoSxQC6 6RPyIPzXViANj8kz4h108kcnSJkRLgwmetPypJfoFP9RKWHwe4Gxlrza9PokDMFPoM+L bhkg== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1716810930; x=1717415730; 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=o8f1LecRbm8L/xewCy9nBx8zEScVJTGvFvFOUtsrm0o=; b=E97rZshibDkrC1ZHn5xVkWmQCPcKkGlGf3YjxtAIh0jvKvt+ihW3fE1uQuNAEyv4d/ lcOst1/pXk7GbmfU1VwVlFp2EVW34IdpXumT49S1/nXnPL1cSIdf68W2XD8LvjT81T5t tClPS+2T7DtJAsof8zij4mAB+5uX6VA9wklb5MHVhsz2eenFucV85t/g9sUx2Oyubprx gQZwhPxGOxRXucoY6o92Svu2j/yctNdS9/A9ok4C3ZOxTlR0cYWIdXKmswgWqsPJhdKV 1W34Gyc3nx+pEoLVqh+v1fxZkg9snDnHm8Od7apTQuG9rECEbdpMWkNM9HCP01D17vwH Afqg== X-Gm-Message-State: AOJu0Yzz277paGC1UHqu0rLdDbuSBnw0CfhHlz8CKbFtjhK0L5UKK0DJ CPQtynQfL+EKKl9+WCMZFbZYTNKiol6b0kUyALH93X3kEFzPOz1iHgRqwizHMwt5Z1NMFGdDjI0 T X-Google-Smtp-Source: AGHT+IGs4bvrMoNexnnDdODmsIkTTNsdgLsrdVEM1cY0SzgW/cf36yZ6GqT05DQf2qC+GNVKMFFCyw== X-Received: by 2002:a2e:9646:0:b0:2e1:bdfd:ce70 with SMTP id 38308e7fff4ca-2e95b096f1fmr51705211fa.6.1716810929537; Mon, 27 May 2024 04:55:29 -0700 (PDT) Received: from localhost.localdomain ([2a0d:3344:2340:ce10::2f3]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-42108966b58sm107636955e9.3.2024.05.27.04.55.28 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Mon, 27 May 2024 04:55:29 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH 1/7] OpenMP: dispatch + adjust_args tree data structures and front-end interfaces Date: Mon, 27 May 2024 13:54:33 +0200 Message-ID: <20240527115439.3967217-2-parras@baylibre.com> X-Mailer: git-send-email 2.45.1 In-Reply-To: <20240527115439.3967217-1-parras@baylibre.com> References: <20240527115439.3967217-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-10.9 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org This patch introduces the OMP_DISPATCH tree node, as well as two new clauses `nocontext` and `novariants`. It defines/exposes interfaces that will be used in subsequent patches that add front-end and middle-end support, but nothing generates these nodes yet. It also adds support for new OpenMP context selectors: `dispatch` as trait selector and `need_device_ptr` as pseudo-trait set selector. The purpose of the latter is for the C++ front-end to store the list of arguments (that need to be converted to device pointers) until the declaration of the variant function becomes available. gcc/ChangeLog: * builtin-types.def (BT_FN_PTR_CONST_PTR_INT): New. * omp-selectors.h (enum omp_tss_code): Add OMP_TRAIT_SET_NEED_DEVICE_PTR. (enum omp_ts_code): Add OMP_TRAIT_CONSTRUCT_DISPATCH. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. * tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (dump_generic_node): Handle OMP_DISPATCH. * tree.cc (omp_clause_num_ops): Add OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (omp_clause_code_name): Add "novariants" and "nocontext". * tree.def (OMP_DISPATCH): New. * tree.h (OMP_DISPATCH_BODY): New macro. (OMP_DISPATCH_CLAUSES): New macro. (OMP_CLAUSE_NOVARIANTS_EXPR): New macro. (OMP_CLAUSE_NOCONTEXT_EXPR): New macro. --- gcc/builtin-types.def | 1 + gcc/omp-selectors.h | 3 +++ gcc/tree-core.h | 7 +++++++ gcc/tree-pretty-print.cc | 21 +++++++++++++++++++++ gcc/tree.cc | 4 ++++ gcc/tree.def | 5 +++++ gcc/tree.h | 7 +++++++ 7 files changed, 48 insertions(+) diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def index c97d6bad1de..ef7aaf67d13 100644 --- a/gcc/builtin-types.def +++ b/gcc/builtin-types.def @@ -677,6 +677,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_INT_FEXCEPT_T_PTR_INT, BT_INT, BT_FEXCEPT_T_PTR, DEF_FUNCTION_TYPE_2 (BT_FN_INT_CONST_FEXCEPT_T_PTR_INT, BT_INT, BT_CONST_FEXCEPT_T_PTR, BT_INT) DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_UINT8, BT_PTR, BT_CONST_PTR, BT_UINT8) +DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_INT, BT_PTR, BT_CONST_PTR, BT_INT) DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR) diff --git a/gcc/omp-selectors.h b/gcc/omp-selectors.h index c61808ec0ad..12bc9e9afa0 100644 --- a/gcc/omp-selectors.h +++ b/gcc/omp-selectors.h @@ -31,6 +31,8 @@ enum omp_tss_code { OMP_TRAIT_SET_TARGET_DEVICE, OMP_TRAIT_SET_IMPLEMENTATION, OMP_TRAIT_SET_USER, + OMP_TRAIT_SET_NEED_DEVICE_PTR, // pseudo-set selector used to convey argument + // list until variant has a decl OMP_TRAIT_SET_LAST, OMP_TRAIT_SET_INVALID = -1 }; @@ -55,6 +57,7 @@ enum omp_ts_code { OMP_TRAIT_CONSTRUCT_PARALLEL, OMP_TRAIT_CONSTRUCT_FOR, OMP_TRAIT_CONSTRUCT_SIMD, + OMP_TRAIT_CONSTRUCT_DISPATCH, OMP_TRAIT_LAST, OMP_TRAIT_INVALID = -1 }; diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 9fa74342919..ed6ffdab87f 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -533,6 +533,13 @@ enum omp_clause_code { /* OpenACC clause: nohost. */ OMP_CLAUSE_NOHOST, + + /* OpenMP clause: novariants (scalar-expression). */ + OMP_CLAUSE_NOVARIANTS, + + /* OpenMP clause: nocontext (scalar-expression). */ + OMP_CLAUSE_NOCONTEXT, + }; #undef DEFTREESTRUCT diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index f9ad8562078..bbae3a98e9a 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -506,6 +506,22 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case OMP_CLAUSE_EXCLUSIVE: name = "exclusive"; goto print_remap; + case OMP_CLAUSE_NOVARIANTS: + pp_string (pp, "novariants"); + pp_left_paren (pp); + gcc_assert (OMP_CLAUSE_NOVARIANTS_EXPR (clause)); + dump_generic_node (pp, OMP_CLAUSE_NOVARIANTS_EXPR (clause), spc, flags, + false); + pp_right_paren (pp); + break; + case OMP_CLAUSE_NOCONTEXT: + pp_string (pp, "nocontext"); + pp_left_paren (pp); + gcc_assert (OMP_CLAUSE_NOCONTEXT_EXPR (clause)); + dump_generic_node (pp, OMP_CLAUSE_NOCONTEXT_EXPR (clause), spc, flags, + false); + pp_right_paren (pp); + break; case OMP_CLAUSE__LOOPTEMP_: name = "_looptemp_"; goto print_remap; @@ -3918,6 +3934,11 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags, dump_omp_clauses (pp, OMP_SECTIONS_CLAUSES (node), spc, flags); goto dump_omp_body; + case OMP_DISPATCH: + pp_string (pp, "#pragma omp dispatch"); + dump_omp_clauses (pp, OMP_DISPATCH_CLAUSES (node), spc, flags); + goto dump_omp_body; + case OMP_SECTION: pp_string (pp, "#pragma omp section"); goto dump_omp_body; diff --git a/gcc/tree.cc b/gcc/tree.cc index 6564b002dc1..f0007e24fc4 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -328,6 +328,8 @@ unsigned const char omp_clause_num_ops[] = 0, /* OMP_CLAUSE_IF_PRESENT */ 0, /* OMP_CLAUSE_FINALIZE */ 0, /* OMP_CLAUSE_NOHOST */ + 1, /* OMP_CLAUSE_NOVARIANTS */ + 1, /* OMP_CLAUSE_NOCONTEXT */ }; const char * const omp_clause_code_name[] = @@ -421,6 +423,8 @@ const char * const omp_clause_code_name[] = "if_present", "finalize", "nohost", + "novariants", + "nocontext", }; /* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric diff --git a/gcc/tree.def b/gcc/tree.def index 24128e1e039..ff7ca96d29c 100644 --- a/gcc/tree.def +++ b/gcc/tree.def @@ -1290,6 +1290,11 @@ DEFTREECODE (OMP_MASKED, "omp_masked", tcc_statement, 2) Operand 1: OMP_SCAN_CLAUSES: List of clauses. */ DEFTREECODE (OMP_SCAN, "omp_scan", tcc_statement, 2) +/* OpenMP - #pragma omp dispatch [clause1 ... clauseN] + Operand 0: OMP_DISPATCH_BODY: Expression statement including a target call. + Operand 1: OMP_DISPATCH_CLAUSES: List of clauses. */ +DEFTREECODE (OMP_DISPATCH, "omp_dispatch", tcc_statement, 2) + /* OpenMP - #pragma omp section Operand 0: OMP_SECTION_BODY: Section body. */ DEFTREECODE (OMP_SECTION, "omp_section", tcc_statement, 1) diff --git a/gcc/tree.h b/gcc/tree.h index ee2aae332a4..f23a7366de0 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1599,6 +1599,9 @@ class auto_suppress_location_wrappers #define OMP_SCAN_BODY(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 0) #define OMP_SCAN_CLAUSES(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 1) +#define OMP_DISPATCH_BODY(NODE) TREE_OPERAND (OMP_DISPATCH_CHECK (NODE), 0) +#define OMP_DISPATCH_CLAUSES(NODE) TREE_OPERAND (OMP_DISPATCH_CHECK (NODE), 1) + #define OMP_CLAUSE_SIZE(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \ OMP_CLAUSE_FROM, \ @@ -1742,6 +1745,10 @@ class auto_suppress_location_wrappers OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_HINT), 0) #define OMP_CLAUSE_FILTER_EXPR(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_FILTER), 0) +#define OMP_CLAUSE_NOVARIANTS_EXPR(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NOVARIANTS), 0) +#define OMP_CLAUSE_NOCONTEXT_EXPR(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NOCONTEXT), 0) #define OMP_CLAUSE_GRAINSIZE_EXPR(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_GRAINSIZE),0) From patchwork Mon May 27 11:54:34 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 1939905 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=pFhO6Mqt; 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 4VnvHd5Fq7z20KL for ; Mon, 27 May 2024 21:56:13 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 055353883036 for ; Mon, 27 May 2024 11:56:12 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x32b.google.com (mail-wm1-x32b.google.com [IPv6:2a00:1450:4864:20::32b]) by sourceware.org (Postfix) with ESMTPS id 101A5382EF03 for ; Mon, 27 May 2024 11:55:33 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 101A5382EF03 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 101A5382EF03 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::32b ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810937; cv=none; b=A7FEbxH1PQ+kJOdj2j9xLRLloZ2kgVfg/13XGXKllpT+yw+CsXue8jP5WTYCO3iSz9j7j7kykoW6Zyp0XSXhSWLIQWdhbSC+NIvpuv8Nkg71LsVDerjGshjEgcXdRUWOGm4Jm/zylVVGV7vbhW9goXXBWG4Y/Rgxn2mFyg893ig= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810937; c=relaxed/simple; bh=oShRI+Mj94bTkbxkItzjUc1WgLBjH0zETdleAz3jtxI=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=wulTOrcvW1o6ZTgtlNOVMe3hz2VK+iFE/uYCreaUiDWhr93grribeOk/Bx1LqMTQjQ87Iso/otlPQEn1+yyVmJhTHU/QJ/0bzHl9xeG0Ef4XQrFnqoiZiTbczOXIFBndJs3Pn0TopH2R2W0hXRiNZUbDxDCKXVLsBbPhe1hDOkw= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x32b.google.com with SMTP id 5b1f17b1804b1-420180b58c3so84401695e9.2 for ; Mon, 27 May 2024 04:55:33 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1716810931; x=1717415731; 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=QHtXc2NSa6J1nO/qNeDenUCLybMZI1/h9rWZO2m7dks=; b=pFhO6MqteiUHEkVvxipLdscAcdMrU4Zje6Yma1zkei/wBYmh1yzvoeTVArBmxuWvGj 22tqeWyE+ls6eCJp/cHlHrehdV2Mg5LPFIApEhKFp1RKU14QCnnKoqqzwucbDVgPVSr5 sRK8PYnv0rG3roW//SL19XIz6KWtIBlFCkk0gyRzy2gQ66kZkrA1XoDVymQYqwli83+K eAVyE+lFhXua/W4Mw72LRegZF9QQwOYFAtaeUlMS7JoQtnzLHj0zpZk/p+HgOnP4WCKK QtYEWuQcpfzbVpgdf6/MfqWElYNgC8vtR1f4P9LWWZjSwcnwFh8QAbmn1qSr2cWZOcqN 2KCw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1716810931; x=1717415731; 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=QHtXc2NSa6J1nO/qNeDenUCLybMZI1/h9rWZO2m7dks=; b=KmxdiEI4l+B8TSaKpd2mjwOqNrUzqprj5i1qu3SnIeeIjwl0YtcRYQo3t0x8YdQs3d F2mOIdBrkZ0XbfVbVXlmB+w5LdTHnUXjU9OYFYIw2WhHIq9Uv3tCAKxn3ljXQXSc6+9G 7bmdMbA/7RTsLiJIlpRZjSCYnypkRdxZZzoeSLhL4GBgqDNI+tORKpI/3E5y+mvy8aMl Py+iVtLJx4i0qwOW0VX2IGIlbvZx+xUGMj/rkiGb5MgeUOhZeU2lWsSYNoBP67w0uX5H B/+7S8re/yawdVsRggtxhUx2wjXU+ejjZ6bbru45DBevmRWqXUVvIpLJs521JkSXhron ojbg== X-Gm-Message-State: AOJu0Yz1QlgLwDlEMQCzrSKiLXy/6jW89loKuYXO1T2rkg3xo4+qHHWG rilWeYQ+5kDhOIJ/ajWfqveaWIuJCQWEocY2gMvxCv38G1+N1cr9deguhN44dynN1ojezrrE231 G X-Google-Smtp-Source: AGHT+IGA6e+SjC6GQeQ8qxCQILqEeUx4SPhv5mThkAs/G53DN+EZ+8tvhgXRMCPOPcQMP+ZdNrrJrQ== X-Received: by 2002:a05:600c:444e:b0:420:14fb:de1f with SMTP id 5b1f17b1804b1-42108a5965fmr59542945e9.14.1716810930865; Mon, 27 May 2024 04:55:30 -0700 (PDT) Received: from localhost.localdomain ([2a0d:3344:2340:ce10::2f3]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-42108966b58sm107636955e9.3.2024.05.27.04.55.29 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Mon, 27 May 2024 04:55:30 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH 2/7] OpenMP: middle-end support for dispatch + adjust_args Date: Mon, 27 May 2024 13:54:34 +0200 Message-ID: <20240527115439.3967217-3-parras@baylibre.com> X-Mailer: git-send-email 2.45.1 In-Reply-To: <20240527115439.3967217-1-parras@baylibre.com> References: <20240527115439.3967217-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds middle-end support for the `dispatch` construct and the `adjust_args` clause. The heavy lifting is done in `gimplify_omp_dispatch` and `gimplify_call_expr` respectively. For `adjust_args`, this mostly consists in emitting a call to `gomp_get_mapped_ptr` for the adequate device. For dispatch, the following steps are performed: * Handle the device clause, if any. This may affect `need_device_ptr` arguments. * Handle novariants and nocontext clauses, if any. Evaluate compile-time constants and select a variant, if possible. Otherwise, emit code to handle all possible cases at run time. * Create an explicit task, as if the `task` construct was used, that wraps the body of the `dispatch` statement. Move relevant clauses to the task. gcc/ChangeLog: * gimple-low.cc (lower_stmt): Handle GIMPLE_OMP_DISPATCH. * gimple-pretty-print.cc (dump_gimple_omp_dispatch): New function. (pp_gimple_stmt_1): Handle GIMPLE_OMP_DISPATCH. * gimple-walk.cc (walk_gimple_stmt): Likewise. * gimple.cc (gimple_build_omp_dispatch): New function. (gimple_copy): Handle GIMPLE_OMP_DISPATCH. * gimple.def (GIMPLE_OMP_DISPATCH): Define. * gimple.h (gimple_build_omp_dispatch): Declare. (gimple_has_substatements): Handle GIMPLE_OMP_DISPATCH. (gimple_omp_dispatch_clauses): New function. (gimple_omp_dispatch_clauses_ptr): Likewise. (gimple_omp_dispatch_set_clauses): Likewise. (gimple_return_set_retval): Handle GIMPLE_OMP_DISPATCH. * gimplify.cc (enum omp_region_type): Add ORT_DISPATCH. (gimplify_call_expr): Handle need_device_ptr arguments. (is_gimple_stmt): Handle OMP_DISPATCH. (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_DEVICE in a dispatch construct. Handle OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (omp_construct_selector_matches): Handle OMP_DISPATCH with nocontext clause. (omp_has_novariants): New function. (omp_has_nocontext): Likewise. (gimplify_omp_dispatch): Likewise. (gimplify_expr): Handle OMP_DISPATCH. * gimplify.h (omp_has_novariants): Declare. (omp_has_nocontext): Declare. * omp-builtins.def (BUILT_IN_OMP_GET_MAPPED_PTR): Define. (BUILT_IN_OMP_GET_DEFAULT_DEVICE): Define. (BUILT_IN_OMP_SET_DEFAULT_DEVICE): Define. * omp-expand.cc (expand_omp_dispatch): New function. (expand_omp): Handle GIMPLE_OMP_DISPATCH. (omp_make_gimple_edges): Likewise. * omp-general.cc (omp_construct_traits_to_codes): Add OMP_DISPATCH. (struct omp_ts_info): Add dispatch. (omp_context_selector_matches): Handle OMP_TRAIT_SET_NEED_DEVICE_PTR. (omp_resolve_declare_variant): Handle novariants. Adjust DECL_ASSEMBLER_NAME. --- gcc/gimple-low.cc | 1 + gcc/gimple-pretty-print.cc | 33 +++ gcc/gimple-walk.cc | 1 + gcc/gimple.cc | 20 ++ gcc/gimple.def | 5 + gcc/gimple.h | 33 ++- gcc/gimplify.cc | 417 ++++++++++++++++++++++++++++++++++++- gcc/gimplify.h | 2 + gcc/omp-builtins.def | 6 + gcc/omp-expand.cc | 18 ++ gcc/omp-general.cc | 16 +- gcc/omp-low.cc | 35 ++++ gcc/tree-inline.cc | 7 + 13 files changed, 583 insertions(+), 11 deletions(-) diff --git a/gcc/gimple-low.cc b/gcc/gimple-low.cc index e0371988705..712a1ebf776 100644 --- a/gcc/gimple-low.cc +++ b/gcc/gimple-low.cc @@ -746,6 +746,7 @@ lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data) case GIMPLE_EH_MUST_NOT_THROW: case GIMPLE_OMP_FOR: case GIMPLE_OMP_SCOPE: + case GIMPLE_OMP_DISPATCH: case GIMPLE_OMP_SECTIONS: case GIMPLE_OMP_SECTIONS_SWITCH: case GIMPLE_OMP_SECTION: diff --git a/gcc/gimple-pretty-print.cc b/gcc/gimple-pretty-print.cc index a71e1e0efc7..d9a24ad2169 100644 --- a/gcc/gimple-pretty-print.cc +++ b/gcc/gimple-pretty-print.cc @@ -1726,6 +1726,35 @@ dump_gimple_omp_scope (pretty_printer *buffer, const gimple *gs, } } +/* Dump a GIMPLE_OMP_DISPATCH tuple on the pretty_printer BUFFER. */ + +static void +dump_gimple_omp_dispatch (pretty_printer *buffer, const gimple *gs, int spc, + dump_flags_t flags) +{ + if (flags & TDF_RAW) + { + dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S>%nCLAUSES <", gs, + gimple_omp_body (gs)); + dump_omp_clauses (buffer, gimple_omp_dispatch_clauses (gs), spc, flags); + dump_gimple_fmt (buffer, spc, flags, " >"); + } + else + { + pp_string (buffer, "#pragma omp dispatch"); + dump_omp_clauses (buffer, gimple_omp_dispatch_clauses (gs), spc, flags); + if (!gimple_seq_empty_p (gimple_omp_body (gs))) + { + newline_and_indent (buffer, spc + 2); + pp_left_brace (buffer); + pp_newline (buffer); + dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags); + newline_and_indent (buffer, spc + 2); + pp_right_brace (buffer); + } + } +} + /* Dump a GIMPLE_OMP_TARGET tuple on the pretty_printer BUFFER. */ static void @@ -2805,6 +2834,10 @@ pp_gimple_stmt_1 (pretty_printer *buffer, const gimple *gs, int spc, dump_gimple_omp_scope (buffer, gs, spc, flags); break; + case GIMPLE_OMP_DISPATCH: + dump_gimple_omp_dispatch (buffer, gs, spc, flags); + break; + case GIMPLE_OMP_MASTER: case GIMPLE_OMP_SECTION: case GIMPLE_OMP_STRUCTURED_BLOCK: diff --git a/gcc/gimple-walk.cc b/gcc/gimple-walk.cc index 9f768ca20fd..1122713a98b 100644 --- a/gcc/gimple-walk.cc +++ b/gcc/gimple-walk.cc @@ -707,6 +707,7 @@ walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt, case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: case GIMPLE_OMP_SCOPE: + case GIMPLE_OMP_DISPATCH: case GIMPLE_OMP_SECTIONS: case GIMPLE_OMP_SINGLE: case GIMPLE_OMP_TARGET: diff --git a/gcc/gimple.cc b/gcc/gimple.cc index a9f968cb038..3a26c74a105 100644 --- a/gcc/gimple.cc +++ b/gcc/gimple.cc @@ -1235,6 +1235,21 @@ gimple_build_omp_scope (gimple_seq body, tree clauses) return p; } +/* Build a GIMPLE_OMP_DISPATCH statement. + + BODY is the target function call to be dispatched. + CLAUSES are any of the OMP dispatch construct's clauses: ... */ + +gimple * +gimple_build_omp_dispatch (gimple_seq body, tree clauses) +{ + gimple *p = gimple_alloc (GIMPLE_OMP_DISPATCH, 0); + gimple_omp_dispatch_set_clauses (p, clauses); + if (body) + gimple_omp_set_body (p, body); + + return p; +} /* Build a GIMPLE_OMP_TARGET statement. @@ -2148,6 +2163,11 @@ gimple_copy (gimple *stmt) gimple_omp_scope_set_clauses (copy, t); goto copy_omp_body; + case GIMPLE_OMP_DISPATCH: + t = unshare_expr (gimple_omp_dispatch_clauses (stmt)); + gimple_omp_dispatch_set_clauses (copy, t); + goto copy_omp_body; + case GIMPLE_OMP_TARGET: { gomp_target *omp_target_stmt = as_a (stmt); diff --git a/gcc/gimple.def b/gcc/gimple.def index fbcd727f945..21c7405875d 100644 --- a/gcc/gimple.def +++ b/gcc/gimple.def @@ -350,6 +350,11 @@ DEFGSCODE(GIMPLE_OMP_SCAN, "gimple_omp_scan", GSS_OMP_SINGLE_LAYOUT) CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */ DEFGSCODE(GIMPLE_OMP_SCOPE, "gimple_omp_scope", GSS_OMP_SINGLE_LAYOUT) +/* GIMPLE_OMP_DISPATCH represents #pragma omp dispatch + BODY is the target function call to be dispatched. + CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */ +DEFGSCODE(GIMPLE_OMP_DISPATCH, "gimple_omp_dispatch", GSS_OMP_SINGLE_LAYOUT) + /* OMP_SECTION represents #pragma omp section. BODY is the sequence of statements in the section body. */ DEFGSCODE(GIMPLE_OMP_SECTION, "gimple_omp_section", GSS_OMP) diff --git a/gcc/gimple.h b/gcc/gimple.h index bd315ffc2dd..25590a22ffb 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -746,7 +746,7 @@ struct GTY((tag("GSS_OMP_CONTINUE"))) }; /* GIMPLE_OMP_SINGLE, GIMPLE_OMP_ORDERED, GIMPLE_OMP_TASKGROUP, - GIMPLE_OMP_SCAN, GIMPLE_OMP_MASKED, GIMPLE_OMP_SCOPE. */ + GIMPLE_OMP_SCAN, GIMPLE_OMP_MASKED, GIMPLE_OMP_SCOPE, GIMPLE_OMP_DISPATCH. */ struct GTY((tag("GSS_OMP_SINGLE_LAYOUT"))) gimple_statement_omp_single_layout : public gimple_statement_omp @@ -1595,6 +1595,7 @@ gomp_task *gimple_build_omp_task (gimple_seq, tree, tree, tree, tree, gimple *gimple_build_omp_section (gimple_seq); gimple *gimple_build_omp_structured_block (gimple_seq); gimple *gimple_build_omp_scope (gimple_seq, tree); +gimple *gimple_build_omp_dispatch (gimple_seq, tree); gimple *gimple_build_omp_master (gimple_seq); gimple *gimple_build_omp_masked (gimple_seq, tree); gimple *gimple_build_omp_taskgroup (gimple_seq, tree); @@ -1886,6 +1887,7 @@ gimple_has_substatements (gimple *g) case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: case GIMPLE_OMP_SCOPE: + case GIMPLE_OMP_DISPATCH: case GIMPLE_OMP_SECTIONS: case GIMPLE_OMP_SINGLE: case GIMPLE_OMP_TARGET: @@ -5437,6 +5439,34 @@ gimple_omp_scope_set_clauses (gimple *gs, tree clauses) = clauses; } +/* Return the clauses associated with OMP_DISPATCH statement GS. */ + +inline tree +gimple_omp_dispatch_clauses (const gimple *gs) +{ + GIMPLE_CHECK (gs, GIMPLE_OMP_DISPATCH); + return static_cast (gs)->clauses; +} + +/* Return a pointer to the clauses associated with OMP dispatch statement + GS. */ + +inline tree * +gimple_omp_dispatch_clauses_ptr (gimple *gs) +{ + GIMPLE_CHECK (gs, GIMPLE_OMP_DISPATCH); + return &static_cast (gs)->clauses; +} + +/* Set CLAUSES to be the clauses associated with OMP dispatch statement + GS. */ + +inline void +gimple_omp_dispatch_set_clauses (gimple *gs, tree clauses) +{ + GIMPLE_CHECK (gs, GIMPLE_OMP_DISPATCH); + static_cast (gs)->clauses = clauses; +} /* Return the kind of the OMP_FOR statemement G. */ @@ -6771,6 +6801,7 @@ gimple_return_set_retval (greturn *gs, tree retval) case GIMPLE_OMP_TARGET: \ case GIMPLE_OMP_TEAMS: \ case GIMPLE_OMP_SCOPE: \ + case GIMPLE_OMP_DISPATCH: \ case GIMPLE_OMP_SECTION: \ case GIMPLE_OMP_STRUCTURED_BLOCK: \ case GIMPLE_OMP_MASTER: \ diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index b0ed58ed0f9..1dd69dbb1de 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -161,7 +161,8 @@ enum omp_region_type { ORT_WORKSHARE = 0x00, ORT_TASKGROUP = 0x01, - ORT_SIMD = 0x04, + ORT_DISPATCH = 0x02, + ORT_SIMD = 0x04, ORT_PARALLEL = 0x08, ORT_COMBINED_PARALLEL = ORT_PARALLEL | 1, @@ -4051,6 +4052,7 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value) /* Gimplify the function arguments. */ if (nargs > 0) { + tree device_num = NULL_TREE; for (i = (PUSH_ARGS_REVERSED ? nargs - 1 : 0); PUSH_ARGS_REVERSED ? i >= 0 : i < nargs; PUSH_ARGS_REVERSED ? i-- : i++) @@ -4061,8 +4063,99 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value) be the plain PARM_DECL. */ if ((i != 1) || !builtin_va_start_p) { - t = gimplify_arg (&CALL_EXPR_ARG (*expr_p, i), pre_p, - EXPR_LOCATION (*expr_p), ! returns_twice); + tree *arg_p = &CALL_EXPR_ARG (*expr_p, i); + if (flag_openmp && EXPR_P (CALL_EXPR_FN (*expr_p)) + && lookup_attribute ("omp declare variant variant", + DECL_ATTRIBUTES (TREE_OPERAND ( + CALL_EXPR_FN (*expr_p), 0))) + != NULL_TREE) + { + tree param + = DECL_ARGUMENTS (TREE_OPERAND (CALL_EXPR_FN (*expr_p), 0)); + + if (param != NULL_TREE) + { + for (int param_idx = 0; param_idx < i; param_idx++) + param = TREE_CHAIN (param); + + bool is_device_ptr = false; + if (gimplify_omp_ctxp != NULL + && gimplify_omp_ctxp->code == OMP_DISPATCH) + { + for (tree c = gimplify_omp_ctxp->clauses; c; + c = TREE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) + == OMP_CLAUSE_IS_DEVICE_PTR) + { + tree decl1 = DECL_NAME (OMP_CLAUSE_DECL (c)); + tree decl2 + = tree_strip_nop_conversions (*arg_p); + if (TREE_CODE (decl2) == ADDR_EXPR) + decl2 = TREE_OPERAND (decl2, 0); + gcc_assert (TREE_CODE (decl2) == VAR_DECL + || TREE_CODE (decl2) + == PARM_DECL); + decl2 = DECL_NAME (decl2); + if (decl1 == decl2) + { + is_device_ptr = true; + break; + } + } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE) + device_num = OMP_CLAUSE_OPERAND (c, 0); + } + } + + if (!is_device_ptr + && lookup_attribute ("omp declare variant " + "adjust_args need_device_ptr", + DECL_ATTRIBUTES (param)) + != NULL_TREE) + { + if (device_num == NULL_TREE) + { + // device_num = omp_get_default_device(); + tree fn = builtin_decl_explicit ( + BUILT_IN_OMP_GET_DEFAULT_DEVICE); + gcall *call = gimple_build_call (fn, 0); + device_num = create_tmp_var ( + gimple_call_return_type (call)); + gimple_call_set_lhs (call, device_num); + gimplify_seq_add_stmt (pre_p, call); + } + + // mapped_arg = omp_get_mapped_ptr(arg, device_num); + tree fn = builtin_decl_explicit ( + BUILT_IN_OMP_GET_MAPPED_PTR); + *arg_p = (TREE_CODE (*arg_p) == NOP_EXPR) + ? TREE_OPERAND (*arg_p, 0) + : *arg_p; + gimplify_arg (arg_p, pre_p, loc); + gimplify_arg (&device_num, pre_p, loc); + call = gimple_build_call (fn, 2, *arg_p, device_num); + tree mapped_arg + = create_tmp_var (gimple_call_return_type (call)); + gimple_call_set_lhs (call, mapped_arg); + gimplify_seq_add_stmt (pre_p, call); + + *arg_p = mapped_arg; + + // Mark mapped argument as device pointer to ensure + // idempotency in gimplification + gcc_assert (gimplify_omp_ctxp->code == OMP_DISPATCH); + tree c = build_omp_clause (input_location, + OMP_CLAUSE_IS_DEVICE_PTR); + OMP_CLAUSE_DECL (c) = *arg_p; + OMP_CLAUSE_CHAIN (c) = gimplify_omp_ctxp->clauses; + gimplify_omp_ctxp->clauses = c; + } + } + } + + t = gimplify_arg (arg_p, pre_p, EXPR_LOCATION (*expr_p), + !returns_twice); if (t == GS_ERROR) ret = GS_ERROR; @@ -6307,6 +6400,7 @@ is_gimple_stmt (tree t) case OACC_LOOP: case OMP_SCAN: case OMP_SCOPE: + case OMP_DISPATCH: case OMP_SECTIONS: case OMP_SECTION: case OMP_STRUCTURED_BLOCK: @@ -13080,6 +13174,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE + && code == OMP_DISPATCH) + { + bool saved_into_ssa = gimplify_ctxp->into_ssa; + gimplify_ctxp->into_ssa = false; + if (gimplify_expr (&OMP_CLAUSE_DEVICE_ID (c), pre_p, NULL, + is_gimple_val, fb_rvalue) + == GS_ERROR) + remove = true; + else if (DECL_P (OMP_CLAUSE_DEVICE_ID (c))) + omp_add_variable (ctx, OMP_CLAUSE_DEVICE_ID (c), + GOVD_SHARED | GOVD_SEEN); + gimplify_ctxp->into_ssa = saved_into_ssa; + break; + } /* Fall through. */ case OMP_CLAUSE_PRIORITY: @@ -13309,6 +13418,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } break; + case OMP_CLAUSE_NOVARIANTS: + { + OMP_CLAUSE_NOVARIANTS_EXPR (c); + tree t = gimple_boolify (OMP_CLAUSE_NOVARIANTS_EXPR (c)); + OMP_CLAUSE_NOVARIANTS_EXPR (c) = t; + } + break; + case OMP_CLAUSE_NOCONTEXT: + { + OMP_CLAUSE_NOCONTEXT_EXPR (c); + tree t = gimple_boolify (OMP_CLAUSE_NOCONTEXT_EXPR (c)); + OMP_CLAUSE_NOCONTEXT_EXPR (c) = t; + } + break; case OMP_CLAUSE_NOHOST: default: gcc_unreachable (); @@ -13763,7 +13886,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, { struct gimplify_omp_ctx *octx; for (octx = ctx; octx; octx = octx->outer_context) - if ((octx->region_type & (ORT_PARALLEL | ORT_TASK | ORT_TEAMS)) != 0) + if ((octx->region_type + & (ORT_DISPATCH | ORT_PARALLEL | ORT_TASK | ORT_TEAMS)) + != 0) break; if (octx) { @@ -14574,6 +14699,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_CLAUSE_FINALIZE: case OMP_CLAUSE_INCLUSIVE: case OMP_CLAUSE_EXCLUSIVE: + case OMP_CLAUSE_NOVARIANTS: + case OMP_CLAUSE_NOCONTEXT: break; case OMP_CLAUSE_NOHOST: @@ -14663,9 +14790,9 @@ omp_construct_selector_matches (enum tree_code *constructs, int nconstructs, == ORT_TARGET && ctx->code == OMP_TARGET) || ((ctx->region_type & ORT_TEAMS) && ctx->code == OMP_TEAMS) || (ctx->region_type == ORT_WORKSHARE && ctx->code == OMP_FOR) - || (ctx->region_type == ORT_SIMD - && ctx->code == OMP_SIMD - && !omp_find_clause (ctx->clauses, OMP_CLAUSE_BIND))) + || (ctx->region_type == ORT_SIMD && ctx->code == OMP_SIMD + && !omp_find_clause (ctx->clauses, OMP_CLAUSE_BIND)) + || (ctx->code == OMP_DISPATCH && omp_has_nocontext () != 1)) { ++cnt; if (scores) @@ -14783,6 +14910,60 @@ omp_construct_selector_matches (enum tree_code *constructs, int nconstructs, return 0; } +/* Try to evaluate a novariants clause. Return 1 if true, 0 if false or absent, + * -1 if run-time evaluation is needed. */ + +int +omp_has_novariants (void) +{ + for (struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; ctx; + ctx = ctx->outer_context) + { + if (ctx->code == OMP_DISPATCH) + { + tree c = omp_find_clause (ctx->clauses, OMP_CLAUSE_NOVARIANTS); + if (c != NULL_TREE) + { + if (integer_nonzerop (OMP_CLAUSE_NOVARIANTS_EXPR (c))) + return 1; + else if (integer_zerop (OMP_CLAUSE_NOVARIANTS_EXPR (c))) + return 0; + else + return -1; + } + return 0; + } + } + return 0; +} + +/* Try to evaluate a nocontext clause. Return 1 if true, 0 if false or absent, + * -1 if run-time evaluation is needed. */ + +int +omp_has_nocontext (void) +{ + for (struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; ctx; + ctx = ctx->outer_context) + { + if (ctx->code == OMP_DISPATCH) + { + tree c = omp_find_clause (ctx->clauses, OMP_CLAUSE_NOCONTEXT); + if (c != NULL_TREE) + { + if (integer_nonzerop (OMP_CLAUSE_NOCONTEXT_EXPR (c))) + return 1; + else if (integer_zerop (OMP_CLAUSE_NOCONTEXT_EXPR (c))) + return 0; + else + return -1; + } + return 0; + } + } + return 0; +} + /* Gimplify OACC_CACHE. */ static void @@ -17614,6 +17795,221 @@ gimplify_omp_ordered (tree expr, gimple_seq body) return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr)); } +/* Gimplify an OMP_DISPATCH construct. */ + +static enum gimplify_status +gimplify_omp_dispatch (tree *expr_p, gimple_seq *pre_p) +{ + tree expr = *expr_p; + gimple_seq body = NULL; + + gimplify_scan_omp_clauses (&OMP_DISPATCH_CLAUSES (expr), pre_p, ORT_DISPATCH, + OMP_DISPATCH); + push_gimplify_context (); + + // If device clause, adjust ICV + tree device + = omp_find_clause (OMP_DISPATCH_CLAUSES (expr), OMP_CLAUSE_DEVICE); + if (device) + { + tree t = builtin_decl_explicit (BUILT_IN_OMP_SET_DEFAULT_DEVICE); + t = build_call_expr_loc (input_location, t, 1, + OMP_CLAUSE_DEVICE_ID (device)); + gimplify_and_add (t, &body); + if (DECL_P (OMP_CLAUSE_DEVICE_ID (device))) + omp_notice_variable (gimplify_omp_ctxp, OMP_CLAUSE_DEVICE_ID (device), + true); + } + + // If the novariants and nocontext clauses are not compile-time constants, + // we need to generate code for all possible cases: + // if (novariants) // implies nocontext + // base() + // else if (nocontext) + // variant1() + // else + // variant2() + tree dispatch_body = OMP_DISPATCH_BODY (expr); + if (TREE_CODE (dispatch_body) == BIND_EXPR) + dispatch_body = BIND_EXPR_BODY (dispatch_body); + if (TREE_CODE (dispatch_body) == STATEMENT_LIST) + { + // Fortran FE may insert some pre-call code, for instance when an + // array is passed as argument. Skip to the actual call. + dispatch_body = expr_last (dispatch_body); + } + gcc_assert (TREE_CODE (dispatch_body) == CALL_EXPR + || TREE_CODE (dispatch_body) == MODIFY_EXPR); + tree base_call_expr = dispatch_body; + tree dst; + if (TREE_CODE (base_call_expr) == MODIFY_EXPR) + { + dst = TREE_OPERAND (base_call_expr, 0); + base_call_expr = TREE_OPERAND (base_call_expr, 1); + while (TREE_CODE (base_call_expr) == FLOAT_EXPR + || TREE_CODE (base_call_expr) == CONVERT_EXPR + || TREE_CODE (base_call_expr) == COMPLEX_EXPR) + base_call_expr = TREE_OPERAND (base_call_expr, 0); + } + + tree base_fndecl = get_callee_fndecl (STRIP_NOPS (base_call_expr)); + if (base_fndecl != NULL_TREE) + { + if (DECL_VIRTUAL_P (base_fndecl)) + { + error_at ( + EXPR_LOCATION (base_call_expr), + "%qD is a virtual function but only a direct call is allowed " + "in a dispatch construct", + DECL_NAME (base_fndecl)); + } + + tree variant_fndecl = omp_resolve_declare_variant (base_fndecl); + if (base_fndecl != variant_fndecl + && (omp_has_novariants () == -1 || omp_has_nocontext () == -1)) + { + tree novariants_clause = NULL_TREE, nocontext_clause = NULL_TREE, + novariants_cond = NULL_TREE, nocontext_cond = NULL_TREE; + for (tree c = OMP_DISPATCH_CLAUSES (expr); c; c = TREE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NOVARIANTS) + { + gcc_assert (novariants_cond == NULL_TREE); + novariants_clause = c; + novariants_cond = OMP_CLAUSE_NOVARIANTS_EXPR (c); + } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NOCONTEXT) + { + gcc_assert (nocontext_cond == NULL_TREE); + nocontext_clause = c; + nocontext_cond = OMP_CLAUSE_NOCONTEXT_EXPR (c); + } + } + gcc_assert (novariants_cond != NULL_TREE + || nocontext_cond != NULL_TREE); + + enum gimplify_status ret + = gimplify_expr (&novariants_cond, &body, NULL, is_gimple_val, + fb_rvalue); + if (ret == GS_ERROR || ret == GS_UNHANDLED) + return ret; + ret = gimplify_expr (&nocontext_cond, &body, NULL, is_gimple_val, + fb_rvalue); + if (ret == GS_ERROR || ret == GS_UNHANDLED) + return ret; + + tree base_label = create_artificial_label (UNKNOWN_LOCATION); + tree variant1_label = create_artificial_label (UNKNOWN_LOCATION); + tree cond_label = create_artificial_label (UNKNOWN_LOCATION); + tree variant2_label = create_artificial_label (UNKNOWN_LOCATION); + tree end_label = create_artificial_label (UNKNOWN_LOCATION); + + if (novariants_cond != NULL_TREE) + { + gcond *novariants_cond_stmt + = gimple_build_cond_from_tree (novariants_cond, base_label, + cond_label); + gimplify_seq_add_stmt (&body, novariants_cond_stmt); + + gimplify_seq_add_stmt (&body, gimple_build_label (base_label)); + tree base_call_expr2 = copy_node (base_call_expr); + if (TREE_CODE (dispatch_body) == MODIFY_EXPR) + { + base_call_expr2 = build2 (MODIFY_EXPR, TREE_TYPE (dst), dst, + base_call_expr2); + } + OMP_CLAUSE_NOVARIANTS_EXPR (novariants_clause) + = boolean_true_node; + gimplify_and_add (base_call_expr2, &body); + gimplify_seq_add_stmt (&body, gimple_build_goto (end_label)); + + OMP_CLAUSE_NOVARIANTS_EXPR (novariants_clause) + = boolean_false_node; + } + + gimplify_seq_add_stmt (&body, gimple_build_label (cond_label)); + if (nocontext_cond != NULL_TREE) + { + gcond *nocontext_cond_stmt + = gimple_build_cond_from_tree (nocontext_cond, variant1_label, + variant2_label); + gimplify_seq_add_stmt (&body, nocontext_cond_stmt); + + gimplify_seq_add_stmt (&body, + gimple_build_label (variant1_label)); + tree variant_call_expr = copy_node (base_call_expr); + if (TREE_CODE (dispatch_body) == MODIFY_EXPR) + { + variant_call_expr = build2 (MODIFY_EXPR, TREE_TYPE (dst), dst, + variant_call_expr); + } + OMP_CLAUSE_NOCONTEXT_EXPR (nocontext_clause) = boolean_true_node; + gimplify_and_add (variant_call_expr, &body); + gimplify_seq_add_stmt (&body, gimple_build_goto (end_label)); + OMP_CLAUSE_NOCONTEXT_EXPR (nocontext_clause) = boolean_false_node; + } + + gimplify_seq_add_stmt (&body, gimple_build_label (variant2_label)); + tree variant_call_expr = copy_node (base_call_expr); + if (TREE_CODE (dispatch_body) == MODIFY_EXPR) + { + variant_call_expr + = build2 (MODIFY_EXPR, TREE_TYPE (dst), dst, variant_call_expr); + } + gimplify_and_add (variant_call_expr, &body); + gimplify_seq_add_stmt (&body, gimple_build_goto (end_label)); + gimplify_seq_add_stmt (&body, gimple_build_label (end_label)); + } + else + gimplify_and_add (OMP_DISPATCH_BODY (expr), &body); + } + else + gimplify_and_add (OMP_DISPATCH_BODY (expr), &body); + + // Wrap dispatch body into a bind + gimple *bind = gimple_build_bind (NULL_TREE, body, NULL_TREE); + pop_gimplify_context (bind); + + gimplify_adjust_omp_clauses (pre_p, bind, &OMP_DISPATCH_CLAUSES (expr), + OMP_DISPATCH); + + // Move relevant clauses to the task construct + tree task_clauses = NULL_TREE; + tree *task_clauses_ptr = &task_clauses; + bool has_nowait = false; + for (tree c = OMP_DISPATCH_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND) + { + *task_clauses_ptr = c; + task_clauses_ptr = &OMP_CLAUSE_CHAIN (c); + } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) + { + *task_clauses_ptr + = build_omp_clause (input_location, OMP_CLAUSE_SHARED); + OMP_CLAUSE_DECL (*task_clauses_ptr) = OMP_CLAUSE_DECL (c); + task_clauses_ptr = &OMP_CLAUSE_CHAIN (*task_clauses_ptr); + } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NOWAIT) + has_nowait = true; + } + *task_clauses_ptr = build_omp_clause (input_location, OMP_CLAUSE_IF); + OMP_CLAUSE_IF_EXPR (*task_clauses_ptr) + = has_nowait ? boolean_true_node : boolean_false_node; + + // Wrap bind into a task + gimple *task + = gimple_build_omp_task (bind, task_clauses, NULL_TREE, NULL_TREE, + NULL_TREE, NULL_TREE, NULL_TREE); + + gimple *stmt = gimple_build_omp_dispatch (task, OMP_DISPATCH_CLAUSES (expr)); + gimplify_seq_add_stmt (pre_p, stmt); + *expr_p = NULL_TREE; + return GS_ALL_DONE; +} + /* Convert the GENERIC expression tree *EXPR_P to GIMPLE. If the expression produces a value to be used as an operand inside a GIMPLE statement, the value will be stored back in *EXPR_P. This value will @@ -18540,6 +18936,10 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, ret = gimplify_omp_atomic (expr_p, pre_p); break; + case OMP_DISPATCH: + ret = gimplify_omp_dispatch (expr_p, pre_p); + break; + case TRANSACTION_EXPR: ret = gimplify_transaction (expr_p, pre_p); break; @@ -18865,7 +19265,8 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, && code != OMP_SECTION && code != OMP_STRUCTURED_BLOCK && code != OMP_SINGLE - && code != OMP_SCOPE); + && code != OMP_SCOPE + && code != OMP_DISPATCH); } #endif diff --git a/gcc/gimplify.h b/gcc/gimplify.h index ac3cc8eb552..55aece2b65b 100644 --- a/gcc/gimplify.h +++ b/gcc/gimplify.h @@ -77,6 +77,8 @@ extern enum gimplify_status gimplify_expr (tree *, gimple_seq *, gimple_seq *, bool (*) (tree), fallback_t); int omp_construct_selector_matches (enum tree_code *, int, int *); +int omp_has_novariants (void); +int omp_has_nocontext (void); extern void gimplify_type_sizes (tree, gimple_seq *); extern void gimplify_one_sizepos (tree *, gimple_seq *); diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index 044d5d087b6..c83edabbcc3 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -76,6 +76,12 @@ DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_TEAM_NUM, "omp_get_team_num", BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST) DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_NUM_TEAMS, "omp_get_num_teams", BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST) +DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_MAPPED_PTR, "omp_get_mapped_ptr", + BT_FN_PTR_CONST_PTR_INT, ATTR_NOTHROW_LEAF_LIST) +DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_DEFAULT_DEVICE, "omp_get_default_device", + BT_FN_INT, ATTR_NOTHROW_LEAF_LIST) +DEF_GOMP_BUILTIN (BUILT_IN_OMP_SET_DEFAULT_DEVICE, "omp_set_default_device", + BT_FN_INT, ATTR_NOTHROW_LEAF_LIST) DEF_GOMP_BUILTIN (BUILT_IN_GOMP_ATOMIC_START, "GOMP_atomic_start", BT_FN_VOID, ATTR_NOTHROW_LEAF_LIST) diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc index 24287826444..6fa372a550b 100644 --- a/gcc/omp-expand.cc +++ b/gcc/omp-expand.cc @@ -8636,6 +8636,19 @@ expand_omp_single (struct omp_region *region) single_succ_edge (exit_bb)->flags = EDGE_FALLTHRU; } +/* Expand code for an OpenMP dispatch directive... */ + +static void +expand_omp_dispatch (struct omp_region *region) +{ + basic_block entry_bb = region->entry; + gimple_stmt_iterator si = gsi_last_nondebug_bb (entry_bb); + enum gimple_code code = gimple_code (gsi_stmt (si)); + gcc_assert (code == GIMPLE_OMP_DISPATCH); + gsi_remove (&si, true); + single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU; +} + /* Generic expansion for OpenMP synchronization directives: master, ordered and critical. All we need to do here is remove the entry and exit markers for REGION. */ @@ -10654,6 +10667,10 @@ expand_omp (struct omp_region *region) expand_omp_single (region); break; + case GIMPLE_OMP_DISPATCH: + expand_omp_dispatch (region); + break; + case GIMPLE_OMP_ORDERED: { gomp_ordered *ord_stmt @@ -11001,6 +11018,7 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region, case GIMPLE_OMP_MASTER: case GIMPLE_OMP_MASKED: case GIMPLE_OMP_SCOPE: + case GIMPLE_OMP_DISPATCH: case GIMPLE_OMP_CRITICAL: case GIMPLE_OMP_SECTION: cur_region = new_omp_region (bb, code, cur_region); diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc index 2c095200d5b..d585df400d1 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -1037,7 +1037,7 @@ omp_construct_traits_to_codes (tree ctx, int nconstructs, /* Order must match the OMP_TRAIT_CONSTRUCT_* enumerators in enum omp_ts_code. */ static enum tree_code code_map[] - = { OMP_TARGET, OMP_TEAMS, OMP_PARALLEL, OMP_FOR, OMP_SIMD }; + = { OMP_TARGET, OMP_TEAMS, OMP_PARALLEL, OMP_FOR, OMP_SIMD, OMP_DISPATCH }; for (tree ts = ctx; ts; ts = TREE_CHAIN (ts), i--) { @@ -1130,6 +1130,7 @@ const char *omp_tss_map[] = "target_device", "implementation", "user", + "need_device_ptr", NULL }; @@ -1236,10 +1237,14 @@ struct omp_ts_info omp_ts_map[] = OMP_TRAIT_PROPERTY_CLAUSE_LIST, false, NULL }, + { "dispatch", + (1 << OMP_TRAIT_SET_CONSTRUCT), + OMP_TRAIT_PROPERTY_NONE, false, + NULL + }, { NULL, 0, OMP_TRAIT_PROPERTY_NONE, false, NULL } /* OMP_TRAIT_LAST */ }; - /* Return a name from PROP, a property in selectors accepting name lists. */ @@ -1445,6 +1450,8 @@ omp_context_selector_matches (tree ctx) for (tree tss = ctx; tss; tss = TREE_CHAIN (tss)) { enum omp_tss_code set = OMP_TSS_CODE (tss); + if (set == OMP_TRAIT_SET_NEED_DEVICE_PTR) + continue; tree selectors = OMP_TSS_TRAIT_SELECTORS (tss); /* Immediately reject the match if there are any ignored @@ -2484,6 +2491,9 @@ omp_resolve_declare_variant (tree base) if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0) return omp_resolve_late_declare_variant (base); + if (omp_has_novariants () == 1) + return base; + auto_vec variants; auto_vec defer; bool any_deferred = false; @@ -2630,6 +2640,8 @@ omp_resolve_declare_variant (tree base) (*slot)->variants = entry.variants; tree alt = build_decl (DECL_SOURCE_LOCATION (base), FUNCTION_DECL, DECL_NAME (base), TREE_TYPE (base)); + if (DECL_ASSEMBLER_NAME_SET_P (base)) + SET_DECL_ASSEMBLER_NAME (alt, DECL_ASSEMBLER_NAME (base)); DECL_ARTIFICIAL (alt) = 1; DECL_IGNORED_P (alt) = 1; TREE_STATIC (alt) = 1; diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 4d003f42098..693d8ca7d8d 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -4185,6 +4185,11 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, scan_omp (gimple_omp_body_ptr (stmt), ctx); break; + case GIMPLE_OMP_DISPATCH: + ctx = new_omp_context (stmt, ctx); + scan_omp (gimple_omp_body_ptr (stmt), ctx); + break; + case GIMPLE_OMP_SECTIONS: scan_omp_sections (as_a (stmt), ctx); break; @@ -8926,6 +8931,31 @@ lower_omp_scope (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (BLOCK_VARS (block)) TREE_USED (block) = 1; } + +/* Lower code for an OMP dispatch directive. */ + +static void +lower_omp_dispatch (gimple_stmt_iterator *gsi_p, omp_context *ctx) +{ + tree block; + gimple *stmt = gsi_stmt (*gsi_p); + gbind *bind; + + push_gimplify_context (); + + block = make_node (BLOCK); + bind = gimple_build_bind (NULL, NULL, block); + gsi_replace (gsi_p, bind, true); + + lower_omp (gimple_omp_body_ptr (stmt), ctx); + gimple_bind_set_body (bind, maybe_catch_exception (gimple_omp_body (stmt))); + + pop_gimplify_context (bind); + + gimple_bind_append_vars (bind, ctx->block_vars); + BLOCK_VARS (block) = ctx->block_vars; +} + /* Expand code for an OpenMP master or masked directive. */ static void @@ -14399,6 +14429,11 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_assert (ctx); lower_omp_scope (gsi_p, ctx); break; + case GIMPLE_OMP_DISPATCH: + ctx = maybe_lookup_ctx (stmt); + gcc_assert (ctx); + lower_omp_dispatch (gsi_p, ctx); + break; case GIMPLE_OMP_SINGLE: ctx = maybe_lookup_ctx (stmt); gcc_assert (ctx); diff --git a/gcc/tree-inline.cc b/gcc/tree-inline.cc index f31a34ac410..2e06b706025 100644 --- a/gcc/tree-inline.cc +++ b/gcc/tree-inline.cc @@ -1679,6 +1679,12 @@ remap_gimple_stmt (gimple *stmt, copy_body_data *id) (s1, gimple_omp_scope_clauses (stmt)); break; + case GIMPLE_OMP_DISPATCH: + s1 = remap_gimple_seq (gimple_omp_body (stmt), id); + copy = gimple_build_omp_dispatch (s1, + gimple_omp_dispatch_clauses (stmt)); + break; + case GIMPLE_OMP_TASKGROUP: s1 = remap_gimple_seq (gimple_omp_body (stmt), id); copy = gimple_build_omp_taskgroup @@ -4609,6 +4615,7 @@ estimate_num_insns (gimple *stmt, eni_weights *weights) case GIMPLE_OMP_MASTER: case GIMPLE_OMP_MASKED: case GIMPLE_OMP_SCOPE: + case GIMPLE_OMP_DISPATCH: case GIMPLE_OMP_TASKGROUP: case GIMPLE_OMP_ORDERED: case GIMPLE_OMP_SCAN: From patchwork Mon May 27 11:54:35 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 1939910 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=g6TIC9sW; 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 4VnvL40F8qz20KL for ; Mon, 27 May 2024 21:58:20 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 401C2388300E for ; Mon, 27 May 2024 11:58:18 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x333.google.com (mail-wm1-x333.google.com [IPv6:2a00:1450:4864:20::333]) by sourceware.org (Postfix) with ESMTPS id 2817D3870C14 for ; Mon, 27 May 2024 11:55:34 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 2817D3870C14 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 2817D3870C14 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::333 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810938; cv=none; b=tJ7kYG4gthc6oEmELdhdXFcoW2gSsPJN8JAkdTNliYCbuVqd6/5hkikTZ7H4SiOtZZEsOgjrGBkTf9CkkrZ1Nm6BB7yGRm0uGqh5zajshey7WcyJwRMzl/MMStHSe+NtY8aiuZFL4Yw5E9+HMaEmW699t4jA19Hv4PKuD9J88Qs= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810938; c=relaxed/simple; bh=yIobTv42+nqN65mAmbJfbV4ZWZeprQIttBN/MYYiqSg=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=IfJ51vhEXqEuHF/sOM7HS25rgA1OxwFt4lio8YmkJzZXnMSW1In1dWVeOe4zvlZtHiXMwsrYD1abTNC053BGbtKo7BKuKZZva/xPIO1qKvMC/ZTJI2lQzyFMC5DcvrfBLe/jgqpQx7wOv+WHe9Rp4hSjOjBBJASLYQr79LYW700= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x333.google.com with SMTP id 5b1f17b1804b1-4211249fbafso8344015e9.3 for ; Mon, 27 May 2024 04:55:34 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1716810932; x=1717415732; 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=jlItdgYQFtBMfxKMxcjUzoARhAhspLAxnx+bXcMcNe0=; b=g6TIC9sWFae9AI7sQjfpczDHd5utxSXuEmyePqq9IUwZ6d0PcMNbijNksKN9muAyaz +aXzsZvy9BmtxS4XrTvh1jia+9KMtkagGHEqjxoiHb4CPFEaI/S4Y93+3/txqFw9hDyF tp3AJYoK4n4XUfnafiWQhPbXl2H6zOQxIBeO/rmErZhlpSq4YFVUY76RQOswui+UR9ba bLL89ip0fex4Uos9VTVYq0Vw7cdKHLg0BmjRUEDQQU5YT/L5d3VCBn13xUca4G/WgAdq i/azkakKMfDlGo3gi41K2jBqXkTo+etOMJKiy1WlBh2xIYHIccAPYvedqiHIpUWLOXxj CVfw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1716810932; x=1717415732; 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=jlItdgYQFtBMfxKMxcjUzoARhAhspLAxnx+bXcMcNe0=; b=Hyl2bkxmMWyDz1Hs2pa5DXToefW++SPLgZ4lXfIgaUHcG2bLBWwCRxlcbbI5ABlerK vZC8QEcdyvNLP+kwIhY6ikPoJW5XdhwMeqffV+jR4U2UbTTDjFpjPgNyyQUinROS8TEK fECuiUaXtJsFeH8gP4Rz4qySwzsgr5PBusgZlBJy/IKxBcF5GByx1I87z9Bw1qsfyDi7 /vQ4X7lh3abkPTCqgS2FInfAg7kcql0GHoA7hDnp1ca4MYsP0fRipeA2nsrqVNHXBqAa SFwcYfnTQtiQpXhsXy3fklhHXCvbDk0rp9sXRapDljK/d94i+EQd76mDI/CrZa4yqICJ MlnQ== X-Gm-Message-State: AOJu0YxyffpVTFT+nQx37yRoWixJyJbuTqc4fhTmSMfUSnn7VoE0WrNE +RVf3Yabrk/FsV3i6knKVP3x9t0MNrtVOFGBTKJkiXFIjWtfrXIscoFx61NMR84Hpe2D8kaozqh a X-Google-Smtp-Source: AGHT+IEc4o8CiCjZmiZlDMX+h+9R6ywO0FnUt1g5L0RCtpx7lYoOA02sxzJrq+9I9gpvCiWWF+eXwA== X-Received: by 2002:a05:600c:19d2:b0:418:2ccf:cbc7 with SMTP id 5b1f17b1804b1-421089ebd3dmr65403295e9.2.1716810932299; Mon, 27 May 2024 04:55:32 -0700 (PDT) Received: from localhost.localdomain ([2a0d:3344:2340:ce10::2f3]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-42108966b58sm107636955e9.3.2024.05.27.04.55.31 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Mon, 27 May 2024 04:55:31 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH 3/7] OpenMP: C front-end support for dispatch + adjust_args Date: Mon, 27 May 2024 13:54:35 +0200 Message-ID: <20240527115439.3967217-4-parras@baylibre.com> X-Mailer: git-send-email 2.45.1 In-Reply-To: <20240527115439.3967217-1-parras@baylibre.com> References: <20240527115439.3967217-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.2 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds support to the C front-end to parse the `dispatch` construct and the `adjust_args` clause. It also includes some common C/C++ bits for pragmas and attributes. Additional common C/C++ testcases are in a later patch in the series. gcc/c-family/ChangeLog: * c-attribs.cc (c_common_gnu_attributes): Add attribute for adjust_args need_device_ptr. * c-omp.cc (c_omp_directives): Uncomment dispatch. * c-pragma.cc (omp_pragmas): Add dispatch. * c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_DISPATCH. (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_NOCONTEXT and PRAGMA_OMP_CLAUSE_NOVARIANTS. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_dispatch): New function. (c_parser_omp_clause_name): Handle nocontext and novariants clauses. (c_parser_omp_clause_novariants): New function. (c_parser_omp_clause_nocontext): Likewise. (c_parser_omp_all_clauses): Handle nocontext and novariants clauses. (c_parser_omp_dispatch_body): New function adapted from c_parser_expr_no_commas. (OMP_DISPATCH_CLAUSE_MASK): Define. (c_parser_omp_dispatch): New function. (c_finish_omp_declare_variant): Parse adjust_args. (c_parser_omp_construct): Handle PRAGMA_OMP_DISPATCH. * c-typeck.cc (c_finish_omp_clauses): Handle OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. gcc/testsuite/ChangeLog: * gcc.dg/gomp/adjust-args-1.c: New test. * gcc.dg/gomp/dispatch-1.c: New test. --- gcc/c-family/c-attribs.cc | 2 + gcc/c-family/c-omp.cc | 4 +- gcc/c-family/c-pragma.cc | 1 + gcc/c-family/c-pragma.h | 3 + gcc/c/c-parser.cc | 483 +++++++++++++++++++--- gcc/c/c-typeck.cc | 2 + gcc/testsuite/gcc.dg/gomp/adjust-args-1.c | 32 ++ gcc/testsuite/gcc.dg/gomp/dispatch-1.c | 53 +++ libgomp/testsuite/libgomp.c/dispatch-1.c | 76 ++++ 9 files changed, 601 insertions(+), 55 deletions(-) create mode 100644 gcc/testsuite/gcc.dg/gomp/adjust-args-1.c create mode 100644 gcc/testsuite/gcc.dg/gomp/dispatch-1.c create mode 100644 libgomp/testsuite/libgomp.c/dispatch-1.c diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc index 04e39b41bdf..860a068d527 100644 --- a/gcc/c-family/c-attribs.cc +++ b/gcc/c-family/c-attribs.cc @@ -556,6 +556,8 @@ const struct attribute_spec c_common_gnu_attributes[] = handle_omp_declare_variant_attribute, NULL }, { "omp declare variant variant", 0, -1, true, false, false, false, handle_omp_declare_variant_attribute, NULL }, + { "omp declare variant adjust_args need_device_ptr", 0, -1, true, false, false, false, + handle_omp_declare_variant_attribute, NULL }, { "simd", 0, 1, true, false, false, false, handle_simd_attribute, NULL }, { "omp declare target", 0, -1, true, false, false, false, diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index c0e02aa422f..e6b42dbd01c 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -4196,8 +4196,8 @@ const struct c_omp_directive c_omp_directives[] = { C_OMP_DIR_DECLARATIVE, false }, { "depobj", nullptr, nullptr, PRAGMA_OMP_DEPOBJ, C_OMP_DIR_STANDALONE, false }, - /* { "dispatch", nullptr, nullptr, PRAGMA_OMP_DISPATCH, - C_OMP_DIR_CONSTRUCT, false }, */ + { "dispatch", nullptr, nullptr, PRAGMA_OMP_DISPATCH, + C_OMP_DIR_DECLARATIVE, false }, { "distribute", nullptr, nullptr, PRAGMA_OMP_DISTRIBUTE, C_OMP_DIR_CONSTRUCT, true }, { "end", "assumes", nullptr, PRAGMA_OMP_END, diff --git a/gcc/c-family/c-pragma.cc b/gcc/c-family/c-pragma.cc index 1237ee6e62b..60fadeee286 100644 --- a/gcc/c-family/c-pragma.cc +++ b/gcc/c-family/c-pragma.cc @@ -1526,6 +1526,7 @@ static const struct omp_pragma_def omp_pragmas[] = { { "cancellation", PRAGMA_OMP_CANCELLATION_POINT }, { "critical", PRAGMA_OMP_CRITICAL }, { "depobj", PRAGMA_OMP_DEPOBJ }, + { "dispatch", PRAGMA_OMP_DISPATCH }, { "error", PRAGMA_OMP_ERROR }, { "end", PRAGMA_OMP_END }, { "flush", PRAGMA_OMP_FLUSH }, diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index ce93a52fa57..061a83d1716 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -55,6 +55,7 @@ enum pragma_kind { PRAGMA_OMP_CRITICAL, PRAGMA_OMP_DECLARE, PRAGMA_OMP_DEPOBJ, + PRAGMA_OMP_DISPATCH, PRAGMA_OMP_DISTRIBUTE, PRAGMA_OMP_ERROR, PRAGMA_OMP_END, @@ -132,9 +133,11 @@ enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_LINK, PRAGMA_OMP_CLAUSE_MAP, PRAGMA_OMP_CLAUSE_MERGEABLE, + PRAGMA_OMP_CLAUSE_NOCONTEXT, PRAGMA_OMP_CLAUSE_NOGROUP, PRAGMA_OMP_CLAUSE_NONTEMPORAL, PRAGMA_OMP_CLAUSE_NOTINBRANCH, + PRAGMA_OMP_CLAUSE_NOVARIANTS, PRAGMA_OMP_CLAUSE_NOWAIT, PRAGMA_OMP_CLAUSE_NUM_TASKS, PRAGMA_OMP_CLAUSE_NUM_TEAMS, diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 00f8bf4376e..c9cd36fb429 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -1733,6 +1733,8 @@ static void c_parser_omp_assumption_clauses (c_parser *, bool); static void c_parser_omp_allocate (c_parser *); static void c_parser_omp_assumes (c_parser *); static bool c_parser_omp_ordered (c_parser *, enum pragma_context, bool *); +static tree +c_parser_omp_dispatch (location_t, c_parser *); static void c_parser_oacc_routine (c_parser *, enum pragma_context); /* These Objective-C parser functions are only ever called when @@ -14972,6 +14974,8 @@ c_parser_omp_clause_name (c_parser *parser) case 'n': if (!strcmp ("no_create", p)) result = PRAGMA_OACC_CLAUSE_NO_CREATE; + else if (!strcmp ("nocontext", p)) + result = PRAGMA_OMP_CLAUSE_NOCONTEXT; else if (!strcmp ("nogroup", p)) result = PRAGMA_OMP_CLAUSE_NOGROUP; else if (!strcmp ("nohost", p)) @@ -14980,6 +14984,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_NONTEMPORAL; else if (!strcmp ("notinbranch", p)) result = PRAGMA_OMP_CLAUSE_NOTINBRANCH; + else if (!strcmp ("novariants", p)) + result = PRAGMA_OMP_CLAUSE_NOVARIANTS; else if (!strcmp ("nowait", p)) result = PRAGMA_OMP_CLAUSE_NOWAIT; else if (!strcmp ("num_gangs", p)) @@ -19190,6 +19196,60 @@ c_parser_omp_clause_uniform (c_parser *parser, tree list) return list; } +/* OpenMP 5.1 + novariants ( scalar-expression ) */ + +static tree +c_parser_omp_clause_novariants (c_parser *parser, tree list) +{ + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + location_t loc = c_parser_peek_token (parser)->location; + c_expr expr = c_parser_expr_no_commas (parser, NULL); + tree t = convert_lvalue_to_rvalue (loc, expr, true, true).value; + t = c_objc_common_truthvalue_conversion (loc, t); + t = c_fully_fold (t, false, NULL); + parens.skip_until_found_close (parser); + + check_no_duplicate_clause (list, OMP_CLAUSE_NOVARIANTS, "novariants"); + + tree c = build_omp_clause (loc, OMP_CLAUSE_NOVARIANTS); + OMP_CLAUSE_NOVARIANTS_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + + return list; +} + +/* OpenMP 5.1 + nocontext ( scalar-expression ) */ + +static tree +c_parser_omp_clause_nocontext (c_parser *parser, tree list) +{ + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + location_t loc = c_parser_peek_token (parser)->location; + c_expr expr = c_parser_expr_no_commas (parser, NULL); + tree t = convert_lvalue_to_rvalue (loc, expr, true, true).value; + t = c_objc_common_truthvalue_conversion (loc, t); + t = c_fully_fold (t, false, NULL); + parens.skip_until_found_close (parser); + + check_no_duplicate_clause (list, OMP_CLAUSE_NOCONTEXT, "nocontext"); + + tree c = build_omp_clause (loc, OMP_CLAUSE_NOCONTEXT); + OMP_CLAUSE_NOCONTEXT_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + + return list; +} + /* OpenMP 5.0: detach ( event-handle ) */ @@ -19801,6 +19861,14 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, clauses); c_name = "enter"; break; + case PRAGMA_OMP_CLAUSE_NOVARIANTS: + c_name = "novariants"; + clauses = c_parser_omp_clause_novariants (parser, clauses); + break; + case PRAGMA_OMP_CLAUSE_NOCONTEXT: + c_name = "nocontext"; + clauses = c_parser_omp_clause_nocontext (parser, clauses); + break; default: c_parser_error (parser, "expected an OpenMP clause"); goto saw_error; @@ -23458,6 +23526,168 @@ c_parser_omp_scope (location_t loc, c_parser *parser, bool *if_p) return add_stmt (stmt); } +// Adapted from c_parser_expr_no_commas +static tree +c_parser_omp_dispatch_body (c_parser *parser) +{ + struct c_expr lhs, rhs, ret; + struct c_expr orig_expr; + location_t expr_loc = c_parser_peek_token (parser)->location; + source_range tok_range = c_parser_peek_token (parser)->get_range (); + location_t sizeof_arg_loc[3]; + tree sizeof_arg[3]; + vec *exprlist; + vec arg_loc = vNULL; + vec *origtypes = NULL; + unsigned int literal_zero_mask; + location_t start; + location_t finish; + + lhs = c_parser_conditional_expression (parser, NULL, NULL); + if (TREE_CODE (lhs.value) == CALL_EXPR) + return lhs.value; + else + { + location_t op_location = c_parser_peek_token (parser)->location; + if (!c_parser_require (parser, CPP_EQ, "expected %<=%>")) + return error_mark_node; + + /* Parse function name*/ + if (!c_parser_next_token_is (parser, CPP_NAME)) + { + c_parser_error (parser, "expected a function name"); + rhs.set_error (); + return rhs.value; + } + expr_loc = c_parser_peek_token (parser)->location; + tree id = c_parser_peek_token (parser)->value; + c_parser_consume_token (parser); + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + return error_mark_node; + + rhs.value = build_external_ref (expr_loc, id, true, &rhs.original_type); + set_c_expr_source_range (&rhs, tok_range); + /* Parse argument list */ + { + for (int i = 0; i < 3; i++) + { + sizeof_arg[i] = NULL_TREE; + sizeof_arg_loc[i] = UNKNOWN_LOCATION; + } + literal_zero_mask = 0; + if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN)) + exprlist = NULL; + else + exprlist = c_parser_expr_list (parser, true, false, &origtypes, + sizeof_arg_loc, sizeof_arg, &arg_loc, + &literal_zero_mask); + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>"); + } + orig_expr = rhs; + mark_exp_read (rhs.value); + if (warn_sizeof_pointer_memaccess) + sizeof_pointer_memaccess_warning (sizeof_arg_loc, rhs.value, exprlist, + sizeof_arg, + sizeof_ptr_memacc_comptypes); + if (TREE_CODE (rhs.value) == FUNCTION_DECL) + { + if (fndecl_built_in_p (rhs.value, BUILT_IN_MEMSET) + && vec_safe_length (exprlist) == 3) + { + tree arg0 = (*exprlist)[0]; + tree arg2 = (*exprlist)[2]; + warn_for_memset (expr_loc, arg0, arg2, literal_zero_mask); + } + if (warn_absolute_value + && fndecl_built_in_p (rhs.value, BUILT_IN_NORMAL) + && vec_safe_length (exprlist) == 1) + warn_for_abs (expr_loc, rhs.value, (*exprlist)[0]); + if (parser->omp_for_parse_state + && parser->omp_for_parse_state->in_intervening_code + && omp_runtime_api_call (rhs.value)) + { + error_at (expr_loc, "calls to the OpenMP runtime API are " + "not permitted in intervening code"); + parser->omp_for_parse_state->fail = true; + } + } + + start = rhs.get_start (); + finish = parser->tokens_buf[0].get_finish (); + rhs.value = c_build_function_call_vec (expr_loc, arg_loc, rhs.value, + exprlist, origtypes); + set_c_expr_source_range (&rhs, start, finish); + rhs.m_decimal = 0; + + rhs.original_code = ERROR_MARK; + if (TREE_CODE (rhs.value) == INTEGER_CST + && TREE_CODE (orig_expr.value) == FUNCTION_DECL + && fndecl_built_in_p (orig_expr.value, BUILT_IN_CONSTANT_P)) + rhs.original_code = C_MAYBE_CONST_EXPR; + rhs.original_type = NULL; + if (exprlist) + { + release_tree_vector (exprlist); + release_tree_vector (origtypes); + } + arg_loc.release (); + + /* Build assignment */ + rhs = convert_lvalue_to_rvalue (expr_loc, rhs, true, true); + ret.value + = build_modify_expr (op_location, lhs.value, lhs.original_type, + NOP_EXPR, expr_loc, rhs.value, rhs.original_type); + ret.m_decimal = 0; + set_c_expr_source_range (&ret, lhs.get_start (), rhs.get_finish ()); + ret.original_code = MODIFY_EXPR; + ret.original_type = NULL; + return ret.value; + } +} + +/* OpenMP 5.1: + # pragma omp dispatch dispatch-clause[optseq] new-line + expression-stmt + + LOC is the location of the #pragma. +*/ + +#define OMP_DISPATCH_CLAUSE_MASK \ + ((OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT)) + +static tree +c_parser_omp_dispatch (location_t loc, c_parser *parser) +{ + tree stmt = make_node (OMP_DISPATCH); + SET_EXPR_LOCATION (stmt, loc); + TREE_TYPE (stmt) = void_type_node; + + OMP_DISPATCH_CLAUSES (stmt) + = c_parser_omp_all_clauses (parser, OMP_DISPATCH_CLAUSE_MASK, + "#pragma omp dispatch"); + + // Parse body as expression statement + loc = c_parser_peek_token (parser)->location; + tree dispatch_body = c_parser_omp_dispatch_body (parser); + if (dispatch_body == error_mark_node) + { + inform (loc, "%<#pragma omp dispatch%> must be followed by a function " + "call with optional assignment"); + c_parser_skip_to_end_of_block_or_statement (parser); + return NULL_TREE; + } + + c_parser_skip_until_found (parser, CPP_SEMICOLON, "expected %<;%>"); + OMP_DISPATCH_BODY (stmt) = dispatch_body; + + return add_stmt (stmt); +} + /* OpenMP 3.0: # pragma omp task task-clause[optseq] new-line @@ -24438,6 +24668,10 @@ check_clauses: OpenMP 5.0: # pragma omp declare variant (identifier) match(context-selector) new-line + + OpenMP 5.1 + # pragma omp declare variant (identifier) match(context-selector) \ + adjust_args(adjust-op:argument-list) new-line */ #define OMP_DECLARE_SIMD_CLAUSE_MASK \ @@ -24901,19 +25135,39 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms) parens.require_close (parser); - if (c_parser_next_token_is (parser, CPP_COMMA) - && c_parser_peek_2nd_token (parser)->type == CPP_NAME) - c_parser_consume_token (parser); + vec adjust_args_list = vNULL; + bool has_match = false, has_adjust_args = false; + location_t adjust_args_loc; - const char *clause = ""; - location_t match_loc = c_parser_peek_token (parser)->location; - if (c_parser_next_token_is (parser, CPP_NAME)) - clause = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); - if (strcmp (clause, "match")) + do { - c_parser_error (parser, "expected %"); - goto fail; - } + if (c_parser_next_token_is (parser, CPP_COMMA) + && c_parser_peek_2nd_token (parser)->type == CPP_NAME) + c_parser_consume_token (parser); + + const char *clause = ""; + location_t match_loc = c_parser_peek_token (parser)->location; + if (c_parser_next_token_is (parser, CPP_NAME)) + clause = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); + + enum clause + { + match, + adjust_args + } ccode; + + if (strcmp (clause, "match") == 0) + ccode = match; + else if (strcmp (clause, "adjust_args") == 0) + { + ccode = adjust_args; + adjust_args_loc = match_loc; + } + else + { + c_parser_error (parser, "expected % or %"); + goto fail; + } c_parser_consume_token (parser); @@ -24923,55 +25177,172 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms) if (parms == NULL_TREE) parms = error_mark_node; - tree ctx = c_parser_omp_context_selector_specification (parser, parms); - if (ctx == error_mark_node) - goto fail; - ctx = omp_check_context_selector (match_loc, ctx); - if (ctx != error_mark_node && variant != error_mark_node) + if (ccode == match) { - if (TREE_CODE (variant) != FUNCTION_DECL) + has_match = true; + tree ctx = c_parser_omp_context_selector_specification (parser, parms); + if (ctx == error_mark_node) + goto fail; + ctx = omp_check_context_selector (match_loc, ctx); + if (ctx != error_mark_node && variant != error_mark_node) { - error_at (token->location, "variant %qD is not a function", variant); - variant = error_mark_node; - } - else if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT, - OMP_TRAIT_CONSTRUCT_SIMD) - && !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant))) - { - error_at (token->location, "variant %qD and base %qD have " - "incompatible types", variant, fndecl); - variant = error_mark_node; - } - else if (fndecl_built_in_p (variant) - && (strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), - "__builtin_", strlen ("__builtin_")) == 0 - || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), - "__sync_", strlen ("__sync_")) == 0 - || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), - "__atomic_", strlen ("__atomic_")) == 0)) - { - error_at (token->location, "variant %qD is a built-in", variant); - variant = error_mark_node; - } - if (variant != error_mark_node) - { - C_DECL_USED (variant) = 1; - tree construct - = omp_get_context_selector_list (ctx, OMP_TRAIT_SET_CONSTRUCT); - omp_mark_declare_variant (match_loc, variant, construct); - if (omp_context_selector_matches (ctx)) + if (TREE_CODE (variant) != FUNCTION_DECL) { - tree attr - = tree_cons (get_identifier ("omp declare variant base"), - build_tree_list (variant, ctx), - DECL_ATTRIBUTES (fndecl)); - DECL_ATTRIBUTES (fndecl) = attr; + error_at (token->location, "variant %qD is not a function", + variant); + variant = error_mark_node; + } + else if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT, + OMP_TRAIT_CONSTRUCT_SIMD) + && !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant))) + { + error_at (token->location, + "variant %qD and base %qD have " + "incompatible types", + variant, fndecl); + variant = error_mark_node; + } + else if (fndecl_built_in_p (variant) + && (strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), + "__builtin_", strlen ("__builtin_")) + == 0 + || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), + "__sync_", strlen ("__sync_")) + == 0 + || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), + "__atomic_", strlen ("__atomic_")) + == 0)) + { + error_at (token->location, "variant %qD is a built-in", variant); + variant = error_mark_node; + } + if (variant != error_mark_node) + { + C_DECL_USED (variant) = 1; + tree construct + = omp_get_context_selector_list (ctx, OMP_TRAIT_SET_CONSTRUCT); + omp_mark_declare_variant (match_loc, variant, construct); + if (omp_context_selector_matches (ctx)) + { + tree attr + = tree_cons (get_identifier ("omp declare variant base"), + build_tree_list (variant, ctx), + DECL_ATTRIBUTES (fndecl)); + DECL_ATTRIBUTES (fndecl) = attr; + } } } } + else if (ccode == adjust_args) + { + has_adjust_args = true; + if (c_parser_next_token_is (parser, CPP_NAME) + && c_parser_peek_2nd_token (parser)->type == CPP_COLON) + { + const char *p + = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); + if (strcmp (p, "need_device_ptr") == 0 + || strcmp (p, "nothing") == 0) + { + c_parser_consume_token (parser); // need_device_ptr + c_parser_consume_token (parser); // : + + location_t loc = c_parser_peek_token (parser)->location; + tree list + = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_ERROR, + NULL_TREE); + + tree arg; + for (tree c = list; c != NULL_TREE; c = TREE_CHAIN (c)) + { + tree decl = TREE_PURPOSE (c); + int idx; + for (arg = parms, idx = 0; arg != NULL; + arg = TREE_CHAIN (arg), idx++) + if (arg == decl) + break; + if (arg == NULL_TREE) + { + error_at (loc, "%qD is not a function argument", + decl); + goto fail; + } + if (adjust_args_list.contains (arg)) + { + error_at (loc, "%qD is specified more than once", + decl); + goto fail; + } + if (strcmp (p, "need_device_ptr") == 0 + && TREE_CODE (TREE_TYPE (arg)) != POINTER_TYPE) + { + error_at (loc, "%qD is not a C pointer", decl); + goto fail; + } + adjust_args_list.safe_push (arg); + if (strcmp (p, "need_device_ptr") == 0) + { + tree variant_decl = (TREE_CODE (variant) == NOP_EXPR) + ? TREE_OPERAND (variant, 0) + : variant; + tree variant_parm = DECL_ARGUMENTS (variant_decl); + for (int i = 0; i < idx; i++) + { + variant_parm = TREE_CHAIN (variant_parm); + gcc_assert (variant_parm != NULL); + } + tree attr = tree_cons ( + get_identifier ("omp declare variant adjust_args " + "need_device_ptr"), + NULL_TREE, DECL_ATTRIBUTES (variant_parm)); + DECL_ATTRIBUTES (variant_parm) = attr; + } + } + } + else + { + error_at (c_parser_peek_token (parser)->location, + "expected % or %"); + goto fail; + } + } + else + { + error_at (c_parser_peek_token (parser)->location, + "expected % or % " + "followed by %<:%>"); + goto fail; + } + } parens.require_close (parser); + } while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL)); c_parser_skip_to_pragma_eol (parser); + + if (has_adjust_args) + { + if (!has_match) + { + error_at ( + adjust_args_loc, + "an % clause can only be specified if the " + "% selector of the construct selector set appears " + "in the % clause"); + } + else + { + tree attr = lookup_attribute ("omp declare variant base", + DECL_ATTRIBUTES (fndecl)); + tree ctx = TREE_VALUE (TREE_VALUE (attr)); + if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT, + OMP_TRAIT_CONSTRUCT_DISPATCH)) + error_at ( + adjust_args_loc, + "an % clause can only be specified if the " + "% selector of the construct selector set appears " + "in the % clause"); + } + } } /* Finalize #pragma omp declare simd or #pragma omp declare variant @@ -25788,7 +26159,6 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context) types.release (); } - /* OpenMP 4.0 #pragma omp declare simd declare-simd-clauses[optseq] new-line #pragma omp declare reduction (reduction-id : typename-list : expression) \ @@ -25796,7 +26166,11 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context) #pragma omp declare target new-line OpenMP 5.0 - #pragma omp declare variant (identifier) match (context-selector) */ + #pragma omp declare variant (identifier) match (context-selector) + + OpenMP 5.1 + #pragma omp declare variant (identifier) match (context-selector) \ + adjust_args(adjust-op:argument-list) */ static bool c_parser_omp_declare (c_parser *parser, enum pragma_context context) @@ -26599,6 +26973,9 @@ c_parser_omp_construct (c_parser *parser, bool *if_p) case PRAGMA_OMP_ASSUME: c_parser_omp_assume (parser, if_p); return; + case PRAGMA_OMP_DISPATCH: + stmt = c_parser_omp_dispatch (loc, parser); + break; default: gcc_unreachable (); } diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 2d092357e0f..48b4bd906e8 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -16054,6 +16054,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_FINALIZE: case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_INDIRECT: + case OMP_CLAUSE_NOVARIANTS: + case OMP_CLAUSE_NOCONTEXT: pc = &OMP_CLAUSE_CHAIN (c); continue; diff --git a/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c b/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c new file mode 100644 index 00000000000..393a44de8e0 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c @@ -0,0 +1,32 @@ +/* Test parsing of OMP clause adjust_args */ +/* { dg-do compile } */ + +int b; + +int f0 (void *a); +int g (void *a); +int f1 (int); + +#pragma omp declare variant (f0) match (construct={target}) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */ +int f2 (void *a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) /* { dg-error "expected 'nothing' or 'need_device_ptr'" } */ +int f3 (int a); +#pragma omp declare variant (f0) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */ +int f4 (void *a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args () /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */ +int f5 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */ +int f6 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) /* { dg-error "expected identifier before '\\)' token" } */ +int f7 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: z) /* { dg-error "'z' undeclared here \\(not in a function\\)" } */ +int f8 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: a) /* { dg-error "'a' is not a C pointer" } */ +int f9 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (nothing: a) /* { dg-error "'a' is specified more than once" } */ +int f10 (int a); +#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: a) /* { dg-error "'a' is specified more than once" } */ +int f11 (void *a); +#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) /* { dg-error "'b' is not a function argument" } */ +int f12 (void *a); + diff --git a/gcc/testsuite/gcc.dg/gomp/dispatch-1.c b/gcc/testsuite/gcc.dg/gomp/dispatch-1.c new file mode 100644 index 00000000000..c8f45c12be6 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/dispatch-1.c @@ -0,0 +1,53 @@ +/* Test parsing of #pragma omp dispatch */ +/* { dg-do compile } */ + +int f0 (int); + +void f1 (void) +{ + int a, b; + double x; + struct {int a; int b;} s; + int arr[1]; + +#pragma omp dispatch + int c = f0 (a); /* { dg-error "expected expression before 'int'" } */ +#pragma omp dispatch + int f2 (int d); /* { dg-error "expected expression before 'int'" } */ +#pragma omp dispatch + a = b; /* { dg-error "expected '\\(' before ';' token" } */ +#pragma omp dispatch + s.a = f0(a) + b; /* { dg-error "expected ';' before '\\+' token" } */ +#pragma omp dispatch + b = !f0(a); /* { dg-error "expected a function name before '!' token" } */ +#pragma omp dispatch + s.b += f0(s.a); /* { dg-error "expected '=' before '\\+=' token" } */ +#pragma omp dispatch +#pragma omp threadprivate(a) /* { dg-error "expected expression before '#pragma'" } */ + a = f0(b); + +#pragma omp dispatch nocontext(s) /* { dg-error "used struct type value where scalar is required" } */ + f0(a); +#pragma omp dispatch nocontext(a, b) /* { dg-error "expected '\\)' before ','" } */ + f0(a); +#pragma omp dispatch nocontext(a) nocontext(b) /* { dg-error "too many 'nocontext' clauses" } */ + f0(a); +#pragma omp dispatch novariants(s) /* { dg-error "used struct type value where scalar is required" } */ + f0(a); +#pragma omp dispatch novariants(a, b) /* { dg-error "expected '\\)' before ','" } */ + f0(a); +#pragma omp dispatch novariants(a) novariants(b) /* { dg-error "too many 'novariants' clauses" } */ + f0(a); +#pragma omp dispatch nowait nowait /* { dg-error "too many 'nowait' clauses" } */ + f0(a); +#pragma omp dispatch device(x) /* { dg-error "expected integer expression before end of line" } */ + f0(a); +#pragma omp dispatch device(arr) /* { dg-error "expected integer expression before end of line" } */ + f0(a); +#pragma omp dispatch is_device_ptr(x) /* { dg-error "'is_device_ptr' variable is neither a pointer nor an array" } */ + f0(a); +#pragma omp dispatch is_device_ptr(&x) /* { dg-error "expected identifier before '&' token" } */ + f0(a); +#pragma omp dispatch depend(inout: f0) /* { dg-error "'f0' is not lvalue expression nor array section in 'depend' clause" } */ + f0(a); +} diff --git a/libgomp/testsuite/libgomp.c/dispatch-1.c b/libgomp/testsuite/libgomp.c/dispatch-1.c new file mode 100644 index 00000000000..0efc075a859 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/dispatch-1.c @@ -0,0 +1,76 @@ +// Adapted from OpenMP examples + +#include +#include +#include + +int baz (double *d_bv, const double *d_av, int n) +{ +#pragma omp distribute parallel for + for (int i = 0; i < n; i++) + d_bv[i] = d_av[i] * i; + return -3; +} + +int bar (double *d_bv, const double *d_av, int n) +{ +#pragma omp target is_device_ptr(d_bv, d_av) + for (int i = 0; i < n; i++) + d_bv[i] = d_av[i] * i; + return -2; +} + +#pragma omp declare variant(bar) match(construct={dispatch}) adjust_args(need_device_ptr: bv, av) +#pragma omp declare variant(baz) match(implementation={vendor(gnu)}) +int foo (double *bv, const double *av, int n) +{ + for (int i = 0; i < n; i++) + bv[i] = av[i] * i; + return -1; +} + +int test (int n) +{ + const double e = 2.71828; + + double *av = (double *) malloc (n * sizeof (*av)); + double *bv = (double *) malloc (n * sizeof (*bv)); + double *d_bv = (double *) malloc (n * sizeof (*d_bv)); + + for (int i = 0; i < n; i++) + { + av[i] = e * i; + bv[i] = 0.0; + d_bv[i] = 0.0; + } + + int f, last_dev = omp_get_num_devices () - 1; +#pragma omp target data map(to: av[:n]) map(from: d_bv[:n]) device(last_dev) if (n == 1024) + { + #pragma omp dispatch nocontext(n > 1024) novariants(n < 1024) device(last_dev) + f = foo (d_bv, av, n); + } + + foo (bv, av, n); + for (int i = 0; i < n; i++) + { + if (d_bv[i] != bv[i]) + { + fprintf (stderr, "ERROR at %d: %lf (act) != %lf (exp)\n", i, d_bv[i], bv[i]); + return 1; + } + } + return f; +} + +int +main (void) +{ + int ret = test(1023); + if (ret != -1) return 1; + ret = test(1024); + if (ret != -2) return 1; + ret = test(1025); + if (ret != -3) return 1; + return 0; +} From patchwork Mon May 27 11:54:36 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 1939907 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=mtpsr2VM; 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 4VnvJH0kmNz20KL for ; Mon, 27 May 2024 21:56:47 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 49C6D3884503 for ; Mon, 27 May 2024 11:56:45 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lj1-x233.google.com (mail-lj1-x233.google.com [IPv6:2a00:1450:4864:20::233]) by sourceware.org (Postfix) with ESMTPS id CFC2F3870881 for ; Mon, 27 May 2024 11:55:35 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org CFC2F3870881 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 CFC2F3870881 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::233 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810941; cv=none; b=rM90am1sF1iZLJOr+D78ghU5p4DDHYOexPZTEAmIYcjbAwtkMqcaE0pfMtzSylE1SqwYZ3gzFObwsFBDGt6e+lLD0TUET20jx8sC8VF5VzL7iMeZOCu0t2n/Dg8TwDrJXihqNYaNfhbSFYGBpknzNutqvf7g0vOmtGeRpaeVCoE= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810941; c=relaxed/simple; bh=Yx2Jcx1l6DFxOwu60oopUe/1K4Tc7HBxFs9XFalNlmE=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=emGsZcksCPtqsnW0tATL6Le4JQ2Ha/Eby9qRNTnQoBdfNNQ1Mkamikr6hXQVnmNIqfKNsedZne+7rwE0dnOwcrp3ble6ChfvR44Zlzl0B1RBOTJM883ZSaYL/kWmoM/exEUgIhQeDU+nEETV1Vt8dwvt3qXlxUS/Bz5tvnQMlGg= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lj1-x233.google.com with SMTP id 38308e7fff4ca-2e95a1f9c53so39017761fa.0 for ; Mon, 27 May 2024 04:55:35 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1716810934; x=1717415734; 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=epc+jH+j5w2Ran4Gh633RXlf2rrRMbTYMxob17hhL0g=; b=mtpsr2VMRN7CWR/07EYAzxi8xa3sTf5kWRmDgIETwXbIn6TLoxEaPiC9oZfU8G0h/a eD1O/25uaH5+xIPbk4Gk++ockn6y8DcTl7coOoMaIzO7SZqWlk7nSlkfVHTrBtwhouxB lNcFmXK9yoHySa9sF5pd1AkpwJrwJcNc/ycYxJnzknW5CpV/X2e41RUhZ5ShisFvud1R MI85hOLt67QhIysaXxx7vT0yIqBqCQU2KxBOK/YpmhVvovGKNUkVuXvmoafgF571yDg/ CbjLdyLlLipmyswvJyZhQN5mIJOZD5LhBMccY+Qt+9QyRQnUz4ti3cFYdeGbQD7rp4Hp xn/w== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1716810934; x=1717415734; 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=epc+jH+j5w2Ran4Gh633RXlf2rrRMbTYMxob17hhL0g=; b=lmSpzZazA6jgt7yVX6BaxcenugJibbmTcGSnQPFhEnORCx3PRs9dlUPY71IALHSjVj SGdIhTuPNZx/LSXzxH81rS19YW0Fd9kx+V+6w8LIzTC1JMst+BmtCFFy6SZ9szfQZc9v 5I/5PPvoY+5uA4BW4SPwz3hIGAqpgTexZby+AJPsTTW/Cv4lqLpK6xKtvdKb83O9MZu3 0Nsn3ULbt1etMjFv999Hc7CTG6iyInWUG6VAwjBgS30spDxtFghEnSXgYIZgf6HmVQhD CfHiFr6M21Xr9rUAMqFnrNCqpJkGxgG5mxnDDL8XYYSPYixJpPXkDiJqjc6Vxi3Bqmdq QvsA== X-Gm-Message-State: AOJu0YwoBR7qc3CAwG/1Ojffo30lYeoA6zH4V7xB8k8GKnWD5FoADhOp Md3ij5Mn0SnV/3pIpxxVEAB3Sak2NvQmnq0OKpjOIMGaOGJvBNBlPBgqWIN0pKvqu2PXptF1Vrf A X-Google-Smtp-Source: AGHT+IEKg3DFiCgxX/bUtGVHjR0TPBkXjABXQI/X0khSpfYkc/+C2AcEBY9vavjdAyBbl7iW/ohACA== X-Received: by 2002:a2e:87c7:0:b0:2e9:794c:19ae with SMTP id 38308e7fff4ca-2e9794c1b09mr13221131fa.23.1716810933561; Mon, 27 May 2024 04:55:33 -0700 (PDT) Received: from localhost.localdomain ([2a0d:3344:2340:ce10::2f3]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-42108966b58sm107636955e9.3.2024.05.27.04.55.32 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Mon, 27 May 2024 04:55:33 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH 4/7] OpenMP: C++ front-end support for dispatch + adjust_args Date: Mon, 27 May 2024 13:54:36 +0200 Message-ID: <20240527115439.3967217-5-parras@baylibre.com> X-Mailer: git-send-email 2.45.1 In-Reply-To: <20240527115439.3967217-1-parras@baylibre.com> References: <20240527115439.3967217-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.6 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, KAM_SHORT, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds C++ support for the `dispatch` construct and the `adjust_args` clause. It relies on the c-family bits comprised in the corresponding C front end patch for pragmas and attributes. Additional C/C++ common testcases are provided in a subsequent patch in the series. gcc/cp/ChangeLog: * decl.cc (omp_declare_variant_finalize_one): Set adjust_args need_device_ptr attribute. * parser.cc (cp_parser_direct_declarator): Update call to cp_parser_late_return_type_opt. (cp_parser_late_return_type_opt): Add parameter. Update call to cp_parser_late_parsing_omp_declare_simd. (cp_parser_omp_clause_name): Handle nocontext and novariants clauses. (cp_parser_omp_clause_novariants): New function. (cp_parser_omp_clause_nocontext): Likewise. (cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_NOVARIANTS and PRAGMA_OMP_CLAUSE_NOCONTEXT. (cp_parser_omp_dispatch_body): New function, inspired from cp_parser_assignment_expression and cp_parser_postfix_expression. (OMP_DISPATCH_CLAUSE_MASK): Define. (cp_parser_omp_dispatch): New function. (cp_finish_omp_declare_variant): Add parameter. Handle adjust_args clause. (cp_parser_late_parsing_omp_declare_simd): Add parameter. Update calls to cp_finish_omp_declare_variant and cp_finish_omp_declare_variant. (cp_parser_omp_construct): Handle PRAGMA_OMP_DISPATCH. (cp_parser_pragma): Likewise. * pt.cc (tsubst_attribute): Skip pseudo-TSS need_device_ptr. * semantics.cc (finish_omp_clauses): Handle OMP_CLAUSE_NOCONTEXT and OMP_CLAUSE_NOVARIANTS. gcc/testsuite/ChangeLog: * g++.dg/gomp/adjust-args-1.C: New test. * g++.dg/gomp/adjust-args-2.C: New test. * g++.dg/gomp/dispatch-1.C: New test. * g++.dg/gomp/dispatch-2.C: New test. --- gcc/cp/decl.cc | 27 + gcc/cp/parser.cc | 613 ++++++++++++++++++++-- gcc/cp/pt.cc | 3 + gcc/cp/semantics.cc | 20 + gcc/testsuite/g++.dg/gomp/adjust-args-1.C | 39 ++ gcc/testsuite/g++.dg/gomp/adjust-args-2.C | 51 ++ gcc/testsuite/g++.dg/gomp/dispatch-1.C | 53 ++ gcc/testsuite/g++.dg/gomp/dispatch-2.C | 62 +++ 8 files changed, 822 insertions(+), 46 deletions(-) create mode 100644 gcc/testsuite/g++.dg/gomp/adjust-args-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/adjust-args-2.C create mode 100644 gcc/testsuite/g++.dg/gomp/dispatch-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/dispatch-2.C diff --git a/gcc/cp/decl.cc b/gcc/cp/decl.cc index a992d54dc8f..fe53e59add1 100644 --- a/gcc/cp/decl.cc +++ b/gcc/cp/decl.cc @@ -8360,6 +8360,33 @@ omp_declare_variant_finalize_one (tree decl, tree attr) if (!omp_context_selector_matches (ctx)) return true; TREE_PURPOSE (TREE_VALUE (attr)) = variant; + + for (tree a = ctx; a != NULL_TREE; a = TREE_CHAIN (a)) + { + if (OMP_TSS_CODE (a) == OMP_TRAIT_SET_NEED_DEVICE_PTR) + { + tree parm_decl = TREE_VALUE (TREE_VALUE (a)); + bool found_arg = false; + for (tree arg = DECL_ARGUMENTS (variant); arg != NULL; + arg = TREE_CHAIN (arg)) + if (DECL_NAME (arg) == DECL_NAME (parm_decl)) + { + DECL_ATTRIBUTES (arg) + = tree_cons (get_identifier ( + "omp declare variant adjust_args " + "need_device_ptr"), + NULL_TREE, DECL_ATTRIBUTES (arg)); + found_arg = true; + } + if (!found_arg) + { + error_at (varid_loc, + "variant %qD does not have a parameter %qD", + variant, parm_decl); + return true; + } + } + } } } else if (!processing_template_decl) diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 779625144db..d61c37729a7 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -19,6 +19,7 @@ along with GCC; see the file COPYING3. If not see . */ #include "config.h" +#include "omp-selectors.h" #define INCLUDE_MEMORY #include "system.h" #include "coretypes.h" @@ -2587,7 +2588,7 @@ static cp_ref_qualifier cp_parser_ref_qualifier_opt static tree cp_parser_tx_qualifier_opt (cp_parser *); static tree cp_parser_late_return_type_opt - (cp_parser *, cp_declarator *, tree &); + (cp_parser *, cp_declarator *, tree &, tree); static tree cp_parser_declarator_id (cp_parser *, bool); static tree cp_parser_type_id @@ -2622,7 +2623,7 @@ static void cp_parser_ctor_initializer_opt_and_function_body (cp_parser *, bool); static tree cp_parser_late_parsing_omp_declare_simd - (cp_parser *, tree); + (cp_parser *, tree, tree); static tree cp_parser_late_parsing_oacc_routine (cp_parser *, tree); @@ -24150,7 +24151,7 @@ cp_parser_direct_declarator (cp_parser* parser, tree requires_clause = NULL_TREE; late_return = cp_parser_late_return_type_opt (parser, declarator, - requires_clause); + requires_clause, params); cp_finalize_omp_declare_simd (parser, &odsd); @@ -25014,8 +25015,8 @@ parsing_function_declarator () function. */ static tree -cp_parser_late_return_type_opt (cp_parser* parser, cp_declarator *declarator, - tree& requires_clause) +cp_parser_late_return_type_opt (cp_parser *parser, cp_declarator *declarator, + tree &requires_clause, tree parms) { cp_token *token; tree type = NULL_TREE; @@ -25051,8 +25052,8 @@ cp_parser_late_return_type_opt (cp_parser* parser, cp_declarator *declarator, if (declare_simd_p) declarator->attributes - = cp_parser_late_parsing_omp_declare_simd (parser, - declarator->attributes); + = cp_parser_late_parsing_omp_declare_simd (parser, declarator->attributes, + parms); if (oacc_routine_p) declarator->attributes = cp_parser_late_parsing_oacc_routine (parser, @@ -38237,6 +38238,8 @@ cp_parser_omp_clause_name (cp_parser *parser) case 'n': if (!strcmp ("no_create", p)) result = PRAGMA_OACC_CLAUSE_NO_CREATE; + else if (!strcmp ("nocontext", p)) + result = PRAGMA_OMP_CLAUSE_NOCONTEXT; else if (!strcmp ("nogroup", p)) result = PRAGMA_OMP_CLAUSE_NOGROUP; else if (!strcmp ("nohost", p)) @@ -38245,6 +38248,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_NONTEMPORAL; else if (!strcmp ("notinbranch", p)) result = PRAGMA_OMP_CLAUSE_NOTINBRANCH; + else if (!strcmp ("novariants", p)) + result = PRAGMA_OMP_CLAUSE_NOVARIANTS; else if (!strcmp ("nowait", p)) result = PRAGMA_OMP_CLAUSE_NOWAIT; else if (!strcmp ("num_gangs", p)) @@ -40647,6 +40652,56 @@ cp_parser_omp_clause_thread_limit (cp_parser *parser, tree list, return c; } +/* OpenMP 5.1 + novariants ( scalar-expression ) */ + +static tree +cp_parser_omp_clause_novariants (cp_parser *parser, tree list, location_t loc) +{ + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + tree t = cp_parser_assignment_expression (parser); + if (t == error_mark_node || !parens.require_close (parser)) + cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + + check_no_duplicate_clause (list, OMP_CLAUSE_NOVARIANTS, "novariants", loc); + + tree c = build_omp_clause (loc, OMP_CLAUSE_NOVARIANTS); + OMP_CLAUSE_NOVARIANTS_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + + return c; +} + +/* OpenMP 5.1 + nocontext ( scalar-expression ) */ + +static tree +cp_parser_omp_clause_nocontext (cp_parser *parser, tree list, location_t loc) +{ + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + tree t = cp_parser_assignment_expression (parser); + if (t == error_mark_node || !parens.require_close (parser)) + cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + + check_no_duplicate_clause (list, OMP_CLAUSE_NOCONTEXT, "nocontext", loc); + + tree c = build_omp_clause (loc, OMP_CLAUSE_NOCONTEXT); + OMP_CLAUSE_NOCONTEXT_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + + return c; +} + /* OpenMP 4.0: aligned ( variable-list ) aligned ( variable-list : constant-expression ) */ @@ -42745,6 +42800,16 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses); c_name = "enter"; break; + case PRAGMA_OMP_CLAUSE_NOVARIANTS: + clauses = cp_parser_omp_clause_novariants (parser, clauses, + token->location); + c_name = "novariants"; + break; + case PRAGMA_OMP_CLAUSE_NOCONTEXT: + clauses + = cp_parser_omp_clause_nocontext (parser, clauses, token->location); + c_name = "nocontext"; + break; default: cp_parser_error (parser, "expected an OpenMP clause"); goto saw_error; @@ -48653,12 +48718,305 @@ cp_parser_omp_assumes (cp_parser *parser, cp_token *pragma_tok) return false; } +/* Parse a function dispatch structured block: + + lvalue-expression = target-call ( [expression-list] ); + or + target-call ( [expression-list] ); + + Inspired from cp_parser_assignment_expression and + cp_parser_postfix_expression. +*/ + +static tree +cp_parser_omp_dispatch_body (cp_parser *parser) +{ + cp_expr expr; + cp_id_kind idk = CP_ID_KIND_NONE; + + /* Parse the binary expressions (lvalue-expression or target-call). */ + expr = cp_parser_binary_expression (parser, false, false, false, + PREC_NOT_OPERATOR, NULL); + if (TREE_CODE (expr) == CALL_EXPR || TREE_CODE (expr) == ERROR_MARK) + return expr; + + /* We have the lvalue, now deal with the assignment. */ + + if (!cp_parser_require (parser, CPP_EQ, RT_EQ)) + return error_mark_node; + + /* Peek at the next token. */ + cp_token *token = cp_lexer_peek_token (parser->lexer); + location_t loc = token->location; + location_t start_loc = get_range_from_loc (line_table, loc).m_start; + + /* Parse function name as primary expression. */ + cp_expr rhs + = cp_parser_primary_expression (parser, false, false, false, false, &idk); + if (TREE_CODE (rhs) == ERROR_MARK) + return rhs; + + /* Keep looping until the postfix-expression is complete. */ + bool parens_found = false; + while (true) + { + if (idk == CP_ID_KIND_UNQUALIFIED && identifier_p (rhs) + && cp_lexer_next_token_is_not (parser->lexer, CPP_OPEN_PAREN)) + /* It is not a Koenig lookup function call. */ + rhs = unqualified_name_lookup_error (rhs); + + /* Peek at the next token. */ + token = cp_lexer_peek_token (parser->lexer); + + switch (token->type) + { + case CPP_OPEN_PAREN: + /* postfix-expression ( expression-list [opt] ) */ + { + if (parens_found) + { + cp_parser_error ( + parser, + "only one function call is allowed in a dispatch construct"); + return error_mark_node; + } + parens_found = true; + + bool koenig_p; + tsubst_flags_t complain = complain_flags (false); + vec *args; + location_t close_paren_loc = UNKNOWN_LOCATION; + location_t combined_loc = UNKNOWN_LOCATION; + + args = (cp_parser_parenthesized_expression_list ( + parser, non_attr, + /*cast_p=*/false, /*allow_expansion_p=*/true, + /*non_constant_p=*/NULL, + /*close_paren_loc=*/&close_paren_loc, + /*wrap_locations_p=*/true)); + + if (args == NULL) + { + rhs = error_mark_node; + break; + } + + koenig_p = false; + if (idk == CP_ID_KIND_UNQUALIFIED || idk == CP_ID_KIND_TEMPLATE_ID) + { + if (identifier_p (rhs) + /* In C++20, we may need to perform ADL for a template + name. */ + || (TREE_CODE (rhs) == TEMPLATE_ID_EXPR + && identifier_p (TREE_OPERAND (rhs, 0)))) + { + if (!args->is_empty ()) + { + koenig_p = true; + if (!any_type_dependent_arguments_p (args)) + rhs = perform_koenig_lookup (rhs, args, complain); + } + else + rhs = unqualified_fn_lookup_error (rhs); + } + /* We do not perform argument-dependent lookup if + normal lookup finds a non-function, in accordance + with the expected resolution of DR 218. */ + else if (!args->is_empty () && is_overloaded_fn (rhs)) + { + /* Do not do argument dependent lookup if regular + lookup finds a member function or a block-scope + function declaration. [basic.lookup.argdep]/3 */ + bool do_adl_p = true; + tree fns = get_fns (rhs); + for (lkp_iterator iter (fns); iter; ++iter) + { + tree fn = STRIP_TEMPLATE (*iter); + if ((TREE_CODE (fn) == USING_DECL + && DECL_DEPENDENT_P (fn)) + || DECL_FUNCTION_MEMBER_P (fn) + || DECL_LOCAL_DECL_P (fn)) + { + do_adl_p = false; + break; + } + } + + if (do_adl_p) + { + koenig_p = true; + if (!any_type_dependent_arguments_p (args)) + rhs = perform_koenig_lookup (rhs, args, complain); + } + } + } + + /* Temporarily set input_location to the combined location + with call expression range, as e.g. build_out_target_exprs + called from convert_default_arg relies on input_location, + so updating it only when the call is fully built results + in inconsistencies between location handling in templates + and outside of templates. */ + if (close_paren_loc != UNKNOWN_LOCATION) + combined_loc + = make_location (token->location, start_loc, close_paren_loc); + iloc_sentinel ils (combined_loc); + + if (TREE_CODE (rhs) == COMPONENT_REF) + { + tree instance = TREE_OPERAND (rhs, 0); + tree fn = TREE_OPERAND (rhs, 1); + + if (processing_template_decl + && (type_dependent_object_expression_p (instance) + || (!BASELINK_P (fn) && TREE_CODE (fn) != FIELD_DECL) + || type_dependent_expression_p (fn) + || any_type_dependent_arguments_p (args))) + { + maybe_generic_this_capture (instance, fn); + rhs = build_min_nt_call_vec (rhs, args); + } + else if (BASELINK_P (fn)) + { + rhs + = (build_new_method_call (instance, fn, &args, NULL_TREE, + (idk == CP_ID_KIND_QUALIFIED + ? LOOKUP_NORMAL + | LOOKUP_NONVIRTUAL + : LOOKUP_NORMAL), + /*fn_p=*/NULL, complain)); + } + else + rhs = finish_call_expr (rhs, &args, + /*disallow_virtual=*/false, + /*koenig_p=*/false, complain); + } + else if (TREE_CODE (rhs) == OFFSET_REF + || TREE_CODE (rhs) == MEMBER_REF + || TREE_CODE (rhs) == DOTSTAR_EXPR) + rhs = (build_offset_ref_call_from_tree (rhs, &args, complain)); + else if (idk == CP_ID_KIND_QUALIFIED) + /* A call to a static class member, or a namespace-scope + function. */ + rhs = finish_call_expr (rhs, &args, + /*disallow_virtual=*/true, koenig_p, + complain); + else + /* All other function calls. */ + { + if (DECL_P (rhs) && parser->omp_for_parse_state + && parser->omp_for_parse_state->in_intervening_code + && omp_runtime_api_call (rhs)) + { + error_at (loc, "calls to the OpenMP runtime API are " + "not permitted in intervening code"); + parser->omp_for_parse_state->fail = true; + } + rhs = finish_call_expr (rhs, &args, + /*disallow_virtual=*/false, koenig_p, + complain); + } + if (close_paren_loc != UNKNOWN_LOCATION) + rhs.set_location (combined_loc); + + /* The expr is certainly no longer an id. */ + idk = CP_ID_KIND_NONE; + + release_tree_vector (args); + } + break; + + case CPP_DOT: + case CPP_DEREF: + /* postfix-expression . template [opt] id-expression + postfix-expression . pseudo-destructor-name + postfix-expression -> template [opt] id-expression + postfix-expression -> pseudo-destructor-name */ + + /* Consume the `.' or `->' operator. */ + cp_lexer_consume_token (parser->lexer); + + rhs = cp_parser_postfix_dot_deref_expression (parser, token->type, + rhs, false, &idk, loc); + + break; + + default: + goto finish; + } + } +finish: + if (!parens_found) + { + cp_parser_error (parser, "expected %<(%>"); + return error_mark_node; + } + + /* Build the assignment expression. Its default + location: + LHS = RHS + ~~~~^~~~~ + is the location of the '=' token as the + caret, ranging from the start of the lhs to the + end of the rhs. */ + loc = make_location (loc, expr.get_start (), rhs.get_finish ()); + expr + = cp_build_modify_expr (loc, expr, NOP_EXPR, rhs, complain_flags (false)); + + return expr; +} + +/* OpenMP 5.1: + # pragma omp dispatch dispatch-clause[optseq] new-line + expression-stmt + + LOC is the location of the #pragma. +*/ + +#define OMP_DISPATCH_CLAUSE_MASK \ + ((OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT)) + +static tree +cp_parser_omp_dispatch (cp_parser *parser, cp_token *pragma_tok) +{ + location_t loc = cp_lexer_peek_token (parser->lexer)->location; + tree stmt = make_node (OMP_DISPATCH); + SET_EXPR_LOCATION (stmt, loc); + TREE_TYPE (stmt) = void_type_node; + + OMP_DISPATCH_CLAUSES (stmt) + = cp_parser_omp_all_clauses (parser, OMP_DISPATCH_CLAUSE_MASK, + "#pragma omp dispatch", pragma_tok); + + // Parse expression statement + loc = cp_lexer_peek_token (parser->lexer)->location; + tree dispatch_body = cp_parser_omp_dispatch_body (parser); + if (dispatch_body == error_mark_node) + { + inform (loc, + "%<#pragma omp dispatch%> must be followed by a direct function " + "call with optional assignment"); + cp_parser_skip_to_end_of_block_or_statement (parser); + return NULL_TREE; + } + + cp_parser_consume_semicolon_at_end_of_statement (parser); + OMP_DISPATCH_BODY (stmt) = dispatch_body; + + return add_stmt (stmt); +} + /* Finalize #pragma omp declare variant after a fndecl has been parsed, and put that into "omp declare variant base" attribute. */ static tree cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok, - tree attrs) + tree attrs, tree parms) { matching_parens parens; if (!parens.require_open (parser)) @@ -48716,45 +49074,197 @@ cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok, location_t finish_loc = get_finish (varid.get_location ()); location_t varid_loc = make_location (caret_loc, start_loc, finish_loc); - if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA) - && cp_lexer_nth_token_is (parser->lexer, 2, CPP_NAME)) - cp_lexer_consume_token (parser->lexer); + vec adjust_args_list = vNULL; + bool has_match = false, has_adjust_args = false; + location_t adjust_args_loc; + tree need_device_ptr_list = NULL_TREE, *need_device_ptr_chain_p = NULL; - const char *clause = ""; - location_t match_loc = cp_lexer_peek_token (parser->lexer)->location; - if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) - clause = IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value); - if (strcmp (clause, "match")) + do { - cp_parser_error (parser, "expected %"); - goto fail; + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA) + && cp_lexer_nth_token_is (parser->lexer, 2, CPP_NAME)) + cp_lexer_consume_token (parser->lexer); + + const char *clause = ""; + location_t match_loc = cp_lexer_peek_token (parser->lexer)->location; + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + clause + = IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value); + + enum clause + { + match, + adjust_args + } ccode; + + if (strcmp (clause, "match") == 0) + ccode = match; + else if (strcmp (clause, "adjust_args") == 0) + { + ccode = adjust_args; + adjust_args_loc = match_loc; + } + else + { + cp_parser_error (parser, "expected % or %"); + goto fail; + } + + cp_lexer_consume_token (parser->lexer); + + if (!parens.require_open (parser)) + goto fail; + + if (ccode == match) + { + has_match = true; + tree ctx + = cp_parser_omp_context_selector_specification (parser, true); + if (ctx == error_mark_node) + goto fail; + ctx = omp_check_context_selector (match_loc, ctx); + if (ctx != error_mark_node && variant != error_mark_node) + { + tree match_loc_node + = maybe_wrap_with_location (integer_zero_node, match_loc); + tree loc_node + = maybe_wrap_with_location (integer_zero_node, varid_loc); + loc_node + = tree_cons (match_loc_node, + build_int_cst (integer_type_node, idk), + build_tree_list (loc_node, integer_zero_node)); + attrs = tree_cons (get_identifier ("omp declare variant base"), + tree_cons (variant, ctx, loc_node), attrs); + if (processing_template_decl) + ATTR_IS_DEPENDENT (attrs) = 1; + } + if (!parens.require_close (parser)) + goto fail; + } + else if (ccode == adjust_args) + { + has_adjust_args = true; + cp_token *adjust_op_tok = cp_lexer_peek_token (parser->lexer); + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME) + && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON)) + { + const char *p = IDENTIFIER_POINTER (adjust_op_tok->u.value); + if (strcmp (p, "need_device_ptr") == 0 + || strcmp (p, "nothing") == 0) + { + cp_lexer_consume_token (parser->lexer); // need_device_ptr + cp_lexer_consume_token (parser->lexer); // : + location_t arg_loc + = cp_lexer_peek_token (parser->lexer)->location; + + tree arg; + tree list + = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_ERROR, + NULL_TREE, NULL); + + for (tree c = list; c != NULL_TREE; c = TREE_CHAIN (c)) + { + tree decl = TREE_PURPOSE (c); + int idx; + for (arg = parms, idx = 0; arg != NULL; + arg = TREE_CHAIN (arg), idx++) + if (TREE_VALUE (arg) == decl) + break; + if (arg == NULL_TREE) + { + error_at (arg_loc, "%qD is not a function argument", + decl); + continue; + } + arg = TREE_VALUE (arg); + if (adjust_args_list.contains (arg)) + { + error_at (arg_loc, "%qD is specified more than once", + decl); + continue; + } + if (strcmp (p, "need_device_ptr") == 0) + { + bool is_ptr_or_template + = TEMPLATE_PARM_P (TREE_TYPE (arg)) + || POINTER_TYPE_P (TREE_TYPE (arg)); + if (!is_ptr_or_template) + { + error_at (arg_loc, "%qD is not a C pointer", + decl); + continue; + } + } + adjust_args_list.safe_push (arg); + if (strcmp (p, "need_device_ptr") == 0) + { + tree attr = tree_cons (NULL_TREE, TREE_PURPOSE (c), + NULL_TREE); + if (need_device_ptr_list == NULL_TREE) + { + gcc_assert (need_device_ptr_chain_p == NULL); + need_device_ptr_list = attr; + } + else + *need_device_ptr_chain_p = attr; + need_device_ptr_chain_p = &TREE_CHAIN (attr); + } + } + } + else + { + error_at (adjust_op_tok->location, + "expected % or %"); + goto fail; + } + } + else + { + error_at (adjust_op_tok->location, + "expected % or % followed " + "by %<:%>"); + goto fail; + } + } + } while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)); + + if (has_adjust_args) + { + if (!has_match) + { + error_at ( + adjust_args_loc, + "an % clause can only be specified if the " + "% selector of the construct selector set appears " + "in the % clause"); + } + else + { + tree ctx = TREE_VALUE (TREE_VALUE (attrs)); + if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT, + OMP_TRAIT_CONSTRUCT_DISPATCH)) + error_at ( + adjust_args_loc, + "an % clause can only be specified if the " + "% selector of the construct selector set appears " + "in the % clause"); + } } - cp_lexer_consume_token (parser->lexer); - - if (!parens.require_open (parser)) - goto fail; - - tree ctx = cp_parser_omp_context_selector_specification (parser, true); - if (ctx == error_mark_node) - goto fail; - ctx = omp_check_context_selector (match_loc, ctx); - if (ctx != error_mark_node && variant != error_mark_node) + if (need_device_ptr_list) { - tree match_loc_node = maybe_wrap_with_location (integer_zero_node, - match_loc); - tree loc_node = maybe_wrap_with_location (integer_zero_node, varid_loc); - loc_node = tree_cons (match_loc_node, - build_int_cst (integer_type_node, idk), - build_tree_list (loc_node, integer_zero_node)); - attrs = tree_cons (get_identifier ("omp declare variant base"), - tree_cons (variant, ctx, loc_node), attrs); - if (processing_template_decl) - ATTR_IS_DEPENDENT (attrs) = 1; + // We might not have DECL_ARGUMENTS for the variant yet. So we store the + // need_device_ptr list in the base function attribute beside the context + // selector. + gcc_assert (TREE_PURPOSE (attrs) + == get_identifier ("omp declare variant base")); + gcc_assert (TREE_PURPOSE (TREE_VALUE (attrs)) == variant); + TREE_VALUE (TREE_VALUE (attrs)) + = make_trait_set_selector (OMP_TRAIT_SET_NEED_DEVICE_PTR, + need_device_ptr_list, + TREE_VALUE (TREE_VALUE (attrs))); } - parens.require_close (parser); - cp_parser_skip_to_pragma_eol (parser, pragma_tok); return attrs; } @@ -48763,7 +49273,8 @@ cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok, been parsed, and put that into "omp declare simd" attribute. */ static tree -cp_parser_late_parsing_omp_declare_simd (cp_parser *parser, tree attrs) +cp_parser_late_parsing_omp_declare_simd (cp_parser *parser, tree attrs, + tree parms) { struct cp_token_cache *ce; cp_omp_declare_simd_data *data = parser->omp_declare_simd; @@ -48807,7 +49318,7 @@ cp_parser_late_parsing_omp_declare_simd (cp_parser *parser, tree attrs) { gcc_assert (strcmp (kind, "variant") == 0); attrs - = cp_finish_omp_declare_variant (parser, pragma_tok, attrs); + = cp_finish_omp_declare_variant (parser, pragma_tok, attrs, parms); } cp_parser_pop_lexer (parser); } @@ -48938,9 +49449,8 @@ cp_parser_late_parsing_omp_declare_simd (cp_parser *parser, tree attrs) else { gcc_assert (strcmp (kind, "variant") == 0); - attrs - = cp_finish_omp_declare_variant (parser, pragma_tok, - attrs); + attrs = cp_finish_omp_declare_variant (parser, pragma_tok, + attrs, parms); } gcc_assert (parser->lexer != lexer); vec_safe_truncate (lexer->buffer, 0); @@ -49790,7 +50300,11 @@ cp_parser_omp_declare_reduction (cp_parser *parser, cp_token *pragma_tok, #pragma omp declare target new-line OpenMP 5.0 - #pragma omp declare variant (identifier) match (context-selector) */ + #pragma omp declare variant (identifier) match (context-selector) + + OpenMP 5.1 + #pragma omp declare variant (identifier) match (context-selector) \ + adjust_args (adjust-op:argument-list) */ static bool cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok, @@ -50645,6 +51159,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok, bool *if_p) case PRAGMA_OMP_ASSUME: cp_parser_omp_assume (parser, pragma_tok, if_p); return; + case PRAGMA_OMP_DISPATCH: + stmt = cp_parser_omp_dispatch (parser, pragma_tok); + break; default: gcc_unreachable (); } @@ -51339,6 +51856,10 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p) "%<#pragma omp sections%> construct"); break; + case PRAGMA_OMP_DISPATCH: + cp_parser_omp_dispatch (parser, pragma_tok); + return true; + case PRAGMA_IVDEP: case PRAGMA_UNROLL: case PRAGMA_NOVECTOR: diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index dfce1b3c359..b8ae4707258 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -12014,6 +12014,9 @@ tsubst_attribute (tree t, tree *decl_p, tree args, for (tree tss = ctx; tss; tss = TREE_CHAIN (tss)) { enum omp_tss_code set = OMP_TSS_CODE (tss); + if (set == OMP_TRAIT_SET_NEED_DEVICE_PTR) + continue; + tree selectors = NULL_TREE; for (tree ts = OMP_TSS_TRAIT_SELECTORS (tss); ts; ts = TREE_CHAIN (ts)) diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index f90c304a65b..248aee4bd5b 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -7613,6 +7613,26 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) OMP_CLAUSE_FINAL_EXPR (c) = t; break; + case OMP_CLAUSE_NOCONTEXT: + t = OMP_CLAUSE_NOCONTEXT_EXPR (c); + t = maybe_convert_cond (t); + if (t == error_mark_node) + remove = true; + else if (!processing_template_decl) + t = fold_build_cleanup_point_expr (TREE_TYPE (t), t); + OMP_CLAUSE_NOCONTEXT_EXPR (c) = t; + break; + + case OMP_CLAUSE_NOVARIANTS: + t = OMP_CLAUSE_NOVARIANTS_EXPR (c); + t = maybe_convert_cond (t); + if (t == error_mark_node) + remove = true; + else if (!processing_template_decl) + t = fold_build_cleanup_point_expr (TREE_TYPE (t), t); + OMP_CLAUSE_NOVARIANTS_EXPR (c) = t; + break; + case OMP_CLAUSE_GANG: /* Operand 1 is the gang static: argument. */ t = OMP_CLAUSE_OPERAND (c, 1); diff --git a/gcc/testsuite/g++.dg/gomp/adjust-args-1.C b/gcc/testsuite/g++.dg/gomp/adjust-args-1.C new file mode 100644 index 00000000000..1c6dd8ac97b --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/adjust-args-1.C @@ -0,0 +1,39 @@ +/* Test parsing of OMP clause adjust_args */ +/* { dg-do compile } */ + +int b; + +int f0 (void *a); +int g (void *a); +int f1 (int); + +#pragma omp declare variant (f0) match (construct={target}) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */ +int f2 (void *a); +#pragma omp declare variant (f0) match (construct={dispatch,target}) adjust_args (need_device_ptr: a) /* { dg-error "'int f0.void..' used as a variant with incompatible 'construct' selector sets" } */ +int f2a (void *a); +#pragma omp declare variant (f0) match (construct={target,dispatch}) adjust_args (need_device_ptr: a) /* { dg-error "'int f0.void..' used as a variant with incompatible 'construct' selector sets" } */ +int f2b (void *a); +#pragma omp declare variant (f0) match (construct={dispatch},device={arch(gcn)}) adjust_args (need_device_ptr: a) /* { dg-error "'int f0.void..' used as a variant with incompatible 'construct' selector sets" } */ +int f2c (void *a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) /* { dg-error "expected 'nothing' or 'need_device_ptr'" } */ +int f3 (int a); +#pragma omp declare variant (f0) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */ +int f4 (void *a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args () /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */ +int f5 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */ +int f6 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) /* { dg-error "expected unqualified-id before '\\)' token" } */ +int f7 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: z) /* { dg-error "'z' has not been declared" } */ +int f8 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: a) /* { dg-error "'a' is not a C pointer" } */ +int f9 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (nothing: a) /* { dg-error "'a' is specified more than once" } */ +int f10 (int a); +#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: a) /* { dg-error "'a' is specified more than once" } */ +int f11 (void *a); +#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) /* { dg-error "'b' is not a function argument" } */ +int f12 (void *a); +#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: this) /* { dg-error "expected unqualified-id before 'this'" } */ +int f13 (void *a); diff --git a/gcc/testsuite/g++.dg/gomp/adjust-args-2.C b/gcc/testsuite/g++.dg/gomp/adjust-args-2.C new file mode 100644 index 00000000000..a78f06ec193 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/adjust-args-2.C @@ -0,0 +1,51 @@ +struct S { + int a; + int g (const void *b); + #pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) + int f0(const void *b); + int operator()() { return a; } + bool operator!() { return !a; } +}; + +template +T f0(T a, T *b); + +#pragma omp declare variant (f0) match (construct={dispatch}) adjust_args (need_device_ptr: a, b) +template +T f1(T a, T *b); + +namespace N { + class C{ + public: + void g(C *c); + #pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: c) + void f0(C *c); + }; + void g(C *c); + #pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: c) + void f0(C *c); +} + +#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: c) +void f3(N::C *c); +void f4(S *&s); +#pragma omp declare variant (f4) match (construct={dispatch}) adjust_args (need_device_ptr: s) +void f5(S *&s); + +void test() { + S s, *sp; + N::C c; + int *a, b; + #pragma omp dispatch + s.f0(a); + #pragma omp dispatch + f1(b, a); + #pragma omp dispatch + c.f0(&c); + #pragma omp dispatch + N::f0(&c); + #pragma omp dispatch + f3(&c); + #pragma omp dispatch + f5(sp); +} diff --git a/gcc/testsuite/g++.dg/gomp/dispatch-1.C b/gcc/testsuite/g++.dg/gomp/dispatch-1.C new file mode 100644 index 00000000000..fb467afcd85 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/dispatch-1.C @@ -0,0 +1,53 @@ +struct S { + int a; + void f0(double); + int operator()() { return a; } + bool operator!() { return !a; } +}; + +int f0(int); +template +T f1(T a, T b); +void (*f2)(void); + +namespace N { + class C{}; + void f0(C); + int a; +} + +int test() { + int result; + double d = 5.0; + N::C c; + S s; + S* sp = &s; + int &r = result; + #pragma omp dispatch + result = f0(5); + #pragma omp dispatch + r = f0(5); + #pragma omp dispatch + N::a = ::f0(5); + #pragma omp dispatch + sp->a = f1(5, 10); + #pragma omp dispatch + s.a = f1(5, 10); + #pragma omp dispatch + f2(); + #pragma omp dispatch + N::f0(c); + #pragma omp dispatch + f0(c); + #pragma omp dispatch + s.f0(d); + #pragma omp dispatch + sp->f0(d); + #pragma omp dispatch + sp->f0(d); + #pragma omp dispatch + s(); + #pragma omp dispatch + !s; + return result; +} diff --git a/gcc/testsuite/g++.dg/gomp/dispatch-2.C b/gcc/testsuite/g++.dg/gomp/dispatch-2.C new file mode 100644 index 00000000000..1bc304e005e --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/dispatch-2.C @@ -0,0 +1,62 @@ +/* Test parsing of #pragma omp dispatch */ +/* { dg-do compile } */ + +struct S { + int a; + int b; + virtual int f (double); +}; + +int f0 (int); + +void f1 (void) +{ + int a, b; + double x; + int arr[1]; + S s; + +#pragma omp dispatch + int c = f0 (a); /* { dg-error "expected primary-expression before 'int'" } */ +#pragma omp dispatch + int f2 (int d); /* { dg-error "expected primary-expression before 'int'" } */ +#pragma omp dispatch + a = b; /* { dg-error "expected '\\(' before ';' token" } */ +#pragma omp dispatch + s.a = f0(a) + b; /* { dg-error "expected ';' before '\\+' token" } */ +#pragma omp dispatch + b = !f0(a); /* { dg-error "expected primary-expression before '!' token" } */ +#pragma omp dispatch + s.b += f0(s.a); /* { dg-error "expected '=' before '\\+=' token" } */ +#pragma omp dispatch +#pragma omp threadprivate(a) /* { dg-error "'#pragma' is not allowed here" } */ + a = f0(b); +#pragma omp dispatch + a = s.f(x); /* { dg-error "'f' is a virtual function but only a direct call is allowed in a dispatch construct" } */ + +#pragma omp dispatch nocontext(s) /* { dg-error "could not convert 's' from 'S' to 'bool'" } */ + f0(a); +#pragma omp dispatch nocontext(a, b) /* { dg-error "expected '\\)' before ','" } */ + f0(a); +#pragma omp dispatch nocontext(a) nocontext(b) /* { dg-error "too many 'nocontext' clauses" } */ + f0(a); +#pragma omp dispatch novariants(s) /* { dg-error "could not convert 's' from 'S' to 'bool'" } */ + f0(a); +#pragma omp dispatch novariants(a, b) /* { dg-error "expected '\\)' before ','" } */ + f0(a); +#pragma omp dispatch novariants(a) novariants(b) /* { dg-error "too many 'novariants' clauses" } */ + f0(a); +#pragma omp dispatch nowait nowait /* { dg-error "too many 'nowait' clauses" } */ + f0(a); +#pragma omp dispatch device(x) /* { dg-error "'device' id must be integral" } */ + f0(a); +#pragma omp dispatch device(arr) /* { dg-error "'device' id must be integral" } */ + f0(a); +#pragma omp dispatch is_device_ptr(x) /* { dg-error "'is_device_ptr' variable is neither a pointer, nor an array nor reference to pointer" } */ + f0(a); +#pragma omp dispatch is_device_ptr(&x) /* { dg-error "expected unqualified-id before '&' token" } */ + f0(a); +#pragma omp dispatch depend(inout: s.f) /* { dg-error "'s.S::f' is not lvalue expression nor array section in 'depend' clause" } */ + f0(a); + +} From patchwork Mon May 27 11:54:37 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 1939906 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=VZTGpbGK; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; 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 [8.43.85.97]) (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 4VnvHv6GbSz20KL for ; Mon, 27 May 2024 21:56:27 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 306303884542 for ; Mon, 27 May 2024 11:56:26 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lj1-x22b.google.com (mail-lj1-x22b.google.com [IPv6:2a00:1450:4864:20::22b]) by sourceware.org (Postfix) with ESMTPS id 1D61D388300E for ; Mon, 27 May 2024 11:55:37 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 1D61D388300E 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 1D61D388300E Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::22b ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810940; cv=none; b=WnhBa+S4pwb3jLhKX4a7DLnC5iq3ub8EytqLTpp46gQCm41gv8qasC+TymmPbjIpzJsPBwxZchKQ31YHXZAHK9rcCbc8wMZqzN6a8kV7R3YXKC6/ihf2dfwK2SichZkupBOr3dZP9n1p3sazOKkF7dkCDqr+8ooOw92oezDlRS4= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810940; c=relaxed/simple; bh=Z3Hw7+uYND64LeXr0pbPYF7/xsB3IPSMYKZKozgPD8E=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=Qp5Io5KYlaPp2fAXVKCGCiPuujIzzYiUBuIrIuD/o/P+Mvl0ST6vFef41hCM9z3zPwCda5ODItUFV+8OMWyNcT5YJwQK+wlhRbhpKQj0xInQTLz4LX5XMhDGYuwH9pz5/4fwIi4VKNfInJOV8M2hrsVYY0vXdqyyjcjuIzvVxZQ= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lj1-x22b.google.com with SMTP id 38308e7fff4ca-2e73359b8fbso89674341fa.2 for ; Mon, 27 May 2024 04:55:37 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1716810935; x=1717415735; 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=VeI2NU4EEn4RduNHEsXiMqT26BPJLeATyM7/uzSo+FI=; b=VZTGpbGKCrC/k4dDBerVO/uualxdy71M/K0gnQhzCw0ENtJ+pxxBwDZzWj32OcdBto x6sqrRUyji8YIHScXyQUPPmJWVIPrFj/fIKp6GSkUbJt3YuskSV2nUGN9cEOLD/WhI6X +Qp9Z38mUMGLXq5dhUOvYDRNtwb5w6MkNf9o7oAtUJTQ1SAulXQ66YS+mz4Ve5bsgUBe U1qqGyCo9/mTIfJa7FAWxYRCT9hagZ54Ep37UuRzlFPFGvTQpTpRr6axUzuiHMbzWYK9 baS4+N63ZnLUSr9uxT/BvH00TNWuZ7fq0qmFLKMoSDPNNmH5pigtEIhCcAszGwuCXf8M xbtg== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1716810935; x=1717415735; 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=VeI2NU4EEn4RduNHEsXiMqT26BPJLeATyM7/uzSo+FI=; b=ggOMmIR3yqxRPduiJmUQe0i06OP1CxjAXpXzymGEGsinl6UrE4nT6IWiZ6Pydh6B+t A/fm+S4W0xtAxXE2TSQFnL/ejTL4YHtWZe3MHk21juh8gIbKlE+ROshg+e0V2/vw93hB TNQhdAmVpkrAygukAPnlw60xnqtHDpH8L3YHHxVW5Juv8ZiWhuYZsRNy8rcOoZVypq9h IegTx0j/omyll9HouWHw+Ic0jOUOtTiC3iAnspb83ygY0QgkYaBRpAPGAKq8gF3seqtu HR8zBdrogwLWyVfRj7bpGqFSl/4x/Ztm3fNKvdogOhIMKyY3MvhnIWkLC5l/XEWXnMDe 5Cdg== X-Gm-Message-State: AOJu0YxzG+5auP+Zy7qZHa+dFlYfa0skrlJdLNEkRi0pEzO1KsBW+oGA +A+C6RkBZjKfezj9zRGYpEnKdGh4lJB/3T/HI9lQYbS7BbxpgHozlSBceJ1RA0XCawds8IjJQmc 0 X-Google-Smtp-Source: AGHT+IG3a61ttST04XozqmMpEiVA1vw5efE5jT2ZY52Pgod+3CTjM3eJSuRAubtfANUzXbnp6jkmpA== X-Received: by 2002:a2e:a165:0:b0:2e6:f7f7:772b with SMTP id 38308e7fff4ca-2e95b256484mr80874761fa.37.1716810934832; Mon, 27 May 2024 04:55:34 -0700 (PDT) Received: from localhost.localdomain ([2a0d:3344:2340:ce10::2f3]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-42108966b58sm107636955e9.3.2024.05.27.04.55.33 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Mon, 27 May 2024 04:55:34 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH 5/7] OpenMP: common C/C++ testcases for dispatch + adjust_args Date: Mon, 27 May 2024 13:54:37 +0200 Message-ID: <20240527115439.3967217-6-parras@baylibre.com> X-Mailer: git-send-email 2.45.1 In-Reply-To: <20240527115439.3967217-1-parras@baylibre.com> References: <20240527115439.3967217-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.4 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org gcc/testsuite/ChangeLog: * c-c++-common/gomp/declare-variant-2.c: Adjust dg-error directives. * c-c++-common/gomp/adjust-args-1.c: New test. * c-c++-common/gomp/adjust-args-2.c: New test. * c-c++-common/gomp/dispatch-1.c: New test. * c-c++-common/gomp/dispatch-2.c: New test. * c-c++-common/gomp/dispatch-3.c: New test. * c-c++-common/gomp/dispatch-4.c: New test. * c-c++-common/gomp/dispatch-5.c: New test. * c-c++-common/gomp/dispatch-6.c: New test. * c-c++-common/gomp/dispatch-7.c: New test. --- .../c-c++-common/gomp/adjust-args-1.c | 30 +++++++++ .../c-c++-common/gomp/adjust-args-2.c | 31 +++++++++ .../c-c++-common/gomp/declare-variant-2.c | 4 +- gcc/testsuite/c-c++-common/gomp/dispatch-1.c | 65 +++++++++++++++++++ gcc/testsuite/c-c++-common/gomp/dispatch-2.c | 28 ++++++++ gcc/testsuite/c-c++-common/gomp/dispatch-3.c | 15 +++++ gcc/testsuite/c-c++-common/gomp/dispatch-4.c | 18 +++++ gcc/testsuite/c-c++-common/gomp/dispatch-5.c | 26 ++++++++ gcc/testsuite/c-c++-common/gomp/dispatch-6.c | 19 ++++++ gcc/testsuite/c-c++-common/gomp/dispatch-7.c | 28 ++++++++ 10 files changed, 262 insertions(+), 2 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/adjust-args-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/adjust-args-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-3.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-4.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-5.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-6.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-7.c diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c b/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c new file mode 100644 index 00000000000..728abe62092 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c @@ -0,0 +1,30 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +int f (int a, void *b, float c[2]); + +#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b, c) +int f0 (int a, void *b, float c[2]); +#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c) +int f1 (int a, void *b, float c[2]); + +int test () { + int a; + void *b; + float c[2]; + struct {int a;} s; + + s.a = f0 (a, b, c); + #pragma omp dispatch + s.a = f0 (a, b, c); + + f1 (a, b, c); + #pragma omp dispatch + s.a = f1 (a, b, c); + + return s.a; +} + +/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, D\.\[0-9]+\\);" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, D\.\[0-9]+\\);" 2 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c b/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c new file mode 100644 index 00000000000..e36d93a01d9 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c @@ -0,0 +1,31 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +int f (int a, void *b, float c[2]); + +#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b, c) +int f0 (int a, void *b, float c[2]); +#pragma omp declare variant (f) adjust_args (need_device_ptr: b, c) match (construct={dispatch}) adjust_args (nothing: a) +int f1 (int a, void *b, float c[2]); + +void test () { + int a; + void *b; + float c[2]; + + #pragma omp dispatch + f0 (a, b, c); + + #pragma omp dispatch device (-4852) + f0 (a, b, c); + + #pragma omp dispatch device (a + a) + f0 (a, b, c); +} + +/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, D\.\[0-9]+\\);" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, D\.\[0-9]+\\);" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, -4852\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, -4852\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp dispatch device\\(-4852\\)" 1 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c b/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c index 05e485ef6a8..50d9b2dcf4b 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c @@ -8,9 +8,9 @@ void f3 (void); void f4 (void); #pragma omp declare variant match(user={condition(0)}) /* { dg-error "expected '\\(' before 'match'" } */ void f5 (void); -#pragma omp declare variant (f1) /* { dg-error "expected 'match' before end of line" } */ +#pragma omp declare variant (f1) /* { dg-error "expected 'match' or 'adjust_args' before end of line" } */ void f6 (void); -#pragma omp declare variant (f1) simd /* { dg-error "expected 'match' before 'simd'" } */ +#pragma omp declare variant (f1) simd /* { dg-error "expected 'match' or 'adjust_args' before 'simd'" } */ void f7 (void); #pragma omp declare variant (f1) match /* { dg-error "expected '\\(' before end of line" } */ void f8 (void); diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-1.c b/gcc/testsuite/c-c++-common/gomp/dispatch-1.c new file mode 100644 index 00000000000..e77b2f3ecf6 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-1.c @@ -0,0 +1,65 @@ +#include + +int f0 (int, long, double); +void f2 (void); +int f3 (void); +void (*f4) (void); + +void f1 (void) +{ + int a, c; + long b; + double x; + struct {int a; float b; short c;} s, *sp; + int arr[3]; + +#pragma omp dispatch + c = f0 (a, b, x); +#pragma omp dispatch + x = f0 (a * 4, 2 - b, x * x); +#pragma omp dispatch + s.a = f0 (a, sp->c, x); +#pragma omp dispatch + sp->c = f0 (s.a - 2, b / 3, x * 5); +#pragma omp dispatch + arr[0] = f0 (arr[1], !b, arr[2]); +#pragma omp dispatch + (*sp).c = f0 (s.a, b, x); +#pragma omp dispatch + sp->b = f0 (s.a++, b % 4, --x); +#pragma omp dispatch + f0 (f3(), b, s.b); +#pragma omp dispatch + f2 (); +#pragma omp dispatch + f4 (); + +#pragma omp dispatch nocontext(sp->a * x + arr[2]) + f2 (); +#pragma omp dispatch nocontext(arr - (intptr_t)(x / s.b)) + f2 (); +#pragma omp dispatch nocontext(x == s.c || b != c) + f2 (); +#pragma omp dispatch novariants(b << sp->c) + f2 (); +#pragma omp dispatch novariants(!arr | s.a) + f2 (); +#pragma omp dispatch novariants(s.c ? f3() : a & c) + f2 (); +#pragma omp dispatch nowait + f2 (); +#pragma omp dispatch device(-25373654) + f2 (); +#pragma omp dispatch device(b * (int)(x - sp->b)) + f2 (); +#pragma omp dispatch is_device_ptr(arr) + f2 (); +#pragma omp dispatch is_device_ptr(sp) + f2 (); +#pragma omp dispatch depend(inout: sp) + f2 (); +#pragma omp dispatch depend(inoutset: arr[:2]) + f2 (); +#pragma omp dispatch depend(mutexinoutset: arr) + f2 (); +} diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-2.c b/gcc/testsuite/c-c++-common/gomp/dispatch-2.c new file mode 100644 index 00000000000..24ab9545b73 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-2.c @@ -0,0 +1,28 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +int f0 (void); +int f1 (void); +#pragma omp declare variant (f0) match (construct={dispatch}) +#pragma omp declare variant (f1) match (implementation={vendor(gnu)}) +int f2 (void); + +int test (void) +{ + int a; +#pragma omp dispatch + a = f2 (); +#pragma omp dispatch novariants(1) + a = f2 (); +#pragma omp dispatch novariants(0) + a = f2 (); +#pragma omp dispatch nocontext(1) + a = f2 (); +#pragma omp dispatch nocontext(0) + a = f2 (); + return a; +} + +/* { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 3 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-3.c b/gcc/testsuite/c-c++-common/gomp/dispatch-3.c new file mode 100644 index 00000000000..319f73a84d8 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-3.c @@ -0,0 +1,15 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void f2 (void); + +void test (void) +{ +#pragma omp dispatch /* { dg-final { scan-tree-dump-times "#pragma omp task if\\(0\\)" 1 "gimple" } } */ + f2 (); +#pragma omp dispatch nowait /* { dg-final { scan-tree-dump-times "#pragma omp task if\\(1\\)" 1 "gimple" } } */ + f2 (); +} + + + diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-4.c b/gcc/testsuite/c-c++-common/gomp/dispatch-4.c new file mode 100644 index 00000000000..aeb5c00507f --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-4.c @@ -0,0 +1,18 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void f2 (int a); + +void test (void) +{ + int a; + +#pragma omp dispatch device(-25373654) +/* { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device \\(-25373654\\);" 1 "gimple" } } */ + f2 (a); +#pragma omp dispatch device(a + a) +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = a \\* 2;.*#pragma omp dispatch device\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\).*#pragma omp task shared\\(D\.\[0-9]+\\).*__builtin_omp_set_default_device \\(D\.\[0-9]+\\);" 1 "gimple" } } */ + f2 (a); +} + + diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-5.c b/gcc/testsuite/c-c++-common/gomp/dispatch-5.c new file mode 100644 index 00000000000..78a37bac59b --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-5.c @@ -0,0 +1,26 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void f1 (void* p, int arr[]); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: p, arr) +void f2 (void* p, int arr[]); + +void test (void) +{ + void *p; + int arr[2]; + +#pragma omp dispatch + f2 (p, arr); +#pragma omp dispatch is_device_ptr(p) +/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p\\) shared\\(arr\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*int D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&arr, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(p, D\.\[0-9]+\\);" 1 "gimple" } } */ + f2 (p, arr); +#pragma omp dispatch is_device_ptr(arr) +/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(arr\\) shared\\(p\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*int D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(p, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(D\.\[0-9]+, &arr\\);" 1 "gimple" } } */ + f2 (p, arr); +#pragma omp dispatch is_device_ptr(p, arr) +/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(arr\\) shared\\(p\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*f1 \\(p, &arr\\);" 1 "gimple" } } */ + f2 (p, arr); +} + + diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-6.c b/gcc/testsuite/c-c++-common/gomp/dispatch-6.c new file mode 100644 index 00000000000..7c495179a13 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-6.c @@ -0,0 +1,19 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-ompexp" } */ + +void f2 (void* p); + +void test (void) +{ + void *p; + +#pragma omp dispatch +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, 0B, .*, .*\\);" 1 "ompexp" } } */ + f2 (p); +#pragma omp dispatch depend(inout: p) +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+\\\[2] = &p;" 1 "ompexp" } } */ +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, &D\.\[0-9]+, .*, .*\\);" 1 "ompexp" } } */ + f2 (p); +} + + diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-7.c b/gcc/testsuite/c-c++-common/gomp/dispatch-7.c new file mode 100644 index 00000000000..8cc4526fea0 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-7.c @@ -0,0 +1,28 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-omplower" } */ + +int f0 (void); +int f1 (void); +#pragma omp declare variant (f0) match (construct={dispatch}) +#pragma omp declare variant (f1) match (implementation={vendor(gnu)}) +int f2 (void); + +int test (void) +{ + int a, n; +#pragma omp dispatch novariants(n < 1024) nocontext(n > 1024) + a = f2 (); + return a; +} + +/* { dg-final { scan-tree-dump-times "#pragma omp dispatch nocontext\\(0\\) novariants\\(0\\) shared\\(n\\) shared\\(a\\)" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(n\\) shared\\(a\\)" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 1 "gimple" } } */ + +/* { dg-final { scan-tree-dump-times ".omp_data_o.1.n = n;" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times ".omp_data_o.1.a = &a;" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "n = .omp_data_i->n;" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = .omp_data_i->a;" 3 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "\\*D\.\[0-9]+ = D\.\[0-9]+;" 3 "omplower" } } */ From patchwork Mon May 27 11:54:38 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 1939911 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=g51JMbtA; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; 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 [8.43.85.97]) (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 4VnvLm1flhz20KL for ; Mon, 27 May 2024 21:58:56 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 85BE9382E680 for ; Mon, 27 May 2024 11:58:54 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lj1-x22f.google.com (mail-lj1-x22f.google.com [IPv6:2a00:1450:4864:20::22f]) by sourceware.org (Postfix) with ESMTPS id D0254388300A for ; Mon, 27 May 2024 11:55:38 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org D0254388300A 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 D0254388300A Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::22f ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810946; cv=none; b=Nf29wXHFFBknItjDPXsAcs2YEOsBpg0qrv8n0EUrvLDQXHHsAhvAH52mBwjbEofYJ0s5uWrs784k4BvvA+kxxnt+jUlfg1615Z1ypgDY56eVlej28iRQT7TveLfZT+P6xRpCPB+yb43awH5AdsEfBEO4waJwAY29YTAGNvgMr6s= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810946; c=relaxed/simple; bh=HrKYjO7ohPk1vdXg3dmIfMIl+8uvnlqbN2Qu/IeuA7w=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=BSXtcSPoUwK2sXHmIAxU/D/zeXjO11CnKazyvYQQyQk6KODLDdUwP8sxDZFfTLEjyOf/duQfeiWPWuJ7RTXP8PxYS7TWnVawRz511Ltofy4DJ5MRalyUv3sldbPCnRyt76tGnXG7HOcdA4SWfvidK+f+I3zpeJ7TfHcgQlfhFuM= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lj1-x22f.google.com with SMTP id 38308e7fff4ca-2e95a7545bdso37992681fa.2 for ; Mon, 27 May 2024 04:55:38 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1716810937; x=1717415737; 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=2u7RGaVdy4wxjdTQnWyTqZAncCc10x2nP+KTyvrKduk=; b=g51JMbtAIDLo0LWgBWyPBQtTq7Ys5wWCtsEjVuVAgXQPEZDapSRjSJxiC4HcTnUamc 9t/JUP9ZoIqqtj70YoYMGwKj/cxJAr35MmG12gPdXbp8SQBQYuMULo6NyLDy/X267uhl rmW+Ijg74+rdxv81GCjRcoCerCmpOoYpJtrddMo60Gi2pE/iduviOu5lmVWmZJtUHjXp sgRVLPRN1vv/+6Qxxq6xT0tXh7Pekvj0rH7bqgdQ4GZG+SMlyrzmiPIuKr+CvR+e2698 wKfwhC81NI2Q14uHwpBeVelANphQ8pS2UjGInVLr/SPRNxJwy7LmTPZDNdYRLjyoons/ NqYQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1716810937; x=1717415737; 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=2u7RGaVdy4wxjdTQnWyTqZAncCc10x2nP+KTyvrKduk=; b=dLMrORB3ivBs2nuw7UtRdkw3+vNAv0gcpGy+liklFufIJQYc9cTKueDeBuzyXTu0wT dAJu/SJqtlwtzx+NaCFwZol+sQt4za6V8NNzlo0HMjCXrHbSWXoTxHvoboy/zjewrOUI cXMD/lvfk/P1R68d0c/O+HUwFhftnNSgAw4JCwafKQYbRXbUZQ7/pYycmafe6w3opZCx tfqudFoLrfvmZAlLI0igreaytK7xQRmvr+Nrlb7+72HaZaof+MuE+WLrM7zv6X0SXc0t JkiZRDUsiEZ+mLS33IJZEd+wjORlQ/BMi22cGjCfXTFskqxvIgFI43NhsrB7ixqluJDc 6/3g== X-Gm-Message-State: AOJu0YwdO3bgVnEKYGcUohwAPO4LnhM0sH7EIaD+jxC9YI9wyq9/dnnZ YCyD8oEImeJYeTHVGxKfvsdM6OIGGKqN038W5OtdY7Yvv97iZXCXoIhgvo8JeZFRVGAD+cXfYDk x X-Google-Smtp-Source: AGHT+IFCQJ4tRmadNk89KFw6NKP1kvl2MIhRQkBCMxdaGIJO936ElWmfmYqeaJR7n9VcmpoRArDIlg== X-Received: by 2002:a2e:88d0:0:b0:2e6:ccfd:fae1 with SMTP id 38308e7fff4ca-2e95b096c66mr54465401fa.17.1716810936313; Mon, 27 May 2024 04:55:36 -0700 (PDT) Received: from localhost.localdomain ([2a0d:3344:2340:ce10::2f3]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-42108966b58sm107636955e9.3.2024.05.27.04.55.35 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Mon, 27 May 2024 04:55:35 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH 6/7] OpenMP: Fortran front-end support for dispatch + adjust_args Date: Mon, 27 May 2024 13:54:38 +0200 Message-ID: <20240527115439.3967217-7-parras@baylibre.com> X-Mailer: git-send-email 2.45.1 In-Reply-To: <20240527115439.3967217-1-parras@baylibre.com> References: <20240527115439.3967217-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE, URIBL_BLACK 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 the `dispatch` construct and the `adjust_args` clause to the Fortran front-end. gcc/fortran/ChangeLog: * dump-parse-tree.cc (show_omp_clauses): Handle novariants and nocontext clauses. (show_omp_node): Handle EXEC_OMP_DISPATCH. (show_code_node): Likewise. * frontend-passes.cc (gfc_code_walker): Handle novariants and nocontext. * gfortran.h (enum gfc_statement): Add ST_OMP_DISPATCH. (symbol_attribute): Add omp_declare_variant_need_device_ptr. (gfc_omp_clauses): Add novariants and nocontext. (gfc_omp_declare_variant): Add need_device_ptr_arg_list. (enum gfc_exec_op): Add EXEC_OMP_DISPATCH. * match.h (gfc_match_omp_dispatch): Declare. * openmp.cc (gfc_free_omp_clauses): Free novariants and nocontext clauses. (gfc_free_omp_declare_variant_list): Free need_device_ptr_arg_list namelist. (enum omp_mask2): Add OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (gfc_match_omp_clauses): Handle OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (OMP_DISPATCH_CLAUSES): Define. (gfc_match_omp_dispatch): New function. (gfc_match_omp_declare_variant): Parse adjust_args. (resolve_omp_clauses): Handle adjust_args, novariants and nocontext. Adjust handling of OMP_LIST_IS_DEVICE_PTR. (icode_code_error_callback): Handle EXEC_OMP_DISPATCH. (omp_code_to_statement): Likewise. (resolve_omp_dispatch): New function. (gfc_resolve_omp_directive): Handle EXEC_OMP_DISPATCH. * parse.cc (decode_omp_directive): Match dispatch. (next_statement): Handle ST_OMP_DISPATCH. (gfc_ascii_statement): Likewise. (parse_omp_dispatch): New function. (parse_executable): Handle ST_OMP_DISPATCH. * resolve.cc (gfc_resolve_blocks): Handle EXEC_OMP_DISPATCH. * st.cc (gfc_free_statement): Likewise. * trans-decl.cc (create_function_arglist): Declare. (gfc_get_extern_function_decl): Call it. * trans-openmp.cc (gfc_trans_omp_clauses): Handle novariants and nocontext. (gfc_trans_omp_dispatch): New function. (gfc_trans_omp_directive): Handle EXEC_OMP_DISPATCH. (gfc_trans_omp_declare_variant): Handle adjust_args. * trans.cc (trans_code): Handle EXEC_OMP_DISPATCH:. * types.def (BT_FN_PTR_CONST_PTR_INT): Declare. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/declare-variant-2.f90: Update dg-error. * gfortran.dg/gomp/adjust-args-1.f90: New test. * gfortran.dg/gomp/adjust-args-2.f90: New test. * gfortran.dg/gomp/adjust-args-3.f90: New test. * gfortran.dg/gomp/adjust-args-4.f90: New test. * gfortran.dg/gomp/adjust-args-5.f90: New test. * gfortran.dg/gomp/dispatch-1.f90: New test. * gfortran.dg/gomp/dispatch-2.f90: New test. * gfortran.dg/gomp/dispatch-3.f90: New test. * gfortran.dg/gomp/dispatch-4.f90: New test. * gfortran.dg/gomp/dispatch-5.f90: New test. * gfortran.dg/gomp/dispatch-6.f90: New test. * gfortran.dg/gomp/dispatch-7.f90: New test. * gfortran.dg/gomp/dispatch-8.f90: New test. --- gcc/fortran/dump-parse-tree.cc | 17 ++ gcc/fortran/frontend-passes.cc | 2 + gcc/fortran/gfortran.h | 11 +- gcc/fortran/match.h | 1 + gcc/fortran/openmp.cc | 193 ++++++++++++++++-- gcc/fortran/parse.cc | 38 ++++ gcc/fortran/resolve.cc | 2 + gcc/fortran/st.cc | 1 + gcc/fortran/trans-decl.cc | 9 +- gcc/fortran/trans-openmp.cc | 161 +++++++++++++++ gcc/fortran/trans.cc | 1 + gcc/fortran/types.def | 1 + .../gfortran.dg/gomp/adjust-args-1.f90 | 54 +++++ .../gfortran.dg/gomp/adjust-args-2.f90 | 18 ++ .../gfortran.dg/gomp/adjust-args-3.f90 | 26 +++ .../gfortran.dg/gomp/adjust-args-4.f90 | 58 ++++++ .../gfortran.dg/gomp/adjust-args-5.f90 | 58 ++++++ .../gfortran.dg/gomp/declare-variant-2.f90 | 6 +- gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90 | 77 +++++++ gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90 | 75 +++++++ gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90 | 39 ++++ gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90 | 19 ++ gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 | 24 +++ gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 | 38 ++++ gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 | 27 +++ gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 | 39 ++++ 26 files changed, 976 insertions(+), 19 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/adjust-args-2.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 87a65036a3d..f64dec63655 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -2119,6 +2119,18 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses) } if (omp_clauses->assume) show_omp_assumes (omp_clauses->assume); + if (omp_clauses->novariants) + { + fputs (" NOVARIANTS(", dumpfile); + show_expr (omp_clauses->novariants); + fputc (')', dumpfile); + } + if (omp_clauses->nocontext) + { + fputs (" NOCONTEXT(", dumpfile); + show_expr (omp_clauses->nocontext); + fputc (')', dumpfile); + } } /* Show a single OpenMP or OpenACC directive node and everything underneath it @@ -2156,6 +2168,9 @@ show_omp_node (int level, gfc_code *c) case EXEC_OMP_CANCEL: name = "CANCEL"; break; case EXEC_OMP_CANCELLATION_POINT: name = "CANCELLATION POINT"; break; case EXEC_OMP_CRITICAL: name = "CRITICAL"; break; + case EXEC_OMP_DISPATCH: + name = "DISPATCH"; + break; case EXEC_OMP_DISTRIBUTE: name = "DISTRIBUTE"; break; case EXEC_OMP_DISTRIBUTE_PARALLEL_DO: name = "DISTRIBUTE PARALLEL DO"; break; @@ -2257,6 +2272,7 @@ show_omp_node (int level, gfc_code *c) case EXEC_OMP_ASSUME: case EXEC_OMP_CANCEL: case EXEC_OMP_CANCELLATION_POINT: + case EXEC_OMP_DISPATCH: case EXEC_OMP_DISTRIBUTE: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: @@ -3498,6 +3514,7 @@ show_code_node (int level, gfc_code *c) case EXEC_OMP_BARRIER: case EXEC_OMP_CRITICAL: case EXEC_OMP_DEPOBJ: + case EXEC_OMP_DISPATCH: case EXEC_OMP_DISTRIBUTE: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: diff --git a/gcc/fortran/frontend-passes.cc b/gcc/fortran/frontend-passes.cc index 3c06018fdbb..1a0ef50b91d 100644 --- a/gcc/fortran/frontend-passes.cc +++ b/gcc/fortran/frontend-passes.cc @@ -5669,6 +5669,8 @@ gfc_code_walker (gfc_code **c, walk_code_fn_t codefn, walk_expr_fn_t exprfn, WALK_SUBEXPR (co->ext.omp_clauses->num_tasks); WALK_SUBEXPR (co->ext.omp_clauses->priority); WALK_SUBEXPR (co->ext.omp_clauses->detach); + WALK_SUBEXPR (co->ext.omp_clauses->novariants); + WALK_SUBEXPR (co->ext.omp_clauses->nocontext); for (idx = 0; idx < ARRAY_SIZE (list_types); idx++) for (n = co->ext.omp_clauses->lists[list_types[idx]]; n; n = n->next) diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index de1a7cd0935..361b4bece15 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -321,7 +321,7 @@ enum gfc_statement ST_OMP_ALLOCATE, ST_OMP_ALLOCATE_EXEC, ST_OMP_ALLOCATORS, ST_OMP_END_ALLOCATORS, /* Note: gfc_match_omp_nothing returns ST_NONE. */ - ST_OMP_NOTHING, ST_NONE + ST_OMP_NOTHING, ST_NONE, ST_OMP_DISPATCH }; /* Types of interfaces that we can have. Assignment interfaces are @@ -1004,6 +1004,9 @@ typedef struct ENUM_BITFIELD (gfc_omp_device_type) omp_device_type:2; unsigned omp_allocate:1; + /* Mentioned in OMP DECLARE VARIANT. */ + unsigned omp_declare_variant_need_device_ptr : 1; + /* Mentioned in OACC DECLARE. */ unsigned oacc_declare_create:1; unsigned oacc_declare_copyin:1; @@ -1431,6 +1434,7 @@ enum OMP_LIST_HAS_DEVICE_ADDR, OMP_LIST_ENTER, OMP_LIST_USES_ALLOCATORS, + OMP_LIST_ADJUST_ARGS, OMP_LIST_NUM /* Must be the last. */ }; @@ -1576,6 +1580,8 @@ typedef struct gfc_omp_clauses struct gfc_expr *depobj; struct gfc_expr *dist_chunk_size; struct gfc_expr *message; + struct gfc_expr *novariants; + struct gfc_expr *nocontext; struct gfc_omp_assumptions *assume; const char *critical_name; enum gfc_omp_default_sharing default_sharing; @@ -1702,6 +1708,7 @@ typedef struct gfc_omp_declare_variant struct gfc_symtree *variant_proc_symtree; gfc_omp_set_selector *set_selectors; + gfc_omp_namelist *need_device_ptr_arg_list; bool checked_p : 1; /* Set if previously checked for errors. */ bool error_p : 1; /* Set if error found in directive. */ @@ -3033,7 +3040,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_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS + EXEC_OMP_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS, EXEC_OMP_DISPATCH }; typedef struct gfc_code diff --git a/gcc/fortran/match.h b/gcc/fortran/match.h index b09921357fd..448f631275e 100644 --- a/gcc/fortran/match.h +++ b/gcc/fortran/match.h @@ -163,6 +163,7 @@ match gfc_match_omp_declare_simd (void); match gfc_match_omp_declare_target (void); match gfc_match_omp_declare_variant (void); match gfc_match_omp_depobj (void); +match gfc_match_omp_dispatch (void); match gfc_match_omp_distribute (void); match gfc_match_omp_distribute_parallel_do (void); match gfc_match_omp_distribute_parallel_do_simd (void); diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 5246647e6f8..b29f39febe6 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -71,7 +71,7 @@ static const struct gfc_omp_directive gfc_omp_directives[] = { {"declare target", GFC_OMP_DIR_DECLARATIVE, ST_OMP_DECLARE_TARGET}, {"declare variant", GFC_OMP_DIR_DECLARATIVE, ST_OMP_DECLARE_VARIANT}, {"depobj", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DEPOBJ}, - /* {"dispatch", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DISPATCH}, */ + {"dispatch", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DISPATCH}, {"distribute", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DISTRIBUTE}, {"do", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DO}, /* "error" becomes GFC_OMP_DIR_EXECUTABLE with at(execution) */ @@ -180,6 +180,8 @@ gfc_free_omp_clauses (gfc_omp_clauses *c) gfc_free_expr (c->num_tasks); gfc_free_expr (c->priority); gfc_free_expr (c->detach); + gfc_free_expr (c->novariants); + gfc_free_expr (c->nocontext); gfc_free_expr (c->async_expr); gfc_free_expr (c->gang_num_expr); gfc_free_expr (c->gang_static_expr); @@ -321,6 +323,8 @@ gfc_free_omp_declare_variant_list (gfc_omp_declare_variant *list) gfc_omp_declare_variant *current = list; list = list->next; gfc_free_omp_set_selector_list (current->set_selectors); + gfc_free_omp_namelist (current->need_device_ptr_arg_list, false, false, + false); free (current); } } @@ -1098,6 +1102,8 @@ enum omp_mask2 OMP_CLAUSE_ASSUMPTIONS, /* OpenMP 5.1. */ OMP_CLAUSE_USES_ALLOCATORS, /* OpenMP 5.0 */ OMP_CLAUSE_INDIRECT, /* OpenMP 5.1 */ + OMP_CLAUSE_NOVARIANTS, /* OpenMP 5.1 */ + OMP_CLAUSE_NOCONTEXT, /* OpenMP 5.1 */ /* This must come last. */ OMP_MASK2_LAST }; @@ -3215,6 +3221,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, c->assume->no_parallelism = needs_space = true; continue; } + + if ((mask & OMP_CLAUSE_NOVARIANTS) + && (m = gfc_match_dupl_check (!c->novariants, "novariants", true, + &c->novariants)) + != MATCH_NO) + { + if (m == MATCH_ERROR) + goto error; + continue; + } + if ((mask & OMP_CLAUSE_NOCONTEXT) + && (m = gfc_match_dupl_check (!c->nocontext, "nocontext", true, + &c->nocontext)) + != MATCH_NO) + { + if (m == MATCH_ERROR) + goto error; + continue; + } if ((mask & OMP_CLAUSE_NOGROUP) && (m = gfc_match_dupl_check (!c->nogroup, "nogroup")) != MATCH_NO) @@ -4530,6 +4555,9 @@ cleanup: omp_mask (OMP_CLAUSE_NOWAIT) #define OMP_ALLOCATORS_CLAUSES \ omp_mask (OMP_CLAUSE_ALLOCATE) +#define OMP_DISPATCH_CLAUSES \ + (omp_mask (OMP_CLAUSE_DEVICE) | OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOVARIANTS \ + | OMP_CLAUSE_NOCONTEXT | OMP_CLAUSE_IS_DEVICE_PTR | OMP_CLAUSE_NOWAIT) static match @@ -4843,6 +4871,12 @@ error: return MATCH_ERROR; } +match +gfc_match_omp_dispatch (void) +{ + return match_omp (EXEC_OMP_DISPATCH, OMP_DISPATCH_CLAUSES); +} + match gfc_match_omp_distribute (void) { @@ -6069,6 +6103,7 @@ gfc_match_omp_declare_variant (void) odv = gfc_get_omp_declare_variant (); odv->where = gfc_current_locus; odv->variant_proc_symtree = variant_proc_st; + odv->need_device_ptr_arg_list = NULL; odv->base_proc_symtree = base_proc_st; odv->next = NULL; odv->error_p = false; @@ -6085,13 +6120,29 @@ gfc_match_omp_declare_variant (void) return MATCH_ERROR; } + bool has_match = false, has_adjust_args = false; + locus adjust_args_loc; + for (;;) { - if (gfc_match (" match") != MATCH_YES) + enum clause + { + match, + adjust_args + } ccode; + + if (gfc_match (" match") == MATCH_YES) + ccode = match; + else if (gfc_match (" adjust_args") == MATCH_YES) + { + ccode = adjust_args; + adjust_args_loc = gfc_current_locus; + } + else { if (first_p) { - gfc_error ("expected % at %C"); + gfc_error ("expected % or % at %C"); return MATCH_ERROR; } else @@ -6104,18 +6155,86 @@ gfc_match_omp_declare_variant (void) return MATCH_ERROR; } - if (gfc_match_omp_context_selector_specification (odv) != MATCH_YES) - return MATCH_ERROR; - - if (gfc_match (" )") != MATCH_YES) + if (ccode == match) { - gfc_error ("expected %<)%> at %C"); - return MATCH_ERROR; + has_match = true; + if (gfc_match_omp_context_selector_specification (odv) + != MATCH_YES) + return MATCH_ERROR; + if (gfc_match (" )") != MATCH_YES) + { + gfc_error ("expected %<)%> at %C"); + return MATCH_ERROR; + } + } + else if (ccode == adjust_args) + { + has_adjust_args = true; + bool need_device_ptr_p; + if (gfc_match (" nothing") == MATCH_YES) + need_device_ptr_p = false; + else if (gfc_match (" need_device_ptr") == MATCH_YES) + need_device_ptr_p = true; + else + { + gfc_error ("expected % or % at %C"); + return MATCH_ERROR; + } + if (need_device_ptr_p) + { + if (gfc_match_omp_variable_list (" :", + &odv->need_device_ptr_arg_list, + false) + != MATCH_YES) + { + gfc_error ("expected argument list at %C"); + return MATCH_ERROR; + } + for (gfc_omp_namelist *n = odv->need_device_ptr_arg_list; + n != NULL; n = n->next) + { + if (!n->sym->attr.dummy) + { + gfc_error ("list item %qs at %L is not a dummy argument", + n->sym->name, &n->where); + return MATCH_ERROR; + } + if (n->sym->ts.type != BT_DERIVED + || !n->sym->ts.u.derived->ts.is_iso_c) + { + gfc_error ("argument list item %qs in " + "% at %L must be of " + "TYPE(C_PTR)", + n->sym->name, &n->where); + return MATCH_ERROR; + } + } + } + else + { + gfc_omp_namelist *nothing_arg_list = NULL; + if (gfc_match_omp_variable_list (" :", ¬hing_arg_list, false) + != MATCH_YES) + { + gfc_error ("expected argument list at %C"); + return MATCH_ERROR; + } + gfc_free_omp_namelist (nothing_arg_list, false, false, false); + } } first_p = false; } + if (has_adjust_args && !has_match) + { + gfc_error ("an % clause at %C can only be specified if the " + "% selector of the construct selector set appears " + "in the % clause", + &adjust_args_loc); + return MATCH_ERROR; + } + return MATCH_YES; } @@ -7544,7 +7663,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, "DEVICE_RESIDENT", "LINK", "USE_DEVICE", "CACHE", "IS_DEVICE_PTR", "USE_DEVICE_PTR", "USE_DEVICE_ADDR", "NONTEMPORAL", "ALLOCATE", "HAS_DEVICE_ADDR", "ENTER", - "USES_ALLOCATORS" }; + "USES_ALLOCATORS", "ADJUST_ARGS" }; STATIC_ASSERT (ARRAY_SIZE (clause_names) == OMP_LIST_NUM); if (omp_clauses == NULL) @@ -7726,6 +7845,26 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, gfc_error ("FINAL clause at %L requires a scalar LOGICAL expression", &expr->where); } + if (omp_clauses->novariants) + { + gfc_expr *expr = omp_clauses->novariants; + if (!gfc_resolve_expr (expr) || expr->ts.type != BT_LOGICAL + || expr->rank != 0) + gfc_error ( + "NOVARIANTS clause at %L requires a scalar LOGICAL expression", + &expr->where); + if_without_mod = true; + } + if (omp_clauses->nocontext) + { + gfc_expr *expr = omp_clauses->nocontext; + if (!gfc_resolve_expr (expr) || expr->ts.type != BT_LOGICAL + || expr->rank != 0) + gfc_error ( + "NOCONTEXT clause at %L requires a scalar LOGICAL expression", + &expr->where); + if_without_mod = true; + } if (omp_clauses->num_threads) resolve_positive_int_expr (omp_clauses->num_threads, "NUM_THREADS"); if (omp_clauses->chunk_size) @@ -8675,9 +8814,9 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, last = NULL; for (n = omp_clauses->lists[list]; n != NULL; ) { - if (n->sym->ts.type == BT_DERIVED - && n->sym->ts.u.derived->ts.is_iso_c - && code->op != EXEC_OMP_TARGET) + if ((n->sym->ts.type != BT_DERIVED + || !n->sym->ts.u.derived->ts.is_iso_c) + && code->op == EXEC_OMP_DISPATCH) /* Non-TARGET (i.e. DISPATCH) requires a C_PTR. */ gfc_error ("List item %qs in %s clause at %L must be of " "TYPE(C_PTR)", n->sym->name, name, &n->where); @@ -10290,6 +10429,7 @@ icode_code_error_callback (gfc_code **codep, case EXEC_OMP_MASKED_TASKLOOP_SIMD: case EXEC_OMP_SCOPE: case EXEC_OMP_ERROR: + case EXEC_OMP_DISPATCH: gfc_error ("%s cannot contain OpenMP directive in intervening code " "at %L", state->name, &code->loc); @@ -11168,6 +11308,8 @@ omp_code_to_statement (gfc_code *code) return ST_OMP_PARALLEL_LOOP; case EXEC_OMP_DEPOBJ: return ST_OMP_DEPOBJ; + case EXEC_OMP_DISPATCH: + return ST_OMP_DISPATCH; default: gcc_unreachable (); } @@ -11583,6 +11725,26 @@ resolve_omp_target (gfc_code *code) #undef GFC_IS_TEAMS_CONSTRUCT } +static void +resolve_omp_dispatch (gfc_code *code) +{ + gfc_code *next = code->block->next; + if (next == NULL) + return; + gfc_exec_op op = next->op; + if (op != EXEC_CALL + && (op != EXEC_ASSIGN || next->expr2->expr_type != EXPR_FUNCTION)) + gfc_error ( + "% directive at %L must be followed by a procedure " + "call with optional assignment", + &code->loc); + + if ((op == EXEC_CALL && next->resolved_sym->attr.proc_pointer) + || (op == EXEC_ASSIGN && gfc_expr_attr (next->expr2).proc_pointer)) + gfc_error ("% directive at %L cannot be followed by a " + "procedure pointer", + &code->loc); +} /* Resolve OpenMP directive clauses and check various requirements of each directive. */ @@ -11696,6 +11858,11 @@ 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_DISPATCH: + if (code->ext.omp_clauses) + resolve_omp_clauses (code, code->ext.omp_clauses, ns); + resolve_omp_dispatch (code); + break; default: break; } diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc index 79c810c86ba..74fc249269d 100644 --- a/gcc/fortran/parse.cc +++ b/gcc/fortran/parse.cc @@ -1050,6 +1050,7 @@ decode_omp_directive (void) break; case 'd': matcho ("depobj", gfc_match_omp_depobj, ST_OMP_DEPOBJ); + matcho ("dispatch", gfc_match_omp_dispatch, ST_OMP_DISPATCH); matchs ("distribute parallel do simd", gfc_match_omp_distribute_parallel_do_simd, ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD); @@ -1916,6 +1917,7 @@ next_statement (void) case ST_OMP_LOOP: case ST_OMP_PARALLEL_LOOP: case ST_OMP_TEAMS_LOOP: \ case ST_OMP_TARGET_PARALLEL_LOOP: case ST_OMP_TARGET_TEAMS_LOOP: \ case ST_OMP_ALLOCATE_EXEC: case ST_OMP_ALLOCATORS: case ST_OMP_ASSUME: \ + case ST_OMP_DISPATCH: \ case ST_CRITICAL: \ case ST_OACC_PARALLEL_LOOP: case ST_OACC_PARALLEL: case ST_OACC_KERNELS: \ case ST_OACC_DATA: case ST_OACC_HOST_DATA: case ST_OACC_LOOP: \ @@ -2597,6 +2599,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel) case ST_OMP_DEPOBJ: p = "!$OMP DEPOBJ"; break; + case ST_OMP_DISPATCH: + p = "!$OMP DISPATCH"; + break; case ST_OMP_DISTRIBUTE: p = "!$OMP DISTRIBUTE"; break; @@ -6183,6 +6188,35 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only) } +static gfc_statement +parse_omp_dispatch (void) +{ + gfc_statement st; + gfc_code *cp, *np; + gfc_state_data s; + + accept_statement (ST_OMP_DISPATCH); + + cp = gfc_state_stack->tail; + push_state (&s, COMP_OMP_STRUCTURED_BLOCK, NULL); + np = new_level (cp); + np->op = cp->op; + np->block = NULL; + + st = next_statement (); + if (st == ST_CALL || st == ST_ASSIGNMENT) + accept_statement (st); + else + { + gfc_error ("% directive must be followed by a procedure " + "call with optional assignment at %C"); + reject_statement (); + } + pop_state (); + st = next_statement (); + return st; +} + /* Accept a series of executable statements. We return the first statement that doesn't fit to the caller. Any block statements are passed on to the correct handler, which usually passes the buck @@ -6383,6 +6417,10 @@ parse_executable (gfc_statement st) st = parse_omp_oacc_atomic (true); continue; + case ST_OMP_DISPATCH: + st = parse_omp_dispatch (); + continue; + default: return st; } diff --git a/gcc/fortran/resolve.cc b/gcc/fortran/resolve.cc index d7a0856fcca..755d1302ce9 100644 --- a/gcc/fortran/resolve.cc +++ b/gcc/fortran/resolve.cc @@ -11378,6 +11378,7 @@ gfc_resolve_blocks (gfc_code *b, gfc_namespace *ns) case EXEC_OMP_ALLOCATORS: case EXEC_OMP_ASSUME: case EXEC_OMP_CRITICAL: + case EXEC_OMP_DISPATCH: case EXEC_OMP_DISTRIBUTE: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: @@ -13054,6 +13055,7 @@ start: case EXEC_OMP_CRITICAL: case EXEC_OMP_FLUSH: case EXEC_OMP_DEPOBJ: + case EXEC_OMP_DISPATCH: case EXEC_OMP_DISTRIBUTE: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: diff --git a/gcc/fortran/st.cc b/gcc/fortran/st.cc index 6a605ad91d4..90ee1352ba4 100644 --- a/gcc/fortran/st.cc +++ b/gcc/fortran/st.cc @@ -222,6 +222,7 @@ gfc_free_statement (gfc_code *p) case EXEC_OMP_CANCELLATION_POINT: case EXEC_OMP_CRITICAL: case EXEC_OMP_DEPOBJ: + case EXEC_OMP_DISPATCH: case EXEC_OMP_DISTRIBUTE: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc index dca7779528b..4390769146a 100644 --- a/gcc/fortran/trans-decl.cc +++ b/gcc/fortran/trans-decl.cc @@ -2124,6 +2124,8 @@ get_proc_pointer_decl (gfc_symbol *sym) return decl; } +static void +create_function_arglist (gfc_symbol *sym); /* Get a basic decl for an external function. */ @@ -2377,7 +2379,12 @@ module_sym: if (sym->formal_ns->omp_declare_simd) gfc_trans_omp_declare_simd (sym->formal_ns); if (flag_openmp) - gfc_trans_omp_declare_variant (sym->formal_ns); + { + // We need DECL_ARGUMENTS to put attributes on, in case some arguments + // need adjustment + create_function_arglist (sym->formal_ns->proc_name); + gfc_trans_omp_declare_variant (sym->formal_ns); + } } return fndecl; diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index f867e2240bf..5e4450184d1 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -4233,6 +4233,36 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, omp_clauses = gfc_trans_add_clause (c, omp_clauses); } + if (clauses->novariants) + { + tree novariants_var; + + gfc_init_se (&se, NULL); + gfc_conv_expr (&se, clauses->novariants); + gfc_add_block_to_block (block, &se.pre); + novariants_var = gfc_evaluate_now (se.expr, block); + gfc_add_block_to_block (block, &se.post); + + c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_NOVARIANTS); + OMP_CLAUSE_NOVARIANTS_EXPR (c) = novariants_var; + omp_clauses = gfc_trans_add_clause (c, omp_clauses); + } + + if (clauses->nocontext) + { + tree nocontext_var; + + gfc_init_se (&se, NULL); + gfc_conv_expr (&se, clauses->nocontext); + gfc_add_block_to_block (block, &se.pre); + nocontext_var = gfc_evaluate_now (se.expr, block); + gfc_add_block_to_block (block, &se.post); + + c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_NOCONTEXT); + OMP_CLAUSE_NOCONTEXT_EXPR (c) = nocontext_var; + omp_clauses = gfc_trans_add_clause (c, omp_clauses); + } + if (clauses->num_threads) { tree num_threads; @@ -6311,6 +6341,30 @@ gfc_trans_omp_depobj (gfc_code *code) return gfc_finish_block (&block); } +static tree +gfc_trans_omp_dispatch (gfc_code *code) +{ + stmtblock_t block; + gfc_code *next = code->block->next; + // assume ill-formed "function dispatch structured + // block" have already been rejected by resolve_omp_dispatch + gcc_assert (next->op == EXEC_CALL || next->op == EXEC_ASSIGN); + + tree body = gfc_trans_code (next); + gfc_start_block (&block); + tree omp_clauses + = gfc_trans_omp_clauses (&block, code->ext.omp_clauses, code->loc); + + tree stmt = make_node (OMP_DISPATCH); + SET_EXPR_LOCATION (stmt, gfc_get_location (&code->loc)); + TREE_TYPE (stmt) = void_type_node; + OMP_DISPATCH_BODY (stmt) = body; + OMP_DISPATCH_CLAUSES (stmt) = omp_clauses; + + gfc_add_expr_to_block (&block, stmt); + return gfc_finish_block (&block); +} + static tree gfc_trans_omp_error (gfc_code *code) { @@ -8221,6 +8275,8 @@ gfc_trans_omp_directive (gfc_code *code) case EXEC_OMP_TASKLOOP: return gfc_trans_omp_do (code, code->op, NULL, code->ext.omp_clauses, NULL); + case EXEC_OMP_DISPATCH: + return gfc_trans_omp_dispatch (code); case EXEC_OMP_DISTRIBUTE_PARALLEL_DO: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: case EXEC_OMP_DISTRIBUTE_SIMD: @@ -8337,6 +8393,7 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns) tree base_fn_decl = ns->proc_name->backend_decl; gfc_namespace *search_ns = ns; gfc_omp_declare_variant *next; + vec adjust_args_list = vNULL; for (gfc_omp_declare_variant *odv = search_ns->omp_declare_variant; search_ns; odv = next) @@ -8532,6 +8589,19 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns) variant_proc_sym = NULL; } } + if (odv->need_device_ptr_arg_list != NULL + && omp_get_context_selector (set_selectors, OMP_TRAIT_SET_CONSTRUCT, + OMP_TRAIT_CONSTRUCT_DISPATCH) + == NULL_TREE) + { + gfc_error ("an % clause can only be " + "specified if the " + "% selector of the construct " + "selector set appears " + "in the % clause at %L", + &odv->where); + variant_proc_sym = NULL; + } if (variant_proc_sym != NULL) { gfc_set_sym_referenced (variant_proc_sym); @@ -8548,6 +8618,97 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns) DECL_ATTRIBUTES (base_fn_decl) = tree_cons (id, build_tree_list (variant, set_selectors), DECL_ATTRIBUTES (base_fn_decl)); + + // Handle adjust_args + for (gfc_omp_namelist *arg_list + = odv->need_device_ptr_arg_list; + arg_list != NULL; arg_list = arg_list->next) + { + if (arg_list->sym->backend_decl == NULL_TREE) + { + gfc_error ( + "%s at %L is not a base function argument", + arg_list->sym->name, &arg_list->where); + continue; + } + + tree base_fn_arg_decl = arg_list->sym->backend_decl; + if (base_fn_arg_decl != error_mark_node) + { + // Is t specified more than once? + if (adjust_args_list.contains (base_fn_arg_decl)) + { + gfc_error ( + "%qD at %L is specified more than once", + base_fn_arg_decl, &arg_list->where); + continue; + } + adjust_args_list.safe_push (base_fn_arg_decl); + + // Handle variant argument + tree variant + = gfc_get_symbol_decl (variant_proc_sym); + tree variant_parm = DECL_ARGUMENTS (variant); + int idx; + tree arg; + for (arg = DECL_ARGUMENTS (base_fn_decl), idx = 0; + arg != NULL; arg = TREE_CHAIN (arg), idx++) + if (arg == base_fn_arg_decl) + break; + gcc_assert (arg != NULL_TREE); + if (variant_parm == NULL_TREE) + { + gfc_formal_arglist *arg + = variant_proc_sym->formal; + for (int i = 0; i < idx; i++) + { + arg = arg->next; + gcc_assert (arg != NULL); + } + + // Check we got the right parameter name + if (strcmp (arg_list->sym->name, arg->sym->name) + != 0) + { + gfc_error ("%s at %L is not a variant " + "function argument", + arg_list->sym->name, + &arg_list->where); + continue; + } + arg->sym->attr + .omp_declare_variant_need_device_ptr + = 1; + } + else + { + for (int i = 0; i < idx; i++) + { + variant_parm = TREE_CHAIN (variant_parm); + gcc_assert (variant_parm != NULL_TREE); + } + // Check we got the right parameter name + if (strcmp (arg_list->sym->name, + IDENTIFIER_POINTER ( + DECL_NAME (variant_parm))) + != 0) + { + gfc_error ("%s at %L is not a variant " + "function argument", + arg_list->sym->name, + &arg_list->where); + continue; + } + + tree attr = tree_cons ( + get_identifier ( + "omp declare variant adjust_args " + "need_device_ptr"), + NULL_TREE, DECL_ATTRIBUTES (variant_parm)); + DECL_ATTRIBUTES (variant_parm) = attr; + } + } + } } } } diff --git a/gcc/fortran/trans.cc b/gcc/fortran/trans.cc index badad6ae892..2795cdf7464 100644 --- a/gcc/fortran/trans.cc +++ b/gcc/fortran/trans.cc @@ -2596,6 +2596,7 @@ trans_code (gfc_code * code, tree cond) case EXEC_OMP_CANCELLATION_POINT: case EXEC_OMP_CRITICAL: case EXEC_OMP_DEPOBJ: + case EXEC_OMP_DISPATCH: case EXEC_OMP_DISTRIBUTE: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def index 390cc9542f7..5047c8f816a 100644 --- a/gcc/fortran/types.def +++ b/gcc/fortran/types.def @@ -120,6 +120,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_INT_BOOL, BT_BOOL, BT_INT, BT_BOOL) DEF_FUNCTION_TYPE_2 (BT_FN_VOID_PTR_PTRMODE, BT_VOID, BT_PTR, BT_PTRMODE) DEF_FUNCTION_TYPE_2 (BT_FN_VOID_CONST_PTR_SIZE, BT_VOID, BT_CONST_PTR, BT_SIZE) +DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_INT, BT_PTR, BT_CONST_PTR, BT_INT) DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR) diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90 new file mode 100644 index 00000000000..982c2c2cb9d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90 @@ -0,0 +1,54 @@ +! Test parsing of OMP clause adjust_args +! { dg-do compile } + +module main + use iso_c_binding, only: c_ptr + implicit none + integer :: b +interface + integer function f0 (a) + import c_ptr + type(c_ptr), intent(inout) :: a + end function + integer function g (a) + import c_ptr + type(c_ptr), intent(inout) :: a + end function + integer function f1 (i) + integer, intent(in) :: i + end function + + integer function f3 (a) + import c_ptr + type(c_ptr), intent(inout) :: a + !$omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) ! { dg-error "expected 'nothing' or 'need_device_ptr' at .1." } +end function +integer function f4 (a) +import c_ptr +type(c_ptr), intent(inout) :: a + !$omp declare variant (f0) adjust_args (nothing: a) ! { dg-error "an 'adjust_args' clause at .1. can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } +end function +integer function f5 (i) +integer, intent(inout) :: i + !$omp declare variant (f1) match (construct={dispatch}) adjust_args () ! { dg-error "expected 'nothing' or 'need_device_ptr' at .1." } +end function +integer function f6 (i) +integer, intent(inout) :: i + !$omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) ! { dg-error "expected argument list at .1." } +end function +integer function f7 (i) +integer, intent(inout) :: i + !$omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) ! { dg-error "expected argument list at .1." } + end function + integer function f9 (i) + integer, intent(inout) :: i + !$omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: i) ! { dg-error "argument list item 'i' in 'need_device_ptr' at .1. must be of TYPE.C_PTR." } +end function + integer function f12 (a) + import c_ptr + type(c_ptr), intent(inout) :: a + !$omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) ! { dg-error "list item 'b' at .1. is not a dummy argument" } +end function + + end interface +end module diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-2.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-2.f90 new file mode 100644 index 00000000000..c65a4839ca5 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-2.f90 @@ -0,0 +1,18 @@ +! Test resolution of OMP clause adjust_args +! { dg-do compile } + +module main + implicit none +interface +subroutine f1 (i) + integer, intent(inout) :: i +end subroutine +end interface +contains + + subroutine f3 (i) + integer, intent(inout) :: i + !$omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: z) ! { dg-error "Symbol 'z' at .1. has no IMPLICIT type" } + end subroutine + +end module diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90 new file mode 100644 index 00000000000..b731cb340c1 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90 @@ -0,0 +1,26 @@ +! Test translation of OMP clause adjust_args +! { dg-do compile } + +module main + use iso_c_binding, only: c_ptr + implicit none + !type(c_ptr) :: a + +contains + subroutine base2 (a) + type(c_ptr), intent(inout) :: a + !$omp declare variant (variant2) match (construct={parallel}) adjust_args (need_device_ptr: a) ! { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause at .1." } + end subroutine + subroutine base3 (a) + type(c_ptr), intent(inout) :: a + !$omp declare variant (variant2) match (construct={dispatch}) adjust_args (need_device_ptr: a) adjust_args (need_device_ptr: a) ! { dg-error "'a' at .1. is specified more than once" } + end subroutine + + subroutine variant2 (a) + type(c_ptr), intent(inout) :: a + end subroutine + subroutine variant3 (i) + integer :: i + end subroutine + +end module diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90 new file mode 100644 index 00000000000..75e884044b2 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90 @@ -0,0 +1,58 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module main + use iso_c_binding, only: c_ptr + implicit none + + type :: struct + integer :: a + real :: b + end type + + interface + integer function f(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + end function + integer function f0(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + !$omp declare variant (f) match (construct={dispatch}) & + !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b, c) + end function + integer function f1(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + !$omp declare variant (f) match (construct={dispatch}) & + !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c) + end function + end interface + +contains +subroutine test + integer :: a + type(c_ptr) :: b + type(c_ptr) :: c(2) + type(struct) :: s + + s%a = f0 (a, b, c) + !$omp dispatch + s%a = f0 (a, b, c) + + s%b = f1 (a, b, c) + !$omp dispatch + s%b = f1 (a, b, c) + +end subroutine +end module + +! { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&parm\.\[0-9]+, D\.\[0-9]+\\);" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&b, D\.\[0-9]+\\);" 2 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90 new file mode 100644 index 00000000000..75e884044b2 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90 @@ -0,0 +1,58 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module main + use iso_c_binding, only: c_ptr + implicit none + + type :: struct + integer :: a + real :: b + end type + + interface + integer function f(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + end function + integer function f0(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + !$omp declare variant (f) match (construct={dispatch}) & + !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b, c) + end function + integer function f1(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + !$omp declare variant (f) match (construct={dispatch}) & + !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c) + end function + end interface + +contains +subroutine test + integer :: a + type(c_ptr) :: b + type(c_ptr) :: c(2) + type(struct) :: s + + s%a = f0 (a, b, c) + !$omp dispatch + s%a = f0 (a, b, c) + + s%b = f1 (a, b, c) + !$omp dispatch + s%b = f1 (a, b, c) + +end subroutine +end module + +! { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&parm\.\[0-9]+, D\.\[0-9]+\\);" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&b, D\.\[0-9]+\\);" 2 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90 index 7fc5071feff..62d2cb96fac 100644 --- a/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90 @@ -18,10 +18,10 @@ contains !$omp declare variant match(user={condition(.false.)}) ! { dg-error "expected '\\(' at .1." } end subroutine subroutine f6 () - !$omp declare variant (f1) ! { dg-error "expected 'match' at .1." } + !$omp declare variant (f1) ! { dg-error "expected 'match' or 'adjust_args' at .1." } end subroutine subroutine f7 () - !$omp declare variant (f1) simd ! { dg-error "expected 'match' at .1." } + !$omp declare variant (f1) simd ! { dg-error "expected 'match' or 'adjust_args' at .1." } end subroutine subroutine f8 () !$omp declare variant (f1) match ! { dg-error "expected '\\(' at .1." } @@ -183,7 +183,7 @@ contains !$omp declare variant (f1) match(construct={requires}) ! { dg-warning "unknown selector 'requires' for context selector set 'construct' at .1." } end subroutine subroutine f75 () - !$omp declare variant (f1),match(construct={parallel}) ! { dg-error "expected 'match' at .1." } + !$omp declare variant (f1),match(construct={parallel}) ! { dg-error "expected 'match' or 'adjust_args' at .1." } end subroutine subroutine f76 () !$omp declare variant (f1) match(implementation={atomic_default_mem_order("relaxed")}) ! { dg-error "expected identifier at .1." } diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90 new file mode 100644 index 00000000000..12c30904131 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90 @@ -0,0 +1,77 @@ +module main + use iso_c_binding, only: c_ptr + implicit none + contains + + subroutine f1 () + integer :: a, b, arr(10) + real :: x + complex :: c + character :: ch + logical :: bool + type :: struct + integer :: a + real :: b + end type + type(struct) :: s + type(c_ptr) :: p + + interface + subroutine f0 (a, c, bool, s) + import :: struct + integer, intent(in) :: a + complex, intent(out) :: c + logical, intent(inout) :: bool + type(struct) :: s + end subroutine + integer function f2 (arr, x, ch, b) + integer, intent(inout) :: arr(:) + real, intent(in) :: x + character, intent(out) :: ch + real :: b + end function + subroutine f3 (p) + import :: c_ptr + type(c_ptr) :: p + end subroutine + integer function f4 () + end function + end interface + + !$omp dispatch + b = f2(arr, x, ch, s%b) + !$omp dispatch + c = f2(arr(:5), x * 2.4, ch, s%b) + !$omp dispatch + arr(1) = f2(arr, x, ch, s%b) + !$omp dispatch + s%a = f2(arr, x, ch, s%b) + !$omp dispatch + x = f2(arr, x, ch, s%b) + !$omp dispatch + call f0(a, c, bool, s) + !$omp dispatch + call f0(f4(), c, bool, s) + + !$omp dispatch nocontext(.TRUE.) + call f0(a, c, bool, s) + !$omp dispatch nocontext(arr(2) < 10) + call f0(a, c, bool, s) + !$omp dispatch novariants(.FALSE.) + call f0(a, c, bool, s) + !$omp dispatch novariants(bool) + call f0(a, c, bool, s) + !$omp dispatch nowait + call f0(a, c, bool, s) + !$omp dispatch device(arr(9)) + call f0(a, c, bool, s) + !$omp dispatch device(a + a) + call f0(a, c, bool, s) + !$omp dispatch device(-25373654) + call f0(a, c, bool, s) + !$omp dispatch is_device_ptr(p) + call f3(p) + !$omp dispatch depend(in: a, c, bool) depend(inout: s, arr(:3)) + call f0(a, c, bool, s) + end subroutine +end module diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90 new file mode 100644 index 00000000000..f52df4446c4 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90 @@ -0,0 +1,75 @@ +module main + implicit none + contains + + subroutine f1 () + integer :: a, b, arr(10) + real :: x + complex :: c + character :: ch + logical :: bool + type :: struct + integer :: a + real :: b + end type + type(struct) :: s + + interface + subroutine f0 (a, c, bool, s) + import :: struct + integer, intent(in) :: a + complex, intent(out) :: c + logical, intent(inout) :: bool + type(struct) :: s + end subroutine + integer function f2 (arr, x, ch, b) + integer, intent(inout) :: arr(:) + real, intent(in) :: x + character, intent(out) :: ch + real :: b + end function + end interface + procedure(f0), pointer:: fp => NULL() + + !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. must be followed by a procedure call with optional assignment" } +50 b = f2(arr, x, ch, s%b) + a + !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. must be followed by a procedure call with optional assignment" } + a = b + !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. must be followed by a procedure call with optional assignment" } + b = Not (2) + !$omp dispatch + !$omp threadprivate(a) !{ dg-error "'OMP DISPATCH' directive must be followed by a procedure call with optional assignment at .1." } + a = f2(arr, x, ch, s%b) + !$omp dispatch + print *, 'This is not allowed here.' !{ dg-error "'OMP DISPATCH' directive must be followed by a procedure call with optional assignment at .1." } + !$omp dispatch + goto 50 !{ dg-error "'OMP DISPATCH' directive must be followed by a procedure call with optional assignment at .1." } + !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. cannot be followed by a procedure pointer" } + call fp(a, c, bool, s) + + !$omp dispatch nocontext(s) !{ dg-error "NOCONTEXT clause at .1. requires a scalar LOGICAL expression" } + call f0(a, c, bool, s) + !$omp dispatch nocontext(a, b) !{ dg-error "Invalid expression after 'nocontext.' at .1." } + call f0(a, c, bool, s) + !$omp dispatch nocontext(a) nocontext(b) !{ dg-error "Duplicated 'nocontext' clause at .1." } + call f0(a, c, bool, s) + !$omp dispatch novariants(s) !{ dg-error "NOVARIANTS clause at .1. requires a scalar LOGICAL expression" } + call f0(a, c, bool, s) + !$omp dispatch novariants(a, b) !{ dg-error "Invalid expression after 'novariants.' at .1." } + call f0(a, c, bool, s) + !$omp dispatch novariants(a) novariants(b) !{ dg-error "Duplicated 'novariants' clause at .1." } + call f0(a, c, bool, s) + !$omp dispatch nowait nowait !{ dg-error "Duplicated 'nowait' clause at .1." } + call f0(a, c, bool, s) + !$omp dispatch device(x) !{ dg-error "DEVICE clause at .1. requires a scalar INTEGER expression" } + call f0(a, c, bool, s) + !$omp dispatch device(arr) !{ dg-error "DEVICE clause at .1. requires a scalar INTEGER expression" } + call f0(a, c, bool, s) + !$omp dispatch is_device_ptr(x) !{ dg-error "List item 'x' in IS_DEVICE_PTR clause at .1. must be of TYPE.C_PTR." } + call f0(a, c, bool, s) + !$omp dispatch is_device_ptr(arr) !{ dg-error "List item 'arr' in IS_DEVICE_PTR clause at .1. must be of TYPE.C_PTR." } + call f0(a, c, bool, s) + !$omp dispatch depend(inout: f0) !{ dg-error "Object 'f0' is not a variable at .1." } + call f0(a, c, bool, s) + end subroutine +end module diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90 new file mode 100644 index 00000000000..84590fd883a --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90 @@ -0,0 +1,39 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module main + implicit none + interface + integer function f0 () + end function + + integer function f1 () + end function + + integer function f2 () + !$omp declare variant (f0) match (construct={dispatch}) + !$omp declare variant (f1) match (implementation={vendor(gnu)}) + end function + end interface + contains + + integer function test () + integer :: a + + !$omp dispatch + a = f2 () + !$omp dispatch novariants(.TRUE.) + a = f2 () + !$omp dispatch novariants(.FALSE.) + a = f2 () + !$omp dispatch nocontext(.TRUE.) + a = f2 () + !$omp dispatch nocontext(.FALSE.) + a = f2 () + end function +end module + + +! { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 3 "gimple" } } +! { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90 new file mode 100644 index 00000000000..149d0613b97 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90 @@ -0,0 +1,19 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module main + implicit none + interface + subroutine f2 () + end subroutine + end interface + contains + + subroutine test () + !$omp dispatch ! { dg-final { scan-tree-dump-times "#pragma omp task if\\(0\\)" 1 "gimple" } } + call f2 () + !$omp dispatch nowait ! { dg-final { scan-tree-dump-times "#pragma omp task if\\(1\\)" 1 "gimple" } } + call f2 () + end subroutine +end module + diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 new file mode 100644 index 00000000000..e45397f3f96 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 @@ -0,0 +1,24 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module main + implicit none + interface + subroutine f2 (a) + integer, intent(in) :: a + end subroutine + end interface + contains + + subroutine test () + integer :: a + + !$omp dispatch device(-25373654) + ! { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device \\(-25373654\\);" 1 "gimple" } } + call f2 (a) + !$omp dispatch device(a + a) + ! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = a.0_1 \\* 2;.*#pragma omp dispatch device\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\).*#pragma omp task shared\\(D\.\[0-9]+\\).*__builtin_omp_set_default_device \\(D\.\[0-9]+\\);" 1 "gimple" } } + call f2 (a) + end subroutine +end module + diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 new file mode 100644 index 00000000000..9f4fa2970ca --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 @@ -0,0 +1,38 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module main + use iso_c_binding, only: c_ptr + implicit none + interface + subroutine f1 (p, p2) + import :: c_ptr + type(c_ptr), intent(out) :: p + type(c_ptr), intent(in) :: p2 + end subroutine + subroutine f2 (p, p2) + import :: c_ptr + type(c_ptr), intent(out) :: p + type(c_ptr), intent(in) :: p2 + !$omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: p, p2) + end subroutine + end interface + contains + + subroutine test () + type(c_ptr) :: p, p2 + + !$omp dispatch + call f2 (p, p2) + !$omp dispatch is_device_ptr(p) + ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p\\) shared\\(p2\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*integer\\(kind=4\\) D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&p2, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(&p, D\.\[0-9]+\\);" 1 "gimple" } } + call f2 (p, p2) + !$omp dispatch is_device_ptr(p2) + ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p2\\) shared\\(p\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*integer\\(kind=4\\) D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&p, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(D\.\[0-9]+, &p2\\);" 1 "gimple" } } + call f2 (p, p2) + !$omp dispatch is_device_ptr(p, p2) + ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p\\) shared\\(p2\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*f1 \\(&p, &p2\\);" 1 "gimple" } } + call f2 (p, p2) + end subroutine +end module + diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 new file mode 100644 index 00000000000..32b6347be67 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 @@ -0,0 +1,27 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-ompexp" } + +module main + use iso_c_binding, only: c_ptr + implicit none + interface + subroutine f2 (p) + import :: c_ptr + type(c_ptr), intent(out) :: p + end subroutine + end interface + contains + + subroutine test () + type(c_ptr) :: p + + !$omp dispatch + ! { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, 0B, .*, .*\\);" 1 "ompexp" } } + call f2 (p) + !$omp dispatch depend(inout: p) + ! { dg-final { scan-tree-dump-times "D\.\[0-9]+\\\[2] = &p;" 1 "ompexp" } } + ! { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, &D\.\[0-9]+, .*, .*\\);" 1 "ompexp" } } + call f2 (p) + end subroutine +end module + diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 new file mode 100644 index 00000000000..6771336aa33 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 @@ -0,0 +1,39 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple -fdump-tree-omplower" } + +module main + use iso_c_binding, only: c_ptr + implicit none + interface + integer function f0 () + end function + integer function f1 () + end function + integer function f2 () + !$omp declare variant (f0) match (construct={dispatch}) + !$omp declare variant (f1) match (implementation={vendor(gnu)}) + end function + end interface + contains + + subroutine test () + integer :: a, n + + !$omp dispatch novariants(n < 1024) nocontext(n > 1024) + a = f2 () + end subroutine +end module + +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = n <= 1023;" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = n > 1024;" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp dispatch novariants\\(0\\) nocontext\\(0\\) shared\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 1 "gimple" } } + +! { dg-final { scan-tree-dump-times ".omp_data_o.1.D\.\[0-9]+ = D\.\[0-9]+;" 2 "omplower" } } +! { dg-final { scan-tree-dump-times ".omp_data_o.1.a = &a;" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = .omp_data_i->D\.\[0-9]+;" 2 "omplower" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = .omp_data_i->a;" 3 "omplower" } } +! { dg-final { scan-tree-dump-times "\\*D\.\[0-9]+ = D\.\[0-9]+;" 3 "omplower" } } From patchwork Mon May 27 11:54:39 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 1939909 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=GqvLkXGM; 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 4VnvKY2Gvwz20KL for ; Mon, 27 May 2024 21:57:53 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 938E33883004 for ; Mon, 27 May 2024 11:57:51 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lj1-x229.google.com (mail-lj1-x229.google.com [IPv6:2a00:1450:4864:20::229]) by sourceware.org (Postfix) with ESMTPS id 3AE333883018 for ; Mon, 27 May 2024 11:55:39 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 3AE333883018 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 3AE333883018 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::229 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810941; cv=none; b=jEG7V9d92zvbBgOrVQGztrlvYrU44umAop2XQBnAt1dXrFzlNpG3zVqbhFYpV7XoMZmwy87mNU1/YjKaagW5GUbvjVtWsDNQGdDfr4C157+mHv4JOvYgjzXDoM5x76e7e9FlGdG8SLpPBqxMZBj/0goLib1hotRgLt3RZF7KX6I= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1716810941; c=relaxed/simple; bh=WGxIk6y/MddPi2+AHDinB2bJSm97DjjAO4cdOdOti2A=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=BVuUmqkQMd03m8ETmaGSqNYcmW9rhQIOZfBu0p7yMBcMufJ/jF56subR66dxeNvwJvEwHizGUEGQ7KwIva91LdFgZNBrpdY8CNMJ61N/nxVXBzDOR8FtsxV3cwS5UpFXznab/872Wdo2eR3IOpuXR9p4GvnnCPU3WhArC4wFK2c= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lj1-x229.google.com with SMTP id 38308e7fff4ca-2e724bc46bfso101147621fa.3 for ; Mon, 27 May 2024 04:55:39 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1716810937; x=1717415737; 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=Yk7qbiKF2T0eq62bu/4rSk1hyMZFJxSA1+59WYXI7Y8=; b=GqvLkXGMY+nj78TmOYUuLQbHE0wjFhpKx1E+hpp8TK6H6wVpeBm/H1H+VidVI6345E e919PBWTM0YnyI4aDrfdnFB2nH0yefrAstC3U6XPGbyKcir49lq/OLAA36Y0fKJk8+vN RvJUxFchcHxEQK8o6G8xW20zEXJRFc3ZZQBS+Pgh7uzEDYJwjbCJP/RHi9dOxK7U7M3i t30rPdx0KlT1wRy+3yMiOix29R0xecTTvtMTCYGfD6vSHTnnv0NtZkq7KRi8ufuUVXWV FnrXwnZ3SKeqOC+5KLg+pfKPlaRxC0ETt0a5WI8dvQ638nqhwzci/lt9w1bedMrMcCEp 4N6w== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1716810937; x=1717415737; 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=Yk7qbiKF2T0eq62bu/4rSk1hyMZFJxSA1+59WYXI7Y8=; b=WcA8WtWIZODuQUE/BcdILnN6KbwO/Bv3ILQrVvIxuvuFHdcp7XwSRGlIXnMVL4dCSh jEspHxBoJfyQOroq9A4FFNr5htpCvstqe/8292ecsrPkwKmjZ+2zcUewZYMZY1fmChKr 67X68tCkyGJOcTQnDzVOqZG9Ae65/bxdyeIU2xlXrsf1dJimcQWOOWkGNaRb5RaaEB6S LENucnDVezUwBpnHeSMEJg71cazag6Tj9oSSBbGcOkPwobqCAnly3hnC50E1A/jU9VGj LIHdHmmVfptwknMKLkxS5OGKo4di8PP+0EWXhH8Z1tnEeUOplCCo8I6knDPlJ1uMSkzz e0PQ== X-Gm-Message-State: AOJu0YzDne6UyB2XimMbDxu+0sPX/G0q2UNcLd9GHQ+pOMX1aWonKPj7 k0KE7nqBmL4SPQshZktEL0m4Yu2p/TOh3QtxZu8zDLe+Lyd2czTucIS0KHkkg09gh830P3d/K6V B X-Google-Smtp-Source: AGHT+IEcIrrWJgxPGQcGaEVyiF2HZTBmtD5UYnhV54hLeOnBk04UZ21DuaH8eNdRtdoIRQIqbfTM8A== X-Received: by 2002:a2e:b003:0:b0:2e7:134d:f79a with SMTP id 38308e7fff4ca-2e95b03e0d8mr57317461fa.10.1716810937257; Mon, 27 May 2024 04:55:37 -0700 (PDT) Received: from localhost.localdomain ([2a0d:3344:2340:ce10::2f3]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-42108966b58sm107636955e9.3.2024.05.27.04.55.36 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Mon, 27 May 2024 04:55:36 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH 7/7] OpenMP: update documentation for dispatch and adjust_args Date: Mon, 27 May 2024 13:54:39 +0200 Message-ID: <20240527115439.3967217-8-parras@baylibre.com> X-Mailer: git-send-email 2.45.1 In-Reply-To: <20240527115439.3967217-1-parras@baylibre.com> References: <20240527115439.3967217-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.7 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org libgomp/ChangeLog: * libgomp.texi: --- libgomp/libgomp.texi | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 71d62105a20..b72accd0d26 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -294,8 +294,8 @@ The OpenMP 4.5 specification is fully supported. @item C/C++'s @code{declare variant} directive: elision support of preprocessed code @tab N @tab @item @code{declare variant}: new clauses @code{adjust_args} and - @code{append_args} @tab N @tab -@item @code{dispatch} construct @tab N @tab + @code{append_args} @tab P @tab Only @code{adjust_args} +@item @code{dispatch} construct @tab Y @tab @item device-specific ICV settings with environment variables @tab Y @tab @item @code{assume} and @code{assumes} directives @tab Y @tab @item @code{nothing} directive @tab Y @tab