From patchwork Sat May 4 21:21:51 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Sandra Loosemore X-Patchwork-Id: 1931428 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=2ZybkorU; 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 4VX10n6HzYz1xnT for ; Sun, 5 May 2024 07:25:13 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 17AB1384474A for ; Sat, 4 May 2024 21:25:12 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-il1-x12f.google.com (mail-il1-x12f.google.com [IPv6:2607:f8b0:4864:20::12f]) by sourceware.org (Postfix) with ESMTPS id 7055A384404C for ; Sat, 4 May 2024 21:22:25 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 7055A384404C 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 7055A384404C Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2607:f8b0:4864:20::12f ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1714857753; cv=none; b=HHSadqHG+mVKzYJYTCwWyz6Yq+MOfQ7JN9snWLEx7CQ6UG9XYrJXVLXdGcXvkDSQGl98quNySCXKdwDSsRw5gvmQC3OdjYz1xkR412QXMAwwjVTpSsMvx6fx52k6cLLkIpe2K8+50H8ICDBIfjUoIr+nHCH6s+d+84Pdi2+9CTY= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1714857753; c=relaxed/simple; bh=1C41AJldPVs/uhGF1F+hoJMPOwpdsmQNSpvnY0IRBzk=; h=DKIM-Signature:From:To:Subject:Date:Message-Id:MIME-Version; b=xEMrdJEAKtloJmzm+sTYBZeju4IP8q8Tsq4Lrp1wVZjjcZu+bd2DRYS65POURkv2yQjs8nNnLrjeh+7TOVa64JvHUNR5H/WBE5iTwCUyvAtjJ4PtJkQvcAKnxosy/6SA0SuZTfbVw5CS9veHG8VuDL50BTxk/VBIr/YqFvw7Tis= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-il1-x12f.google.com with SMTP id e9e14a558f8ab-36c5f81dc4bso4040905ab.0 for ; Sat, 04 May 2024 14:22:25 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1714857743; x=1715462543; 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=ShN4k7hmX+49NGPR3EnU4I6Ow25+gWb9DIJsBvN2KR4=; b=2ZybkorUTsZkq0hINX3TTR8jxD0rYMVTnhPMBYYnTt2gngZkGdBet7eB7nXPq8Qg2V Yf8MpEvp+qeM8LyOHKPn/9R1kMU6yFX1qEkA3ivoXebwQbvNVvNJczdcGxlc+EInWFIG 87RSLyf2usTr2krlCIVJ0FggjZ2LAg9u9C3Bh8TAsbALofFje0pWF2L/4RRTeIrW4azw CcW4LFpQ/8/Vp14GeynuhcP0IBTjZb9bL9Z97YhsbyxET/hkWBiUCp9RfSx08TeNke+z eEFpXHXf4uIcG/mE0eDw4DYuKb0ACEBhy3u+N4h0T+93P3OqKT7AQLeJC+yg4qt2wsgi CXbw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1714857743; x=1715462543; 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=ShN4k7hmX+49NGPR3EnU4I6Ow25+gWb9DIJsBvN2KR4=; b=tPTTiiGFO39VgUqEA1RtB798zXpH7WLVCY7/AoHa3Gx9xT5WI4PQeBL2UC6cCvthhu FrI//M3BPLsq+lTzH2k/lSo5xIDgtLHqo3xhsuRWcIiVvEvEZ1wvskg7Mh1iAZAF7rFV jmxMXvzK3WgVjLdHj8GMlqnMOEs4YMFRn6OBbTVyqqUFOvHRNkNXuA7YzkfZJWFrG5R3 rBwfgyB4NvrFLwBTnFZ/ELUWmEvU8fSXvci01AlCIUOFPBnuBgZZT0ClkgPM/yGAelpN EnPMp4UFx2DLY6LQr7k0qTb+ZySfAurdxyNMA6Q7ck79+IaAUGHlmJywknStIHGpQMCf POfQ== X-Gm-Message-State: AOJu0Yzb+Zz63kppAVSBkQr3UwhqsH2y/oCB/hJnIyjPCCVujSpDJoLA VN/kljQGeB3Ocubc/nKloWWn7utO79hJ2jR8Ui8v5KLgMSrkIYqrIjdZIFJaEa2JBjchSACdyfb h X-Google-Smtp-Source: AGHT+IE6iu5zh+mPAt/+lbUWbmebgZL9vsfgrts84AAgSKgIH+NkhegVZ32i/MUMMU8I1OSe9jJpTw== X-Received: by 2002:a92:cd82:0:b0:36c:504f:587e with SMTP id r2-20020a92cd82000000b0036c504f587emr8192331ilb.3.1714857743242; Sat, 04 May 2024 14:22:23 -0700 (PDT) Received: from pondscum.hsd1.co.comcast.net ([2601:281:d901:5620:3e29:4728:ec99:5098]) by smtp.gmail.com with ESMTPSA id ez3-20020a056638614300b004877be21febsm1559468jab.62.2024.05.04.14.22.22 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Sat, 04 May 2024 14:22:22 -0700 (PDT) From: Sandra Loosemore To: gcc-patches@gcc.gnu.org Cc: jakub@redhat.com, tburnus@baylibre.com Subject: [PATCH 11/12] OpenMP: Update "declare target"/OpenMP context interaction Date: Sat, 4 May 2024 15:21:51 -0600 Message-Id: <20240504212153.3561429-12-sloosemore@baylibre.com> X-Mailer: git-send-email 2.34.1 In-Reply-To: <20240504212153.3561429-1-sloosemore@baylibre.com> References: <20240504212153.3561429-1-sloosemore@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-11.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 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 04e39b41bdf..582d99ada1b 100644 --- a/gcc/c-family/c-attribs.cc +++ b/gcc/c-family/c-attribs.cc @@ -570,8 +570,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 52af8f32998..4ab7cd86030 100644 --- a/gcc/c/c-decl.cc +++ b/gcc/c/c-decl.cc @@ -5414,12 +5414,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 806a2a4bc69..028105a5b26 100644 --- a/gcc/cp/decl2.cc +++ b/gcc/cp/decl2.cc @@ -1777,13 +1777,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 b7eca439ad9..986f9e4f558 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -2609,7 +2609,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 9a3fa5230f8..6abaddcd5f4 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" } } */ /* { dg-prune-output "'reverse_offload' clause on 'requires' directive not supported yet" } */ 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