aboutsummaryrefslogtreecommitdiff
path: root/gcc
diff options
context:
space:
mode:
authorSandra Loosemore <sloosemore@baylibre.com>2024-12-26 18:16:00 +0000
committerSandra Loosemore <sloosemore@baylibre.com>2025-01-16 18:12:25 +0000
commit3c8df3693c40c7a29dde3915b66d29fad7134868 (patch)
tree2a7c9811010c9076caf846a639a69851321ddf53 /gcc
parentfdeceba59bee60040fd58203b6fe0239d789eade (diff)
downloadgcc-3c8df3693c40c7a29dde3915b66d29fad7134868.zip
gcc-3c8df3693c40c7a29dde3915b66d29fad7134868.tar.gz
gcc-3c8df3693c40c7a29dde3915b66d29fad7134868.tar.bz2
OpenMP: Update "declare target"/OpenMP context interaction
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/ChangeLog * 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.
Diffstat (limited to 'gcc')
-rw-r--r--gcc/c-family/c-attribs.cc2
-rw-r--r--gcc/c/c-decl.cc8
-rw-r--r--gcc/cp/decl2.cc9
-rw-r--r--gcc/omp-general.cc2
-rw-r--r--gcc/testsuite/c-c++-common/gomp/declare-target-indirect-2.c10
-rw-r--r--gcc/testsuite/c-c++-common/gomp/declare-variant-8.c4
-rw-r--r--gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c2
-rw-r--r--gcc/testsuite/gfortran.dg/gomp/declare-variant-8.f9012
8 files changed, 15 insertions, 34 deletions
diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc
index eb76430..f3181e7 100644
--- a/gcc/c-family/c-attribs.cc
+++ b/gcc/c-family/c-attribs.cc
@@ -592,8 +592,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 56d1388..314b118 100644
--- a/gcc/c/c-decl.cc
+++ b/gcc/c/c-decl.cc
@@ -5476,12 +5476,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 869701b..f64aa84 100644
--- a/gcc/cp/decl2.cc
+++ b/gcc/cp/decl2.cc
@@ -1786,13 +1786,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 c38905b..2ce79bf 100644
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
@@ -3026,7 +3026,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 6ba278b..75a205f 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 a7a3ba4..9cd706e 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 ddc3c2c..8d31905 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 d69e552..e393576 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