Hi Tobias,
this implements 'link' clause behavior for 'enter' maps, when
detected that 'requires self_maps' is in effect.
This basically does it in the front-ends, changing the tagged
attribute to "target link" under the supposed right conditions.
In the Fortran FE, including tree.h/omp-general.h in openmp.cc
is too problematic due to the _MASK symbol clash, so exporting
the self_maps detection from parse.cc via a function.
Tested under x86_64-linux with nvptx offloading, some new testcases
plus existing libgomp.c/omp_target_is_present-2.c adjusted.
Okay for mainline?
Thanks,
Chung-Lin
2026-07-08 Chung-Lin Tang <[email protected]>
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_declare_target): Adjust to tag
"omp declare target link" attribute for enter clause under
'requires self_maps'.
gcc/cp/ChangeLog:
* parser.cc (handle_omp_declare_target_clause): Adjust to tag
"omp declare target link" attribute for enter clause under
'requires self_maps'.
gcc/fortran/ChangeLog:
* openmp.cc (gfc_match_omp_declare_target): Set
n->u.common->omp_declare_target_link = 1 instead of ->omp_declare_target
for enter clause under 'requires self_maps'.
* parse.cc (is_omp_requires_self_maps): New function.
* parse.h (is_omp_requires_self_maps): New declaration.
libgomp/ChangeLog:
* testsuite/libgomp.c/omp_target_is_present.c: Adjust tests for
self_maps.
* testsuite/libgomp.c-c++-common/requires-8.c: New test.
* testsuite/libgomp.fortran/requires-2.f90: New test.diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 4eba2999727..63f3d3db5b2 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -28927,7 +28927,9 @@ c_parser_omp_declare_target (c_parser *parser)
tree at2 = lookup_attribute ("omp declare target link",
DECL_ATTRIBUTES (t));
only_device_type_or_indirect = false;
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK)
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK
+ || (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ENTER
+ && (omp_requires_mask & OMP_REQUIRES_SELF_MAPS) != 0))
{
id = get_identifier ("omp declare target link");
std::swap (at1, at2);
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index aeff4535337..8045abe8ec2 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -54385,7 +54385,9 @@ handle_omp_declare_target_clause (tree c, tree t, int
device_type,
tree at1 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (t));
tree at2 = lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (t));
tree id;
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK)
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK
+ || (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ENTER
+ && (omp_requires_mask & OMP_REQUIRES_SELF_MAPS) != 0))
{
id = get_identifier ("omp declare target link");
std::swap (at1, at2);
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 64355971d9e..20187e66be9 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -6767,7 +6767,12 @@ gfc_match_omp_declare_target (void)
n->u.common->name, &n->where);
if (list == OMP_LIST_TO || list == OMP_LIST_ENTER)
- n->u.common->omp_declare_target = 1;
+ {
+ if (is_omp_requires_self_maps ())
+ n->u.common->omp_declare_target_link = 1;
+ else
+ n->u.common->omp_declare_target = 1;
+ }
if (list == OMP_LIST_LINK)
n->u.common->omp_declare_target_link = 1;
if (list == OMP_LIST_LOCAL)
diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc
index 7b6cf525a80..d5209139a5b 100644
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -8058,3 +8058,10 @@ is_omp_declarative_stmt (gfc_statement st)
return false;
}
}
+
+/* Return true if requires self_maps used. */
+bool
+is_omp_requires_self_maps (void)
+{
+ return (omp_requires_mask & OMP_REQUIRES_SELF_MAPS) != 0;
+}
diff --git a/gcc/fortran/parse.h b/gcc/fortran/parse.h
index 52f82af880f..e65a9d35e08 100644
--- a/gcc/fortran/parse.h
+++ b/gcc/fortran/parse.h
@@ -82,4 +82,5 @@ extern vec<int> gfc_omp_metadirective_region_stack;
match gfc_match_prefix (gfc_typespec *);
bool is_oacc (gfc_state_data *);
+bool is_omp_requires_self_maps (void);
#endif /* GFC_PARSE_H */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-8.c
b/libgomp/testsuite/libgomp.c-c++-common/requires-8.c
new file mode 100644
index 00000000000..949518c490e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-8.c
@@ -0,0 +1,18 @@
+double data [12];
+#pragma omp requires self_maps
+
+#pragma omp declare target enter(data)
+
+int main (void)
+{
+ void *devaddr;
+ #pragma omp target map(from: devaddr)
+ {
+ devaddr = (void *) data;
+ }
+
+ if (devaddr != (void *) data)
+ __builtin_abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/omp_target_is_present.c
b/libgomp/testsuite/libgomp.c/omp_target_is_present.c
index 2e5e7acab05..21994ee2ca4 100644
--- a/libgomp/testsuite/libgomp.c/omp_target_is_present.c
+++ b/libgomp/testsuite/libgomp.c/omp_target_is_present.c
@@ -51,14 +51,13 @@ void check_routines (int dev)
if (omp_target_is_present (&A, dev) != initial_dev)
__builtin_abort ();
- // For link, it points to the pointer var - FIXME: update for self_maps
implying 'link'
- if (omp_target_is_present (&gEnter, dev) != !invalid_dev)
+ if (omp_target_is_present (&gEnter, dev) != (!invalid_dev && (initial_dev ||
!REQ_SELF_MAPS)))
__builtin_abort ();
if (omp_target_is_present (&gLink, dev) != (!invalid_dev && (initial_dev ||
!REQ_SELF_MAPS)))
__builtin_abort ();
- if (omp_target_is_present (&gLocal, dev) != !invalid_dev)
+ if (omp_target_is_present (&gLocal, dev) != (!invalid_dev && (initial_dev ||
!REQ_SELF_MAPS)))
__builtin_abort ();
int *ptr = (int*) 0xDEEDBEEF;
@@ -100,14 +99,15 @@ void check_routines (int dev)
__builtin_abort ();
if (!!omp_target_is_present (&gEnter, dev)
- != (initial_dev || (/* !self_mapping && */ !invalid_dev)))
+ != (initial_dev || (!self_mapping && !invalid_dev)))
__builtin_abort ();
if (!!omp_target_is_present (&gLink, dev)
!= (initial_dev || (!self_mapping && !invalid_dev)))
__builtin_abort ();
- if (!!omp_target_is_present (&gLocal, dev) != !invalid_dev)
+ if (!!omp_target_is_present (&gLocal, dev)
+ != (initial_dev || (!self_mapping && !invalid_dev)))
__builtin_abort ();
int *ptr2 = omp_get_mapped_ptr (&A, dev);
@@ -144,7 +144,7 @@ void check_routines (int dev)
if (ptr2 != &gEnter)
__builtin_abort ();
}
- else if (invalid_dev /* FIXME: || self_mapping */)
+ else if (invalid_dev || self_mapping)
{
if (ptr2 != nullptr)
__builtin_abort ();
@@ -172,7 +172,7 @@ void check_routines (int dev)
if (ptr2 != &gLocal)
__builtin_abort ();
}
- else if (invalid_dev)
+ else if (invalid_dev || self_mapping)
{
if (ptr2 != nullptr)
__builtin_abort ();
diff --git a/libgomp/testsuite/libgomp.fortran/requires-2.f90
b/libgomp/testsuite/libgomp.fortran/requires-2.f90
new file mode 100644
index 00000000000..85d49323b28
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/requires-2.f90
@@ -0,0 +1,17 @@
+program main
+ use iso_c_binding
+ type (c_ptr) :: devaddr
+ real(c_double), target :: data(10)
+ integer(c_intptr_t) :: ip
+ !$omp requires self_maps
+
+ !$omp target enter data map(data)
+
+ !$omp target map(from: devaddr)
+ block
+ devaddr = c_loc (data)
+ end block
+
+ if (transfer(devaddr, ip) /= transfer(c_loc(data), ip)) stop 1
+
+end program main
