From patchwork Wed Aug 7 11:50:06 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: 1970000 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=tYlmxEPM; 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 4Wf7n51TZzz1yXs for ; Wed, 7 Aug 2024 21:51:36 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 482613858428 for ; Wed, 7 Aug 2024 11:51:34 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x331.google.com (mail-wm1-x331.google.com [IPv6:2a00:1450:4864:20::331]) by sourceware.org (Postfix) with ESMTPS id 905D43858C56 for ; Wed, 7 Aug 2024 11:51:00 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 905D43858C56 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 905D43858C56 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::331 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031463; cv=none; b=aBB9ixGdcSwFhb15We2rCZSsfzWThl37JY4ztkvuKK5r9PXLswZiDbK9SEADl3FWRUSQkn7F8qSFrkBh2SDRM3wg/JuzL/mO9nw1olNDF7EaF1N4m2tgA50SEVP+te881SCr24drPs/TCPUIDZIXzXsVz2haPHm8DD3EktUcwFI= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031463; c=relaxed/simple; bh=zbYHm0zugWabsoYZCphDtIvlFtX1KawHgsiIYwg/fLI=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=SGLcvnPAGArWQ4v/mnAjpvMOghiBQVktYj+aaVbytfQGCxiDujBPapeb6Vs53qy/+CJZHZXtgfhDh7SZNnljEP3BIqZ71iBNVbkMWSbj4FN0bAIdyVVaXnDbISQCDYLNS86blfKQgAV4iFbV0TY6uSrQhC9/ASCedXzbOnJ07mc= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x331.google.com with SMTP id 5b1f17b1804b1-4280ee5f1e3so11613765e9.0 for ; Wed, 07 Aug 2024 04:51:00 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1723031459; x=1723636259; 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=Rg/cnZUdz3kJT8UrDYflkGxHgHZN5qSqsKipZzGgwBk=; b=tYlmxEPMivkg1cz7eOWsaim/oyyLEtSNX3aDu66UPy/h0Owy2wwnzQF+VCsFebjS9H IdbdSxh2Ctn4veLb5jjX14hrIGLsHcm2GS+n1ufx8/4tqiQxBmIS5YWo7CmbPvF1Isui HKXRFrzZNn5blC/OVhAAYZqC1UREjohaX39t/rqFPZDiAV0Unt4yD99NlBXrmkCiPQ8G 3ZImvR2psf8QotBi+jPp6SuOtwqBzJU4I1fiLuU7uW+T5vteccYJ66S6IkZo9NakP47K tuAXnhSeY1fMJkEllvfbg9ZZMf6jEZMq7HtjaYMhZ6gweOIlBuzn2trWaNadFuqUtr0/ brww== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1723031459; x=1723636259; 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=Rg/cnZUdz3kJT8UrDYflkGxHgHZN5qSqsKipZzGgwBk=; b=AMooIqSt4OdNC/eckeC3OE7GPLrMXDs5sl2Yh9n57nYclP6Lx+XTTltMMp/QDjNeZf PqpIDu0PTOgkueqme7aodzdmTBD2ssuk4uqP8A8ZwtgcVdGvhRYCYd/wLWYRRQlnl2gm nvcyQcGDwSY71LKGaV8kZF8uGctOM2rSt2l6IG96wq63XOSYD2X3TxxSSrAwfQxlDZi2 L0xWmFUgWeYUoaIln/kbSxRRtZULoMdcs43Ryd254XpZ50+zjNjFeXe4qfU8yWzLQzMx 85lsXbhv+UXm+N/eAnOGa1qBos0/8kc4ZTnOHzXK6SDFrWetjR0eEjg8hY3neP3xxM56 jEkw== X-Gm-Message-State: AOJu0YwT8RqAf+tyXw33AIUS+hHgoEQhM2+BlklP0wbRz/uGSKxjxo9z 4eRKgGmpqb5C/vSkQ6iObhZlPS15vANIclpEvkq111dxG4oL7kJSarbRctZLldi2PKBItN6CiSV J0ew= X-Google-Smtp-Source: AGHT+IFcU0MONwT3RhFcH1TBavYtq8CdlVeDEHpx2TvVoEN4+b9SZ/AIufKD6mnlHudoURfWp2DRPg== X-Received: by 2002:a05:600c:3ba9:b0:426:554a:e0bf with SMTP id 5b1f17b1804b1-428e6b037fcmr117793445e9.16.1723031458661; Wed, 07 Aug 2024 04:50:58 -0700 (PDT) Received: from localhost.localdomain (vlan-185-velingrad-59.comnet.bg. [84.54.185.59]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4290580c970sm25430125e9.46.2024.08.07.04.50.57 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 07 Aug 2024 04:50:58 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v3 1/7] OpenMP: dispatch + adjust_args tree data structures and front-end interfaces Date: Wed, 7 Aug 2024 14:50:06 +0300 Message-ID: <20240807115012.3632947-2-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240807115012.3632947-1-parras@baylibre.com> References: <20240807115012.3632947-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-11.1 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_BARRACUDACENTRAL, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch 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. gcc/ChangeLog: * builtin-types.def (BT_FN_PTR_CONST_PTR_INT): New. * omp-selectors.h (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/fortran/ChangeLog: * types.def (BT_FN_PTR_CONST_PTR_INT): Declare. --- gcc/builtin-types.def | 1 + gcc/fortran/types.def | 1 + gcc/omp-selectors.h | 1 + gcc/tree-core.h | 7 +++++++ gcc/tree-pretty-print.cc | 21 +++++++++++++++++++++ gcc/tree.cc | 4 ++++ gcc/tree.def | 5 +++++ gcc/tree.h | 7 +++++++ 8 files changed, 47 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/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/omp-selectors.h b/gcc/omp-selectors.h index c61808ec0ad..ef3ce9a449a 100644 --- a/gcc/omp-selectors.h +++ b/gcc/omp-selectors.h @@ -55,6 +55,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 27c569c7702..508f5c580d4 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -542,6 +542,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 4bb946bb0e8..752a402e0d0 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; @@ -3947,6 +3963,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 a2d431662bd..4d60143c3ca 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -331,6 +331,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[] = @@ -427,6 +429,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 85ab182c6f5..1a6c9512cfc 100644 --- a/gcc/tree.def +++ b/gcc/tree.def @@ -1298,6 +1298,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 5dcbb2fb5dd..612c2a1b868 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1603,6 +1603,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, \ @@ -1750,6 +1753,10 @@ class auto_suppress_location_wrappers OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_PARTIAL), 0) #define OMP_CLAUSE_SIZES_LIST(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_SIZES), 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 Wed Aug 7 11:50:07 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: 1970002 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=Ezxp7FdK; 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 4Wf7nC4CJpz1yXs for ; Wed, 7 Aug 2024 21:51:43 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 6DB8A385E003 for ; Wed, 7 Aug 2024 11:51:41 +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 6636B3858C33 for ; Wed, 7 Aug 2024 11:51:02 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 6636B3858C33 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 6636B3858C33 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=1723031468; cv=none; b=sNwZQ/fHKkaoSfHyYd9YaKBy8lQjTZZt/g+omCzg0N+YRrVFF80InHUb8npMYN9HLs2peca8xQWvbGxDa5yU2QCg/zJvji7xfediNLFVnA/UDrAZ1E/cuQsy6OimNPOE+evHxddN49XfUIvZEhIXJ6pCBzXgWJlfBw7Q7CS/ZkE= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031468; c=relaxed/simple; bh=+GuNes8HZOSbHy3PkU29OiwU592v4KoKULkWeG+ucew=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=FSPJO+UUcrDZ6reTubf8fGiqEm2V9KzN7eKrOgEEvtcvoz4StzqschEEpxr9XtxPw+GS9OB5TmXCgZQXw99tDhJFLrS+dH5iAA9ORYSABNQJPhFrTuCoPec6OhgAMcNvl7Gw3pY2L7h1s6BAOeiHARkDDGZPc+fh2+ZNyJoDfis= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lj1-x22f.google.com with SMTP id 38308e7fff4ca-2f0271b0ae9so17063061fa.1 for ; Wed, 07 Aug 2024 04:51:02 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1723031460; x=1723636260; 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=YUUC0qLg8Vs919P4fFam+JUY3bZ8wHAFv5pSd20tBmc=; b=Ezxp7FdKi4niVMvxnHWm3bXamjXebjeukMe/IKLO8ilR/5dbRDqKGem+JcQwwNWWa7 /6tcBslOzEO+NC0c389ZBN2vvOgHxkpD3XgEXfmPFTVEg45zp42Jswsd68ZKrLflfaNv yxn0LMP8LzP8ziwJP0SnspXQ/UXaHflYRkpA8J/j5jcUrv9cKcIOQah2BhY5SqReDRlB gCBMPGteOqlmQM6kpHSBVZ9LiI8bEhPdcdluXpfU5hplfUgGl3GRTEtSrAGxsCkL/Scv 7gVokroN0ichBmomUs+FvKW2dZzWozgRIcJvgOOHIO2YHulnJWUWp0O3NO8y3h8ZddyS T5Jw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1723031460; x=1723636260; 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=YUUC0qLg8Vs919P4fFam+JUY3bZ8wHAFv5pSd20tBmc=; b=NBhHHLvGBZgSdJSB0Qo7b668XA5cNHwPb0MkfxKYUCilnjZUafp/N59i5tUxlyYx/K VePh7N6oGuAgYatweoZJU3Yxw+QXCXGWLJVYHC/QGfGMRzPPNi4aqXPUB0Hfgm2Ro0QM 66Yt14+3MLLMi2cJacHhJoV3K4W2nyo6Ug7tgWEdOO7KjvZm6xMLZc5eLeX2SSmBPRzP DiS+WEk9zIqpSzq0TGmUgZqhjUlnvXvEE2HCmlfoH2JRrQP9MPySPDCaI4YYFhZtXLPC sUPjyxhPR6yT7lFfyaDduxVDPvIWTEJvqtutqt5/44KyTfK5NLR1tT/dwiNYDVPLNNoO 8wlw== X-Gm-Message-State: AOJu0Yx2yTB9/TB5om8jVGxJrxsjmQn+7NPEqY6IwdkuvvkS7hVxMi8p fv/tn1EMlDO1vGE4v/Dxd6lcFLk3RvS5kOiAt9O4ttSPayFJbrBmqtYT0VnSyVdgTgxjJfyloHa fl+4= X-Google-Smtp-Source: AGHT+IGKUtG+Igdmd6fVK/JRHTy1zEMcYrBnMcRvXtq5sXRj67XVJz/nZlAC/X+5E66e076e1F6Z/g== X-Received: by 2002:a2e:87ce:0:b0:2ef:2855:533f with SMTP id 38308e7fff4ca-2f15aa8744dmr115559131fa.16.1723031459796; Wed, 07 Aug 2024 04:50:59 -0700 (PDT) Received: from localhost.localdomain (vlan-185-velingrad-59.comnet.bg. [84.54.185.59]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4290580c970sm25430125e9.46.2024.08.07.04.50.58 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 07 Aug 2024 04:50:59 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v3 2/7] OpenMP: middle-end support for dispatch + adjust_args Date: Wed, 7 Aug 2024 14:50:07 +0300 Message-ID: <20240807115012.3632947-3-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240807115012.3632947-1-parras@baylibre.com> References: <20240807115012.3632947-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_BARRACUDACENTRAL, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds 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: set the default-device ICV at the top of the dispatch region and restore its previous value at the end. * 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. * If depend clauses are present, add a taskwait construct before the dispatch region and move them there. 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 | 421 ++++++++++++++++++++++++++++++++++++- gcc/gimplify.h | 2 + gcc/omp-builtins.def | 6 + gcc/omp-expand.cc | 18 ++ gcc/omp-general.cc | 14 +- gcc/omp-low.cc | 35 +++ gcc/tree-inline.cc | 7 + 13 files changed, 585 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 08b823c84ef..e7b2df9a0ef 100644 --- a/gcc/gimple-pretty-print.cc +++ b/gcc/gimple-pretty-print.cc @@ -1726,6 +1726,35 @@ dump_gimple_omp_scope (pretty_printer *pp, 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 PP. */ static void @@ -2805,6 +2834,10 @@ pp_gimple_stmt_1 (pretty_printer *pp, const gimple *gs, int spc, dump_gimple_omp_scope (pp, gs, spc, flags); break; + case GIMPLE_OMP_DISPATCH: + dump_gimple_omp_dispatch(pp, 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 30bfecf67e5..cf41fa20e57 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, @@ -4052,6 +4053,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++) @@ -4062,8 +4064,111 @@ 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); + tree adjust_args_list; + if (flag_openmp && EXPR_P (CALL_EXPR_FN (*expr_p)) + && DECL_P (TREE_OPERAND (CALL_EXPR_FN (*expr_p), 0)) + && (adjust_args_list = lookup_attribute ( + "omp declare variant variant adjust_args", + 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); + } + } + + bool need_device_ptr = false; + for (tree arg + = TREE_PURPOSE (TREE_VALUE (adjust_args_list)); + arg != NULL; arg = TREE_CHAIN (arg)) + { + if (build_int_cst (integer_type_node, i) + == TREE_VALUE (arg)) + { + need_device_ptr = true; + break; + } + } + + if (need_device_ptr && !is_device_ptr) + { + 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; @@ -6310,6 +6415,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: @@ -13130,6 +13236,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: @@ -13359,6 +13480,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } break; + case OMP_CLAUSE_NOVARIANTS: + OMP_CLAUSE_NOVARIANTS_EXPR (c) + = gimple_boolify (OMP_CLAUSE_NOVARIANTS_EXPR (c)); + break; + case OMP_CLAUSE_NOCONTEXT: + OMP_CLAUSE_NOCONTEXT_EXPR (c) + = gimple_boolify (OMP_CLAUSE_NOCONTEXT_EXPR (c)); + break; case OMP_CLAUSE_NOHOST: default: gcc_unreachable (); @@ -13813,7 +13942,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) { @@ -14624,6 +14755,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: @@ -14713,9 +14846,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) @@ -14833,6 +14966,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 @@ -17826,6 +18013,219 @@ 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); + tree saved_device_icv; + if (device) + { + // Save current default-device-var ICV + saved_device_icv = create_tmp_var (integer_type_node); + tree fn = builtin_decl_explicit (BUILT_IN_OMP_GET_DEFAULT_DEVICE); + gcall *call = gimple_build_call (fn, 0); + gimple_call_set_lhs (call, saved_device_icv); + gimplify_seq_add_stmt (&body, call); + + // Set default device + fn = builtin_decl_explicit (BUILT_IN_OMP_SET_DEFAULT_DEVICE); + call = gimple_build_call (fn, 1, OMP_CLAUSE_DEVICE_ID (device)); + gimplify_seq_add_stmt (&body, call); + } + + // 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 = base_call_expr; + 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); + + // Restore default-device-var ICV + if (device) + { + tree fn = builtin_decl_explicit (BUILT_IN_OMP_SET_DEFAULT_DEVICE); + gcall *call = gimple_build_call (fn, 1, saved_device_icv); + gimplify_seq_add_stmt (&body, call); + } + + // 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); + + // Remove nowait as it has no effect on dispatch (OpenMP 5.2), and depend and + // device clauses as they have already been handled + tree *dispatch_clauses_ptr = &OMP_DISPATCH_CLAUSES (expr); + for (tree c = *dispatch_clauses_ptr; c; c = *dispatch_clauses_ptr) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NOWAIT + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE) + { + *dispatch_clauses_ptr = OMP_CLAUSE_CHAIN (c); + break; + } + else + dispatch_clauses_ptr = &OMP_CLAUSE_CHAIN (c); + } + + gimple *stmt = gimple_build_omp_dispatch (bind, 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 @@ -18764,6 +19164,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; @@ -19089,7 +19493,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 0b61335dba4..2604ec4cd2b 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -1048,7 +1048,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--) { @@ -1141,6 +1141,7 @@ const char *omp_tss_map[] = "target_device", "implementation", "user", + "need_device_ptr", NULL }; @@ -1247,10 +1248,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. */ @@ -2495,6 +2500,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; @@ -2641,6 +2649,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 Wed Aug 7 11:50:08 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: 1970006 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=jFRNZ7XV; 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 4Wf7ph1Lm5z1yXs for ; Wed, 7 Aug 2024 21:53:00 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 0B9913857022 for ; Wed, 7 Aug 2024 11:52:58 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x331.google.com (mail-wm1-x331.google.com [IPv6:2a00:1450:4864:20::331]) by sourceware.org (Postfix) with ESMTPS id 068BB3858C3A for ; Wed, 7 Aug 2024 11:51:03 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 068BB3858C3A 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 068BB3858C3A Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::331 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031468; cv=none; b=tAx38xxpfjODA4cg/xL5wW4Q6LNiAgegF10jotP1S6Mm9ftEtf0Yx+bIhBQ5RNmpSPfWndhFXoc7UM1+aPjnaIpt35CQeRpnypRgfmMSVSQQYRFbhpn9r/2JWMUcIIt29HyCNa3PFNOORWZVnMuQVlGyfiu7n4B6xaeka9nDa3c= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031468; c=relaxed/simple; bh=pav5gR6wQ+fV/IzDXr1VGoztT4k6IAQufc3s0AAGt3I=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=HJ1geysWm8B8rjJieM1KabX49NbeotFJKAlL4P7B3pBXTq/1IrklMWylJT0O1Zvx48zhiMY7ezRDOyPil5y6rNXApX9bFMQbiFFuD29xq9lBnJ+72KLsbtw8DbKufIL0RJ39v5JDAKoChVGhSFfoMVCPUg9S6udF6UHeNErfCBY= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x331.google.com with SMTP id 5b1f17b1804b1-42816ca797fso11593575e9.2 for ; Wed, 07 Aug 2024 04:51:02 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1723031461; x=1723636261; 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=k7KxE/1RgVwakz9UCo4BK9ppuocLWAAgwVXLpnCHYD0=; b=jFRNZ7XVMApgnILpLj2TZxC3vZYUSCsQrsiTVj2zIqd3geqsQKTlQMGkOEfnPQ4W4T eSkjdsjKCF9A4CqquTwBEyI6mfYcC11x7Hpa9ApVDbrw4t/SPok1Zv5e/rNbGJuoXMbM pH3L/Tw6nn/B0whgB/M7M16Da6dzTkcAREbTlvBMvLyGjZ+byQVWNjKmhmYtK0qgwDuR hqPJWFiW2pF4arO0ZZRCuEpG7Ui6EF9AW49ZXd4dyEb0F5WtLVNlbd/L8XVEpx0lPTGt 72iiZC0lqDq31d8UWFMH1+QgqP3+hP7ACLxZHiRj1a9wurmfB3WwGA4+zPLIcbLvX71w ssUA== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1723031461; x=1723636261; 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=k7KxE/1RgVwakz9UCo4BK9ppuocLWAAgwVXLpnCHYD0=; b=LpI+dYL4MYY4ZhMsZEbpFdJtvTLbCvgz0TmesC7hIlAHzDKYSQ8kxNxeagVPK9NVfc 4dHoufSkGUhqlYb/DL9hO1M0iOtW+069Z6NvLA500MCTJ4ZQL62tL7z8jjslgCJMkjJC LjyDNna9Rf2sTZyPoxsO+7LZSNVdQ7oJZQG/EOVD5MGNl8vLim8yce8e5/dF5M7mT+eV Dm5B8HQ8yUHwz7YTT6TdwTjlvgX/KCtfTQh39QbYttD0Ie4g+/Op6j9gXn9dAopCNt/h 3JtxvVS0UA4Z6oUeMrLNgU9sH/AD9hQIFX0MWgWNy3m1gQfcM2YWqfb4cWheIW6PwgFp MEMQ== X-Gm-Message-State: AOJu0YyMoWJU487ZuyASluYtz7TdsjR8n/AIfki74WYyu5bTkfYppjjY Q5t5pPFCLOsDJbWsxj83nBK37kXRyTaEcADi8TpR6hS3ggH6WdYaw3DZJLzrgCJUZjNLt6QNIEb uams= X-Google-Smtp-Source: AGHT+IEkZ4PWMy/udo5MM+2ckKB3a/GphG1VVaABE6EezS4z8Ak46VlLNnou1RNbNfXgzYhiLPLQfg== X-Received: by 2002:a05:600c:580a:b0:428:f0c2:ef4a with SMTP id 5b1f17b1804b1-428f0c2f460mr95527425e9.13.1723031460892; Wed, 07 Aug 2024 04:51:00 -0700 (PDT) Received: from localhost.localdomain (vlan-185-velingrad-59.comnet.bg. [84.54.185.59]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4290580c970sm25430125e9.46.2024.08.07.04.50.59 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 07 Aug 2024 04:51:00 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v3 3/7] OpenMP: C front-end support for dispatch + adjust_args Date: Wed, 7 Aug 2024 14:50:08 +0300 Message-ID: <20240807115012.3632947-4-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240807115012.3632947-1-parras@baylibre.com> References: <20240807115012.3632947-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_BARRACUDACENTRAL, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds support 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 | 522 +++++++++++++++++++--- 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 ++++ libgomp/testsuite/libgomp.c/dispatch-2.c | 84 ++++ 10 files changed, 719 insertions(+), 60 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 create mode 100644 libgomp/testsuite/libgomp.c/dispatch-2.c diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc index 685f212683f..91a5356796d 100644 --- a/gcc/c-family/c-attribs.cc +++ b/gcc/c-family/c-attribs.cc @@ -562,6 +562,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 b5ce1466e5d..c74a9fb2691 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -4299,8 +4299,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 25251c2b69f..b956819c0a5 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 2ebde06c471..6b6826b2426 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, @@ -135,9 +136,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 9b9284b1ba4..ef9e5a04c36 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -1747,6 +1747,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 @@ -15089,6 +15091,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)) @@ -15097,6 +15101,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)) @@ -19364,6 +19370,60 @@ c_parser_omp_clause_partial (c_parser *parser, tree list) return c; } +/* 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 ) */ @@ -19983,6 +20043,14 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "partial"; clauses = c_parser_omp_clause_partial (parser, clauses); 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; @@ -23793,6 +23861,189 @@ 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"); + + // Extract depend clauses and create taskwait + tree depend_clauses = NULL_TREE; + tree *depend_clauses_ptr = &depend_clauses; + for (tree c = OMP_DISPATCH_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND) + { + *depend_clauses_ptr = c; + depend_clauses_ptr = &OMP_CLAUSE_CHAIN (c); + } + } + if (depend_clauses != NULL_TREE) + { + tree stmt = make_node (OMP_TASK); + TREE_TYPE (stmt) = void_node; + OMP_TASK_CLAUSES (stmt) = depend_clauses; + OMP_TASK_BODY (stmt) = NULL_TREE; + SET_EXPR_LOCATION (stmt, loc); + add_stmt (stmt); + } + + // 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 @@ -24773,6 +25024,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 \ @@ -25236,77 +25491,222 @@ 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 = UNKNOWN_LOCATION; + tree need_device_ptr_list = make_node (TREE_LIST); - 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); - 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); - if (!parens.require_open (parser)) - goto fail; + enum clause + { + match, + adjust_args + } ccode; - 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 (TREE_CODE (variant) != FUNCTION_DECL) + if (strcmp (clause, "match") == 0) + ccode = match; + else if (strcmp (clause, "adjust_args") == 0) { - error_at (token->location, "variant %qD is not a function", variant); - variant = error_mark_node; + ccode = adjust_args; + adjust_args_loc = match_loc; } - else if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT, - OMP_TRAIT_CONSTRUCT_SIMD) - && !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant))) + else { - error_at (token->location, "variant %qD and base %qD have " - "incompatible types", variant, fndecl); - variant = error_mark_node; + c_parser_error (parser, "expected % or %"); + goto fail; } - 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)) + + c_parser_consume_token (parser); + + if (!parens.require_open (parser)) + goto fail; + + if (parms == NULL_TREE) + parms = error_mark_node; + + if (ccode == match) { - 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)) + 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) { - tree attr - = tree_cons (get_identifier ("omp declare variant base"), - build_tree_list (variant, ctx), - DECL_ATTRIBUTES (fndecl)); - DECL_ATTRIBUTES (fndecl) = attr; + if (TREE_CODE (variant) != FUNCTION_DECL) + { + 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) + { + need_device_ptr_list = chainon ( + need_device_ptr_list, + build_tree_list ( + NULL_TREE, + build_int_cst ( + integer_type_node, + idx))); // Store 0-based argument index, + // as in gimplify_call_expr + } + } + } + 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"); + } } - parens.require_close (parser); - c_parser_skip_to_pragma_eol (parser); + if (need_device_ptr_list && variant != error_mark_node) + { + tree variant_decl = tree_strip_nop_conversions (variant); + DECL_ATTRIBUTES (variant_decl) + = tree_cons (get_identifier ("omp declare variant variant adjust_args"), + build_tree_list (need_device_ptr_list, + NULL_TREE /*need_device_addr */), + DECL_ATTRIBUTES (variant_decl)); + } } /* Finalize #pragma omp declare simd or #pragma omp declare variant @@ -26123,7 +26523,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) \ @@ -26131,7 +26530,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) @@ -27046,6 +27449,9 @@ c_parser_omp_construct (c_parser *parser, bool *if_p) case PRAGMA_OMP_UNROLL: stmt = c_parser_omp_unroll (loc, parser, if_p); break; + 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 094e41fa202..beaa37f8729 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -16330,6 +16330,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; +} diff --git a/libgomp/testsuite/libgomp.c/dispatch-2.c b/libgomp/testsuite/libgomp.c/dispatch-2.c new file mode 100644 index 00000000000..faa0d8a1d1c --- /dev/null +++ b/libgomp/testsuite/libgomp.c/dispatch-2.c @@ -0,0 +1,84 @@ +// Adapted from OpenMP examples + +#include +#include +#include + +int baz (double *d_bv, const double *d_av, int n); +int bar (double *d_bv, const double *d_av, int n); + +#pragma omp declare variant(bar) match(construct={dispatch}) adjust_args(need_device_ptr: f_bv, f_av) +#pragma omp declare variant(baz) match(implementation={vendor(gnu)}) +int foo (double *f_bv, const double *f_av, int n); + +int baz (double *bv, const double *av, int n); +int bar (double *bv, const double *av, int n); + +int foo (double *bv, const double *av, int n) +{ + for (int i = 0; i < n; i++) + bv[i] = av[i] * i; + return -1; +} + +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; +} + +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 Wed Aug 7 11:50:09 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: 1970004 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=KEK9KvkN; 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 4Wf7pJ45gTz1yXs for ; Wed, 7 Aug 2024 21:52:40 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 1067C385828B for ; Wed, 7 Aug 2024 11:52:38 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x32c.google.com (mail-wm1-x32c.google.com [IPv6:2a00:1450:4864:20::32c]) by sourceware.org (Postfix) with ESMTPS id 6D0063858402 for ; Wed, 7 Aug 2024 11:51:04 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 6D0063858402 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 6D0063858402 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::32c ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031469; cv=none; b=FG/ufgJUpToSsQFUc02C7AQ8HoqslNEeviGcnxOcLBtj6yvVQWwLla7FN/8pZXO7PgDTv3ebGQ6eUXeq0EjTP70rSwRa7BXyfqlGTQv0JF4ogqRxaL+W6EwE6q9jPz2jkvdUcz/nw5gnsqLOJKYCX8FfwfsCX+7R5HNdwgfAzzw= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031469; c=relaxed/simple; bh=WdYMQYMVtgQ5u48VmHE38Jo5UX5r8v+tZO0fLNUSpK0=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=c+QXFhlywdfemXOhytqwLlYiuWRw4uZ/g4v5aAXvsAkWXgvf4rTRiWkutoHBwShjr8EOSmie++ZrjBhyBD3zbYcXe7xWL8QfBixe83VrJwNk6bv8dq2Ud8M1oTYwFpqdwO4eKt1zcNaL3lA8w+cnRMwrp5cF4Os0Evc7Wjl4d0k= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x32c.google.com with SMTP id 5b1f17b1804b1-428fb103724so4637505e9.1 for ; Wed, 07 Aug 2024 04:51:04 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1723031463; x=1723636263; 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=p4kXvftxanoCbxa/yoLWrgKflzOYTMaIlJNcCcmuv2U=; b=KEK9KvkNvZWne6HhW0cCdgToHO7vS1SgIt+8mAGDqKs+hMupez0/kC480QwnIGruNq FJapAdLhEmVzg3zfB5PhBDoUc85dlolVO+MK7JUG6ClAmd2pVA4QX+xDDm21dJdyazPg 1jS/taUhqgtz/AUfshkUUTCydZdMZ7neeEWY2gdBIaBJxGvJoGQWqib+nWpPQrxiX7J+ mDBVVmgTwrVHNjGLMv8qBeYoFUq3253hB6XhNJHNKqtfN2+5z3GlWebiZ2my6fLhgCoP gNDIuYdY7r9Hwv0cHJheOQLB2HDh3sDEqN9M9cLnVEzdj4xtH2r4HA/wZlgM/2GxryBn oTFw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1723031463; x=1723636263; 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=p4kXvftxanoCbxa/yoLWrgKflzOYTMaIlJNcCcmuv2U=; b=t47yj7eavarzA+K+WJBAayzG6yG4GC0rXIWx9BFRXn5mNhWP+6FJWYNOcA00OuMfUG fK5LwaDmJ7SG5gDc4mF61W9mT/c2qSYQRwW0TI9NyrTpkwRvZ93mDriCcB72GbHO0Xe9 +Yyllk7+R/8V33vgEV+vUBgBn7YE0Y/mYP6MmrJO7tl5DBog9kVZw32iYynYzYOHh3iR bvU6HkXL0SpV9/GYfvo4Zz08c5zNomyNn1yMz25dOS2prcRFWGuMID9GULJAbj4ASzKh fQD3ke3ga1QCyDeCEQPMh/tT+02XDizGjT7ihf11vNjK3y4yoFHZ8v56uvL0RmgYcVsk kNnw== X-Gm-Message-State: AOJu0Yx+JsuJLcqXgoq9VaOYI6kP8gMCfluQgzKvslmCBjvO4+OUkw0f JNpaobeDD1nbWqTV4Muhxyas6iLHRpgZPY4Ulx6e/pp/ceAxkcBY1SLEVN7aCec1bmeT2xX/By5 UIP0= X-Google-Smtp-Source: AGHT+IH5kCEbpJzIxzOTNhB2YjUR1xkCmcSpQ2OjqoQouIudhIvrmD4rprKKEaPS55wmua6sThObVA== X-Received: by 2002:a05:600c:3b83:b0:426:62a2:34fc with SMTP id 5b1f17b1804b1-429052cbc65mr13828655e9.11.1723031462315; Wed, 07 Aug 2024 04:51:02 -0700 (PDT) Received: from localhost.localdomain (vlan-185-velingrad-59.comnet.bg. [84.54.185.59]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4290580c970sm25430125e9.46.2024.08.07.04.51.01 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 07 Aug 2024 04:51:01 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v3 4/7] OpenMP: C++ front-end support for dispatch + adjust_args Date: Wed, 7 Aug 2024 14:50:09 +0300 Message-ID: <20240807115012.3632947-5-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240807115012.3632947-1-parras@baylibre.com> References: <20240807115012.3632947-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, KAM_SHORT, RCVD_IN_BARRACUDACENTRAL, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds 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. * 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 | 7 + gcc/cp/parser.cc | 631 ++++++++++++++++++++-- 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 +++ 7 files changed, 818 insertions(+), 45 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 279af21eed0..00056d37ecf 100644 --- a/gcc/cp/decl.cc +++ b/gcc/cp/decl.cc @@ -8383,6 +8383,13 @@ omp_declare_variant_finalize_one (tree decl, tree attr) if (!omp_context_selector_matches (ctx)) return true; TREE_PURPOSE (TREE_VALUE (attr)) = variant; + + // Prepend adjust_args list to variant attributes + tree adjust_args_list = TREE_CHAIN (TREE_CHAIN (chain)); + if (adjust_args_list != NULL_TREE) + DECL_ATTRIBUTES (variant) = tree_cons ( + get_identifier ("omp declare variant variant adjust_args"), + TREE_VALUE (adjust_args_list), DECL_ATTRIBUTES (variant)); } } else if (!processing_template_decl) diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index eb102dea829..b43eabb0cff 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); @@ -24193,7 +24194,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); @@ -25058,8 +25059,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; @@ -25095,8 +25096,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, @@ -38180,6 +38181,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)) @@ -38188,6 +38191,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)) @@ -40634,6 +40639,56 @@ cp_parser_omp_clause_partial (cp_parser *parser, tree list, location_t loc) 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 ) */ @@ -42741,6 +42796,16 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_omp_clause_full (clauses, token->location); c_name = "full"; 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; @@ -48946,12 +49011,326 @@ 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); + + // Extract depend clauses and create taskwait + tree depend_clauses = NULL_TREE; + tree *depend_clauses_ptr = &depend_clauses; + for (tree c = OMP_DISPATCH_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND) + { + *depend_clauses_ptr = c; + depend_clauses_ptr = &OMP_CLAUSE_CHAIN (c); + } + } + if (depend_clauses != NULL_TREE) + { + tree stmt = make_node (OMP_TASK); + TREE_TYPE (stmt) = void_node; + OMP_TASK_CLAUSES (stmt) = depend_clauses; + OMP_TASK_BODY (stmt) = NULL_TREE; + SET_EXPR_LOCATION (stmt, loc); + add_stmt (stmt); + } + + // 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)) @@ -49009,44 +49388,195 @@ 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 = UNKNOWN_LOCATION; + tree need_device_ptr_list = make_node (TREE_LIST); - 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) + { + need_device_ptr_list = chainon ( + need_device_ptr_list, + build_tree_list ( + NULL_TREE, + build_int_cst ( + integer_type_node, + idx))); // Store 0-based argument index, + // as in gimplify_call_expr + } + } + } + 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 (TREE_CHAIN (need_device_ptr_list) != NULL_TREE) { - 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 a DECL for the variant yet. So we store the + // need_device_ptr list in the base function attribute, after loc nodes. + gcc_assert (TREE_PURPOSE (attrs) + == get_identifier ("omp declare variant base")); + gcc_assert (TREE_PURPOSE (TREE_VALUE (attrs)) == variant); + TREE_VALUE (attrs) = chainon ( + TREE_VALUE (attrs), + build_tree_list (NULL_TREE, + build_tree_list (need_device_ptr_list, + NULL_TREE /*need_device_addr */))); } - parens.require_close (parser); cp_parser_skip_to_pragma_eol (parser, pragma_tok); return attrs; } @@ -49056,7 +49586,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; @@ -49100,7 +49631,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); } @@ -49231,9 +49762,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); @@ -50083,7 +50613,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, @@ -50944,6 +51478,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok, bool *if_p) case PRAGMA_OMP_UNROLL: stmt = cp_parser_omp_unroll (parser, pragma_tok, if_p); break; + case PRAGMA_OMP_DISPATCH: + stmt = cp_parser_omp_dispatch (parser, pragma_tok); + break; default: gcc_unreachable (); } @@ -51640,6 +52177,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/semantics.cc b/gcc/cp/semantics.cc index 669da4ad969..61ec1cd7a3c 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -7634,6 +7634,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 Wed Aug 7 11:50:10 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: 1970003 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=kW7+/3In; 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 4Wf7nT2cwRz1yXs for ; Wed, 7 Aug 2024 21:51:57 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 307B0385DDCC for ; Wed, 7 Aug 2024 11:51:55 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x32c.google.com (mail-wm1-x32c.google.com [IPv6:2a00:1450:4864:20::32c]) by sourceware.org (Postfix) with ESMTPS id BB3ED385841C for ; Wed, 7 Aug 2024 11:51:04 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org BB3ED385841C 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 BB3ED385841C Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::32c ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031468; cv=none; b=i6PcCNp7O0IH61RY45wMsPgqP4yZzhO5VbukznK2Edf0V1Tx87Ap1By1Z/qotfO3wwPUrMJhxZ+Y2sHZQsSeoCyyQsZiKKrysc7aI5L+qu1gCIBvWEO1Qn/KBxfFLQ15GQsWXqSbbYnSkQmpkDB688nwMx3qph5qo4w0MXG2p0U= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031468; c=relaxed/simple; bh=gB3QDVIIX+HQS+G9aGkwrsvCods8jhjcVKceesmD4N0=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=YhS7dmAzqlHItRTgS8WxloBJsILSPktD5yIosPOp9A7muX2SGHync69RNS349aH3Q4+6eHOhzn0F8JCTthRIuJJ8R2uFsCxoSFoCgo6AZMaAwnmcySROaCi9B2GKoDqbIS4qXZgknXPBYEjzs90AGdMoeTCB81eenhLL8mmyu9g= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x32c.google.com with SMTP id 5b1f17b1804b1-428f5c0833bso4763255e9.0 for ; Wed, 07 Aug 2024 04:51:04 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1723031463; x=1723636263; 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=hp4tnxquIii8L/nVq5n3yl2vfaS0z0s+K/Pm2L+f32s=; b=kW7+/3IntSQMNNYqBsX7KzU5mypjMU9af8ArNmYw7VCfWiE83dKWIWq6i6DWzCxFw1 XM+/GsREHnV70rvuIStHPL+PTcPDj5/wnBNiqD7zO0E/Kfh+a9DAX0dpNXbqA5RXp9FV TXcnxqscY398QKp7zIrxu8r41gV2k/tPci8q1wNeXTt1h4q5hpkAkeJS+f35UUeNYVUO sDcw9YyVNa6tCx54mt9FCT79w51UF6cwNnDsH8r5mx42Lzi4Yqvv+ZnvmdkCIpala0eX R6FG4xQxJQsMcQj5Nw750zMTbg/6BWJJjh4ShLb0vLPfghr1UHa4sCqeaR+6g3p2ivUw fknQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1723031463; x=1723636263; 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=hp4tnxquIii8L/nVq5n3yl2vfaS0z0s+K/Pm2L+f32s=; b=hrl2Tp/YMK0dcr2w6WpI769FNJhaAtxaHYSUT2vRKlZh/GqTrVNEr+xrS7MXJGu45T JRVaLo8fBWPfq7VWpZbFCVI5ukzm8O1tWa0hV3v+j39N39OfrQMwQdSlEWVsIFTjGTdL Gk+Lj4AM+et0u1MEy4xJb63sqAj6uVTJ90IWg6ek8AjPO7IK6iat3LwnLsORkiSa4C6w CTUOFC3KIaZUzmVqAIc7tPYtxKWc5knWZ7eu9z3qQR+7LjrhgDGS/3+2Gv8GFjyvmQ0B hGcFl0tB0ccB2KAXOQ8AJagaODldtajy4YyF8/Mx9jhar5FKT1K+Uy8z0gLAj3tM0WAs GtRQ== X-Gm-Message-State: AOJu0Yyp8F41skcW2VuoW2b126fuVkOa3TdsGGTe6fEcaXO7AzANCcrp ttHomexBQwYPlXXRedMuq3h2BbVu0C8n9y1MCwrv7VADfPtYgrS//ahg/vFdPLxicMXxTa9Hl1t DdEs= X-Google-Smtp-Source: AGHT+IGUCLMhFQs/Jzr5hh78COUq/GZIXTIs9VW+oNkXTeDgxFnJ5g7+hjgdpsbCFETWBiL3D0zflg== X-Received: by 2002:a05:600c:5129:b0:426:5cc7:82f with SMTP id 5b1f17b1804b1-429052cfb3dmr14259225e9.13.1723031463100; Wed, 07 Aug 2024 04:51:03 -0700 (PDT) Received: from localhost.localdomain (vlan-185-velingrad-59.comnet.bg. [84.54.185.59]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4290580c970sm25430125e9.46.2024.08.07.04.51.02 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 07 Aug 2024 04:51:02 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v3 5/7] OpenMP: common C/C++ testcases for dispatch + adjust_args Date: Wed, 7 Aug 2024 14:50:10 +0300 Message-ID: <20240807115012.3632947-6-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240807115012.3632947-1-parras@baylibre.com> References: <20240807115012.3632947-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_BARRACUDACENTRAL, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org 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 | 12 ++++ gcc/testsuite/c-c++-common/gomp/dispatch-4.c | 18 +++++ gcc/testsuite/c-c++-common/gomp/dispatch-5.c | 27 ++++++++ gcc/testsuite/c-c++-common/gomp/dispatch-6.c | 18 +++++ gcc/testsuite/c-c++-common/gomp/dispatch-7.c | 21 ++++++ .../dispatch-1.c | 0 .../dispatch-2.c | 0 12 files changed, 252 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 rename libgomp/testsuite/{libgomp.c => libgomp.c-c++-common}/dispatch-1.c (100%) rename libgomp/testsuite/{libgomp.c => libgomp.c-c++-common}/dispatch-2.c (100%) 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..d2a4a5f4ec4 --- /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 \\(\\);" 3 "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-not "#pragma omp dispatch device" "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..7b020c6da79 --- /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(out: 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..dba1d4ffdc8 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-3.c @@ -0,0 +1,12 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void f2 (void); + +void test (void) +{ +#pragma omp dispatch /* { dg-final { scan-tree-dump-not "#pragma omp task" "gimple" } } */ + f2 (); +#pragma omp dispatch nowait /* { dg-final { scan-tree-dump-not "nowait" "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..b8a4f8457aa --- /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 shared\\(D\.\[0-9]+\\) shared\\(a\\).*__builtin_omp_set_default_device \\(\\1\\);.*f2 \\(a\\)" 2 "gimple" } } */ + f2 (a); +} + +/* { dg-final { scan-tree-dump-times "(D\.\[0-9]+) = __builtin_omp_get_default_device \\(\\);.*__builtin_omp_set_default_device \\(\\1\\);" 4 "gimple" } } */ 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..c4d8f07bb6f --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-5.c @@ -0,0 +1,27 @@ +/* { 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 +/* { dg-final { scan-tree-dump-times "#pragma omp dispatch shared\\(arr\\) shared\\(p\\)\[ \t\n\r\{]*int (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(&arr, \\1\\);\[ \t\n\r]*\\3 = __builtin_omp_get_mapped_ptr \\(p, \\1\\);\[ \t\n\r]*f1 \\(\\3, \\2\\);" 4 "gimple" } } */ + f2 (p, arr); +#pragma omp dispatch is_device_ptr(p) +/* { dg-final { scan-tree-dump-times "#pragma omp dispatch is_device_ptr\\(p\\) shared\\(arr\\)\[ \t\n\r\{]*int (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(&arr, \\1\\);\[ \t\n\r]*f1 \\(p, \\2\\);" 3 "gimple" } } */ + f2 (p, arr); +#pragma omp dispatch is_device_ptr(arr) +/* { dg-final { scan-tree-dump-times "#pragma omp dispatch is_device_ptr\\(arr\\) shared\\(p\\)\[ \t\n\r\{]*int (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(p, \\1\\);\[ \t\n\r]*f1 \\(\\2, &arr\\);" 3 "gimple" } } */ + f2 (p, arr); +#pragma omp dispatch is_device_ptr(p, arr) +/* { dg-final { scan-tree-dump-times "#pragma omp dispatch is_device_ptr\\(arr\\) is_device_ptr\\(p\\)\[ \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..4e900bb5e8d --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-6.c @@ -0,0 +1,18 @@ +/* { 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-not "__builtin_GOMP_task " "ompexp" } } */ + f2 (p); +#pragma omp dispatch depend(inout: p) +/* { dg-final { scan-tree-dump-times "(D\.\[0-9]+)\\\[2] = &p;\[ \n]*__builtin_GOMP_taskwait_depend \\(&\\1\\);" 2 "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..a8c85559c2a --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-7.c @@ -0,0 +1,21 @@ +/* { 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, 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 "a = f2 \\\(\\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 1 "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.c/dispatch-1.c b/libgomp/testsuite/libgomp.c-c++-common/dispatch-1.c similarity index 100% rename from libgomp/testsuite/libgomp.c/dispatch-1.c rename to libgomp/testsuite/libgomp.c-c++-common/dispatch-1.c diff --git a/libgomp/testsuite/libgomp.c/dispatch-2.c b/libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c similarity index 100% rename from libgomp/testsuite/libgomp.c/dispatch-2.c rename to libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c From patchwork Wed Aug 7 11:50:11 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: 1970007 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=HH0VFFL9; 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 4Wf7qV4RbJz1yXs for ; Wed, 7 Aug 2024 21:53:42 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 74D4D385C6C3 for ; Wed, 7 Aug 2024 11:53:40 +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 71ED3385828B for ; Wed, 7 Aug 2024 11:51:06 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 71ED3385828B 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 71ED3385828B 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=1723031475; cv=none; b=vNJ/zhRqicomM65O/Gi8+3p+O2CbmhlLV5Ypi8syK5cwiTHyQOeI1C8qkxPQNH1XGlcWA3h7bEGmIC/m3cR/3zBE48sFBQ3wEZ6WadxSet4byiATy0cywdM1QZ6Rz+hl4hYeM8pvcpbKHyR0aWTsfC3nI6YwFaebQvUsItZWVS8= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031475; c=relaxed/simple; bh=cEWA8oxVwS27S+HepqiSPEmgzOVSaLJ711i3IEPMBzQ=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=LD0qk4DazYqxykp1OQip5xjOuHTOk9N19SaiHBWOSXTtOL1TxWs8hMRndbszQvdMk4h/S2rVTl9w3dpSnnMWOVJz76+qu8UXouVuFzeivNuIuah2tV/lCAPXCTbPYSL6xPQ+KM0cA2wp7kg6QS7RjJXxqy6xRXF9yNd/n20DI0c= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x32b.google.com with SMTP id 5b1f17b1804b1-42812945633so12479625e9.0 for ; Wed, 07 Aug 2024 04:51:06 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1723031465; x=1723636265; 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=FUSCjkb1Hou5TbSgGhI3I2vK8pKr0WeGjn+qiL3v0/4=; b=HH0VFFL9ovD7MJUP5fVToWrDG5C5YcRG50fC/eMJLZX9RYbsscGLTkK8AyMVQ3wd7M HEtIcgPSwtU/vrw1GCq+lpwVUg6QwKko+wryPxQUhFfnefJ+pzkiIlQG3v8CswbQTICZ +I5sgrNVT5LIJ55oy5tek4mC625yM1ONndRKrzHEf60hdGxuHEHp9vg5atuaXEByFSF3 j2qqd+yYTeqHxU5x2WEzfYYSFBZXvhNzCOpdefFAHqYlnYpFrcZ0lmpo23chp9XNyPab Pe1qKf4eWrSqjc37SGJ/OO9CFjWQp37t71+bvxNZZgBS1x68B07pUMfs8L/S9KkRPpja BnkQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1723031465; x=1723636265; 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=FUSCjkb1Hou5TbSgGhI3I2vK8pKr0WeGjn+qiL3v0/4=; b=DfsuICI257mXb5oVmlmtRzz995IfpHsLdy+gn5fNaBAY1cVL2CIuYvuNBkBKRHSejc d+/2CUlwx6o92aTCO9WjB3NOfjIN1DvpDlaiF9gZ2SkyJoQJdUhgdyansOjP7zitU4Iy M2qnos38mURwxYD6ub7rMdQnzoPpXOLXm+5a4V2vzz6srEpvoVyv0TU+SrErEbdTNerv p4KNdlp435D3rBCGrRVfcPTd3xLEA5p7YpgSKqIcajDxZRdmbSBOO3Kg9LksxHYt1ZxE YmVSB9fMrBNQGTbrHdGcXc0hs7UhFMqFYBtXJ7DlGLTPQm1nbR27sORZcEDx/fIRWtQa htEQ== X-Gm-Message-State: AOJu0YyJp5L1AK0G83BWUsw7/V95bF2bcTtePyMX+WIKPYrvcJLHsKJx Pwr2VxBUc2UosWJ3vcFb0kaIywDChXy4f36d3KwhCU8Bhn558StcIYsn2IIoNdojS0d3k8eRZtz Cy08= X-Google-Smtp-Source: AGHT+IFs0IVK77xhTiXGnCFiPKly5OusJrUzU4HZQQQ9m+zLqrg2mw2cSXm4l2qaTdHUjs+1zOf/XA== X-Received: by 2002:a05:600c:3108:b0:426:55c4:ff34 with SMTP id 5b1f17b1804b1-428e6aff8fcmr124951575e9.15.1723031464122; Wed, 07 Aug 2024 04:51:04 -0700 (PDT) Received: from localhost.localdomain (vlan-185-velingrad-59.comnet.bg. [84.54.185.59]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4290580c970sm25430125e9.46.2024.08.07.04.51.03 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 07 Aug 2024 04:51:03 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v3 6/7] OpenMP: Fortran front-end support for dispatch + adjust_args Date: Wed, 7 Aug 2024 14:50:11 +0300 Message-ID: <20240807115012.3632947-7-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240807115012.3632947-1-parras@baylibre.com> References: <20240807115012.3632947-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, KAM_SHORT, RCVD_IN_BARRACUDACENTRAL, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds support for the `dispatch` construct and the `adjust_args` clause to the Fortran front-end. Handling of `adjust_args` across translation units is missing due to PR115271. 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:. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/declare-variant-2.f90: Update dg-error. * gfortran.dg/gomp/declare-variant-21.f90: New test (xfail). * gfortran.dg/gomp/declare-variant-21-aux.f90: New test. * 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 | 192 ++++++++++++++++-- gcc/fortran/parse.cc | 39 +++- gcc/fortran/resolve.cc | 2 + gcc/fortran/st.cc | 1 + gcc/fortran/trans-decl.cc | 9 +- gcc/fortran/trans-openmp.cc | 135 ++++++++++++ gcc/fortran/trans.cc | 1 + .../gfortran.dg/gomp/adjust-args-1.f90 | 58 ++++++ .../gfortran.dg/gomp/adjust-args-2.f90 | 18 ++ .../gfortran.dg/gomp/adjust-args-3.f90 | 27 +++ .../gfortran.dg/gomp/adjust-args-4.f90 | 58 ++++++ .../gfortran.dg/gomp/adjust-args-5.f90 | 58 ++++++ .../gfortran.dg/gomp/declare-variant-2.f90 | 6 +- .../gomp/declare-variant-21-aux.f90 | 25 +++ .../gfortran.dg/gomp/declare-variant-21.f90 | 22 ++ gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90 | 77 +++++++ gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90 | 79 +++++++ 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 | 25 +++ gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 | 39 ++++ gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 | 26 +++ gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 | 33 +++ 27 files changed, 998 insertions(+), 21 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/declare-variant-21-aux.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/declare-variant-21.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 80aa8ef84e7..a15a17c086c 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -2139,6 +2139,18 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses) } fputc (')', dumpfile); } + 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 @@ -2176,6 +2188,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; @@ -2279,6 +2294,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: @@ -3522,6 +3538,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 6207ad6ed1d..ee8eb5483d0 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -323,7 +323,7 @@ enum gfc_statement /* Note: gfc_match_omp_nothing returns ST_NONE. */ ST_OMP_NOTHING, ST_NONE, ST_OMP_UNROLL, ST_OMP_END_UNROLL, - ST_OMP_TILE, ST_OMP_END_TILE + ST_OMP_TILE, ST_OMP_END_TILE, ST_OMP_DISPATCH }; /* Types of interfaces that we can have. Assignment interfaces are @@ -1006,6 +1006,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; @@ -1433,6 +1436,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. */ }; @@ -1578,6 +1582,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; struct gfc_expr_list *sizes_list; const char *critical_name; @@ -1707,6 +1713,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. */ @@ -3044,7 +3051,7 @@ enum gfc_exec_op EXEC_OMP_PARALLEL_MASKED_TASKLOOP, EXEC_OMP_PARALLEL_MASKED_TASKLOOP_SIMD, EXEC_OMP_MASKED_TASKLOOP, EXEC_OMP_MASKED_TASKLOOP_SIMD, EXEC_OMP_SCOPE, EXEC_OMP_UNROLL, EXEC_OMP_TILE, - EXEC_OMP_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 c2b7d69c37c..31280ba15ad 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 333f0c7fe7f..f1ecdc0fdaf 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -72,7 +72,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) */ @@ -181,6 +181,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); @@ -323,6 +325,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); } } @@ -1106,6 +1110,8 @@ enum omp_mask2 OMP_CLAUSE_FULL, /* OpenMP 5.1. */ OMP_CLAUSE_PARTIAL, /* OpenMP 5.1. */ OMP_CLAUSE_SIZES, /* OpenMP 5.1. */ + OMP_CLAUSE_NOVARIANTS, /* OpenMP 5.1 */ + OMP_CLAUSE_NOCONTEXT, /* OpenMP 5.1 */ /* This must come last. */ OMP_MASK2_LAST }; @@ -3231,6 +3237,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) @@ -4590,6 +4615,9 @@ cleanup: (omp_mask (OMP_CLAUSE_SIZES)) #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 @@ -4903,6 +4931,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) { @@ -6129,6 +6163,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; @@ -6145,13 +6180,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 @@ -6164,18 +6215,79 @@ 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->ts.type != BT_DERIVED + || !n->sym->ts.u.derived->ts.is_iso_c + || (n->sym->ts.u.derived->intmod_sym_id != ISOCBINDING_PTR)) + { + 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 %L 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; } @@ -7618,7 +7730,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) @@ -7800,6 +7912,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) @@ -8749,14 +8881,18 @@ 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 + || (n->sym->ts.u.derived->intmod_sym_id + != ISOCBINDING_PTR)) + && 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); else if (n->sym->ts.type != BT_DERIVED - || !n->sym->ts.u.derived->ts.is_iso_c) + || !n->sym->ts.u.derived->ts.is_iso_c + || (n->sym->ts.u.derived->intmod_sym_id + != ISOCBINDING_PTR)) { /* For TARGET, non-C_PTR are deprecated and handled as has_device_addr. */ @@ -10391,6 +10527,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); @@ -11365,6 +11502,8 @@ omp_code_to_statement (gfc_code *code) return ST_OMP_TILE; case EXEC_OMP_UNROLL: return ST_OMP_UNROLL; + case EXEC_OMP_DISPATCH: + return ST_OMP_DISPATCH; default: gcc_unreachable (); } @@ -11780,6 +11919,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. */ @@ -11895,6 +12054,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 b28c8a94547..67e1157be93 100644 --- a/gcc/fortran/parse.cc +++ b/gcc/fortran/parse.cc @@ -1058,6 +1058,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); @@ -1924,7 +1925,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_TILE: case ST_OMP_UNROLL: \ + case ST_OMP_TILE: case ST_OMP_UNROLL: 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: \ @@ -2606,6 +2607,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; @@ -6214,6 +6218,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 @@ -6416,6 +6449,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 503029364c1..48d29c3f578 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: @@ -13058,6 +13059,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 0218d290782..3d0a40a4b41 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 0fdc41b1784..69fac5cfae9 100644 --- a/gcc/fortran/trans-decl.cc +++ b/gcc/fortran/trans-decl.cc @@ -2189,6 +2189,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. */ @@ -2440,7 +2442,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 df1bf144e23..20dab559741 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; @@ -6360,6 +6390,51 @@ 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); + + // Extract depend clauses and create taskwait + tree depend_clauses = NULL_TREE; + tree *depend_clauses_ptr = &depend_clauses; + for (tree c = omp_clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND) + { + *depend_clauses_ptr = c; + depend_clauses_ptr = &OMP_CLAUSE_CHAIN (c); + } + } + if (depend_clauses != NULL_TREE) + { + tree stmt = make_node (OMP_TASK); + TREE_TYPE (stmt) = void_node; + OMP_TASK_CLAUSES (stmt) = depend_clauses; + OMP_TASK_BODY (stmt) = NULL_TREE; + SET_EXPR_LOCATION (stmt, gfc_get_location (&code->loc)); + gfc_add_expr_to_block (&block, stmt); + } + + 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) { @@ -8272,6 +8347,8 @@ gfc_trans_omp_directive (gfc_code *code) case EXEC_OMP_UNROLL: 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: @@ -8388,6 +8465,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) @@ -8583,6 +8661,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); @@ -8599,6 +8690,50 @@ 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 + tree need_device_ptr_list = make_node (TREE_LIST); + vec adjust_args_list = vNULL; + for (gfc_omp_namelist *arg_list + = odv->need_device_ptr_arg_list; + arg_list != NULL; arg_list = arg_list->next) + { + if (!arg_list->sym->attr.dummy) + { + gfc_error ( + "list item %qs at %L is not a dummy argument", + arg_list->sym->name, &arg_list->where); + continue; + } + if (adjust_args_list.contains (arg_list->sym)) + { + gfc_error ("%qs at %L is specified more than once", + arg_list->sym->name, &arg_list->where); + continue; + } + adjust_args_list.safe_push (arg_list->sym); + int idx; + gfc_formal_arglist *arg; + for (arg = ns->proc_name->formal, idx = 0; arg != NULL; + arg = arg->next, idx++) + if (arg->sym == arg_list->sym) + break; + gcc_assert (arg != NULL); + need_device_ptr_list + = chainon (need_device_ptr_list, + build_tree_list ( + NULL_TREE, + build_int_cst ( + integer_type_node, + idx))); // Store 0-based argument index, + // as in gimplify_call_expr + } + + DECL_ATTRIBUTES (variant) = tree_cons ( + get_identifier ("omp declare variant variant adjust_args"), + build_tree_list (need_device_ptr_list, + NULL_TREE /*need_device_addr */), + DECL_ATTRIBUTES (variant)); } } } diff --git a/gcc/fortran/trans.cc b/gcc/fortran/trans.cc index d4c54093cbc..97827f6c7a1 100644 --- a/gcc/fortran/trans.cc +++ b/gcc/fortran/trans.cc @@ -2599,6 +2599,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/testsuite/gfortran.dg/gomp/adjust-args-1.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90 new file mode 100644 index 00000000000..ae8c3afa73d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90 @@ -0,0 +1,58 @@ +! Test parsing of OMP clause adjust_args +! { dg-do compile } + +module main + use iso_c_binding, only: c_ptr, c_funptr + 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 h (a) + import c_funptr + type(c_funptr), 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 f13 (a) + import c_funptr + type(c_funptr), intent(inout) :: a + !$omp declare variant (h) match (construct={dispatch}) adjust_args (need_device_ptr: a) ! { dg-error "argument list item 'a' in 'need_device_ptr' at .1. must be of TYPE.C_PTR." } + 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..291bb47aaa2 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90 @@ -0,0 +1,27 @@ +! Test translation of OMP clause adjust_args +! { dg-do compile } + +module main + use iso_c_binding, only: c_ptr + implicit none + type(c_ptr) :: b + +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 base4 (a) + type(c_ptr), intent(inout) :: a + !$omp declare variant (variant2) match (construct={dispatch}) adjust_args (need_device_ptr: b) ! { dg-error "list item 'b' at .1. is not a dummy argument" } + end subroutine + + subroutine variant2 (a) + type(c_ptr), intent(inout) :: a + 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/declare-variant-21-aux.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-variant-21-aux.f90 new file mode 100644 index 00000000000..59b55e0bb85 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/declare-variant-21-aux.f90 @@ -0,0 +1,25 @@ +! { dg-do compile { target skip-all-targets } } + +! Test XFAILed due to https://gcc.gnu.org/PR115271 + + +subroutine base_proc (a) + use iso_c_binding, only: c_ptr + type(c_ptr), intent(inout) :: a +end subroutine + +program main + use iso_c_binding, only: c_ptr + use my_mod + implicit none + + type(c_ptr) :: a + + + call base_proc(a) + !call variant_proc(a) + + !$omp dispatch + call base_proc(a) + +end program main diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-variant-21.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-variant-21.f90 new file mode 100644 index 00000000000..0a89c8ff231 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/declare-variant-21.f90 @@ -0,0 +1,22 @@ +! { dg-do run } +! { dg-additional-sources declare-variant-21-aux.f90 } +! { dg-additional-options "-fdump-tree-gimple" } + +module my_mod + use iso_c_binding, only: c_ptr + implicit none + interface + subroutine base_proc (a) + use iso_c_binding, only: c_ptr + type(c_ptr), intent(inout) :: a + end subroutine + end interface + +contains + subroutine variant_proc (a) + type(c_ptr), intent(inout) :: a + !$omp declare variant (base_proc) match (construct={dispatch}) adjust_args(need_device_ptr: a) + end subroutine +end module + +! { dg-final { scan-tree-dump "variant_proc \\(&a\\)" "gimple" { xfail *-*-* } } } 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..d2d555b5932 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90 @@ -0,0 +1,79 @@ +module main + use iso_c_binding, only: c_funptr + 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_funptr) :: 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 + 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 is_device_ptr(p) !{ dg-error "List item 'p' 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..edcd799a718 --- /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-not "#pragma omp task" "gimple" } } + call f2 () + !$omp dispatch nowait ! { dg-final { scan-tree-dump-not "nowait" "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..f96dce64d44 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 @@ -0,0 +1,25 @@ +! { 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-9_]+ \\* 2;.*#pragma omp dispatch shared\\(D\.\[0-9]+\\) shared\\(a\\).*__builtin_omp_set_default_device \\(\\1\\);.*f2 \\(&a\\)" 2 "gimple" } } + call f2 (a) + end subroutine +end module + +! { dg-final { scan-tree-dump-times "(D\.\[0-9]+) = __builtin_omp_get_default_device \\(\\);.*__builtin_omp_set_default_device \\(\\1\\);" 4 "gimple" } } 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..29e06de151a --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 @@ -0,0 +1,39 @@ +! { 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 + ! { dg-final { scan-tree-dump-times "#pragma omp dispatch shared\\(p2\\) shared\\(p\\)\[ \t\n\r\{]*integer\\(kind=4\\) (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(&p2, \\1\\);\[ \t\n\r]*\\3 = __builtin_omp_get_mapped_ptr \\(&p, \\1\\);\[ \t\n\r]*f1 \\(\\3, \\2\\);" 4 "gimple" } } + call f2 (p, p2) + !$omp dispatch is_device_ptr(p) + ! { dg-final { scan-tree-dump-times "#pragma omp dispatch is_device_ptr\\(p\\) shared\\(p2\\)\[ \t\n\r\{]*integer\\(kind=4\\) (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(&p2, \\1\\);\[ \t\n\r]*f1 \\(&p, \\2\\);" 3 "gimple" } } + call f2 (p, p2) + !$omp dispatch is_device_ptr(p2) + ! { dg-final { scan-tree-dump-times "#pragma omp dispatch is_device_ptr\\(p2\\) shared\\(p\\)\[ \t\n\r\{]*integer\\(kind=4\\) (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(&p, \\1\\);\[ \t\n\r]*f1 \\(\\2, &p2\\);" 3 "gimple" } } + call f2 (p, p2) + !$omp dispatch is_device_ptr(p, p2) + ! { dg-final { scan-tree-dump-times "#pragma omp dispatch is_device_ptr\\(p\\) is_device_ptr\\(p2\\)\[ \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..6b40af6f315 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 @@ -0,0 +1,26 @@ +! { 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-not "__builtin_GOMP_task " "ompexp" } } + call f2 (p) + !$omp dispatch depend(inout: p) + ! { dg-final { scan-tree-dump-times "(D\.\[0-9]+)\\\[2] = &p;\[ \n]*__builtin_GOMP_taskwait_depend \\(&\\1\\);" 2 "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..62f6da76b8f --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 @@ -0,0 +1,33 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +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 "a = f2 \\\(\\\);" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 1 "gimple" } } + From patchwork Wed Aug 7 11:50:12 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: 1970005 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=LFvy2taf; 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 4Wf7pR0Bczz1yXs for ; Wed, 7 Aug 2024 21:52:47 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id DA32F385B508 for ; Wed, 7 Aug 2024 11:52:44 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x335.google.com (mail-wm1-x335.google.com [IPv6:2a00:1450:4864:20::335]) by sourceware.org (Postfix) with ESMTPS id C28D4385B508 for ; Wed, 7 Aug 2024 11:51:06 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org C28D4385B508 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 C28D4385B508 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::335 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031469; cv=none; b=J+bVbV0wkJrFiThxz1vF4Vcx5VHGDIviYKyGaWS6c267Cz6RcuFinQ/4JJBaPqRzfn/PFr9Ufmqq08S8/naCWW6IX9temOhYPqCCO93Xru/O+uU8gY3c14QAL/LKA7MUQoNpsy9jyDlVmfI8GWUZBoM/kCGefi92p2qbmBP/Lmo= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1723031469; c=relaxed/simple; bh=Y/KcfE67Xb5QVtsun75YDR7zX/5OAXNinNVFXQODwog=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=ckoitgkz4mRmQpGszBFetEuOtJk/AeTHeaAAm9NmnyTvmSO6Cn9FWIPTHQ6lhEopVUTYfIMYOljwn8kDokn5l3u+w41TfqXYyMK4xFDQZZuIbVbSXXu7xG+JFgy/jaMSdYvYtH/86ACKIUNgT4mm1NfGy/d2EMcouI6yfauYBCU= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x335.google.com with SMTP id 5b1f17b1804b1-428141be2ddso11679515e9.2 for ; Wed, 07 Aug 2024 04:51:06 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1723031465; x=1723636265; 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=8Lt++dV3JV7xSZJM7gmNDzDB1aUua8xnFDESRhpP5h0=; b=LFvy2taf/L+c5iPAOlqrahb3mQQxkpZmEXzrm5UPISfaZkFm3FUfZg+dfSqmBsdqep JCaBIOgYy5TsOUYyQZMsTaYmeGCSAiUVRaQaqXpDFv2ArN2Y/D/qieh5JzX3X5uDBzFB YbfC5iMZlw8TRABs/3a7vTigAh8u+F4aAUcfL2GOlxsMGsbj3NF0sWwgI7KNOGx5Ypc0 rYVwjBo+Q+eWMmZbfjhG0eY4zqXR+KOTSRjdHtiMB4siQfkCIC4EuBnWgYJy+u89bAWU s/Rk5eqPHmn9c8SXrbowIjxRgBvSUGYOUNtYm9E6bqmcZqVc9LRRJq9Fp4WJ7hJv6T/C H0Bw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1723031465; x=1723636265; 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=8Lt++dV3JV7xSZJM7gmNDzDB1aUua8xnFDESRhpP5h0=; b=p75hvVpwyzBlTKfVyaV60PZS8XxfHPhpXd+DARoLBzrfbk1hckmp/TgqDTLy6NLZP0 hf0qIG6IK1wJSUmU5+cESOrwbBGp/5nfq4XQtGZnBuxypfo4QGulLbRRX1VjaqdS8l8A 1LcLFWrBR4+TCaWCoz7ki54fYPXV1bLIuitvW4CG28zuX8ouxV8GB60sMPtSXx+1rbTC 8cMuPz4bm6ZtBm5gdWyPD7AOIEV5Zj/cLeZy9GQ1TsXEZA8+MlwnsKBfLX02wyOauYSu M28r4vXwd8YFC18GHhtTdIKnEpky9X7vEWOQVL9RWnNSfrjRGB8kwu9/KJGyU7igmkoP AwAg== X-Gm-Message-State: AOJu0YxsxY2ZwfP4xWcFhQwmLsdi7Zqn+tFM4kMjWr8/ZjHpQRpCUMCT M4aYKP49Glr6G2Lw1U30XSZT2B9rv1qVqAEwyodqM5/Bxz2U4H8jlUElr9sacMtoj1/5iGA1iTl PsMI= X-Google-Smtp-Source: AGHT+IGNCZzYt8LVmz8mZI6LHx2VBlYUEtGCnbpHb7Q0Ymh9aNZ6H4a91hf8iv3Gz4l50KZZNqRjIw== X-Received: by 2002:a05:600c:458c:b0:428:1991:eda with SMTP id 5b1f17b1804b1-428e6b32f32mr113730555e9.21.1723031464942; Wed, 07 Aug 2024 04:51:04 -0700 (PDT) Received: from localhost.localdomain (vlan-185-velingrad-59.comnet.bg. [84.54.185.59]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4290580c970sm25430125e9.46.2024.08.07.04.51.04 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 07 Aug 2024 04:51:04 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v3 7/7] OpenMP: update documentation for dispatch and adjust_args Date: Wed, 7 Aug 2024 14:50:12 +0300 Message-ID: <20240807115012.3632947-8-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240807115012.3632947-1-parras@baylibre.com> References: <20240807115012.3632947-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_BARRACUDACENTRAL, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org 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 07cd75124b0..b35424c047a 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