https://gcc.gnu.org/g:e39cea88b2ed5518ff622b6ffab414854b234a2d
commit r17-3987-ge39cea88b2ed5518ff622b6ffab414854b234a2d Author: Chung-Lin Tang <[email protected]> Date: Tue Sep 8 10:09:23 2026 +0000 OpenMP: Implement 'link' behavior for maps under 'requires self_maps' [PR122205] This patch implements 'link' clause behavior for 'enter' maps, when detected that 'requires self_maps' or 'unified_shared_memory' is in effect. This is done during implicit declare-target discovery, across all global variables, changing any "omp declare target" attributes into "omp declare target link". PR middle-end/122205 gcc/ChangeLog: * omp-offload.cc (omp_discover_implicit_declare_target): For variables, under 'OMP_REQUIRES_SELF_MAPS | OMP_REQUIRES_UNIFIED_SHARED_MEMORY', remove "omp declare target" attribute and add "omp declare target link". 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: --- gcc/omp-offload.cc | 20 +++++++++++++ .../testsuite/libgomp.c-c++-common/requires-8.c | 35 ++++++++++++++++++++++ .../testsuite/libgomp.c/omp_target_is_present.c | 14 ++++----- libgomp/testsuite/libgomp.fortran/requires-2.f90 | 27 +++++++++++++++++ 4 files changed, 89 insertions(+), 7 deletions(-) diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc index 3c37956a28c5..c9b4a94255cf 100644 --- a/gcc/omp-offload.cc +++ b/gcc/omp-offload.cc @@ -395,6 +395,26 @@ omp_discover_implicit_declare_target (void) &worklist); } + if (omp_requires_mask + & (OMP_REQUIRES_SELF_MAPS | OMP_REQUIRES_UNIFIED_SHARED_MEMORY)) + FOR_EACH_VARIABLE (vnode) + { + /* If 'self_maps' or 'unified_shared_memory' is enabled, + remove 'enter/to' and add 'link'. */ + if (lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (vnode->decl))) + { + DECL_ATTRIBUTES (vnode->decl) + = remove_attribute ("omp declare target", + DECL_ATTRIBUTES (vnode->decl)); + if (!lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (vnode->decl))) + DECL_ATTRIBUTES (vnode->decl) + = tree_cons (get_identifier ("omp declare target link"), + NULL_TREE, DECL_ATTRIBUTES (vnode->decl)); + } + } + lang_hooks.decls.omp_finish_decl_inits (); } 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 000000000000..6efbe45efcd8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-8.c @@ -0,0 +1,35 @@ +#include <stdint.h> + +double data [12]; +#pragma omp requires self_maps + +#pragma omp declare target enter(data) + +void * __attribute__((noinline)) +get_addr_from_function (void) +{ + return (void *) &data; +} + +int main (void) +{ + void *devaddr; + #pragma omp target map(from: devaddr) + { + devaddr = (void *) data; + } + + if (devaddr != (void *) data) + __builtin_abort (); + + #pragma omp target map(from: devaddr) + { + devaddr = get_addr_from_function (); + } + + 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 2e5e7acab05c..21994ee2ca4e 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 000000000000..82d481330f9b --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/requires-2.f90 @@ -0,0 +1,27 @@ +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 + + !$omp target map(from: devaddr) + block + devaddr = get_addr_from_function () + end block + if (transfer(devaddr, ip) /= transfer(c_loc(data), ip)) stop 2 + +contains + type (c_ptr) function get_addr_from_function () + get_addr_from_function = c_loc (data) + end function get_addr_from_function + +end program main
