From patchwork Fri Jul 12 14:11:48 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: 1959895 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=2wthKW2g; 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 4WLD8s1VwQz1xqj for ; Sat, 13 Jul 2024 00:13:33 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id C65C238319EF for ; Fri, 12 Jul 2024 14:13:30 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lj1-x231.google.com (mail-lj1-x231.google.com [IPv6:2a00:1450:4864:20::231]) by sourceware.org (Postfix) with ESMTPS id 88DE838323D1 for ; Fri, 12 Jul 2024 14:12:54 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 88DE838323D1 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 88DE838323D1 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::231 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793577; cv=none; b=Kx3RnKHdP2TrdNJ1PVuJ+znB8awVwcl6/JeBNjv435OBCHSjgflz4fH4tgnOLRzP8yADQn7VT8ZWQwfB22n9I6rAJ20W/K3bt1ez2M3N/lR45cRLxILakEg0Zy/6biZuA4taO6OMFwCcjcwfQ9qYt7bHqNHJBqYAnVgpWF+eJUI= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793577; c=relaxed/simple; bh=EVxwIzUzyYNOh2WgbBrc8m8+9uDcMQANyaChZZ41wHU=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=dQ5ndvO40f4vHWZweGlbpnw8Egmwgl4MXGb3WDijEpooG/vvxJvZzpozj4ufgrLmGEX9Agr9ZkS34FrPLq4+XIWfP7EEbb9UsH1bpnDD6sywP1dUmsTmwPxIhC4x8meDqGOKow2TJNNrzfG0pHe2AR2WAfyEqoE3iHldmvJzDSY= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lj1-x231.google.com with SMTP id 38308e7fff4ca-2eabd22d3f4so23180281fa.1 for ; Fri, 12 Jul 2024 07:12:54 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1720793572; x=1721398372; 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=IqNaBUfYPX2aUtbDRS0nM7EhKX6Mq0/fd0eiY6ARppY=; b=2wthKW2gKcbLJd5y5a2YplOtDjWZ9uE21dL51utnBsC5AW7Md6btJK8y88VVO+WyDP M5Nib/3XIf+nhiTh1RRSCcEP0EZ4zIGRuda5HmCzQDppf2s+ZgW18+XPFsH2QP9dOoln /7EiYTNLqq6rVjE5n3f83PDCqS+5jk/ezCjFWYHKQBoxFIGfxZaSsFzeuf4Qf2Sn8mn/ aDOM3VhDD1/Z0CEY8rXO4yED/E7dntz3WDbjqngQhV2+xzFn91HcoH06Hs0w/siQKSnC EAvzkcjNfcqosre1PNrpX9dAqdJpGqQCtLZrFE+E13JWZ/UJZ1EqQUNzRR/+lFi0VDfM FHBQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1720793572; x=1721398372; 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=IqNaBUfYPX2aUtbDRS0nM7EhKX6Mq0/fd0eiY6ARppY=; b=DML06Zm/bW5y8513ZFvZsYKLFpMPqn2W7qPtzyrAp+xZr/ey3ZjHLnqrDhx1c1qPSe PAd4dAxg/FfazOKzTguHvX75IK/9Bjpf5sHtGD4p85fVKuYrfKfpU1nmikyx4YJM81AR 8lI/WwzxMFGmtLjqzSDQ6uxvKRXGXROA1DEOfHTKBbQHh8SDPpEfTiOOT5BP85yyphFU 8vE0j30gucArA+fPDhf0gt7JFw9Rm//Ipge4mdB1iau/4LB/6hnWLPLJ2AAggvz5N57F leU06TzaZhEKD7IwrvvqnNyajHjhJPFCNFMStkWMkUrorX9OmC0jP1b47DMAD3VH/4ZP 4Gzw== X-Gm-Message-State: AOJu0YxzjWo7ao2+GNdmkqAo4oacwco9d6MMJADqOfDwa1blokwg3R6S GNXemAc0nSQJ5LxjwTQWAFCR29TFNtjpe4zxQkutY0r1zps0DZHsqMR9UY7mI/bCIQBqIWjjxMK I X-Google-Smtp-Source: AGHT+IFG0lH0pvj9EmjVFuIqyyBtKKXGa21Z3JelZ4qPpwtjHZjqfIMi1kDyoxkVVR2gPRGMotYhyQ== X-Received: by 2002:a2e:bc13:0:b0:2ec:568e:336e with SMTP id 38308e7fff4ca-2eeb30ba026mr88618641fa.1.1720793572263; Fri, 12 Jul 2024 07:12:52 -0700 (PDT) Received: from localhost.localdomain ([169.155.255.128]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4279f25b946sm24680325e9.19.2024.07.12.07.12.51 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 12 Jul 2024 07:12:51 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v2 1/8] Fix warnings for tree formats in gfc_error Date: Fri, 12 Jul 2024 16:11:48 +0200 Message-ID: <20240712141155.255186-2-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240712141155.255186-1-parras@baylibre.com> References: <20240712141155.255186-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-11.9 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This enables proper warnings for formats like %qD. gcc/c-family/ChangeLog: * c-format.cc (gcc_gfc_char_table): Add formats for tree objects. --- gcc/c-family/c-format.cc | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/gcc/c-family/c-format.cc b/gcc/c-family/c-format.cc index 5bfd2fc4469..f4163c9cbc0 100644 --- a/gcc/c-family/c-format.cc +++ b/gcc/c-family/c-format.cc @@ -847,6 +847,10 @@ static const format_char_info gcc_gfc_char_table[] = /* This will require a "locus" at runtime. */ { "L", 0, STD_C89, { T89_V, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN }, "", "R", NULL }, + /* These will require a "tree" at runtime. */ + { "DFTV", 1, STD_C89, { T89_T, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN }, "q+", "'", NULL }, + { "E", 1, STD_C89, { T89_T, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN, BADLEN }, "q+", "", NULL }, + /* These will require nothing. */ { "<>",0, STD_C89, NOARGUMENTS, "", "", NULL }, { NULL, 0, STD_C89, NOLENGTHS, NULL, NULL, NULL } From patchwork Fri Jul 12 14:11:49 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: 1959899 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=MgBRhieC; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4WLDBx3srWz1xpd for ; Sat, 13 Jul 2024 00:15:21 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id D11BA3831E1D for ; Fri, 12 Jul 2024 14:15:19 +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 F3D4938323D6 for ; Fri, 12 Jul 2024 14:12:55 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org F3D4938323D6 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 F3D4938323D6 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=1720793578; cv=none; b=vf924UWxuecQg3M8IaCrH00Ydk7weopG1yF1G2fre26p6fP0OIY6IGhy/L5ffD8oTpc0VnimMPR00b8dwDZpIioCIZ5Hy4GvxCnE031H8Lw6hCjuCDeMLK+jD/SF0//8QMCQea/GUat0S2PdNjqeugS9+A0S0+oIOetQw2TWXeM= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793578; c=relaxed/simple; bh=y0ZBvi8T/2JYfxZnMeUF+InKrLf/xpK0OmQq8VOMDRk=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=W8UvImkmXKRjVoCCfVycOaJ6kk+yd9d6Oh3wGtUwgXXkvfRjPqYxdxKhdq5/8H7jZQvytXOi0thJMQwJym67ZfBOjidMiZGy+E8bsOF9E6dRcXjR8fFSl4a+mGdrfA29yIyT5XlTh1txdg68qD26JJHszz+P/rC79Qi+8s1PrYg= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x32c.google.com with SMTP id 5b1f17b1804b1-42565670e20so18845565e9.0 for ; Fri, 12 Jul 2024 07:12:55 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1720793574; x=1721398374; 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=PODGL7t6MAykTH3cOhsXntwKfhXJKfvOTWIwWucEqk8=; b=MgBRhieCYuT4IxmjZ33DZH+HQX86UmmLoWWdZChJ7jNrn2DAmQseQZJ78UN5G3e8bC RW3/Z0MALoxTzE+CUnMKeKHduE2tyMV+2vM+N+1xXOiNuwjimTyMbavJYZaUAPTfC1br E9gS4akCUAr65uD7GbBLlziyFIO5MZ6Oz6rLvMHAG2xxbLT2v7ifAR9I7VZccShO1sOE PDUOyL14Fx50s5Oz1GiIth60FMBZx8tkwhpEAXE7buXs/CIRxr4tTHDAAWciy/Bs2Xsg w0Brff9ZSgtoG4D8m34Yvgq2AWeuS5vWYfLIPGriby7J3QnKw2lYARO4EYF1vPrmW4Lt OqbQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1720793574; x=1721398374; 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=PODGL7t6MAykTH3cOhsXntwKfhXJKfvOTWIwWucEqk8=; b=B+0SK7dlHqieEe9H9ez7pktuUfrdDxFE8+F3EujqZ0uKqwPSCBHejx0FsM6KMCnZHg ILxHw5gaZtKpXZm5veDPxPlCp7/sk1Ho/fMuOfE5/JLZhRX5sqAp6ClUyHxOUbvhECLF ROh4kmq+vydhb5tXxt8CxH+EVPY8sPtRT+VxmnoVeRG3GggpP8nh8//CiFUQXbzGq4/h Qanf1ACJtFCyAGXyJtjI/wpq0O/YJ9GR2LO/FO2eAslb33i5d6tgiCZErqn/N/bxcU4Q v5zqgAahOaTK1Jv4Z37TX5NCjYa3+NdkwuPgBQ1DvUcmz1ecMtN/6uZ4rEOexRiQ1afO YVOQ== X-Gm-Message-State: AOJu0YyUJMkueBJeFGTGFQ/gD7ftZrAsNTJVD5T3WXYohp1TC8enwU92 xCrFaSJzkqFktB8FpPFEGLg0fFnJwTlR+3A9Q7tebxZTHsJydQe/efeDAx/JajF6P0D8e2/dJBd s X-Google-Smtp-Source: AGHT+IFbdb+yTkzhjfNY6n9aWplNLk+NmsVy4E+DZ5xrgQLls9CcJBUhbBsedf9Uf9eti4Qz2dt/oA== X-Received: by 2002:a05:600c:15d6:b0:426:686f:7ad with SMTP id 5b1f17b1804b1-4279dae54b3mr23862185e9.10.1720793573607; Fri, 12 Jul 2024 07:12:53 -0700 (PDT) Received: from localhost.localdomain ([169.155.255.128]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4279f25b946sm24680325e9.19.2024.07.12.07.12.52 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 12 Jul 2024 07:12:52 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v2 2/8] OpenMP: dispatch + adjust_args tree data structures and front-end interfaces Date: Fri, 12 Jul 2024 16:11:49 +0200 Message-ID: <20240712141155.255186-3-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240712141155.255186-1-parras@baylibre.com> References: <20240712141155.255186-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.3 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch introduces the OMP_DISPATCH tree node, as well as two new clauses `nocontext` and `novariants`. It defines/exposes interfaces that will be used in subsequent patches that add front-end and middle-end support, but nothing generates these nodes yet. It also adds support for new OpenMP context selectors: `dispatch` as trait selector and `need_device_ptr` as pseudo-trait set selector. The purpose of the latter is for the C++ front-end to store the list of arguments (that need to be converted to device pointers) until the declaration of the variant function becomes available. gcc/ChangeLog: * builtin-types.def (BT_FN_PTR_CONST_PTR_INT): New. * omp-selectors.h (enum omp_tss_code): Add OMP_TRAIT_SET_NEED_DEVICE_PTR. (enum omp_ts_code): Add OMP_TRAIT_CONSTRUCT_DISPATCH. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. * tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (dump_generic_node): Handle OMP_DISPATCH. * tree.cc (omp_clause_num_ops): Add OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (omp_clause_code_name): Add "novariants" and "nocontext". * tree.def (OMP_DISPATCH): New. * tree.h (OMP_DISPATCH_BODY): New macro. (OMP_DISPATCH_CLAUSES): New macro. (OMP_CLAUSE_NOVARIANTS_EXPR): New macro. (OMP_CLAUSE_NOCONTEXT_EXPR): New macro. --- gcc/builtin-types.def | 1 + gcc/omp-selectors.h | 3 +++ gcc/tree-core.h | 7 +++++++ gcc/tree-pretty-print.cc | 21 +++++++++++++++++++++ gcc/tree.cc | 4 ++++ gcc/tree.def | 5 +++++ gcc/tree.h | 7 +++++++ 7 files changed, 48 insertions(+) diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def index c97d6bad1de..ef7aaf67d13 100644 --- a/gcc/builtin-types.def +++ b/gcc/builtin-types.def @@ -677,6 +677,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_INT_FEXCEPT_T_PTR_INT, BT_INT, BT_FEXCEPT_T_PTR, DEF_FUNCTION_TYPE_2 (BT_FN_INT_CONST_FEXCEPT_T_PTR_INT, BT_INT, BT_CONST_FEXCEPT_T_PTR, BT_INT) DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_UINT8, BT_PTR, BT_CONST_PTR, BT_UINT8) +DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_INT, BT_PTR, BT_CONST_PTR, BT_INT) DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR) diff --git a/gcc/omp-selectors.h b/gcc/omp-selectors.h index c61808ec0ad..12bc9e9afa0 100644 --- a/gcc/omp-selectors.h +++ b/gcc/omp-selectors.h @@ -31,6 +31,8 @@ enum omp_tss_code { OMP_TRAIT_SET_TARGET_DEVICE, OMP_TRAIT_SET_IMPLEMENTATION, OMP_TRAIT_SET_USER, + OMP_TRAIT_SET_NEED_DEVICE_PTR, // pseudo-set selector used to convey argument + // list until variant has a decl OMP_TRAIT_SET_LAST, OMP_TRAIT_SET_INVALID = -1 }; @@ -55,6 +57,7 @@ enum omp_ts_code { OMP_TRAIT_CONSTRUCT_PARALLEL, OMP_TRAIT_CONSTRUCT_FOR, OMP_TRAIT_CONSTRUCT_SIMD, + OMP_TRAIT_CONSTRUCT_DISPATCH, OMP_TRAIT_LAST, OMP_TRAIT_INVALID = -1 }; diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 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 2d2d5b6db6e..9da9630199b 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 28e8e71b036..961615a6030 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 Fri Jul 12 14:11:50 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: 1959898 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=g0iEpheU; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4WLDBY4g9Nz1xqx for ; Sat, 13 Jul 2024 00:15:01 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id CA69038323EA for ; Fri, 12 Jul 2024 14:14:59 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x330.google.com (mail-wm1-x330.google.com [IPv6:2a00:1450:4864:20::330]) by sourceware.org (Postfix) with ESMTPS id 0CAC33860768 for ; Fri, 12 Jul 2024 14:12:57 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 0CAC33860768 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 0CAC33860768 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::330 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793583; cv=none; b=bmkZ1aY9FveW/5B0MttUa9uQY2DBnwlzDPJb9Fm03ua6wGoKoI6OnIVb45GzLjiWQOIH92aBJkFjPNpnH9uY4lMDoFyNbRzGl0Wu+YPr5AkDVWcYoXdQSzCvXsG6MWNAQ/oMujKNh88i7R+82YwdWklhU8syoMjvRC/KH0nP3kg= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793583; c=relaxed/simple; bh=+z8lHlpgxWU+sFkyabDhbjX1u2Hcnv2m8KHJwv91+44=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=QCywhZmIGaVdK1g21z66BfC+ajQtLvHidC+ntkSOTInAEtyIyTFzgq+lGD3j6KSmF+7FB00llLZvsFTI/P3sSZ4ZmezYcxvN5jvdIv51q+6C7CEN/ONeZxYaYWK4Cv5FYKGB0ijKXq6t8LJNYJ3/ze6tvrOlPZqPOLdGSNmYIYc= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x330.google.com with SMTP id 5b1f17b1804b1-4265b7514fcso18763065e9.1 for ; Fri, 12 Jul 2024 07:12:56 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1720793575; x=1721398375; 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=cf1aVYxWXeRsESAykqV+VkA86RN3ZhTsbClAdlVOYRU=; b=g0iEpheU7GjOzunvlnlQGQTRe78mfZhjf5M3bjPUDRL4oVXH8geEZ0QsoDvUeSqc86 lhPjRP6dnpTqhbyfXWhm/IdDI8GHF8TDw6xdp5Uhr8i/h4zmWS+ubd6BGuyWceBAZZvW TlizFVsPKqsInDw+WHH2RyZen1VdsedpjS5JNGDslp0zXmjD+3dXL3TyTgkk9xC6cO8i khPOkc+0ocP1mU1DsPf31/5utTOf/c4CLsJNCeG5XaUV3M4NVyuitnGDMBrPFU3uDGbH mGXsQriR94AJ7kHbq+kEXL7e9c1pRQDlWDMxIPLxjMlCXVXJ/k9iBhc0jFGBDDCzXcDB /mZw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1720793575; x=1721398375; 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=cf1aVYxWXeRsESAykqV+VkA86RN3ZhTsbClAdlVOYRU=; b=OHiPffnDlqre7carPRGNFxM/S+aKGbM/exLZZxT39VhW/VWMRcYPXFqvTgrXfZupr8 3ZLEi2kdLtKQY01KbZ6YLjHdTK6xbdb2pn0OlUtLnP6zNbVCrz39QOGtPQdK4HB6zWBg FWenSarq//0j7WbL21UVocnD52Z7e7Ytikp+931BTmjD7N1jWOkL+VXACo5yMlFPOFcd A59WY2JamqA5oyYKzaXd2IlMjyWfI8fGOyG5+K8MlAIAtTNc+q3T3cBmaoaTNcRCVUZ3 P8ORQ62HuciDL+XlNOx5VLWyoTSKfa6VhigUwVR08DuZjGQOIB7Z4fTB7Ik9IekKoiAT wC2w== X-Gm-Message-State: AOJu0YwelgBpJ68/Kjuaq7Oly2x1v2ul5CvCCvjDNlObsjJzccjNU8nJ A/lWlgEUlZ+bNUshG+kAs3XkiR5sdL9nAri10y8fidMvxRUqoZwLInkG80lmEt2oZJzeFF3SEyg 2 X-Google-Smtp-Source: AGHT+IHvlsMjnsLEAH5TQbgHFdR90ng8fmrfoiQLSh6s4FePPVb04awR2cEjhx2RBOZNojZ4t6/i0w== X-Received: by 2002:a05:600c:6c8c:b0:426:62a2:34fc with SMTP id 5b1f17b1804b1-4279dae54e5mr23924175e9.11.1720793575133; Fri, 12 Jul 2024 07:12:55 -0700 (PDT) Received: from localhost.localdomain ([169.155.255.128]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4279f25b946sm24680325e9.19.2024.07.12.07.12.53 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 12 Jul 2024 07:12:54 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v2 3/8] OpenMP: middle-end support for dispatch + adjust_args Date: Fri, 12 Jul 2024 16:11:50 +0200 Message-ID: <20240712141155.255186-4-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240712141155.255186-1-parras@baylibre.com> References: <20240712141155.255186-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.4 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds middle-end support for the `dispatch` construct and the `adjust_args` clause. The heavy lifting is done in `gimplify_omp_dispatch` and `gimplify_call_expr` respectively. For `adjust_args`, this mostly consists in emitting a call to `gomp_get_mapped_ptr` for the adequate device. For dispatch, the following steps are performed: * Handle the device clause, if any. This may affect `need_device_ptr` arguments. * Handle novariants and nocontext clauses, if any. Evaluate compile-time constants and select a variant, if possible. Otherwise, emit code to handle all possible cases at run time. * Create an explicit task, as if the `task` construct was used, that wraps the body of the `dispatch` statement. Move relevant clauses to the task. gcc/ChangeLog: * gimple-low.cc (lower_stmt): Handle GIMPLE_OMP_DISPATCH. * gimple-pretty-print.cc (dump_gimple_omp_dispatch): New function. (pp_gimple_stmt_1): Handle GIMPLE_OMP_DISPATCH. * gimple-walk.cc (walk_gimple_stmt): Likewise. * gimple.cc (gimple_build_omp_dispatch): New function. (gimple_copy): Handle GIMPLE_OMP_DISPATCH. * gimple.def (GIMPLE_OMP_DISPATCH): Define. * gimple.h (gimple_build_omp_dispatch): Declare. (gimple_has_substatements): Handle GIMPLE_OMP_DISPATCH. (gimple_omp_dispatch_clauses): New function. (gimple_omp_dispatch_clauses_ptr): Likewise. (gimple_omp_dispatch_set_clauses): Likewise. (gimple_return_set_retval): Handle GIMPLE_OMP_DISPATCH. * gimplify.cc (enum omp_region_type): Add ORT_DISPATCH. (gimplify_call_expr): Handle need_device_ptr arguments. (is_gimple_stmt): Handle OMP_DISPATCH. (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_DEVICE in a dispatch construct. Handle OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (omp_construct_selector_matches): Handle OMP_DISPATCH with nocontext clause. (omp_has_novariants): New function. (omp_has_nocontext): Likewise. (gimplify_omp_dispatch): Likewise. (gimplify_expr): Handle OMP_DISPATCH. * gimplify.h (omp_has_novariants): Declare. (omp_has_nocontext): Declare. * omp-builtins.def (BUILT_IN_OMP_GET_MAPPED_PTR): Define. (BUILT_IN_OMP_GET_DEFAULT_DEVICE): Define. (BUILT_IN_OMP_SET_DEFAULT_DEVICE): Define. * omp-expand.cc (expand_omp_dispatch): New function. (expand_omp): Handle GIMPLE_OMP_DISPATCH. (omp_make_gimple_edges): Likewise. * omp-general.cc (omp_construct_traits_to_codes): Add OMP_DISPATCH. (struct omp_ts_info): Add dispatch. (omp_context_selector_matches): Handle OMP_TRAIT_SET_NEED_DEVICE_PTR. (omp_resolve_declare_variant): Handle novariants. Adjust DECL_ASSEMBLER_NAME. --- gcc/gimple-low.cc | 1 + gcc/gimple-pretty-print.cc | 33 +++ gcc/gimple-walk.cc | 1 + gcc/gimple.cc | 20 ++ gcc/gimple.def | 5 + gcc/gimple.h | 33 ++- gcc/gimplify.cc | 412 ++++++++++++++++++++++++++++++++++++- gcc/gimplify.h | 2 + gcc/omp-builtins.def | 6 + gcc/omp-expand.cc | 18 ++ gcc/omp-general.cc | 16 +- gcc/omp-low.cc | 35 ++++ gcc/tree-inline.cc | 7 + 13 files changed, 578 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 02faaf7114c..8e4329d2b42 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -161,7 +161,8 @@ enum omp_region_type { ORT_WORKSHARE = 0x00, ORT_TASKGROUP = 0x01, - ORT_SIMD = 0x04, + ORT_DISPATCH = 0x02, + ORT_SIMD = 0x04, ORT_PARALLEL = 0x08, ORT_COMBINED_PARALLEL = ORT_PARALLEL | 1, @@ -4051,6 +4052,7 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value) /* Gimplify the function arguments. */ if (nargs > 0) { + tree device_num = NULL_TREE; for (i = (PUSH_ARGS_REVERSED ? nargs - 1 : 0); PUSH_ARGS_REVERSED ? i >= 0 : i < nargs; PUSH_ARGS_REVERSED ? i-- : i++) @@ -4061,8 +4063,100 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value) be the plain PARM_DECL. */ if ((i != 1) || !builtin_va_start_p) { - t = gimplify_arg (&CALL_EXPR_ARG (*expr_p, i), pre_p, - EXPR_LOCATION (*expr_p), ! returns_twice); + tree *arg_p = &CALL_EXPR_ARG (*expr_p, i); + if (flag_openmp && EXPR_P (CALL_EXPR_FN (*expr_p)) + && DECL_P (TREE_OPERAND (CALL_EXPR_FN (*expr_p), 0)) + && lookup_attribute ("omp declare variant variant", + DECL_ATTRIBUTES (TREE_OPERAND ( + CALL_EXPR_FN (*expr_p), 0))) + != NULL_TREE) + { + tree param + = DECL_ARGUMENTS (TREE_OPERAND (CALL_EXPR_FN (*expr_p), 0)); + + if (param != NULL_TREE) + { + for (int param_idx = 0; param_idx < i; param_idx++) + param = TREE_CHAIN (param); + + bool is_device_ptr = false; + if (gimplify_omp_ctxp != NULL + && gimplify_omp_ctxp->code == OMP_DISPATCH) + { + for (tree c = gimplify_omp_ctxp->clauses; c; + c = TREE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) + == OMP_CLAUSE_IS_DEVICE_PTR) + { + tree decl1 = DECL_NAME (OMP_CLAUSE_DECL (c)); + tree decl2 + = tree_strip_nop_conversions (*arg_p); + if (TREE_CODE (decl2) == ADDR_EXPR) + decl2 = TREE_OPERAND (decl2, 0); + gcc_assert (TREE_CODE (decl2) == VAR_DECL + || TREE_CODE (decl2) + == PARM_DECL); + decl2 = DECL_NAME (decl2); + if (decl1 == decl2) + { + is_device_ptr = true; + break; + } + } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE) + device_num = OMP_CLAUSE_OPERAND (c, 0); + } + } + + if (!is_device_ptr + && lookup_attribute ("omp declare variant " + "adjust_args need_device_ptr", + DECL_ATTRIBUTES (param)) + != NULL_TREE) + { + if (device_num == NULL_TREE) + { + // device_num = omp_get_default_device(); + tree fn = builtin_decl_explicit ( + BUILT_IN_OMP_GET_DEFAULT_DEVICE); + gcall *call = gimple_build_call (fn, 0); + device_num = create_tmp_var ( + gimple_call_return_type (call)); + gimple_call_set_lhs (call, device_num); + gimplify_seq_add_stmt (pre_p, call); + } + + // mapped_arg = omp_get_mapped_ptr(arg, device_num); + tree fn = builtin_decl_explicit ( + BUILT_IN_OMP_GET_MAPPED_PTR); + *arg_p = (TREE_CODE (*arg_p) == NOP_EXPR) + ? TREE_OPERAND (*arg_p, 0) + : *arg_p; + gimplify_arg (arg_p, pre_p, loc); + gimplify_arg (&device_num, pre_p, loc); + call = gimple_build_call (fn, 2, *arg_p, device_num); + tree mapped_arg + = create_tmp_var (gimple_call_return_type (call)); + gimple_call_set_lhs (call, mapped_arg); + gimplify_seq_add_stmt (pre_p, call); + + *arg_p = mapped_arg; + + // Mark mapped argument as device pointer to ensure + // idempotency in gimplification + gcc_assert (gimplify_omp_ctxp->code == OMP_DISPATCH); + tree c = build_omp_clause (input_location, + OMP_CLAUSE_IS_DEVICE_PTR); + OMP_CLAUSE_DECL (c) = *arg_p; + OMP_CLAUSE_CHAIN (c) = gimplify_omp_ctxp->clauses; + gimplify_omp_ctxp->clauses = c; + } + } + } + + t = gimplify_arg (arg_p, pre_p, EXPR_LOCATION (*expr_p), + !returns_twice); if (t == GS_ERROR) ret = GS_ERROR; @@ -6309,6 +6403,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: @@ -13128,6 +13223,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: @@ -13357,6 +13467,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 (); @@ -13811,7 +13929,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) { @@ -14622,6 +14742,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: @@ -14711,9 +14833,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) @@ -14831,6 +14953,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 @@ -17824,6 +18000,221 @@ gimplify_omp_ordered (tree expr, gimple_seq body) return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr)); } +/* Gimplify an OMP_DISPATCH construct. */ + +static enum gimplify_status +gimplify_omp_dispatch (tree *expr_p, gimple_seq *pre_p) +{ + tree expr = *expr_p; + gimple_seq body = NULL; + + gimplify_scan_omp_clauses (&OMP_DISPATCH_CLAUSES (expr), pre_p, ORT_DISPATCH, + OMP_DISPATCH); + push_gimplify_context (); + + // If device clause, adjust ICV + tree device + = omp_find_clause (OMP_DISPATCH_CLAUSES (expr), OMP_CLAUSE_DEVICE); + if (device) + { + tree t = builtin_decl_explicit (BUILT_IN_OMP_SET_DEFAULT_DEVICE); + t = build_call_expr_loc (input_location, t, 1, + OMP_CLAUSE_DEVICE_ID (device)); + gimplify_and_add (t, &body); + if (DECL_P (OMP_CLAUSE_DEVICE_ID (device))) + omp_notice_variable (gimplify_omp_ctxp, OMP_CLAUSE_DEVICE_ID (device), + true); + } + + // If the novariants and nocontext clauses are not compile-time constants, + // we need to generate code for all possible cases: + // if (novariants) // implies nocontext + // base() + // else if (nocontext) + // variant1() + // else + // variant2() + tree dispatch_body = OMP_DISPATCH_BODY (expr); + if (TREE_CODE (dispatch_body) == BIND_EXPR) + dispatch_body = BIND_EXPR_BODY (dispatch_body); + if (TREE_CODE (dispatch_body) == STATEMENT_LIST) + { + // Fortran FE may insert some pre-call code, for instance when an + // array is passed as argument. Skip to the actual call. + dispatch_body = expr_last (dispatch_body); + } + gcc_assert (TREE_CODE (dispatch_body) == CALL_EXPR + || TREE_CODE (dispatch_body) == MODIFY_EXPR); + tree base_call_expr = dispatch_body; + tree dst = 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); + + // Wrap dispatch body into a bind + gimple *bind = gimple_build_bind (NULL_TREE, body, NULL_TREE); + pop_gimplify_context (bind); + + gimplify_adjust_omp_clauses (pre_p, bind, &OMP_DISPATCH_CLAUSES (expr), + OMP_DISPATCH); + + // Move relevant clauses to the task construct + tree task_clauses = NULL_TREE; + tree *task_clauses_ptr = &task_clauses; + bool has_nowait = false; + for (tree c = OMP_DISPATCH_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND) + { + *task_clauses_ptr = c; + task_clauses_ptr = &OMP_CLAUSE_CHAIN (c); + } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) + { + *task_clauses_ptr + = build_omp_clause (input_location, OMP_CLAUSE_SHARED); + OMP_CLAUSE_DECL (*task_clauses_ptr) = OMP_CLAUSE_DECL (c); + task_clauses_ptr = &OMP_CLAUSE_CHAIN (*task_clauses_ptr); + } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NOWAIT) + has_nowait = true; + } + *task_clauses_ptr = build_omp_clause (input_location, OMP_CLAUSE_IF); + OMP_CLAUSE_IF_EXPR (*task_clauses_ptr) + = has_nowait ? boolean_true_node : boolean_false_node; + + // Wrap bind into a task + gimple *task + = gimple_build_omp_task (bind, task_clauses, NULL_TREE, NULL_TREE, + NULL_TREE, NULL_TREE, NULL_TREE); + + gimple *stmt = gimple_build_omp_dispatch (task, OMP_DISPATCH_CLAUSES (expr)); + gimplify_seq_add_stmt (pre_p, stmt); + *expr_p = NULL_TREE; + return GS_ALL_DONE; +} + /* Convert the GENERIC expression tree *EXPR_P to GIMPLE. If the expression produces a value to be used as an operand inside a GIMPLE statement, the value will be stored back in *EXPR_P. This value will @@ -18752,6 +19143,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; @@ -19077,7 +19472,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..42a3091fd00 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. */ @@ -1456,6 +1461,8 @@ omp_context_selector_matches (tree ctx) for (tree tss = ctx; tss; tss = TREE_CHAIN (tss)) { enum omp_tss_code set = OMP_TSS_CODE (tss); + if (set == OMP_TRAIT_SET_NEED_DEVICE_PTR) + continue; tree selectors = OMP_TSS_TRAIT_SELECTORS (tss); /* Immediately reject the match if there are any ignored @@ -2495,6 +2502,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 +2651,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 Fri Jul 12 14:11:51 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: 1959901 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=JpyZzQvg; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4WLDCy0qsFz1xpd for ; Sat, 13 Jul 2024 00:16:14 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 4E03138323F5 for ; Fri, 12 Jul 2024 14:16:12 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lf1-x130.google.com (mail-lf1-x130.google.com [IPv6:2a00:1450:4864:20::130]) by sourceware.org (Postfix) with ESMTPS id CA128385DDE9 for ; Fri, 12 Jul 2024 14:12:58 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org CA128385DDE9 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 CA128385DDE9 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::130 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793583; cv=none; b=ICwxXSBZ7pNxA7CbCwdz0f2fn6j7MSx8Hepz76k1282X8gF8E618VkxGCaWyLPe2ZSsbfIek95aYjiURBSsH/DFxJXWi2zvWe6la53AP486HozmiOFruwUwOmpvItmxY4kKf/UJvNwiTPiJ/59VKFgBk0olxo+GGz6306O1yriI= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793583; c=relaxed/simple; bh=MQnBsEmyR+kix6dYpQL77iF63/+MnECZUhna83nGtzA=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=l2aD8hffTagfW5E0mCMSZIAHRR75aREVvE8N4vKOY1mSmPaIVcBU4SwDwkk0RDy83TmzkD4RLNC0WUwqSCi5VcVDg4nx31vKq5m8t3BR/f2/ReJYDNS3FPt9dJiUkDMMCp+8phhU8eRMwsgogCU7Qn3Yv522XZRPDo0y9G36e/0= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lf1-x130.google.com with SMTP id 2adb3069b0e04-52e99060b41so2189260e87.2 for ; Fri, 12 Jul 2024 07:12:58 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1720793576; x=1721398376; 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=tLNfdHs718tFAuvaMDy34TpGxYu1pgZ0ZRdWrjF4aX8=; b=JpyZzQvgrXCt11s/AhFJuYViody/XzlRbte+UrdQCOErWZqb9Rthf/N04K2xv6zH16 jLKLe0in74Lto9CvEakcS9yFU7rKynmFsG1J+qc1k9VrxLtr0YNNKxsvbvZbbvHBUFmT 26h0rNoiJ3U3RnntpfGJUl0vLvVK+0XOliO5oDoCEAyutxTQOVlI/E6T3eBqT9P5QPCb 7iGEVtqHVZjo2j8Bp2/Sw44wAMR6UAZWHxhh/YDbD2YX3AFbxfGESVAclQRRgdbBQkdY 72R4Vogiq9bnuOBF2doEB+z+2ySGXltGx6NVWv09pexSR4D4hSSNJ59DpN7i6wUU0RP1 VTIw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1720793576; x=1721398376; 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=tLNfdHs718tFAuvaMDy34TpGxYu1pgZ0ZRdWrjF4aX8=; b=A4R2lzJ8DSde0ol8RRF5AmoRFeMjmHXEnETyXaKcOS1Xnp+Hn34SvFmMK0tCJsIWZN rsG+OKK6Tb6wEX2EbZqeiaM9R9uTaR+6a5V9vJVsRhv2H5Lmly6DRKlMTQPrSE1DzpxN deLmfYzB4hGt493BQGwxMOKQ6F4LDukE71fmrW60q9hTjq3kdnk6Rtc1jq+rNsngO7rT Pr7V16SxuisDquaB8URQ3TqRgVCUH4IK+ovVdNMnDEfAtxcix+wnAiRCdGlJu53mnfUZ 4SlWdI0c+X9bDux+5WS28n00TEIyCmNdftTs2ggqQxqvNOgtL/ZpOAFfOhkeU2ZJ7Ceu WbGQ== X-Gm-Message-State: AOJu0YwWpEf1YHZgf51f/j1b+6PcJCG0ofrkLo0zrca6RqDn9gsf6ako NFgBOuuUHvKcaxo40+q099Ii5xFI+lvA09FKPvDXeDoi10iZfqhWg4D3rmRZOxvcEbuKolazGsV z X-Google-Smtp-Source: AGHT+IE3TzwEmq9W9wvwF0OS1DBWTgvh2ZW3LbYMi3j7qUS5Rjr9e6PtGmcTROGIl8qYfmq+LSlPTg== X-Received: by 2002:a05:6512:3b84:b0:52e:9ecd:3465 with SMTP id 2adb3069b0e04-52eb99d4e22mr8299286e87.57.1720793575992; Fri, 12 Jul 2024 07:12:55 -0700 (PDT) Received: from localhost.localdomain ([169.155.255.128]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4279f25b946sm24680325e9.19.2024.07.12.07.12.55 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 12 Jul 2024 07:12:55 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v2 4/8] OpenMP: C front-end support for dispatch + adjust_args Date: Fri, 12 Jul 2024 16:11:51 +0200 Message-ID: <20240712141155.255186-5-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240712141155.255186-1-parras@baylibre.com> References: <20240712141155.255186-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.9 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds support 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 | 496 +++++++++++++++++++--- gcc/c/c-typeck.cc | 2 + gcc/testsuite/gcc.dg/gomp/adjust-args-1.c | 32 ++ gcc/testsuite/gcc.dg/gomp/dispatch-1.c | 53 +++ libgomp/testsuite/libgomp.c/dispatch-1.c | 76 ++++ 9 files changed, 609 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 diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc index f9b229aba7f..1cb49d7b911 100644 --- a/gcc/c-family/c-attribs.cc +++ b/gcc/c-family/c-attribs.cc @@ -560,6 +560,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 12c5ed5d92c..8a6653057ef 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -1740,6 +1740,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 @@ -15044,6 +15046,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)) @@ -15052,6 +15056,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)) @@ -19319,6 +19325,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 ) */ @@ -19938,6 +19998,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; @@ -23748,6 +23816,168 @@ c_parser_omp_scope (location_t loc, c_parser *parser, bool *if_p) return add_stmt (stmt); } +// Adapted from c_parser_expr_no_commas +static tree +c_parser_omp_dispatch_body (c_parser *parser) +{ + struct c_expr lhs, rhs, ret; + struct c_expr orig_expr; + location_t expr_loc = c_parser_peek_token (parser)->location; + source_range tok_range = c_parser_peek_token (parser)->get_range (); + location_t sizeof_arg_loc[3]; + tree sizeof_arg[3]; + vec *exprlist; + vec arg_loc = vNULL; + vec *origtypes = NULL; + unsigned int literal_zero_mask; + location_t start; + location_t finish; + + lhs = c_parser_conditional_expression (parser, NULL, NULL); + if (TREE_CODE (lhs.value) == CALL_EXPR) + return lhs.value; + else + { + location_t op_location = c_parser_peek_token (parser)->location; + if (!c_parser_require (parser, CPP_EQ, "expected %<=%>")) + return error_mark_node; + + /* Parse function name*/ + if (!c_parser_next_token_is (parser, CPP_NAME)) + { + c_parser_error (parser, "expected a function name"); + rhs.set_error (); + return rhs.value; + } + expr_loc = c_parser_peek_token (parser)->location; + tree id = c_parser_peek_token (parser)->value; + c_parser_consume_token (parser); + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + return error_mark_node; + + rhs.value = build_external_ref (expr_loc, id, true, &rhs.original_type); + set_c_expr_source_range (&rhs, tok_range); + /* Parse argument list */ + { + for (int i = 0; i < 3; i++) + { + sizeof_arg[i] = NULL_TREE; + sizeof_arg_loc[i] = UNKNOWN_LOCATION; + } + literal_zero_mask = 0; + if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN)) + exprlist = NULL; + else + exprlist = c_parser_expr_list (parser, true, false, &origtypes, + sizeof_arg_loc, sizeof_arg, &arg_loc, + &literal_zero_mask); + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>"); + } + orig_expr = rhs; + mark_exp_read (rhs.value); + if (warn_sizeof_pointer_memaccess) + sizeof_pointer_memaccess_warning (sizeof_arg_loc, rhs.value, exprlist, + sizeof_arg, + sizeof_ptr_memacc_comptypes); + if (TREE_CODE (rhs.value) == FUNCTION_DECL) + { + if (fndecl_built_in_p (rhs.value, BUILT_IN_MEMSET) + && vec_safe_length (exprlist) == 3) + { + tree arg0 = (*exprlist)[0]; + tree arg2 = (*exprlist)[2]; + warn_for_memset (expr_loc, arg0, arg2, literal_zero_mask); + } + if (warn_absolute_value + && fndecl_built_in_p (rhs.value, BUILT_IN_NORMAL) + && vec_safe_length (exprlist) == 1) + warn_for_abs (expr_loc, rhs.value, (*exprlist)[0]); + if (parser->omp_for_parse_state + && parser->omp_for_parse_state->in_intervening_code + && omp_runtime_api_call (rhs.value)) + { + error_at (expr_loc, "calls to the OpenMP runtime API are " + "not permitted in intervening code"); + parser->omp_for_parse_state->fail = true; + } + } + + start = rhs.get_start (); + finish = parser->tokens_buf[0].get_finish (); + rhs.value = c_build_function_call_vec (expr_loc, arg_loc, rhs.value, + exprlist, origtypes); + set_c_expr_source_range (&rhs, start, finish); + rhs.m_decimal = 0; + + rhs.original_code = ERROR_MARK; + if (TREE_CODE (rhs.value) == INTEGER_CST + && TREE_CODE (orig_expr.value) == FUNCTION_DECL + && fndecl_built_in_p (orig_expr.value, BUILT_IN_CONSTANT_P)) + rhs.original_code = C_MAYBE_CONST_EXPR; + rhs.original_type = NULL; + if (exprlist) + { + release_tree_vector (exprlist); + release_tree_vector (origtypes); + } + arg_loc.release (); + + /* Build assignment */ + rhs = convert_lvalue_to_rvalue (expr_loc, rhs, true, true); + ret.value + = build_modify_expr (op_location, lhs.value, lhs.original_type, + NOP_EXPR, expr_loc, rhs.value, rhs.original_type); + ret.m_decimal = 0; + set_c_expr_source_range (&ret, lhs.get_start (), rhs.get_finish ()); + ret.original_code = MODIFY_EXPR; + ret.original_type = NULL; + return ret.value; + } +} + +/* OpenMP 5.1: + # pragma omp dispatch dispatch-clause[optseq] new-line + expression-stmt + + LOC is the location of the #pragma. +*/ + +#define OMP_DISPATCH_CLAUSE_MASK \ + ((OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT)) + +static tree +c_parser_omp_dispatch (location_t loc, c_parser *parser) +{ + tree stmt = make_node (OMP_DISPATCH); + SET_EXPR_LOCATION (stmt, loc); + TREE_TYPE (stmt) = void_type_node; + + OMP_DISPATCH_CLAUSES (stmt) + = c_parser_omp_all_clauses (parser, OMP_DISPATCH_CLAUSE_MASK, + "#pragma omp dispatch"); + + // Parse body as expression statement + loc = c_parser_peek_token (parser)->location; + tree dispatch_body = c_parser_omp_dispatch_body (parser); + if (dispatch_body == error_mark_node) + { + inform (loc, "%<#pragma omp dispatch%> must be followed by a function " + "call with optional assignment"); + c_parser_skip_to_end_of_block_or_statement (parser); + return NULL_TREE; + } + + c_parser_skip_until_found (parser, CPP_SEMICOLON, "expected %<;%>"); + OMP_DISPATCH_BODY (stmt) = dispatch_body; + + return add_stmt (stmt); +} + /* OpenMP 3.0: # pragma omp task task-clause[optseq] new-line @@ -24728,6 +24958,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 \ @@ -25191,77 +25425,217 @@ 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; - 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); // : - parens.require_close (parser); + location_t loc = c_parser_peek_token (parser)->location; + tree list + = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_ERROR, + NULL_TREE); + + tree arg; + for (tree c = list; c != NULL_TREE; c = TREE_CHAIN (c)) + { + tree decl = TREE_PURPOSE (c); + int idx; + for (arg = parms, idx = 0; arg != NULL; + arg = TREE_CHAIN (arg), idx++) + if (arg == decl) + break; + if (arg == NULL_TREE) + { + error_at (loc, "%qD is not a function argument", + decl); + goto fail; + } + if (adjust_args_list.contains (arg)) + { + error_at (loc, "%qD is specified more than once", + decl); + goto fail; + } + if (strcmp (p, "need_device_ptr") == 0 + && TREE_CODE (TREE_TYPE (arg)) != POINTER_TYPE) + { + error_at (loc, "%qD is not a C pointer", decl); + goto fail; + } + adjust_args_list.safe_push (arg); + if (strcmp (p, "need_device_ptr") == 0) + { + tree variant_decl = (TREE_CODE (variant) == NOP_EXPR) + ? TREE_OPERAND (variant, 0) + : variant; + tree variant_parm = DECL_ARGUMENTS (variant_decl); + for (int i = 0; i < idx; i++) + { + variant_parm = TREE_CHAIN (variant_parm); + gcc_assert (variant_parm != NULL); + } + tree attr = tree_cons ( + get_identifier ("omp declare variant adjust_args " + "need_device_ptr"), + NULL_TREE, DECL_ATTRIBUTES (variant_parm)); + DECL_ATTRIBUTES (variant_parm) = attr; + } + } + } + else + { + error_at (c_parser_peek_token (parser)->location, + "expected % or %"); + goto fail; + } + } + else + { + error_at (c_parser_peek_token (parser)->location, + "expected % or % " + "followed by %<:%>"); + goto fail; + } + } + + parens.require_close (parser); + } while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL)); c_parser_skip_to_pragma_eol (parser); + + if (has_adjust_args) + { + if (!has_match) + { + error_at ( + adjust_args_loc, + "an % clause can only be specified if the " + "% selector of the construct selector set appears " + "in the % clause"); + } + else + { + tree attr = lookup_attribute ("omp declare variant base", + DECL_ATTRIBUTES (fndecl)); + tree ctx = TREE_VALUE (TREE_VALUE (attr)); + if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT, + OMP_TRAIT_CONSTRUCT_DISPATCH)) + error_at ( + adjust_args_loc, + "an % clause can only be specified if the " + "% selector of the construct selector set appears " + "in the % clause"); + } + } } /* Finalize #pragma omp declare simd or #pragma omp declare variant @@ -26078,7 +26452,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) \ @@ -26086,7 +26459,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) @@ -27001,6 +27378,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 36f88fcd03d..5745d0cca47 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -16305,6 +16305,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_FINALIZE: case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_INDIRECT: + case OMP_CLAUSE_NOVARIANTS: + case OMP_CLAUSE_NOCONTEXT: pc = &OMP_CLAUSE_CHAIN (c); continue; diff --git a/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c b/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c new file mode 100644 index 00000000000..393a44de8e0 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c @@ -0,0 +1,32 @@ +/* Test parsing of OMP clause adjust_args */ +/* { dg-do compile } */ + +int b; + +int f0 (void *a); +int g (void *a); +int f1 (int); + +#pragma omp declare variant (f0) match (construct={target}) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */ +int f2 (void *a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) /* { dg-error "expected 'nothing' or 'need_device_ptr'" } */ +int f3 (int a); +#pragma omp declare variant (f0) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */ +int f4 (void *a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args () /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */ +int f5 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */ +int f6 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) /* { dg-error "expected identifier before '\\)' token" } */ +int f7 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: z) /* { dg-error "'z' undeclared here \\(not in a function\\)" } */ +int f8 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: a) /* { dg-error "'a' is not a C pointer" } */ +int f9 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (nothing: a) /* { dg-error "'a' is specified more than once" } */ +int f10 (int a); +#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: a) /* { dg-error "'a' is specified more than once" } */ +int f11 (void *a); +#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) /* { dg-error "'b' is not a function argument" } */ +int f12 (void *a); + diff --git a/gcc/testsuite/gcc.dg/gomp/dispatch-1.c b/gcc/testsuite/gcc.dg/gomp/dispatch-1.c new file mode 100644 index 00000000000..c8f45c12be6 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/dispatch-1.c @@ -0,0 +1,53 @@ +/* Test parsing of #pragma omp dispatch */ +/* { dg-do compile } */ + +int f0 (int); + +void f1 (void) +{ + int a, b; + double x; + struct {int a; int b;} s; + int arr[1]; + +#pragma omp dispatch + int c = f0 (a); /* { dg-error "expected expression before 'int'" } */ +#pragma omp dispatch + int f2 (int d); /* { dg-error "expected expression before 'int'" } */ +#pragma omp dispatch + a = b; /* { dg-error "expected '\\(' before ';' token" } */ +#pragma omp dispatch + s.a = f0(a) + b; /* { dg-error "expected ';' before '\\+' token" } */ +#pragma omp dispatch + b = !f0(a); /* { dg-error "expected a function name before '!' token" } */ +#pragma omp dispatch + s.b += f0(s.a); /* { dg-error "expected '=' before '\\+=' token" } */ +#pragma omp dispatch +#pragma omp threadprivate(a) /* { dg-error "expected expression before '#pragma'" } */ + a = f0(b); + +#pragma omp dispatch nocontext(s) /* { dg-error "used struct type value where scalar is required" } */ + f0(a); +#pragma omp dispatch nocontext(a, b) /* { dg-error "expected '\\)' before ','" } */ + f0(a); +#pragma omp dispatch nocontext(a) nocontext(b) /* { dg-error "too many 'nocontext' clauses" } */ + f0(a); +#pragma omp dispatch novariants(s) /* { dg-error "used struct type value where scalar is required" } */ + f0(a); +#pragma omp dispatch novariants(a, b) /* { dg-error "expected '\\)' before ','" } */ + f0(a); +#pragma omp dispatch novariants(a) novariants(b) /* { dg-error "too many 'novariants' clauses" } */ + f0(a); +#pragma omp dispatch nowait nowait /* { dg-error "too many 'nowait' clauses" } */ + f0(a); +#pragma omp dispatch device(x) /* { dg-error "expected integer expression before end of line" } */ + f0(a); +#pragma omp dispatch device(arr) /* { dg-error "expected integer expression before end of line" } */ + f0(a); +#pragma omp dispatch is_device_ptr(x) /* { dg-error "'is_device_ptr' variable is neither a pointer nor an array" } */ + f0(a); +#pragma omp dispatch is_device_ptr(&x) /* { dg-error "expected identifier before '&' token" } */ + f0(a); +#pragma omp dispatch depend(inout: f0) /* { dg-error "'f0' is not lvalue expression nor array section in 'depend' clause" } */ + f0(a); +} diff --git a/libgomp/testsuite/libgomp.c/dispatch-1.c b/libgomp/testsuite/libgomp.c/dispatch-1.c new file mode 100644 index 00000000000..0efc075a859 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/dispatch-1.c @@ -0,0 +1,76 @@ +// Adapted from OpenMP examples + +#include +#include +#include + +int baz (double *d_bv, const double *d_av, int n) +{ +#pragma omp distribute parallel for + for (int i = 0; i < n; i++) + d_bv[i] = d_av[i] * i; + return -3; +} + +int bar (double *d_bv, const double *d_av, int n) +{ +#pragma omp target is_device_ptr(d_bv, d_av) + for (int i = 0; i < n; i++) + d_bv[i] = d_av[i] * i; + return -2; +} + +#pragma omp declare variant(bar) match(construct={dispatch}) adjust_args(need_device_ptr: bv, av) +#pragma omp declare variant(baz) match(implementation={vendor(gnu)}) +int foo (double *bv, const double *av, int n) +{ + for (int i = 0; i < n; i++) + bv[i] = av[i] * i; + return -1; +} + +int test (int n) +{ + const double e = 2.71828; + + double *av = (double *) malloc (n * sizeof (*av)); + double *bv = (double *) malloc (n * sizeof (*bv)); + double *d_bv = (double *) malloc (n * sizeof (*d_bv)); + + for (int i = 0; i < n; i++) + { + av[i] = e * i; + bv[i] = 0.0; + d_bv[i] = 0.0; + } + + int f, last_dev = omp_get_num_devices () - 1; +#pragma omp target data map(to: av[:n]) map(from: d_bv[:n]) device(last_dev) if (n == 1024) + { + #pragma omp dispatch nocontext(n > 1024) novariants(n < 1024) device(last_dev) + f = foo (d_bv, av, n); + } + + foo (bv, av, n); + for (int i = 0; i < n; i++) + { + if (d_bv[i] != bv[i]) + { + fprintf (stderr, "ERROR at %d: %lf (act) != %lf (exp)\n", i, d_bv[i], bv[i]); + return 1; + } + } + return f; +} + +int +main (void) +{ + int ret = test(1023); + if (ret != -1) return 1; + ret = test(1024); + if (ret != -2) return 1; + ret = test(1025); + if (ret != -3) return 1; + return 0; +} From patchwork Fri Jul 12 14:11:52 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: 1959900 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=IJymFUV7; 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 4WLDC53VS8z1xpd for ; Sat, 13 Jul 2024 00:15:29 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 814AB3838A1D for ; Fri, 12 Jul 2024 14:15:27 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x333.google.com (mail-wm1-x333.google.com [IPv6:2a00:1450:4864:20::333]) by sourceware.org (Postfix) with ESMTPS id 8B3BF3831E29 for ; Fri, 12 Jul 2024 14:12:59 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 8B3BF3831E29 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 8B3BF3831E29 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::333 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793587; cv=none; b=agaxKEfR/tUqD4arnj+MXlM0CMIMfyFmqqWALSd0tGtOFhzGYF/0xqE3UuVFKlt27Kc97RIVK0bCVOGQLqhjhtup0EGN0674qF6rN8USB5jn8X36oRTwwE+HKwRhiRSB7LoKCW+OhCDmVF71u8n6YM4ZNTUQsYlf+wFOmQ5h70Y= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793587; c=relaxed/simple; bh=BTVt3q6ez6//oq+urcHPeloh/NQJHBqG3b/NWWnaKgo=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=AAiUzo3cYWM7d7lQFA3X0pW3KZNjC2yWpAJIua5FuN2hk2MKLYezR2CWJA4VfwYuWnSz+wkf/leWRI9otPcnfzOjYVhVhEZh5Tzpz33teFYcbduU5rXK+UDuxxUlqALDaauBojtXeE5s7MoHtoa2CYfszq8N+505G68uOrBQaZk= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x333.google.com with SMTP id 5b1f17b1804b1-4266fd395eeso14308115e9.3 for ; Fri, 12 Jul 2024 07:12:59 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1720793577; x=1721398377; 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=iyaMITeLRS7KFk7QMWY5ucfwCfzrKkvtj2S0aDI3zDM=; b=IJymFUV7lA3ifIWdcOtd4Rqagcxv4WVfiLgUHJ4W1lXVxLRHFK6IYT49oVOCGMlCqW dUPTQMmDD0gjneYUQgJH8tKnPZUHkZMAz9DnC+NWLPvzlxbpBzG0T7o+0jpb0tpAM+dx qA3XoKza3WrfSIsTLxohWsaqN6WCPdR31RZlhawYPdYdytffKWytM0i6nMCyErZNBjBq +drIaVDQ19lDf27mQHq4GNG7k++0/fFxPDy55O+mmDj2ZecGtlvesKyruyZW1TXrze6E VfxBDjWldYSpiwdbJ/6gI/T1KejBp+pVpB4Q4wyCnyq2vt7x3IBM+uysa92AxJxK22ay T+8Q== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1720793577; x=1721398377; 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=iyaMITeLRS7KFk7QMWY5ucfwCfzrKkvtj2S0aDI3zDM=; b=cGRBgS4ot1BXG418yWudZtbyxUfLL9+F0CYSqkvGtjl8OWpTYRhbA/KXv/gTSIbd6C bpq1o5qKKRNJS33NX0YA63HBmjIXrlkMzeUywruTwkUejHprqmBRMvSaPEHKQ3PKVUBM nX4pUm8tMx/zSG2GWL9sOgsqGdd502yiGuiAV8tfP5r1xql3ATtSzxhfDaYj5MSavuxO nxDix6iC0tKLjtjg8SpY2N5rMqDXb7PG/ce/YvzbNf7b43TUDF7UPSaUYhMYCzapa0Uo 0tr6T47BVbiY0Et1JxLwH5xKLnPqCOSBZBLE1/s51cp12uSFjv1AL9+YRJvXuhXFrp3m cbdQ== X-Gm-Message-State: AOJu0YzRk8FX3pFhkV44UhXTAQKi65ELh+6p5kSitLzlZNI2gOJiFaJV luUubpnBvhYYg5RAEZKuAZpis7wNU8Ke7E8DcsLY0B+p108w18nB+5KqW4/zbxoVr892nWT917p D X-Google-Smtp-Source: AGHT+IGuTApTfA+Cfvc9znKwbckNNYMp5NqmrOIr3aoUjrSHAF4edEPtzuoB3QlL992gvd5i7CNj0w== X-Received: by 2002:a05:600c:54ca:b0:426:6902:7053 with SMTP id 5b1f17b1804b1-426707cc0d7mr85614275e9.15.1720793576991; Fri, 12 Jul 2024 07:12:56 -0700 (PDT) Received: from localhost.localdomain ([169.155.255.128]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4279f25b946sm24680325e9.19.2024.07.12.07.12.56 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 12 Jul 2024 07:12:56 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v2 5/8] OpenMP: C++ front-end support for dispatch + adjust_args Date: Fri, 12 Jul 2024 16:11:52 +0200 Message-ID: <20240712141155.255186-6-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240712141155.255186-1-parras@baylibre.com> References: <20240712141155.255186-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-13.0 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, KAM_SHORT, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds C++ support for the `dispatch` construct and the `adjust_args` clause. It relies on the c-family bits comprised in the corresponding C front end patch for pragmas and attributes. Additional C/C++ common testcases are provided in a subsequent patch in the series. gcc/cp/ChangeLog: * decl.cc (omp_declare_variant_finalize_one): Set adjust_args need_device_ptr attribute. * parser.cc (cp_parser_direct_declarator): Update call to cp_parser_late_return_type_opt. (cp_parser_late_return_type_opt): Add parameter. Update call to cp_parser_late_parsing_omp_declare_simd. (cp_parser_omp_clause_name): Handle nocontext and novariants clauses. (cp_parser_omp_clause_novariants): New function. (cp_parser_omp_clause_nocontext): Likewise. (cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_NOVARIANTS and PRAGMA_OMP_CLAUSE_NOCONTEXT. (cp_parser_omp_dispatch_body): New function, inspired from cp_parser_assignment_expression and cp_parser_postfix_expression. (OMP_DISPATCH_CLAUSE_MASK): Define. (cp_parser_omp_dispatch): New function. (cp_finish_omp_declare_variant): Add parameter. Handle adjust_args clause. (cp_parser_late_parsing_omp_declare_simd): Add parameter. Update calls to cp_finish_omp_declare_variant and cp_finish_omp_declare_variant. (cp_parser_omp_construct): Handle PRAGMA_OMP_DISPATCH. (cp_parser_pragma): Likewise. * pt.cc (tsubst_attribute): Skip pseudo-TSS need_device_ptr. * semantics.cc (finish_omp_clauses): Handle OMP_CLAUSE_NOCONTEXT and OMP_CLAUSE_NOVARIANTS. gcc/testsuite/ChangeLog: * g++.dg/gomp/adjust-args-1.C: New test. * g++.dg/gomp/adjust-args-2.C: New test. * g++.dg/gomp/dispatch-1.C: New test. * g++.dg/gomp/dispatch-2.C: New test. --- gcc/cp/decl.cc | 33 ++ gcc/cp/parser.cc | 612 ++++++++++++++++++++-- gcc/cp/pt.cc | 3 + gcc/cp/semantics.cc | 20 + gcc/testsuite/g++.dg/gomp/adjust-args-1.C | 39 ++ gcc/testsuite/g++.dg/gomp/adjust-args-2.C | 51 ++ gcc/testsuite/g++.dg/gomp/dispatch-1.C | 53 ++ gcc/testsuite/g++.dg/gomp/dispatch-2.C | 62 +++ 8 files changed, 828 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 edf4c155bf7..fae37d508b0 100644 --- a/gcc/cp/decl.cc +++ b/gcc/cp/decl.cc @@ -8375,6 +8375,39 @@ omp_declare_variant_finalize_one (tree decl, tree attr) if (!omp_context_selector_matches (ctx)) return true; TREE_PURPOSE (TREE_VALUE (attr)) = variant; + + for (tree a = ctx; a != NULL_TREE; a = TREE_CHAIN (a)) + { + if (OMP_TSS_CODE (a) == OMP_TRAIT_SET_NEED_DEVICE_PTR) + { + for (tree need_device_ptr_list = TREE_VALUE (a); + need_device_ptr_list != NULL_TREE; + need_device_ptr_list = TREE_CHAIN (need_device_ptr_list)) + { + tree parm_decl = TREE_VALUE (need_device_ptr_list); + bool found_arg = false; + for (tree arg = DECL_ARGUMENTS (variant); arg != NULL; + arg = TREE_CHAIN (arg)) + if (DECL_NAME (arg) == DECL_NAME (parm_decl)) + { + DECL_ATTRIBUTES (arg) + = tree_cons (get_identifier ( + "omp declare variant adjust_args " + "need_device_ptr"), + NULL_TREE, DECL_ATTRIBUTES (arg)); + found_arg = true; + break; + } + if (!found_arg) + { + error_at (varid_loc, + "variant %qD does not have a parameter %qD", + variant, parm_decl); + return true; + } + } + } + } } } else if (!processing_template_decl) diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 6bf3f52a059..b85c9c387fb 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); @@ -24154,7 +24155,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); @@ -25019,8 +25020,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; @@ -25056,8 +25057,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, @@ -38129,6 +38130,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)) @@ -38137,6 +38140,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)) @@ -40583,6 +40588,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 ) */ @@ -42690,6 +42745,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; @@ -48895,12 +48960,305 @@ cp_parser_omp_assumes (cp_parser *parser, cp_token *pragma_tok) return false; } +/* Parse a function dispatch structured block: + + lvalue-expression = target-call ( [expression-list] ); + or + target-call ( [expression-list] ); + + Inspired from cp_parser_assignment_expression and + cp_parser_postfix_expression. +*/ + +static tree +cp_parser_omp_dispatch_body (cp_parser *parser) +{ + cp_expr expr; + cp_id_kind idk = CP_ID_KIND_NONE; + + /* Parse the binary expressions (lvalue-expression or target-call). */ + expr = cp_parser_binary_expression (parser, false, false, false, + PREC_NOT_OPERATOR, NULL); + if (TREE_CODE (expr) == CALL_EXPR || TREE_CODE (expr) == ERROR_MARK) + return expr; + + /* We have the lvalue, now deal with the assignment. */ + + if (!cp_parser_require (parser, CPP_EQ, RT_EQ)) + return error_mark_node; + + /* Peek at the next token. */ + cp_token *token = cp_lexer_peek_token (parser->lexer); + location_t loc = token->location; + location_t start_loc = get_range_from_loc (line_table, loc).m_start; + + /* Parse function name as primary expression. */ + cp_expr rhs + = cp_parser_primary_expression (parser, false, false, false, false, &idk); + if (TREE_CODE (rhs) == ERROR_MARK) + return rhs; + + /* Keep looping until the postfix-expression is complete. */ + bool parens_found = false; + while (true) + { + if (idk == CP_ID_KIND_UNQUALIFIED && identifier_p (rhs) + && cp_lexer_next_token_is_not (parser->lexer, CPP_OPEN_PAREN)) + /* It is not a Koenig lookup function call. */ + rhs = unqualified_name_lookup_error (rhs); + + /* Peek at the next token. */ + token = cp_lexer_peek_token (parser->lexer); + + switch (token->type) + { + case CPP_OPEN_PAREN: + /* postfix-expression ( expression-list [opt] ) */ + { + if (parens_found) + { + cp_parser_error ( + parser, + "only one function call is allowed in a dispatch construct"); + return error_mark_node; + } + parens_found = true; + + bool koenig_p; + tsubst_flags_t complain = complain_flags (false); + vec *args; + location_t close_paren_loc = UNKNOWN_LOCATION; + location_t combined_loc = UNKNOWN_LOCATION; + + args = (cp_parser_parenthesized_expression_list ( + parser, non_attr, + /*cast_p=*/false, /*allow_expansion_p=*/true, + /*non_constant_p=*/NULL, + /*close_paren_loc=*/&close_paren_loc, + /*wrap_locations_p=*/true)); + + if (args == NULL) + { + rhs = error_mark_node; + break; + } + + koenig_p = false; + if (idk == CP_ID_KIND_UNQUALIFIED || idk == CP_ID_KIND_TEMPLATE_ID) + { + if (identifier_p (rhs) + /* In C++20, we may need to perform ADL for a template + name. */ + || (TREE_CODE (rhs) == TEMPLATE_ID_EXPR + && identifier_p (TREE_OPERAND (rhs, 0)))) + { + if (!args->is_empty ()) + { + koenig_p = true; + if (!any_type_dependent_arguments_p (args)) + rhs = perform_koenig_lookup (rhs, args, complain); + } + else + rhs = unqualified_fn_lookup_error (rhs); + } + /* We do not perform argument-dependent lookup if + normal lookup finds a non-function, in accordance + with the expected resolution of DR 218. */ + else if (!args->is_empty () && is_overloaded_fn (rhs)) + { + /* Do not do argument dependent lookup if regular + lookup finds a member function or a block-scope + function declaration. [basic.lookup.argdep]/3 */ + bool do_adl_p = true; + tree fns = get_fns (rhs); + for (lkp_iterator iter (fns); iter; ++iter) + { + tree fn = STRIP_TEMPLATE (*iter); + if ((TREE_CODE (fn) == USING_DECL + && DECL_DEPENDENT_P (fn)) + || DECL_FUNCTION_MEMBER_P (fn) + || DECL_LOCAL_DECL_P (fn)) + { + do_adl_p = false; + break; + } + } + + if (do_adl_p) + { + koenig_p = true; + if (!any_type_dependent_arguments_p (args)) + rhs = perform_koenig_lookup (rhs, args, complain); + } + } + } + + /* Temporarily set input_location to the combined location + with call expression range, as e.g. build_out_target_exprs + called from convert_default_arg relies on input_location, + so updating it only when the call is fully built results + in inconsistencies between location handling in templates + and outside of templates. */ + if (close_paren_loc != UNKNOWN_LOCATION) + combined_loc + = make_location (token->location, start_loc, close_paren_loc); + iloc_sentinel ils (combined_loc); + + if (TREE_CODE (rhs) == COMPONENT_REF) + { + tree instance = TREE_OPERAND (rhs, 0); + tree fn = TREE_OPERAND (rhs, 1); + + if (processing_template_decl + && (type_dependent_object_expression_p (instance) + || (!BASELINK_P (fn) && TREE_CODE (fn) != FIELD_DECL) + || type_dependent_expression_p (fn) + || any_type_dependent_arguments_p (args))) + { + maybe_generic_this_capture (instance, fn); + rhs = build_min_nt_call_vec (rhs, args); + } + else if (BASELINK_P (fn)) + { + rhs + = (build_new_method_call (instance, fn, &args, NULL_TREE, + (idk == CP_ID_KIND_QUALIFIED + ? LOOKUP_NORMAL + | LOOKUP_NONVIRTUAL + : LOOKUP_NORMAL), + /*fn_p=*/NULL, complain)); + } + else + rhs = finish_call_expr (rhs, &args, + /*disallow_virtual=*/false, + /*koenig_p=*/false, complain); + } + else if (TREE_CODE (rhs) == OFFSET_REF + || TREE_CODE (rhs) == MEMBER_REF + || TREE_CODE (rhs) == DOTSTAR_EXPR) + rhs = (build_offset_ref_call_from_tree (rhs, &args, complain)); + else if (idk == CP_ID_KIND_QUALIFIED) + /* A call to a static class member, or a namespace-scope + function. */ + rhs = finish_call_expr (rhs, &args, + /*disallow_virtual=*/true, koenig_p, + complain); + else + /* All other function calls. */ + { + if (DECL_P (rhs) && parser->omp_for_parse_state + && parser->omp_for_parse_state->in_intervening_code + && omp_runtime_api_call (rhs)) + { + error_at (loc, "calls to the OpenMP runtime API are " + "not permitted in intervening code"); + parser->omp_for_parse_state->fail = true; + } + rhs = finish_call_expr (rhs, &args, + /*disallow_virtual=*/false, koenig_p, + complain); + } + if (close_paren_loc != UNKNOWN_LOCATION) + rhs.set_location (combined_loc); + + /* The expr is certainly no longer an id. */ + idk = CP_ID_KIND_NONE; + + release_tree_vector (args); + } + break; + + case CPP_DOT: + case CPP_DEREF: + /* postfix-expression . template [opt] id-expression + postfix-expression . pseudo-destructor-name + postfix-expression -> template [opt] id-expression + postfix-expression -> pseudo-destructor-name */ + + /* Consume the `.' or `->' operator. */ + cp_lexer_consume_token (parser->lexer); + + rhs = cp_parser_postfix_dot_deref_expression (parser, token->type, + rhs, false, &idk, loc); + + break; + + default: + goto finish; + } + } +finish: + if (!parens_found) + { + cp_parser_error (parser, "expected %<(%>"); + return error_mark_node; + } + + /* Build the assignment expression. Its default + location: + LHS = RHS + ~~~~^~~~~ + is the location of the '=' token as the + caret, ranging from the start of the lhs to the + end of the rhs. */ + loc = make_location (loc, expr.get_start (), rhs.get_finish ()); + expr + = cp_build_modify_expr (loc, expr, NOP_EXPR, rhs, complain_flags (false)); + + return expr; +} + +/* OpenMP 5.1: + # pragma omp dispatch dispatch-clause[optseq] new-line + expression-stmt + + LOC is the location of the #pragma. +*/ + +#define OMP_DISPATCH_CLAUSE_MASK \ + ((OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT)) + +static tree +cp_parser_omp_dispatch (cp_parser *parser, cp_token *pragma_tok) +{ + location_t loc = cp_lexer_peek_token (parser->lexer)->location; + tree stmt = make_node (OMP_DISPATCH); + SET_EXPR_LOCATION (stmt, loc); + TREE_TYPE (stmt) = void_type_node; + + OMP_DISPATCH_CLAUSES (stmt) + = cp_parser_omp_all_clauses (parser, OMP_DISPATCH_CLAUSE_MASK, + "#pragma omp dispatch", pragma_tok); + + // Parse expression statement + loc = cp_lexer_peek_token (parser->lexer)->location; + tree dispatch_body = cp_parser_omp_dispatch_body (parser); + if (dispatch_body == error_mark_node) + { + inform (loc, + "%<#pragma omp dispatch%> must be followed by a direct function " + "call with optional assignment"); + cp_parser_skip_to_end_of_block_or_statement (parser); + return NULL_TREE; + } + + cp_parser_consume_semicolon_at_end_of_statement (parser); + OMP_DISPATCH_BODY (stmt) = dispatch_body; + + return add_stmt (stmt); +} + /* Finalize #pragma omp declare variant after a fndecl has been parsed, and put that into "omp declare variant base" attribute. */ static tree cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok, - tree attrs) + tree attrs, tree parms) { matching_parens parens; if (!parens.require_open (parser)) @@ -48958,44 +49316,197 @@ cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok, location_t finish_loc = get_finish (varid.get_location ()); location_t varid_loc = make_location (caret_loc, start_loc, finish_loc); - if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA) - && cp_lexer_nth_token_is (parser->lexer, 2, CPP_NAME)) - cp_lexer_consume_token (parser->lexer); + vec adjust_args_list = vNULL; + bool has_match = false, has_adjust_args = false; + location_t adjust_args_loc = UNKNOWN_LOCATION; + tree need_device_ptr_list = NULL_TREE, *need_device_ptr_chain_p = NULL; - const char *clause = ""; - location_t match_loc = cp_lexer_peek_token (parser->lexer)->location; - if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) - clause = IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value); - if (strcmp (clause, "match")) + do { - cp_parser_error (parser, "expected %"); - goto fail; + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA) + && cp_lexer_nth_token_is (parser->lexer, 2, CPP_NAME)) + cp_lexer_consume_token (parser->lexer); + + const char *clause = ""; + location_t match_loc = cp_lexer_peek_token (parser->lexer)->location; + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + clause + = IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value); + + enum clause + { + match, + adjust_args + } ccode; + + if (strcmp (clause, "match") == 0) + ccode = match; + else if (strcmp (clause, "adjust_args") == 0) + { + ccode = adjust_args; + adjust_args_loc = match_loc; + } + else + { + cp_parser_error (parser, "expected % or %"); + goto fail; + } + + cp_lexer_consume_token (parser->lexer); + + if (!parens.require_open (parser)) + goto fail; + + if (ccode == match) + { + has_match = true; + tree ctx + = cp_parser_omp_context_selector_specification (parser, true); + if (ctx == error_mark_node) + goto fail; + ctx = omp_check_context_selector (match_loc, ctx); + if (ctx != error_mark_node && variant != error_mark_node) + { + tree match_loc_node + = maybe_wrap_with_location (integer_zero_node, match_loc); + tree loc_node + = maybe_wrap_with_location (integer_zero_node, varid_loc); + loc_node + = tree_cons (match_loc_node, + build_int_cst (integer_type_node, idk), + build_tree_list (loc_node, integer_zero_node)); + attrs = tree_cons (get_identifier ("omp declare variant base"), + tree_cons (variant, ctx, loc_node), attrs); + if (processing_template_decl) + ATTR_IS_DEPENDENT (attrs) = 1; + } + if (!parens.require_close (parser)) + goto fail; + } + else if (ccode == adjust_args) + { + has_adjust_args = true; + cp_token *adjust_op_tok = cp_lexer_peek_token (parser->lexer); + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME) + && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON)) + { + const char *p = IDENTIFIER_POINTER (adjust_op_tok->u.value); + if (strcmp (p, "need_device_ptr") == 0 + || strcmp (p, "nothing") == 0) + { + cp_lexer_consume_token (parser->lexer); // need_device_ptr + cp_lexer_consume_token (parser->lexer); // : + location_t arg_loc + = cp_lexer_peek_token (parser->lexer)->location; + + tree arg; + tree list + = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_ERROR, + NULL_TREE, NULL); + + for (tree c = list; c != NULL_TREE; c = TREE_CHAIN (c)) + { + tree decl = TREE_PURPOSE (c); + int idx; + for (arg = parms, idx = 0; arg != NULL; + arg = TREE_CHAIN (arg), idx++) + if (TREE_VALUE (arg) == decl) + break; + if (arg == NULL_TREE) + { + error_at (arg_loc, "%qD is not a function argument", + decl); + continue; + } + arg = TREE_VALUE (arg); + if (adjust_args_list.contains (arg)) + { + error_at (arg_loc, "%qD is specified more than once", + decl); + continue; + } + if (strcmp (p, "need_device_ptr") == 0) + { + bool is_ptr_or_template + = TEMPLATE_PARM_P (TREE_TYPE (arg)) + || POINTER_TYPE_P (TREE_TYPE (arg)); + if (!is_ptr_or_template) + { + error_at (arg_loc, "%qD is not a C pointer", + decl); + continue; + } + } + adjust_args_list.safe_push (arg); + if (strcmp (p, "need_device_ptr") == 0) + { + tree attr = tree_cons (NULL_TREE, TREE_PURPOSE (c), + NULL_TREE); + if (need_device_ptr_list == NULL_TREE) + { + gcc_assert (need_device_ptr_chain_p == NULL); + need_device_ptr_list = attr; + } + else + *need_device_ptr_chain_p = attr; + need_device_ptr_chain_p = &TREE_CHAIN (attr); + } + } + } + else + { + error_at (adjust_op_tok->location, + "expected % or %"); + goto fail; + } + } + else + { + error_at (adjust_op_tok->location, + "expected % or % followed " + "by %<:%>"); + goto fail; + } + } + } while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)); + + if (has_adjust_args) + { + if (!has_match) + { + error_at ( + adjust_args_loc, + "an % clause can only be specified if the " + "% selector of the construct selector set appears " + "in the % clause"); + } + else + { + tree ctx = TREE_VALUE (TREE_VALUE (attrs)); + if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT, + OMP_TRAIT_CONSTRUCT_DISPATCH)) + error_at ( + adjust_args_loc, + "an % clause can only be specified if the " + "% selector of the construct selector set appears " + "in the % clause"); + } } - cp_lexer_consume_token (parser->lexer); - - if (!parens.require_open (parser)) - goto fail; - - tree ctx = cp_parser_omp_context_selector_specification (parser, true); - if (ctx == error_mark_node) - goto fail; - ctx = omp_check_context_selector (match_loc, ctx); - if (ctx != error_mark_node && variant != error_mark_node) + if (need_device_ptr_list) { - tree match_loc_node = maybe_wrap_with_location (integer_zero_node, - match_loc); - tree loc_node = maybe_wrap_with_location (integer_zero_node, varid_loc); - loc_node = tree_cons (match_loc_node, - build_int_cst (integer_type_node, idk), - build_tree_list (loc_node, integer_zero_node)); - attrs = tree_cons (get_identifier ("omp declare variant base"), - tree_cons (variant, ctx, loc_node), attrs); - if (processing_template_decl) - ATTR_IS_DEPENDENT (attrs) = 1; + // We might not have DECL_ARGUMENTS for the variant yet. So we store the + // need_device_ptr list in the base function attribute beside the context + // selector. + gcc_assert (TREE_PURPOSE (attrs) + == get_identifier ("omp declare variant base")); + gcc_assert (TREE_PURPOSE (TREE_VALUE (attrs)) == variant); + TREE_VALUE (TREE_VALUE (attrs)) + = make_trait_set_selector (OMP_TRAIT_SET_NEED_DEVICE_PTR, + need_device_ptr_list, + TREE_VALUE (TREE_VALUE (attrs))); } - parens.require_close (parser); cp_parser_skip_to_pragma_eol (parser, pragma_tok); return attrs; } @@ -49005,7 +49516,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; @@ -49049,7 +49561,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); } @@ -49180,9 +49692,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); @@ -50032,7 +50543,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, @@ -50893,6 +51408,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 (); } @@ -51589,6 +52107,10 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p) "%<#pragma omp sections%> construct"); break; + case PRAGMA_OMP_DISPATCH: + cp_parser_omp_dispatch (parser, pragma_tok); + return true; + case PRAGMA_IVDEP: case PRAGMA_UNROLL: case PRAGMA_NOVECTOR: diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index e38e02488be..af6b112d463 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -11987,6 +11987,9 @@ tsubst_attribute (tree t, tree *decl_p, tree args, for (tree tss = ctx; tss; tss = TREE_CHAIN (tss)) { enum omp_tss_code set = OMP_TSS_CODE (tss); + if (set == OMP_TRAIT_SET_NEED_DEVICE_PTR) + continue; + tree selectors = NULL_TREE; for (tree ts = OMP_TSS_TRAIT_SELECTORS (tss); ts; ts = TREE_CHAIN (ts)) diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index cd3df13772d..76a5d3b23f3 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -7631,6 +7631,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 Fri Jul 12 14:11:53 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: 1959896 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=jkjxNvBB; 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 4WLD9L0XJtz1xqj for ; Sat, 13 Jul 2024 00:13:58 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id B31F43831E31 for ; Fri, 12 Jul 2024 14:13:55 +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 AA92838323F1 for ; Fri, 12 Jul 2024 14:12:59 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org AA92838323F1 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 AA92838323F1 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=1720793583; cv=none; b=pOVTr8D1X1X7gDzU/wmuCU+P+gthDfwR0Qi9Lh4MBU5yncbChyD+6dKM91JAgYeb5IvKTvbZaTqLH7AM9kaNLBESGGsJi6xlqzDK0ZmyLL+0yorQ4UFGqdM9mqNuWIRL8DoT15RW20ALJ1mOBnZa7grekduoC6rAkwgDujUy6e8= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793583; c=relaxed/simple; bh=EsCyckCXJYg3rv/fPwj14oU4hjj6S0aU41jhGKIfxB0=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=LY80MCBWwnxMldRo9S8Z64WSTw0Fl1u6+oZ0pphDKtkTydPncFBiXHbVa2E6ZaKg73lY8XuXTgu+F/KjstuzLv27C+vGX41EaYanZJS6dGgDMMK94uD/GS9e5pN1HW3xLU+pETplHUQBMARw057G7X/qB79hKmBDeIftYQDlcPM= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x32b.google.com with SMTP id 5b1f17b1804b1-42797289c8bso15183555e9.0 for ; Fri, 12 Jul 2024 07:12:59 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1720793578; x=1721398378; 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=VI7TNZ4Ew9jutPGIhjqw0PfdJ9f2t2Qjw0EJ5LxDBxA=; b=jkjxNvBBGs5MaYw565WA/3732qoZw/hMde/4qEw8pqgnvkVXerqPQ+uyxUQsm3Ml7t ip6Ck/B8kl0+5Plu3qtmT072V53vkMqPeOoi6ruYq7GdB4UC1jNhyfRHF2i8vPaF4fNo Q+ks19bxSaZrrEeZFKpPLK4UzQ/EROd6XhspIfHn3bVPa4gC0TpeOabUmOntggS0RfKU aHjg7YXollbWACs0YjurJpORA42qcoxTNNxKcUyYb9N7cr8rXhBaEIGWVp13X4eLycgi ZPW1NkJRjdhd2djTwAKJaG9Iotx3cy8lfXPxQp+tOawNhOXo6Tj4KFWUo1cHfPdj5qFA DH6A== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1720793578; x=1721398378; 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=VI7TNZ4Ew9jutPGIhjqw0PfdJ9f2t2Qjw0EJ5LxDBxA=; b=esO/84nIeA3lKNvlh7Ai4WB1zppQUqKQGRSwmfdG7PFCH2eGPgsVig0lSXHIYVUOuq Nfr9z72fMsJwdPYlUXXb7EtXFRcwsbKgvXVnBMaJ/qY+KPPuRBxgmg8YItggewNlNOfy LOcbEdBEFwXxQ12MxmglmBjQOhywO235zHB1wXhlNDUdlzQDepI2o+yR0Frq9/zCZC1a bdf9dCZ7CnQXX5+D2t50VBHSCFAtxPRtPP5FW4qeY2gNL4NRFI3TFopT0V/c3/35IneR TeUc5hLuekRS6W3TSlYIDPdT/KJa63gsU8kTs+nZCwkttrYZkbRhukJ2kAHQcg+AWLI+ XnTg== X-Gm-Message-State: AOJu0Yxmj3GmMp7t+Vu24Z1ryenjSCMCzWeZTu8rCQx9kRoReAERSzSL CtnGctgZ+AsPamcDN01Nm9Dk048KwXns2L53/6SDBDxoEbZXveR+FvgfIRmkkXDF9nngMaZTmcV k X-Google-Smtp-Source: AGHT+IGkWzwjt6nM7z1jQkh4qN3L174sY6aBsSl5fnAugGHoeMUjgix4wLYDw7cjZWfD32ETY7GX3g== X-Received: by 2002:a05:600c:3644:b0:426:5cef:ee41 with SMTP id 5b1f17b1804b1-426708f1928mr77042075e9.38.1720793577913; Fri, 12 Jul 2024 07:12:57 -0700 (PDT) Received: from localhost.localdomain ([169.155.255.128]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4279f25b946sm24680325e9.19.2024.07.12.07.12.57 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 12 Jul 2024 07:12:57 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v2 6/8] OpenMP: common C/C++ testcases for dispatch + adjust_args Date: Fri, 12 Jul 2024 16:11:53 +0200 Message-ID: <20240712141155.255186-7-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240712141155.255186-1-parras@baylibre.com> References: <20240712141155.255186-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.5 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org gcc/testsuite/ChangeLog: * c-c++-common/gomp/declare-variant-2.c: Adjust dg-error directives. * c-c++-common/gomp/adjust-args-1.c: New test. * c-c++-common/gomp/adjust-args-2.c: New test. * c-c++-common/gomp/dispatch-1.c: New test. * c-c++-common/gomp/dispatch-2.c: New test. * c-c++-common/gomp/dispatch-3.c: New test. * c-c++-common/gomp/dispatch-4.c: New test. * c-c++-common/gomp/dispatch-5.c: New test. * c-c++-common/gomp/dispatch-6.c: New test. * c-c++-common/gomp/dispatch-7.c: New test. --- .../c-c++-common/gomp/adjust-args-1.c | 30 +++++++++ .../c-c++-common/gomp/adjust-args-2.c | 31 +++++++++ .../c-c++-common/gomp/declare-variant-2.c | 4 +- gcc/testsuite/c-c++-common/gomp/dispatch-1.c | 65 +++++++++++++++++++ gcc/testsuite/c-c++-common/gomp/dispatch-2.c | 28 ++++++++ gcc/testsuite/c-c++-common/gomp/dispatch-3.c | 15 +++++ gcc/testsuite/c-c++-common/gomp/dispatch-4.c | 18 +++++ gcc/testsuite/c-c++-common/gomp/dispatch-5.c | 26 ++++++++ gcc/testsuite/c-c++-common/gomp/dispatch-6.c | 19 ++++++ gcc/testsuite/c-c++-common/gomp/dispatch-7.c | 28 ++++++++ 10 files changed, 262 insertions(+), 2 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/adjust-args-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/adjust-args-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-3.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-4.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-5.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-6.c create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-7.c diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c b/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c new file mode 100644 index 00000000000..728abe62092 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c @@ -0,0 +1,30 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +int f (int a, void *b, float c[2]); + +#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b, c) +int f0 (int a, void *b, float c[2]); +#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c) +int f1 (int a, void *b, float c[2]); + +int test () { + int a; + void *b; + float c[2]; + struct {int a;} s; + + s.a = f0 (a, b, c); + #pragma omp dispatch + s.a = f0 (a, b, c); + + f1 (a, b, c); + #pragma omp dispatch + s.a = f1 (a, b, c); + + return s.a; +} + +/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, D\.\[0-9]+\\);" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, D\.\[0-9]+\\);" 2 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c b/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c new file mode 100644 index 00000000000..e36d93a01d9 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c @@ -0,0 +1,31 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +int f (int a, void *b, float c[2]); + +#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b, c) +int f0 (int a, void *b, float c[2]); +#pragma omp declare variant (f) adjust_args (need_device_ptr: b, c) match (construct={dispatch}) adjust_args (nothing: a) +int f1 (int a, void *b, float c[2]); + +void test () { + int a; + void *b; + float c[2]; + + #pragma omp dispatch + f0 (a, b, c); + + #pragma omp dispatch device (-4852) + f0 (a, b, c); + + #pragma omp dispatch device (a + a) + f0 (a, b, c); +} + +/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, D\.\[0-9]+\\);" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, D\.\[0-9]+\\);" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, -4852\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, -4852\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp dispatch device\\(-4852\\)" 1 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c b/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c index 05e485ef6a8..50d9b2dcf4b 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c @@ -8,9 +8,9 @@ void f3 (void); void f4 (void); #pragma omp declare variant match(user={condition(0)}) /* { dg-error "expected '\\(' before 'match'" } */ void f5 (void); -#pragma omp declare variant (f1) /* { dg-error "expected 'match' before end of line" } */ +#pragma omp declare variant (f1) /* { dg-error "expected 'match' or 'adjust_args' before end of line" } */ void f6 (void); -#pragma omp declare variant (f1) simd /* { dg-error "expected 'match' before 'simd'" } */ +#pragma omp declare variant (f1) simd /* { dg-error "expected 'match' or 'adjust_args' before 'simd'" } */ void f7 (void); #pragma omp declare variant (f1) match /* { dg-error "expected '\\(' before end of line" } */ void f8 (void); diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-1.c b/gcc/testsuite/c-c++-common/gomp/dispatch-1.c new file mode 100644 index 00000000000..e77b2f3ecf6 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-1.c @@ -0,0 +1,65 @@ +#include + +int f0 (int, long, double); +void f2 (void); +int f3 (void); +void (*f4) (void); + +void f1 (void) +{ + int a, c; + long b; + double x; + struct {int a; float b; short c;} s, *sp; + int arr[3]; + +#pragma omp dispatch + c = f0 (a, b, x); +#pragma omp dispatch + x = f0 (a * 4, 2 - b, x * x); +#pragma omp dispatch + s.a = f0 (a, sp->c, x); +#pragma omp dispatch + sp->c = f0 (s.a - 2, b / 3, x * 5); +#pragma omp dispatch + arr[0] = f0 (arr[1], !b, arr[2]); +#pragma omp dispatch + (*sp).c = f0 (s.a, b, x); +#pragma omp dispatch + sp->b = f0 (s.a++, b % 4, --x); +#pragma omp dispatch + f0 (f3(), b, s.b); +#pragma omp dispatch + f2 (); +#pragma omp dispatch + f4 (); + +#pragma omp dispatch nocontext(sp->a * x + arr[2]) + f2 (); +#pragma omp dispatch nocontext(arr - (intptr_t)(x / s.b)) + f2 (); +#pragma omp dispatch nocontext(x == s.c || b != c) + f2 (); +#pragma omp dispatch novariants(b << sp->c) + f2 (); +#pragma omp dispatch novariants(!arr | s.a) + f2 (); +#pragma omp dispatch novariants(s.c ? f3() : a & c) + f2 (); +#pragma omp dispatch nowait + f2 (); +#pragma omp dispatch device(-25373654) + f2 (); +#pragma omp dispatch device(b * (int)(x - sp->b)) + f2 (); +#pragma omp dispatch is_device_ptr(arr) + f2 (); +#pragma omp dispatch is_device_ptr(sp) + f2 (); +#pragma omp dispatch depend(inout: sp) + f2 (); +#pragma omp dispatch depend(inoutset: arr[:2]) + f2 (); +#pragma omp dispatch depend(mutexinoutset: arr) + f2 (); +} diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-2.c b/gcc/testsuite/c-c++-common/gomp/dispatch-2.c new file mode 100644 index 00000000000..24ab9545b73 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-2.c @@ -0,0 +1,28 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +int f0 (void); +int f1 (void); +#pragma omp declare variant (f0) match (construct={dispatch}) +#pragma omp declare variant (f1) match (implementation={vendor(gnu)}) +int f2 (void); + +int test (void) +{ + int a; +#pragma omp dispatch + a = f2 (); +#pragma omp dispatch novariants(1) + a = f2 (); +#pragma omp dispatch novariants(0) + a = f2 (); +#pragma omp dispatch nocontext(1) + a = f2 (); +#pragma omp dispatch nocontext(0) + a = f2 (); + return a; +} + +/* { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 3 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-3.c b/gcc/testsuite/c-c++-common/gomp/dispatch-3.c new file mode 100644 index 00000000000..319f73a84d8 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-3.c @@ -0,0 +1,15 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void f2 (void); + +void test (void) +{ +#pragma omp dispatch /* { dg-final { scan-tree-dump-times "#pragma omp task if\\(0\\)" 1 "gimple" } } */ + f2 (); +#pragma omp dispatch nowait /* { dg-final { scan-tree-dump-times "#pragma omp task if\\(1\\)" 1 "gimple" } } */ + f2 (); +} + + + diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-4.c b/gcc/testsuite/c-c++-common/gomp/dispatch-4.c new file mode 100644 index 00000000000..aeb5c00507f --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-4.c @@ -0,0 +1,18 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void f2 (int a); + +void test (void) +{ + int a; + +#pragma omp dispatch device(-25373654) +/* { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device \\(-25373654\\);" 1 "gimple" } } */ + f2 (a); +#pragma omp dispatch device(a + a) +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = a \\* 2;.*#pragma omp dispatch device\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\).*#pragma omp task shared\\(D\.\[0-9]+\\).*__builtin_omp_set_default_device \\(D\.\[0-9]+\\);" 1 "gimple" } } */ + f2 (a); +} + + diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-5.c b/gcc/testsuite/c-c++-common/gomp/dispatch-5.c new file mode 100644 index 00000000000..78a37bac59b --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-5.c @@ -0,0 +1,26 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void f1 (void* p, int arr[]); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: p, arr) +void f2 (void* p, int arr[]); + +void test (void) +{ + void *p; + int arr[2]; + +#pragma omp dispatch + f2 (p, arr); +#pragma omp dispatch is_device_ptr(p) +/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p\\) shared\\(arr\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*int D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&arr, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(p, D\.\[0-9]+\\);" 1 "gimple" } } */ + f2 (p, arr); +#pragma omp dispatch is_device_ptr(arr) +/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(arr\\) shared\\(p\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*int D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(p, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(D\.\[0-9]+, &arr\\);" 1 "gimple" } } */ + f2 (p, arr); +#pragma omp dispatch is_device_ptr(p, arr) +/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(arr\\) shared\\(p\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*f1 \\(p, &arr\\);" 1 "gimple" } } */ + f2 (p, arr); +} + + diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-6.c b/gcc/testsuite/c-c++-common/gomp/dispatch-6.c new file mode 100644 index 00000000000..7c495179a13 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-6.c @@ -0,0 +1,19 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-ompexp" } */ + +void f2 (void* p); + +void test (void) +{ + void *p; + +#pragma omp dispatch +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, 0B, .*, .*\\);" 1 "ompexp" } } */ + f2 (p); +#pragma omp dispatch depend(inout: p) +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+\\\[2] = &p;" 1 "ompexp" } } */ +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, &D\.\[0-9]+, .*, .*\\);" 1 "ompexp" } } */ + f2 (p); +} + + diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-7.c b/gcc/testsuite/c-c++-common/gomp/dispatch-7.c new file mode 100644 index 00000000000..8cc4526fea0 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/dispatch-7.c @@ -0,0 +1,28 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-omplower" } */ + +int f0 (void); +int f1 (void); +#pragma omp declare variant (f0) match (construct={dispatch}) +#pragma omp declare variant (f1) match (implementation={vendor(gnu)}) +int f2 (void); + +int test (void) +{ + int a, n; +#pragma omp dispatch novariants(n < 1024) nocontext(n > 1024) + a = f2 (); + return a; +} + +/* { dg-final { scan-tree-dump-times "#pragma omp dispatch nocontext\\(0\\) novariants\\(0\\) shared\\(n\\) shared\\(a\\)" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(n\\) shared\\(a\\)" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 1 "gimple" } } */ + +/* { dg-final { scan-tree-dump-times ".omp_data_o.1.n = n;" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times ".omp_data_o.1.a = &a;" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "n = .omp_data_i->n;" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = .omp_data_i->a;" 3 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "\\*D\.\[0-9]+ = D\.\[0-9]+;" 3 "omplower" } } */ From patchwork Fri Jul 12 14:11:54 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: 1959902 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=dQZWgozi; 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 4WLDDx2xJMz1xpd for ; Sat, 13 Jul 2024 00:17:05 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 92AD138323D4 for ; Fri, 12 Jul 2024 14:17:03 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lf1-x12a.google.com (mail-lf1-x12a.google.com [IPv6:2a00:1450:4864:20::12a]) by sourceware.org (Postfix) with ESMTPS id 25F8E38323FC for ; Fri, 12 Jul 2024 14:13:03 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 25F8E38323FC 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 25F8E38323FC Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::12a ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793590; cv=none; b=Qq2B5ceIy+pTBLz7a+BWsNa06lmVRJ54DG8gk/ZDEiD0aKuBgGZnQcLZ/3QwBEo0+Pk02BvSSbHYJwjQR3FibDF6X1ULksjV0TPyaSmQyzl3aBob1oOTJQOe0DhtV94BBnRmdROVmldwkn48v/pBkyXp/9no0NSB1zouWutkDF8= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793590; c=relaxed/simple; bh=aOvEudm6YX9WCMXe/ke9iMQx8CKzuSGsxifiONyKseo=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=DxTwu4rFSBpLMXgS32pzs4vyc1Fqu1osVU5pg7PVkdfI1jQSM5SK/bgV/dQl+WP3OPYSHcIe+aoVn48p3lJn41ocoN6QpU5At4Nw6+wD1KlSo0oixjaiy9PzmKg7XX2Fbva8b2vnCRpNedgA8OV2b569rd2s+MyTaJqk8U8/LBw= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lf1-x12a.google.com with SMTP id 2adb3069b0e04-52e98087e32so2393351e87.2 for ; Fri, 12 Jul 2024 07:13:03 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1720793581; x=1721398381; 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=AVJ0DpRWNKmAMgg6avEq7cyjdmU7ivqtDZ8Q8AqmB1s=; b=dQZWgoziUHNU9w2vDbPGSi2Jyhq2tKiDXb2IWCvy4spNRfP9DeORFaN7L6NuLXlUMB wZq8FVkWlwu8aRT0DApPPFSCRoHKnHLc6M9uGSsTVbeJBX4iLn2f6pSba+aG/27H6oma AMOof5GYd6DUKyNxbmRXMu6Ct4CkAk/fArKEZV06XM8jLfYoCC+2f3VUUnXl3/fIMx1h aGZF5vFqNuaIA1kj+DVhJFYCU9YS7KktqmifLs0aqZAuSPVtVRzyNoxHfhC7Tt+i6qJO 4gz418KXQuy50NVihYfiVzRxDZDib4iLphyICMsGSfmuSbpADkQKurXlinMLJAFdDdwL aJ3Q== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1720793581; x=1721398381; 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=AVJ0DpRWNKmAMgg6avEq7cyjdmU7ivqtDZ8Q8AqmB1s=; b=se1++dMvHGIyTWTPTn2bZ/nO7J3hdqb3jtwftezaExbd6+Kdnj+A2LJtXjoUyiSdIU FmMZ4rW2qgbrTEtQFGjJPHFccS0m3vXmuFwng4bTrl8/3VFbFOaZPsUW1DHAsFqghH9W Ehf9+rqmzvcOXXYHOnWQ98u5Yrj2Yvob+zlQy3mhYStmoM/HcU8HlaMjbEzkL5fEp3uc O1ImOe552XMGkEGi1b+MyA+Rl4gjLEzdM5Tokd9gVHsrLm3h/zWh2/5M2HkBQZdu47dE 0qxU7NnH8CBjvo//+vjID0NCvPbY7gulfDPBgzDzWioYcBBM2err4wCZKuWT+O8uhoF8 1jow== X-Gm-Message-State: AOJu0Yw9c/onkMi3X2lcdH/p9DFS23Iw8Y5NgwtQHkndKB9D2XcqfAHS m15wX6dcFj4/NCR17NYVH1g3B+QRbOd5YgcBByFyDHQ+AHYbCcdQ71eBD/Ok+vwu5EMJQSvkkIu f X-Google-Smtp-Source: AGHT+IEd1CI6FWvCOZ4pamvW02IPinQb4L6mAfe7XoHrcuAsRGZ5PePvhj9K3FlHCB3vbO5Z6//6rA== X-Received: by 2002:a05:651c:198b:b0:2ec:40cf:fa9 with SMTP id 38308e7fff4ca-2eeb3103c4dmr107205831fa.29.1720793579039; Fri, 12 Jul 2024 07:12:59 -0700 (PDT) Received: from localhost.localdomain ([169.155.255.128]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4279f25b946sm24680325e9.19.2024.07.12.07.12.58 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 12 Jul 2024 07:12:58 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v2 7/8] OpenMP: Fortran front-end support for dispatch + adjust_args Date: Fri, 12 Jul 2024 16:11:54 +0200 Message-ID: <20240712141155.255186-8-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240712141155.255186-1-parras@baylibre.com> References: <20240712141155.255186-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.2 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, KAM_SHORT, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP, URIBL_BLACK autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds support for the `dispatch` construct and the `adjust_args` clause to the Fortran front-end. 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:. * types.def (BT_FN_PTR_CONST_PTR_INT): Declare. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/declare-variant-2.f90: Update dg-error. * gfortran.dg/gomp/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 | 201 ++++++++++++++++-- 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 | 161 ++++++++++++++ gcc/fortran/trans.cc | 1 + gcc/fortran/types.def | 1 + .../gfortran.dg/gomp/adjust-args-1.f90 | 63 ++++++ .../gfortran.dg/gomp/adjust-args-2.f90 | 18 ++ .../gfortran.dg/gomp/adjust-args-3.f90 | 26 +++ .../gfortran.dg/gomp/adjust-args-4.f90 | 58 +++++ .../gfortran.dg/gomp/adjust-args-5.f90 | 58 +++++ .../gfortran.dg/gomp/declare-variant-2.f90 | 6 +- .../gomp/declare-variant-21-aux.f90 | 18 ++ .../gfortran.dg/gomp/declare-variant-21.f90 | 28 +++ 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 | 24 +++ gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 | 38 ++++ gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 | 27 +++ gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 | 39 ++++ 28 files changed, 1042 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 ed1213a41cb..c06f69588e1 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. */ @@ -3037,7 +3044,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..c7a89924b78 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,88 @@ gfc_match_omp_declare_variant (void) return MATCH_ERROR; } - if (gfc_match_omp_context_selector_specification (odv) != MATCH_YES) - return MATCH_ERROR; - - if (gfc_match (" )") != MATCH_YES) + if (ccode == match) { - gfc_error ("expected %<)%> at %C"); - return MATCH_ERROR; + has_match = true; + if (gfc_match_omp_context_selector_specification (odv) + != MATCH_YES) + return MATCH_ERROR; + if (gfc_match (" )") != MATCH_YES) + { + gfc_error ("expected %<)%> at %C"); + return MATCH_ERROR; + } + } + else if (ccode == adjust_args) + { + has_adjust_args = true; + bool need_device_ptr_p; + if (gfc_match (" nothing") == MATCH_YES) + need_device_ptr_p = false; + else if (gfc_match (" need_device_ptr") == MATCH_YES) + need_device_ptr_p = true; + else + { + gfc_error ("expected % or % at %C"); + return MATCH_ERROR; + } + if (need_device_ptr_p) + { + if (gfc_match_omp_variable_list (" :", + &odv->need_device_ptr_arg_list, + false) + != MATCH_YES) + { + gfc_error ("expected argument list at %C"); + return MATCH_ERROR; + } + for (gfc_omp_namelist *n = odv->need_device_ptr_arg_list; + n != NULL; n = n->next) + { + if (!n->sym->attr.dummy) + { + gfc_error ("list item %qs at %L is not a dummy argument", + n->sym->name, &n->where); + return MATCH_ERROR; + } + if (n->sym->ts.type != BT_DERIVED + || !n->sym->ts.u.derived->ts.is_iso_c + || (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 +7739,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 +7921,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 +8890,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 +10536,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 +11511,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 +11928,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 +12063,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 4f4fafa4217..3ad44b0dde7 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 54ab60b4935..ad9fedc3452 100644 --- a/gcc/fortran/trans-decl.cc +++ b/gcc/fortran/trans-decl.cc @@ -2156,6 +2156,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. */ @@ -2409,7 +2411,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..a4d32811663 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,30 @@ gfc_trans_omp_depobj (gfc_code *code) return gfc_finish_block (&block); } +static tree +gfc_trans_omp_dispatch (gfc_code *code) +{ + stmtblock_t block; + gfc_code *next = code->block->next; + // assume ill-formed "function dispatch structured + // block" have already been rejected by resolve_omp_dispatch + gcc_assert (next->op == EXEC_CALL || next->op == EXEC_ASSIGN); + + tree body = gfc_trans_code (next); + gfc_start_block (&block); + tree omp_clauses + = gfc_trans_omp_clauses (&block, code->ext.omp_clauses, code->loc); + + tree stmt = make_node (OMP_DISPATCH); + SET_EXPR_LOCATION (stmt, gfc_get_location (&code->loc)); + TREE_TYPE (stmt) = void_type_node; + OMP_DISPATCH_BODY (stmt) = body; + OMP_DISPATCH_CLAUSES (stmt) = omp_clauses; + + gfc_add_expr_to_block (&block, stmt); + return gfc_finish_block (&block); +} + static tree gfc_trans_omp_error (gfc_code *code) { @@ -8272,6 +8326,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 +8444,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 +8640,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 +8669,97 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns) DECL_ATTRIBUTES (base_fn_decl) = tree_cons (id, build_tree_list (variant, set_selectors), DECL_ATTRIBUTES (base_fn_decl)); + + // Handle adjust_args + for (gfc_omp_namelist *arg_list + = odv->need_device_ptr_arg_list; + arg_list != NULL; arg_list = arg_list->next) + { + if (arg_list->sym->backend_decl == NULL_TREE) + { + gfc_error ( + "%s at %L is not a base function argument", + arg_list->sym->name, &arg_list->where); + continue; + } + + tree base_fn_arg_decl = arg_list->sym->backend_decl; + if (base_fn_arg_decl != error_mark_node) + { + // Is t specified more than once? + if (adjust_args_list.contains (base_fn_arg_decl)) + { + gfc_error ( + "%qD at %L is specified more than once", + base_fn_arg_decl, &arg_list->where); + continue; + } + adjust_args_list.safe_push (base_fn_arg_decl); + + // Handle variant argument + tree variant + = gfc_get_symbol_decl (variant_proc_sym); + tree variant_parm = DECL_ARGUMENTS (variant); + int idx; + tree arg; + for (arg = DECL_ARGUMENTS (base_fn_decl), idx = 0; + arg != NULL; arg = TREE_CHAIN (arg), idx++) + if (arg == base_fn_arg_decl) + break; + gcc_assert (arg != NULL_TREE); + if (variant_parm == NULL_TREE) + { + gfc_formal_arglist *arg + = variant_proc_sym->formal; + for (int i = 0; i < idx; i++) + { + arg = arg->next; + gcc_assert (arg != NULL); + } + + // Check we got the right parameter name + if (strcmp (arg_list->sym->name, arg->sym->name) + != 0) + { + gfc_error ("%s at %L is not a variant " + "function argument", + arg_list->sym->name, + &arg_list->where); + continue; + } + arg->sym->attr + .omp_declare_variant_need_device_ptr + = 1; + } + else + { + for (int i = 0; i < idx; i++) + { + variant_parm = TREE_CHAIN (variant_parm); + gcc_assert (variant_parm != NULL_TREE); + } + // Check we got the right parameter name + if (strcmp (arg_list->sym->name, + IDENTIFIER_POINTER ( + DECL_NAME (variant_parm))) + != 0) + { + gfc_error ("%s at %L is not a variant " + "function argument", + arg_list->sym->name, + &arg_list->where); + continue; + } + + tree attr = tree_cons ( + get_identifier ( + "omp declare variant adjust_args " + "need_device_ptr"), + NULL_TREE, DECL_ATTRIBUTES (variant_parm)); + DECL_ATTRIBUTES (variant_parm) = attr; + } + } + } } } } diff --git a/gcc/fortran/trans.cc b/gcc/fortran/trans.cc index 1067e032621..882d205b183 100644 --- a/gcc/fortran/trans.cc +++ b/gcc/fortran/trans.cc @@ -2597,6 +2597,7 @@ trans_code (gfc_code * code, tree cond) case EXEC_OMP_CANCELLATION_POINT: case EXEC_OMP_CRITICAL: case EXEC_OMP_DEPOBJ: + case EXEC_OMP_DISPATCH: case EXEC_OMP_DISTRIBUTE: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def index 390cc9542f7..5047c8f816a 100644 --- a/gcc/fortran/types.def +++ b/gcc/fortran/types.def @@ -120,6 +120,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_INT_BOOL, BT_BOOL, BT_INT, BT_BOOL) DEF_FUNCTION_TYPE_2 (BT_FN_VOID_PTR_PTRMODE, BT_VOID, BT_PTR, BT_PTRMODE) DEF_FUNCTION_TYPE_2 (BT_FN_VOID_CONST_PTR_SIZE, BT_VOID, BT_CONST_PTR, BT_SIZE) +DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_INT, BT_PTR, BT_CONST_PTR, BT_INT) DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR) diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90 new file mode 100644 index 00000000000..68adb60a397 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90 @@ -0,0 +1,63 @@ +! 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 f12 (a) + import c_ptr + type(c_ptr), intent(inout) :: a + !$omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) ! { dg-error "list item 'b' at .1. is not a dummy argument" } + end function + 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..b731cb340c1 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90 @@ -0,0 +1,26 @@ +! Test translation of OMP clause adjust_args +! { dg-do compile } + +module main + use iso_c_binding, only: c_ptr + implicit none + !type(c_ptr) :: a + +contains + subroutine base2 (a) + type(c_ptr), intent(inout) :: a + !$omp declare variant (variant2) match (construct={parallel}) adjust_args (need_device_ptr: a) ! { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause at .1." } + end subroutine + subroutine base3 (a) + type(c_ptr), intent(inout) :: a + !$omp declare variant (variant2) match (construct={dispatch}) adjust_args (need_device_ptr: a) adjust_args (need_device_ptr: a) ! { dg-error "'a' at .1. is specified more than once" } + end subroutine + + subroutine variant2 (a) + type(c_ptr), intent(inout) :: a + end subroutine + subroutine variant3 (i) + integer :: i + end subroutine + +end module diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90 new file mode 100644 index 00000000000..75e884044b2 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90 @@ -0,0 +1,58 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module main + use iso_c_binding, only: c_ptr + implicit none + + type :: struct + integer :: a + real :: b + end type + + interface + integer function f(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + end function + integer function f0(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + !$omp declare variant (f) match (construct={dispatch}) & + !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b, c) + end function + integer function f1(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + !$omp declare variant (f) match (construct={dispatch}) & + !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c) + end function + end interface + +contains +subroutine test + integer :: a + type(c_ptr) :: b + type(c_ptr) :: c(2) + type(struct) :: s + + s%a = f0 (a, b, c) + !$omp dispatch + s%a = f0 (a, b, c) + + s%b = f1 (a, b, c) + !$omp dispatch + s%b = f1 (a, b, c) + +end subroutine +end module + +! { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&parm\.\[0-9]+, D\.\[0-9]+\\);" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&b, D\.\[0-9]+\\);" 2 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90 new file mode 100644 index 00000000000..75e884044b2 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90 @@ -0,0 +1,58 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module main + use iso_c_binding, only: c_ptr + implicit none + + type :: struct + integer :: a + real :: b + end type + + interface + integer function f(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + end function + integer function f0(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + !$omp declare variant (f) match (construct={dispatch}) & + !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b, c) + end function + integer function f1(a, b, c) + import c_ptr + integer, intent(in) :: a + type(c_ptr), intent(inout) :: b + type(c_ptr), intent(out) :: c(:) + !$omp declare variant (f) match (construct={dispatch}) & + !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c) + end function + end interface + +contains +subroutine test + integer :: a + type(c_ptr) :: b + type(c_ptr) :: c(2) + type(struct) :: s + + s%a = f0 (a, b, c) + !$omp dispatch + s%a = f0 (a, b, c) + + s%b = f1 (a, b, c) + !$omp dispatch + s%b = f1 (a, b, c) + +end subroutine +end module + +! { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&parm\.\[0-9]+, D\.\[0-9]+\\);" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&b, D\.\[0-9]+\\);" 2 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90 index 7fc5071feff..62d2cb96fac 100644 --- a/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90 @@ -18,10 +18,10 @@ contains !$omp declare variant match(user={condition(.false.)}) ! { dg-error "expected '\\(' at .1." } end subroutine subroutine f6 () - !$omp declare variant (f1) ! { dg-error "expected 'match' at .1." } + !$omp declare variant (f1) ! { dg-error "expected 'match' or 'adjust_args' at .1." } end subroutine subroutine f7 () - !$omp declare variant (f1) simd ! { dg-error "expected 'match' at .1." } + !$omp declare variant (f1) simd ! { dg-error "expected 'match' or 'adjust_args' at .1." } end subroutine subroutine f8 () !$omp declare variant (f1) match ! { dg-error "expected '\\(' at .1." } @@ -183,7 +183,7 @@ contains !$omp declare variant (f1) match(construct={requires}) ! { dg-warning "unknown selector 'requires' for context selector set 'construct' at .1." } end subroutine subroutine f75 () - !$omp declare variant (f1),match(construct={parallel}) ! { dg-error "expected 'match' at .1." } + !$omp declare variant (f1),match(construct={parallel}) ! { dg-error "expected 'match' or 'adjust_args' at .1." } end subroutine subroutine f76 () !$omp declare variant (f1) match(implementation={atomic_default_mem_order("relaxed")}) ! { dg-error "expected identifier at .1." } diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-variant-21-aux.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-variant-21-aux.f90 new file mode 100644 index 00000000000..4e8bb129d40 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/declare-variant-21-aux.f90 @@ -0,0 +1,18 @@ +! { dg-do compile { target skip-all-targets } } + +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 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..022ae04dac0 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/declare-variant-21.f90 @@ -0,0 +1,28 @@ +! { dg-do run } +! { dg-additional-sources declare-variant-21-aux.f90 } +! { dg-additional-options "-fdump-tree-gimple" } + +! 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) +! { dg-final { scan-tree-dump "variant_proc \\(&a\\)" "gimple" { xfail *-*-* } } } + +end program main 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..149d0613b97 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90 @@ -0,0 +1,19 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module main + implicit none + interface + subroutine f2 () + end subroutine + end interface + contains + + subroutine test () + !$omp dispatch ! { dg-final { scan-tree-dump-times "#pragma omp task if\\(0\\)" 1 "gimple" } } + call f2 () + !$omp dispatch nowait ! { dg-final { scan-tree-dump-times "#pragma omp task if\\(1\\)" 1 "gimple" } } + call f2 () + end subroutine +end module + diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 new file mode 100644 index 00000000000..e45397f3f96 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 @@ -0,0 +1,24 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module main + implicit none + interface + subroutine f2 (a) + integer, intent(in) :: a + end subroutine + end interface + contains + + subroutine test () + integer :: a + + !$omp dispatch device(-25373654) + ! { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device \\(-25373654\\);" 1 "gimple" } } + call f2 (a) + !$omp dispatch device(a + a) + ! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = a.0_1 \\* 2;.*#pragma omp dispatch device\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\).*#pragma omp task shared\\(D\.\[0-9]+\\).*__builtin_omp_set_default_device \\(D\.\[0-9]+\\);" 1 "gimple" } } + call f2 (a) + end subroutine +end module + diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 new file mode 100644 index 00000000000..9f4fa2970ca --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 @@ -0,0 +1,38 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module main + use iso_c_binding, only: c_ptr + implicit none + interface + subroutine f1 (p, p2) + import :: c_ptr + type(c_ptr), intent(out) :: p + type(c_ptr), intent(in) :: p2 + end subroutine + subroutine f2 (p, p2) + import :: c_ptr + type(c_ptr), intent(out) :: p + type(c_ptr), intent(in) :: p2 + !$omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: p, p2) + end subroutine + end interface + contains + + subroutine test () + type(c_ptr) :: p, p2 + + !$omp dispatch + call f2 (p, p2) + !$omp dispatch is_device_ptr(p) + ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p\\) shared\\(p2\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*integer\\(kind=4\\) D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&p2, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(&p, D\.\[0-9]+\\);" 1 "gimple" } } + call f2 (p, p2) + !$omp dispatch is_device_ptr(p2) + ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p2\\) shared\\(p\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*integer\\(kind=4\\) D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&p, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(D\.\[0-9]+, &p2\\);" 1 "gimple" } } + call f2 (p, p2) + !$omp dispatch is_device_ptr(p, p2) + ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p\\) shared\\(p2\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*f1 \\(&p, &p2\\);" 1 "gimple" } } + call f2 (p, p2) + end subroutine +end module + diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 new file mode 100644 index 00000000000..32b6347be67 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 @@ -0,0 +1,27 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-ompexp" } + +module main + use iso_c_binding, only: c_ptr + implicit none + interface + subroutine f2 (p) + import :: c_ptr + type(c_ptr), intent(out) :: p + end subroutine + end interface + contains + + subroutine test () + type(c_ptr) :: p + + !$omp dispatch + ! { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, 0B, .*, .*\\);" 1 "ompexp" } } + call f2 (p) + !$omp dispatch depend(inout: p) + ! { dg-final { scan-tree-dump-times "D\.\[0-9]+\\\[2] = &p;" 1 "ompexp" } } + ! { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, &D\.\[0-9]+, .*, .*\\);" 1 "ompexp" } } + call f2 (p) + end subroutine +end module + diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 new file mode 100644 index 00000000000..6771336aa33 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 @@ -0,0 +1,39 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple -fdump-tree-omplower" } + +module main + use iso_c_binding, only: c_ptr + implicit none + interface + integer function f0 () + end function + integer function f1 () + end function + integer function f2 () + !$omp declare variant (f0) match (construct={dispatch}) + !$omp declare variant (f1) match (implementation={vendor(gnu)}) + end function + end interface + contains + + subroutine test () + integer :: a, n + + !$omp dispatch novariants(n < 1024) nocontext(n > 1024) + a = f2 () + end subroutine +end module + +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = n <= 1023;" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = n > 1024;" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp dispatch novariants\\(0\\) nocontext\\(0\\) shared\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 1 "gimple" } } + +! { dg-final { scan-tree-dump-times ".omp_data_o.1.D\.\[0-9]+ = D\.\[0-9]+;" 2 "omplower" } } +! { dg-final { scan-tree-dump-times ".omp_data_o.1.a = &a;" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = .omp_data_i->D\.\[0-9]+;" 2 "omplower" } } +! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = .omp_data_i->a;" 3 "omplower" } } +! { dg-final { scan-tree-dump-times "\\*D\.\[0-9]+ = D\.\[0-9]+;" 3 "omplower" } } From patchwork Fri Jul 12 14:11:55 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: 1959903 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=MZy/2WXZ; 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 4WLDF34p1Wz1xpd for ; Sat, 13 Jul 2024 00:17:11 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id E26FD3858C48 for ; Fri, 12 Jul 2024 14:17:09 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x332.google.com (mail-wm1-x332.google.com [IPv6:2a00:1450:4864:20::332]) by sourceware.org (Postfix) with ESMTPS id CC9DF38323F5 for ; Fri, 12 Jul 2024 14:13:02 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org CC9DF38323F5 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 CC9DF38323F5 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::332 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793586; cv=none; b=FeMVGWT/wkJ11iVakeAOsPl0xRehKpWnQFYYiqLCndPuXFrtfAtiqbS5o03ELki96cHxoDXJZRR2dD29JEKFdoB8xFmJvgPpnN3Q4AG0Xfks4/2yFhSgLBnpHIVIVTn2WxKe0PARJOwDwomzgW8Vc//WX/AQF5/KqccpYnoKBwI= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793586; c=relaxed/simple; bh=GmI4XuEmx2TjSE8cD3DAICVYwKVM+ui+SuKWLe10mXk=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=xOBNxSGDfcW6Ig6q8VISEng9os7CTxb3WHWRBA6bHS353z9D9udPH4hbPnm9nlS2/F57tnncZV3DZ9r9Vyooqbqk2lx2LhLA3OAkfUYoxonxTi1invpUw+vikKumzxwkbF3UiL9uAxLEwJ8Ze8rBb3StWuHQdFR7J5w5dr+jRwM= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x332.google.com with SMTP id 5b1f17b1804b1-4266f3e0df8so13697405e9.2 for ; Fri, 12 Jul 2024 07:13:02 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1720793581; x=1721398381; 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=rPdByaLT8UvZR3wQE4zyP5lyzCSBxBHBqC33r5u1mCo=; b=MZy/2WXZWsddrWw67iCZ7bXLZTrNY/Nc0xrXCpuTdvitPqWAyRwJWZvhJ+sVBpD/JH erb9pMAFCkqWZR1PjJa71uLeyK6YHvml8zxBaLgz6iAwtfZe+YtEeffy6sBYxQ1KRU71 BcvU3kwNQK83U9llLUF8O11pxckr3CWeiECfJIHSr+Z2e0QNUN7GRtsCD/Jb2UKTayyH vAqLeIqZP4FVV8X1bUSuSmWKJW+EG0zuYVeTfiACUT5UpUGs/CgqTxQGpS/NIN4Ns3vh nYsaB1KYmYzBUrpBDJ28UhULORYI7lL/G8MAFo7B1ftv7jf2ZzVHBicbwjljf3/9b9T0 4wyg== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1720793581; x=1721398381; 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=rPdByaLT8UvZR3wQE4zyP5lyzCSBxBHBqC33r5u1mCo=; b=IbwKut+VGO/ZQtuujKz7hb+R38aEZ1Csk4Xf5SlsZYTx7CiIypKdMa9UElqyPTvntO bysKRyBpEvFFXTZe4h76MsUt+P7/8z6lB00IgZKCbMKvDC/MH9XSuQ+XAG/iWxPmpPA7 8N3TzKdQNZvPtSLwYvgSoTsy7aXOQfd78kWwwCYiK0LlHVjm1vmFeiRXgHrvMMDEF1uc xtv0k4uyCb9yzeUuOQItFm8RY4V7NCz+eqlOyq7RcdLwy7ca6vmaDMHj3jrkl3Q/MBwi 1gIfyi7x3weBy2JRX4kDIkzI/nDCtZTCU1gez0vl8TCb8rlIwY4g4Dg5pEHPx9EJMRtN pC8Q== X-Gm-Message-State: AOJu0YzSpEjfZmq4sYKBGBWXh6NJ1BLs1a9fsWmTY+feYCznA4MKwlL3 EfAtP7qb/aPRfxd4yflZzMQ8izU31GzlAz6e5tKM1cpMiZK7eWuiqXTh8+RJ0ecYLg9mUL54hKX 3 X-Google-Smtp-Source: AGHT+IFSov6H1faQPXzQohHDeEB5Ar+640IhKcDICwxV5Q6LyceQ689UmXeGvEQB4quWGUoIVHPOEw== X-Received: by 2002:a05:600c:41d1:b0:426:59fc:cdec with SMTP id 5b1f17b1804b1-426707e1fd1mr77479285e9.21.1720793580737; Fri, 12 Jul 2024 07:13:00 -0700 (PDT) Received: from localhost.localdomain ([169.155.255.128]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4279f25b946sm24680325e9.19.2024.07.12.07.13.00 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 12 Jul 2024 07:13:00 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v2 8/8] OpenMP: update documentation for dispatch and adjust_args Date: Fri, 12 Jul 2024 16:11:55 +0200 Message-ID: <20240712141155.255186-9-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240712141155.255186-1-parras@baylibre.com> References: <20240712141155.255186-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.9 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org 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 50da248b74d..a2f5897463a 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