From patchwork Sat Jul 20 20:42:30 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Sandra Loosemore X-Patchwork-Id: 94275 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id AF15B385E457 for ; Sat, 20 Jul 2024 20:46:19 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-il1-x133.google.com (mail-il1-x133.google.com [IPv6:2607:f8b0:4864:20::133]) by sourceware.org (Postfix) with ESMTPS id 075C9385DDFB for ; Sat, 20 Jul 2024 20:42:58 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 075C9385DDFB 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 075C9385DDFB Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2607:f8b0:4864:20::133 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1721508181; cv=none; b=wIXEsqsZQpmSNlWsgmHE3vBrvFQHOA9SVwNUQk74EZjJNCV6Xq4LZBA4gJSnaKmh5Kq/YaWEB8hFsd/cr6YCnTToOY7bVZWLz/RTkDxc8BImRuEPqVrznMZtHqORsF3phuw/bdEZ2N6iQwKRaxwqgoqqA948gq747KPS/Dc9if8= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1721508181; c=relaxed/simple; bh=hThLZnmgxnfCJsydldzYbxwmQysL+Zn6RZDbBCyQsF4=; h=DKIM-Signature:From:To:Subject:Date:Message-Id:MIME-Version; b=eOGDSJ2H22DmDTF9ns4XrSzbACSgi8wZ7O5uXTZhGYqLcVyK1nB+gzAOkCVGHK1Y4B/8PGbY4rgcj2rAAzgL5HLULyuwWnq9RX0sAIcmOwztXbL7GPI6vkkTQOFNeAUK03KzxYAl1dDubPKdzjcHEm9URhmG2sed1hRfx0z0UmI= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-il1-x133.google.com with SMTP id e9e14a558f8ab-397a363cb13so7895935ab.2 for ; Sat, 20 Jul 2024 13:42:58 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1721508177; x=1722112977; 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=yFcECitphx8qRmj+Z3Y8CghPRvlhRvcbOrun1pfG+LY=; b=E+y1d0TxnvJlSYouv+xV1kmLz514qqS2SBaACymvnV+lo7gEVwcn0wE402DeUxEIE6 s5ssNpeQZF17lP+Pva2EZWp85ilKHqxv62INGMHu/X0fXk8Vu5G/iFgUkEKSPAV76ppw WH959xmAFdgLb5mH8nVjZ7SGvj297pT3mxCKRc5hr18RFJZUO3GEAyfOwDdLB4qMBy1P OHCKy/xx3WSvTeqxVGXjN/ST/qiq9klA+iRVD23eViiBtTV/LExJMef2q7z9rvi1fOOr wgpDt8NqXk9B7rArM71DGtbPMhXG05ee0wdthDyP0/nOpk68yDL1bJnTyJ2hGszO7TKf /yIQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1721508177; x=1722112977; 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=yFcECitphx8qRmj+Z3Y8CghPRvlhRvcbOrun1pfG+LY=; b=IZDot2/FD/iDCuh+vo3jWtEzttFQgoy8BBGMZiSYGVsEdARk43G95sY5J+mCylc7a4 HOQHh4mLxF6OhjFbtaXMa43qnfDgy+eos0oaIpB1BXYH9iFGu61c6ZHKmrTrqfHMBtCh J8/gh6mtxS1vn1E1VcmAr+hY581c5QP+KKZ/LnzDre8O0TYvXMuHFAicgv+Edg6pxq+b dAysc1C/Qv7DuOB4RUAoEst4WJ1Q7eaRYi8rYT8Fj0BlmGlwEWvcbfL2r7pDRpYv/L2z Pq1Fv9yXmI9wO4AGh1zqVL9ILH9vvyh7jcxIzAo+1M92e17PtK1YX/F//0fB0KnYYO0r M0Dg== X-Gm-Message-State: AOJu0YwAyezDqwPVNj+1OmsrY+GjyFgTB1Tstyn3aY6Z9+tOaJwp51z7 BeyBGy26FCN3W7Is+GteipHWYtYkdkWKwmstWNlx+r6x1Z4esSdb+iZKDJNx4ly5Ldw1j4KHwKZ 2WDo= X-Google-Smtp-Source: AGHT+IF6fhSPD76pM9HXQryCSe6mPIldHguxzn6tgTxiS6avYTriAW3p2cDH/NqdTBK3qtG/+SdEAg== X-Received: by 2002:a05:6e02:1aad:b0:397:12e5:5b7c with SMTP id e9e14a558f8ab-398e7447373mr52030585ab.27.1721508176912; Sat, 20 Jul 2024 13:42:56 -0700 (PDT) Received: from pondscum.hsd1.co.comcast.net ([2601:281:d901:5620:ee50:43b7:d7db:ddf2]) by smtp.gmail.com with ESMTPSA id e9e14a558f8ab-397f5c2b74esm13133105ab.35.2024.07.20.13.42.56 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Sat, 20 Jul 2024 13:42:56 -0700 (PDT) From: Sandra Loosemore To: gcc-patches@gcc.gnu.org Cc: jakub@redhat.com, tburnus@baylibre.com Subject: [PATCH v3 11/12] OpenMP: Update "declare target"/OpenMP context interaction Date: Sat, 20 Jul 2024 14:42:30 -0600 Message-Id: <20240720204231.2229891-12-sloosemore@baylibre.com> X-Mailer: git-send-email 2.34.1 In-Reply-To: <20240720204231.2229891-1-sloosemore@baylibre.com> References: <20240720204231.2229891-1-sloosemore@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.0 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~patchwork=sourceware.org@gcc.gnu.org The code and test case previously implemented the OpenMP 5.0 spec, which said in section 2.3.1: "For functions within a declare target block, the target trait is added to the beginning of the set..." In OpenMP 5.1, this was changed to "For device routines, the target trait is added to the beginning of the set..." In OpenMP 5.2 and TR12, it says: "For procedures that are determined to be target function variants by a declare target directive..." The definition of "device routine" in OpenMP 5.1 is confusing, but certainly the intent of the later versions of the spec is clear that it doesn't just apply to functions within a begin declare target/end declare target block. The only use of the "omp declare target block" function attribute was to support the 5.0 language, so it can be removed. This patch changes the context augmentation to use the "omp declare target" attribute instead. gcc/c-family/ChangeLog * c-attribs.cc (c_common_gnu_attributes): Delete "omp declare target block". gcc/c/ChangeLog * c-decl.cc (c_decl_attributes): Don't add "omp declare target block". gcc/cp/decl2.cc * decl2.cc (cplus_decl_attributes): Don't add "omp declare target block". gcc/ChangeLog * omp-general.cc (omp_complete_construct_context): Check "omp declare target" attribute, not "omp declare target block". gcc/testsuite/ChangeLog * c-c++-common/gomp/declare-target-indirect-2.c : Adjust expected output for removal of "omp declare target block". * c-c++-common/gomp/declare-variant-8.c: Likewise, the variant call to f20 is now resolved differently. * c-c++-common/gomp/reverse-offload-1.c: Adjust expected output. * gfortran.dg/gomp/declare-variant-8.f90: Likewise, both f18 and f20 now resolve to the variant. Delete obsolete comments. --- gcc/c-family/c-attribs.cc | 2 -- gcc/c/c-decl.cc | 8 ++------ gcc/cp/decl2.cc | 9 ++------- gcc/omp-general.cc | 2 +- .../c-c++-common/gomp/declare-target-indirect-2.c | 10 +++++----- gcc/testsuite/c-c++-common/gomp/declare-variant-8.c | 4 ++-- gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c | 2 +- gcc/testsuite/gfortran.dg/gomp/declare-variant-8.f90 | 12 ++---------- 8 files changed, 15 insertions(+), 34 deletions(-) diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc index f9b229aba7f..7f74c25d3d7 100644 --- a/gcc/c-family/c-attribs.cc +++ b/gcc/c-family/c-attribs.cc @@ -574,8 +574,6 @@ const struct attribute_spec c_common_gnu_attributes[] = handle_omp_declare_target_attribute, NULL }, { "omp declare target nohost", 0, 0, true, false, false, false, handle_omp_declare_target_attribute, NULL }, - { "omp declare target block", 0, 0, true, false, false, false, - handle_omp_declare_target_attribute, NULL }, { "non overlapping", 0, 0, true, false, false, false, handle_non_overlapping_attribute, NULL }, { "alloc_align", 1, 1, false, true, true, false, diff --git a/gcc/c/c-decl.cc b/gcc/c/c-decl.cc index 97f1d346835..e2a62bf2d56 100644 --- a/gcc/c/c-decl.cc +++ b/gcc/c/c-decl.cc @@ -5418,12 +5418,8 @@ c_decl_attributes (tree *node, tree attributes, int flags) attributes = tree_cons (get_identifier ("omp declare target implicit"), NULL_TREE, attributes); else - { - attributes = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, attributes); - attributes = tree_cons (get_identifier ("omp declare target block"), - NULL_TREE, attributes); - } + attributes = tree_cons (get_identifier ("omp declare target"), + NULL_TREE, attributes); if (TREE_CODE (*node) == FUNCTION_DECL) { int device_type diff --git a/gcc/cp/decl2.cc b/gcc/cp/decl2.cc index 6d674684931..da881e9e797 100644 --- a/gcc/cp/decl2.cc +++ b/gcc/cp/decl2.cc @@ -1781,13 +1781,8 @@ cplus_decl_attributes (tree *decl, tree attributes, int flags) = tree_cons (get_identifier ("omp declare target implicit"), NULL_TREE, attributes); else - { - attributes = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, attributes); - attributes - = tree_cons (get_identifier ("omp declare target block"), - NULL_TREE, attributes); - } + attributes = tree_cons (get_identifier ("omp declare target"), + NULL_TREE, attributes); if (TREE_CODE (*decl) == FUNCTION_DECL) { cp_omp_declare_target_attr &last diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc index 2cbb907fb54..c9cc0a9b81a 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -2631,7 +2631,7 @@ omp_complete_construct_context (tree construct_context, bool *completep) } /* Add target trait when in a target variant. */ - if (lookup_attribute ("omp declare target block", attributes)) + if (lookup_attribute ("omp declare target", attributes)) construct_context = make_trait_selector (OMP_TRAIT_CONSTRUCT_TARGET, NULL_TREE, NULL_TREE, construct_context); diff --git a/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-2.c b/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-2.c index 6ba278b3ef0..75a205feb95 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-2.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-2.c @@ -4,12 +4,12 @@ #pragma omp begin declare target indirect void fn1 (void) { } #pragma omp end declare target -/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target, omp declare target block, omp declare target indirect\\\)\\\)\\\nvoid fn1" "gimple" } } */ +/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target, omp declare target indirect\\\)\\\)\\\nvoid fn1" "gimple" } } */ #pragma omp begin declare target indirect (0) void fn2 (void) { } #pragma omp end declare target -/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target, omp declare target block\\\)\\\)\\\nvoid fn2" "gimple" } } */ +/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nvoid fn2" "gimple" } } */ void fn3 (void) { } #pragma omp declare target indirect to (fn3) @@ -27,6 +27,6 @@ void fn4 (void) { } #pragma omp declare target indirect enter(baz) #pragma omp end declare target #pragma omp end declare target -/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target, omp declare target block, omp declare target indirect\\\)\\\)\\\nint foo" "gimple" } } */ -/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target, omp declare target block\\\)\\\)\\\nint bar" "gimple" } } */ -/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target indirect, omp declare target, omp declare target block\\\)\\\)\\\nint baz" "gimple" } } */ +/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target, omp declare target indirect\\\)\\\)\\\nint foo" "gimple" } } */ +/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nint bar" "gimple" } } */ +/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target indirect, omp declare target\\\)\\\)\\\nint baz" "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-8.c b/gcc/testsuite/c-c++-common/gomp/declare-variant-8.c index a7a3ba41b97..9cd706e896f 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-variant-8.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-8.c @@ -1,4 +1,4 @@ -/* { dg-do compile { target c } } */ +/* { dg-do compile } */ /* { dg-additional-options "-fdump-tree-gimple" } */ void f01 (void); @@ -102,7 +102,7 @@ void test3 (void) { #pragma omp parallel - f20 (); /* { dg-final { scan-tree-dump-times "f20 \\\(\\\);" 1 "gimple" } } */ + f20 (); /* { dg-final { scan-tree-dump-times "f19 \\\(\\\);" 1 "gimple" } } */ } void diff --git a/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c index ddc3c2c6be1..8d319058ce1 100644 --- a/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c +++ b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c @@ -4,7 +4,7 @@ /* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target\\)\\)\[\n\r\]*int called_in_target1" 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target\\)\\)\[\n\r\]*int called_in_target2" 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target, omp declare target block\\)\\)\[\n\r\]*void tg_fn" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target\\)\\)\[\n\r\]*void tg_fn" 1 "omplower" } } */ #pragma omp requires reverse_offload diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-variant-8.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-variant-8.f90 index d69e552eeb7..e3935768bc4 100644 --- a/gcc/testsuite/gfortran.dg/gomp/declare-variant-8.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/declare-variant-8.f90 @@ -167,23 +167,15 @@ contains end subroutine subroutine test2 () - ! OpenMP 5.0 specifies that the 'target' trait should be added for - ! functions within a declare target block, but Fortran does not have - ! the notion of a declare target _block_, so the variant is not used here. - ! This may change in later versions of OpenMP. - !$omp declare target !$omp parallel - call f18 () ! { dg-final { scan-tree-dump-times "f18 \\\(\\\);" 1 "gimple" } } + call f18 () ! { dg-final { scan-tree-dump-times "f17 \\\(\\\);" 1 "gimple" } } !$omp end parallel end subroutine subroutine test3 () - ! In the C version, this test was used to check that the - ! 'declare target to' form of the directive did not result in the variant - ! being used. !$omp parallel - call f20 () ! { dg-final { scan-tree-dump-times "f20 \\\(\\\);" 1 "gimple" } } + call f20 () ! { dg-final { scan-tree-dump-times "f19 \\\(\\\);" 1 "gimple" } } !$omp end parallel end subroutine