commit 38944ec2a6fa108d24e5cfbb24c52020f9aa3015 Author: Tobias Burnus Date: Mon Jun 12 18:15:28 2023 +0200 OpenMP: Cleanups related to the 'present' modifier Reduce number of enum values passed to libgomp as GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} have the same semantic as GOMP_MAP_FORCE_PRESENT (i.e. abort if not present, otherwise ignore); that's different to GOMP_MAP_ALWAYS_PRESENT_{TO,TOFROM,FROM} which also abort if not present but copy data when present. This is is a follow-up to the commit r14-1579-g4ede915d5dde93 done 6 days ago. Additionally, the commit improves a libgomp run-time and a C/C++ compile-time error wording and extends testcases a tiny bit. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_clause_map): Reword error message for clearness especially with 'omp target (enter/exit) data.' gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_clause_map): Reword error message for clearness especially with 'omp target (enter/exit) data.' * semantics.cc (handle_omp_array_sections): Handle GOMP_MAP_{ALWAYS_,}PRESENT_{TO,TOFROM,FROM,ALLOC} enum values. gcc/ChangeLog: * gimplify.cc (gimplify_adjust_omp_clauses_1): Use GOMP_MAP_FORCE_PRESENT for 'present alloc' implicit mapping. (gimplify_adjust_omp_clauses): Change GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} to the equivalent GOMP_MAP_FORCE_PRESENT. * omp-low.cc (lower_omp_target): Remove handling of no-longer valid GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC}; update map kinds used for to/from clauses with present modifier. include/ChangeLog: * gomp-constants.h (enum gomp_map_kind): Change the enum values GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} to be compiler only. (GOMP_MAP_PRESENT_P): Update to include also GOMP_MAP_FORCE_PRESENT. libgomp/ChangeLog: * target.c (gomp_to_device_kind_p, gomp_map_vars_internal): Replace GOMP_MAP_PRESENT_{FROM,TO,TOFROM,ACLLOC} by GOMP_MAP_FORCE_PRESENT. (gomp_map_vars_internal, gomp_update): Likewise; unify and improve error message. * testsuite/libgomp.c-c++-common/target-present-2.c: Update for changed error message. * testsuite/libgomp.fortran/target-present-1.f90: Likewise. * testsuite/libgomp.fortran/target-present-2.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/present-1.c: Likewise. * testsuite/libgomp.c-c++-common/target-present-1.c: Likewise and extend testcase to check that data is copied when needed. * testsuite/libgomp.c-c++-common/target-present-3.c: Likewise. * testsuite/libgomp.fortran/target-present-3.f90: Likewise. gcc/testsuite/ChangeLog: * c-c++-common/gomp/defaultmap-4.c: Update scan-tree-dump. * c-c++-common/gomp/map-9.c: Likewise. * gfortran.dg/gomp/defaultmap-8.f90: Likewise. * gfortran.dg/gomp/map-11.f90: Likewise. * gfortran.dg/gomp/target-update-1.f90: Likewise. * gfortran.dg/gomp/map-12.f90: Likewise; also check original dump. * c-c++-common/gomp/map-6.c: Update dg-error and also check clause error with 'target (enter/exit) data'. --- gcc/c/c-parser.cc | 5 +- gcc/cp/parser.cc | 5 +- gcc/cp/semantics.cc | 7 +++ gcc/gimplify.cc | 13 ++++- gcc/omp-low.cc | 14 ++---- gcc/testsuite/c-c++-common/gomp/defaultmap-4.c | 4 +- gcc/testsuite/c-c++-common/gomp/map-6.c | 14 +++++- gcc/testsuite/c-c++-common/gomp/map-9.c | 8 ++-- gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90 | 4 +- gcc/testsuite/gfortran.dg/gomp/map-11.f90 | 8 ++-- gcc/testsuite/gfortran.dg/gomp/map-12.f90 | 35 +++++++------- gcc/testsuite/gfortran.dg/gomp/target-update-1.f90 | 2 +- include/gomp-constants.h | 19 ++++---- libgomp/target.c | 55 ++++++++++------------ .../libgomp.c-c++-common/target-present-1.c | 20 ++++++-- .../libgomp.c-c++-common/target-present-2.c | 2 +- .../libgomp.c-c++-common/target-present-3.c | 15 +++++- .../testsuite/libgomp.fortran/target-present-1.f90 | 2 +- .../testsuite/libgomp.fortran/target-present-2.f90 | 2 +- .../testsuite/libgomp.fortran/target-present-3.f90 | 19 ++++++-- .../libgomp.oacc-c-c++-common/present-1.c | 2 +- 21 files changed, 153 insertions(+), 102 deletions(-) diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 72f6fbae6a6..61487cc7481 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -17160,9 +17160,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list) } else { - c_parser_error (parser, "%<#pragma omp target%> with " - "modifier other than %, % " - "or % on % clause"); + c_parser_error (parser, "% clause with map-type modifier other " + "than %, % or %"); parens.skip_until_found_close (parser); return list; } diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 09cba713437..f5f9f5a4510 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -40625,9 +40625,8 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) } else { - cp_parser_error (parser, "%<#pragma omp target%> with " - "modifier other than %, % " - "or % on % clause"); + cp_parser_error (parser, "% clause with map-type modifier other" + " than %, % or %"); cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, /*or_comma=*/false, diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index a2e74a5d2c7..8fb47fd179e 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -5819,6 +5819,13 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_PRESENT_ALLOC: + case GOMP_MAP_PRESENT_TO: + case GOMP_MAP_PRESENT_FROM: + case GOMP_MAP_PRESENT_TOFROM: + case GOMP_MAP_ALWAYS_PRESENT_TO: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: case GOMP_MAP_FORCE_TO: diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 91640deecea..0e24b915b8f 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -12479,7 +12479,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) kind = GOMP_MAP_FORCE_PRESENT; break; case GOVD_MAP_FORCE_PRESENT | GOVD_MAP_ALLOC_ONLY: - kind = GOMP_MAP_PRESENT_ALLOC; + kind = GOMP_MAP_FORCE_PRESENT; break; default: gcc_unreachable (); @@ -12797,6 +12797,17 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, break; case OMP_CLAUSE_MAP: + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_PRESENT_ALLOC: + case GOMP_MAP_PRESENT_TO: + case GOMP_MAP_PRESENT_FROM: + case GOMP_MAP_PRESENT_TOFROM: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_PRESENT); + break; + default: + break; + } if (code == OMP_TARGET_EXIT_DATA && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER) { diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 1857b5b960f..b882df048ef 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -12800,10 +12800,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_PRESENT_ALLOC: - case GOMP_MAP_PRESENT_FROM: - case GOMP_MAP_PRESENT_TO: - case GOMP_MAP_PRESENT_TOFROM: + case GOMP_MAP_FORCE_PRESENT: case GOMP_MAP_ALWAYS_PRESENT_FROM: case GOMP_MAP_ALWAYS_PRESENT_TO: case GOMP_MAP_ALWAYS_PRESENT_TOFROM: @@ -12822,7 +12819,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_FORCE_TO: case GOMP_MAP_FORCE_FROM: case GOMP_MAP_FORCE_TOFROM: - case GOMP_MAP_FORCE_PRESENT: case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: case GOMP_MAP_LINK: @@ -13349,10 +13345,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_PRESENT_ALLOC: - case GOMP_MAP_PRESENT_TO: - case GOMP_MAP_PRESENT_FROM: - case GOMP_MAP_PRESENT_TOFROM: case GOMP_MAP_ALWAYS_PRESENT_TO: case GOMP_MAP_ALWAYS_PRESENT_FROM: case GOMP_MAP_ALWAYS_PRESENT_TOFROM: @@ -13397,13 +13389,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case OMP_CLAUSE_TO: tkind = (OMP_CLAUSE_MOTION_PRESENT (c) - ? GOMP_MAP_PRESENT_TO : GOMP_MAP_TO); + ? GOMP_MAP_ALWAYS_PRESENT_TO : GOMP_MAP_TO); tkind_zero = tkind; break; case OMP_CLAUSE_FROM: tkind = (OMP_CLAUSE_MOTION_PRESENT (c) - ? GOMP_MAP_PRESENT_FROM : GOMP_MAP_FROM); + ? GOMP_MAP_ALWAYS_PRESENT_FROM : GOMP_MAP_FROM); tkind_zero = tkind; break; default: diff --git a/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c b/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c index 1afff7ea38f..b84f89b0c7c 100644 --- a/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c +++ b/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c @@ -20,5 +20,5 @@ foo (void) c[i] = a[i] + b[i]; } -/* { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\) map\\(present,alloc:c \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" "gimple" } } */ -/* { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\) map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\) map\\(force_present:c \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c index 8c5b7f7cca1..5152d9d7c21 100644 --- a/gcc/testsuite/c-c++-common/gomp/map-6.c +++ b/gcc/testsuite/c-c++-common/gomp/map-6.c @@ -13,12 +13,22 @@ foo (void) #pragma omp target map (to:a) ; - #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ ; - #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ ; + #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + ; + + #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + ; + + #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + ; + + #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ /* { dg-error "'close' has not been declared" "" { target c++ } .-1 } */ /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */ diff --git a/gcc/testsuite/c-c++-common/gomp/map-9.c b/gcc/testsuite/c-c++-common/gomp/map-9.c index 4b4bd6d2aa3..fcf3125ceee 100644 --- a/gcc/testsuite/c-c++-common/gomp/map-9.c +++ b/gcc/testsuite/c-c++-common/gomp/map-9.c @@ -25,8 +25,8 @@ foo (void) c[i] = a[i] + b[i]; } -/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,to:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */ -/* { dg-final { scan-tree-dump "pragma omp target data map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */ -/* { dg-final { scan-tree-dump "pragma omp target.*map\\(present,from:c \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,to:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target data map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\) map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target.*map\\(force_present:c \\\[len: \[0-9\]+\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "pragma omp target exit data map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */ -/* { dg-final { scan-tree-dump "pragma omp target.*map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\) map\\(to:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target.*map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\) map\\(to:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */ diff --git a/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90 b/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90 index 669a623f746..e26d1e004b1 100644 --- a/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90 @@ -22,5 +22,5 @@ program main !$omp end target end program -! { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\).*map\\(present,alloc:c \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" "gimple" } } -! { dg-final { scan-tree-dump "pragma omp target.*map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\) defaultmap\\(present:aggregate\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\).*map\\(force_present:c \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target.*map\\(force_present:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\) defaultmap\\(present:aggregate\\)" "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/map-11.f90 b/gcc/testsuite/gfortran.dg/gomp/map-11.f90 index 9eb956f1aa1..7ef9d46f2f8 100644 --- a/gcc/testsuite/gfortran.dg/gomp/map-11.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/map-11.f90 @@ -27,8 +27,8 @@ program main !$omp end target end program -! { dg-final { scan-tree-dump "pragma omp target enter data map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\)" "gimple" } } -! { dg-final { scan-tree-dump "pragma omp target data map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } -! { dg-final { scan-tree-dump "pragma omp target.*map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target enter data map\\(force_present:a \\\[len: \[0-9\]+\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target data map\\(force_present:a \\\[len: \[0-9\]+\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target.*map\\(force_present:a \\\[len: \[0-9\]+\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(force_present:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } ! { dg-final { scan-tree-dump "pragma omp target exit data map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } -! { dg-final { scan-tree-dump "pragma omp target.*map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(to:a \\\[len: \[0-9\]+\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target.*map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(to:a \\\[len: \[0-9\]+\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/map-12.f90 b/gcc/testsuite/gfortran.dg/gomp/map-12.f90 index 74bd01f5f5a..ac9a0f8aae0 100644 --- a/gcc/testsuite/gfortran.dg/gomp/map-12.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/map-12.f90 @@ -1,4 +1,4 @@ -! { dg-additional-options "-fdump-tree-omplower" } +! { dg-additional-options "-fdump-tree-omplower -fdump-tree-original" } subroutine foo implicit none @@ -40,28 +40,29 @@ subroutine foo !$omp target exit data map(present,always,close,delete: a) map(close,present,always,release:b) map(always close present,from:b1) end subroutine +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:b1\\)\[\r\n\]" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,tofrom:b1\\)\[\r\n\]" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(present,tofrom:b1\\)\[\r\n\]" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,present,tofrom:b1\\)\[\r\n\]" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(alloc:a\\) map\\(to:b\\) map\\(to:b1\\)\[\r\n\]" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(alloc:a\\) map\\(always,to:b\\) map\\(always,to:b1\\)\[\r\n\]" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a\\) map\\(present,to:b\\) map\\(present,to:b1\\)\[\r\n\]" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a\\) map\\(always,present,to:b\\) map\\(always,present,to:b1\\)\[\r\n\]" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(delete:a\\) map\\(release:b\\) map\\(from:b1\\)\[\r\n\]" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(delete:a\\) map\\(release:b\\) map\\(always,from:b1\\)\[\r\n\]" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(delete:a\\) map\\(release:b\\) map\\(present,from:b1\\)\[\r\n\]" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(delete:a\\) map\\(release:b\\) map\\(always,present,from:b1\\)\[\r\n\]" 2 "original" } } + ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(present,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(present,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(force_present:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,present,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,present,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(to:b \\\[len: 4\\\]\\) map\\(to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(to:b \\\[len: 4\\\]\\) map\\(to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(always,to:b \\\[len: 4\\\]\\) map\\(always,to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(always,to:b \\\[len: 4\\\]\\) map\\(always,to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a \\\[len: 4\\\]\\) map\\(present,to:b \\\[len: 4\\\]\\) map\\(present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a \\\[len: 4\\\]\\) map\\(present,to:b \\\[len: 4\\\]\\) map\\(present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a \\\[len: 4\\\]\\) map\\(always,present,to:b \\\[len: 4\\\]\\) map\\(always,present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a \\\[len: 4\\\]\\) map\\(always,present,to:b \\\[len: 4\\\]\\) map\\(always,present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } +! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(force_present:a \\\[len: 4\\\]\\) map\\(force_present:b \\\[len: 4\\\]\\) map\\(force_present:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } +! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(force_present:a \\\[len: 4\\\]\\) map\\(always,present,to:b \\\[len: 4\\\]\\) map\\(always,present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(present,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(present,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,present,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } +! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(force_present:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,present,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90 index f99bffe2e0b..a9db2f1a39f 100644 --- a/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90 @@ -10,4 +10,4 @@ program main !$omp target update to(c) to(present: a) from(d) from(present: b) to(e) end program -! { dg-final { scan-tree-dump "#pragma omp target update to\\(c \\\[len: \[0-9\]+\\\]\\) to\\(present:a \\\[len: \[0-9\]+\\\]\\) to\\(e \\\[len: \[0-9\]+\\\]\\) from\\(present:b \\\[len: \[0-9\]+\\\]\\) from\\(d \\\[len: \[0-9\]+\\\]\\)" "gimple" } } +! { dg-final { scan-tree-dump "#pragma omp target update to\\(c \\\[len: \[0-9\]+\\\]\\) to\\(present:a \\\[len: \[0-9\]+\\\]\\) to\\(e \\\[len: \[0-9\]+\\\]\\) from\\(d \\\[len: \[0-9\]+\\\]\\) from\\(present:b \\\[len: \[0-9\]+\\\]\\)" "gimple" } } diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 49b7dd86ff5..8d4e8e81303 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -136,14 +136,6 @@ enum gomp_map_kind device. */ GOMP_MAP_ALWAYS_TOFROM = (GOMP_MAP_FLAG_SPECIAL_2 | GOMP_MAP_TOFROM), - /* Must already be present. */ - GOMP_MAP_PRESENT_ALLOC = (GOMP_MAP_FLAG_PRESENT | GOMP_MAP_ALLOC), - /* Must already be present, copy to device. */ - GOMP_MAP_PRESENT_TO = (GOMP_MAP_FLAG_PRESENT | GOMP_MAP_TO), - /* Must already be present, copy from device. */ - GOMP_MAP_PRESENT_FROM = (GOMP_MAP_FLAG_PRESENT | GOMP_MAP_FROM), - /* Must already be present, copy to and from device. */ - GOMP_MAP_PRESENT_TOFROM = (GOMP_MAP_FLAG_PRESENT | GOMP_MAP_TOFROM), /* Must already be present, unconditionally copy to device. */ GOMP_MAP_ALWAYS_PRESENT_TO = (GOMP_MAP_FLAG_ALWAYS_PRESENT | GOMP_MAP_TO), @@ -205,7 +197,13 @@ enum gomp_map_kind /* An attach or detach operation. Rewritten to the appropriate type during gimplification, depending on directive (i.e. "enter data" or parallel/kernels region vs. "exit data"). */ - GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3) + GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3), + /* Must already be present - all of following map to GOMP_MAP_FORCE_PRESENT + as no data transfer is needed. */ + GOMP_MAP_PRESENT_ALLOC = (GOMP_MAP_LAST | 4), + GOMP_MAP_PRESENT_TO = (GOMP_MAP_LAST | 5), + GOMP_MAP_PRESENT_FROM = (GOMP_MAP_LAST | 6), + GOMP_MAP_PRESENT_TOFROM = (GOMP_MAP_LAST | 7) }; #define GOMP_MAP_COPY_TO_P(X) \ @@ -243,7 +241,8 @@ enum gomp_map_kind (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_FLAG_FORCE) #define GOMP_MAP_PRESENT_P(X) \ - (((X) & GOMP_MAP_FLAG_PRESENT) == GOMP_MAP_FLAG_PRESENT) + (((X) & GOMP_MAP_FLAG_PRESENT) == GOMP_MAP_FLAG_PRESENT \ + || (X) == GOMP_MAP_FORCE_PRESENT) /* Asynchronous behavior. Keep in sync with diff --git a/libgomp/target.c b/libgomp/target.c index a9e8005c588..e3c4121a09f 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -358,8 +358,8 @@ gomp_to_device_kind_p (int kind) case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_FROM: case GOMP_MAP_ALWAYS_FROM: - case GOMP_MAP_PRESENT_FROM: case GOMP_MAP_ALWAYS_PRESENT_FROM: + case GOMP_MAP_FORCE_PRESENT: return false; default: return true; @@ -1699,37 +1699,29 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, i = j - 1; break; case GOMP_MAP_FORCE_PRESENT: + case GOMP_MAP_ALWAYS_PRESENT_TO: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: { /* We already looked up the memory region above and it was missing. */ size_t size = k->host_end - k->host_start; gomp_mutex_unlock (&devicep->lock); #ifdef HAVE_INTTYPES_H - gomp_fatal ("present clause: !acc_is_present (%p, " - "%"PRIu64" (0x%"PRIx64"))", - (void *) k->host_start, - (uint64_t) size, (uint64_t) size); + gomp_fatal ("present clause: not present on the device " + "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), " + "dev: %d)", (void *) k->host_start, + (uint64_t) size, (uint64_t) size, + devicep->target_id); #else - gomp_fatal ("present clause: !acc_is_present (%p, " - "%lu (0x%lx))", (void *) k->host_start, - (unsigned long) size, (unsigned long) size); + gomp_fatal ("present clause: not present on the device " + "(addr: %p, size: %lu (0x%lx), dev: %d)", + (void *) k->host_start, + (unsigned long) size, (unsigned long) size, + devicep->target_id); #endif } break; - case GOMP_MAP_PRESENT_ALLOC: - case GOMP_MAP_PRESENT_TO: - case GOMP_MAP_PRESENT_FROM: - case GOMP_MAP_PRESENT_TOFROM: - case GOMP_MAP_ALWAYS_PRESENT_TO: - case GOMP_MAP_ALWAYS_PRESENT_FROM: - case GOMP_MAP_ALWAYS_PRESENT_TOFROM: - /* We already looked up the memory region above and it - was missing. */ - gomp_mutex_unlock (&devicep->lock); - gomp_fatal ("present clause: not present on the device " - "(%p, %d)", - (void *) k->host_start, devicep->target_id); - break; case GOMP_MAP_FORCE_DEVICEPTR: assert (k->host_end - k->host_start == sizeof (void *)); gomp_copy_host2dev (devicep, aq, @@ -2149,9 +2141,18 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, /* We already looked up the memory region above and it was missing. */ gomp_mutex_unlock (&devicep->lock); +#ifdef HAVE_INTTYPES_H gomp_fatal ("present clause: not present on the device " - "(%p, %d)", - (void *) hostaddrs[i], devicep->target_id); + "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), " + "dev: %d)", (void *) hostaddrs[i], + (uint64_t) sizes[i], (uint64_t) sizes[i], + devicep->target_id); +#else + gomp_fatal ("present clause: not present on the device " + "(addr: %p, size: %lu (0x%lx), dev: %d)", + (void *) hostaddrs[i], (unsigned long) sizes[i], + (unsigned long) sizes[i], devicep->target_id); +#endif } } } @@ -3465,9 +3466,7 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, case GOMP_MAP_FORCE_TOFROM: case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_PRESENT_FROM: - case GOMP_MAP_PRESENT_TO: - case GOMP_MAP_PRESENT_TOFROM: + case GOMP_MAP_FORCE_PRESENT: case GOMP_MAP_ALWAYS_PRESENT_FROM: case GOMP_MAP_ALWAYS_PRESENT_TO: case GOMP_MAP_ALWAYS_PRESENT_TOFROM: @@ -3710,8 +3709,6 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, case GOMP_MAP_FORCE_TOFROM: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_PRESENT_FROM: - case GOMP_MAP_PRESENT_TOFROM: case GOMP_MAP_ALWAYS_PRESENT_FROM: case GOMP_MAP_ALWAYS_PRESENT_TOFROM: copy = true; diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c index aa343197e35..5eaa9cd1ad3 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c +++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c @@ -4,24 +4,34 @@ int main (void) { - int a[N], b[N], c[N]; + int a[N], b[N], c[N], d[N]; for (int i = 0; i < N; i++) { a[i] = i * 2; b[i] = i * 3 + 1; + d[i] = i * 5; } - #pragma omp target enter data map (alloc: a, c) - /* a has already been allocated, so this should be okay. */ - #pragma omp target map (present, to: a) + #pragma omp target enter data map (alloc: c, d) map(to: a) + #pragma omp target map (present, always, to: d) + for (int i = 0; i < N; i++) + if (d[i] != i * 5) + __builtin_abort (); + + /* a has already been mapped and 'c' allocated so this should be okay. */ + #pragma omp target map (present, to: a) map(present, always, from: c) for (int i = 0; i < N; i++) c[i] = a[i]; + for (int i = 0; i < N; i++) + if (c[i] != i * 2) + __builtin_abort (); + fprintf (stderr, "CheCKpOInT\n"); /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ /* b has not been allocated, so this should result in an error. */ - /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } */ + /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */ /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */ #pragma omp target map (present, to: b) for (int i = 0; i < N; i++) diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c index ad11023b2d6..07ae90b5aaa 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c +++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c @@ -21,7 +21,7 @@ int main (void) /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ /* b has not been allocated, so this should result in an error. */ - /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } */ + /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */ /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */ #pragma omp target defaultmap (present) for (int i = 0; i < N; i++) diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c index 455519af405..582247dc05e 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c +++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c @@ -16,11 +16,24 @@ int main (void) /* This should work as a has already been allocated. */ #pragma omp target update to (present: a) + #pragma omp target map(present,alloc: a, c) + for (int i = 0; i < N; i++) { + if (a[i] != i * 2) + __builtin_abort (); + c[i] = 23*i; + } + + #pragma omp target update from(present : c) + for (int i = 0; i < N; i++) { + if (c[i] != 23*i) + __builtin_abort (); + } + fprintf (stderr, "CheCKpOInT\n"); /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ /* This should fail as b has not been allocated. */ - /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } */ + /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */ /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */ #pragma omp target update to (present: b) diff --git a/libgomp/testsuite/libgomp.fortran/target-present-1.f90 b/libgomp/testsuite/libgomp.fortran/target-present-1.f90 index 768166fcff7..fc13609d528 100644 --- a/libgomp/testsuite/libgomp.fortran/target-present-1.f90 +++ b/libgomp/testsuite/libgomp.fortran/target-present-1.f90 @@ -20,7 +20,7 @@ program main ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } ! b has not been allocated, so this should result in an error. - ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } + ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } !$omp target map (present, to: b) do i = 1, N diff --git a/libgomp/testsuite/libgomp.fortran/target-present-2.f90 b/libgomp/testsuite/libgomp.fortran/target-present-2.f90 index 8f2c24ef5f2..524d01d9465 100644 --- a/libgomp/testsuite/libgomp.fortran/target-present-2.f90 +++ b/libgomp/testsuite/libgomp.fortran/target-present-2.f90 @@ -20,7 +20,7 @@ program main ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } ! b has not been allocated, so this should result in an error. - ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } + ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } !$omp target defaultmap (present) do i = 1, N diff --git a/libgomp/testsuite/libgomp.fortran/target-present-3.f90 b/libgomp/testsuite/libgomp.fortran/target-present-3.f90 index eb29c907624..dd4af4c1651 100644 --- a/libgomp/testsuite/libgomp.fortran/target-present-3.f90 +++ b/libgomp/testsuite/libgomp.fortran/target-present-3.f90 @@ -9,14 +9,27 @@ program main end do !$omp target enter data map (alloc: a, c) - ! This should work as a has already been allocated. - !$omp target update to (present: a) + + ! This should work as a has already been allocated. + !$omp target update to (present: a) + + !$omp target map(present, alloc: a, c) + do i = 1, N + if (a(i) /= i * 2) stop 1 + c(i) = 23 * i + end do + !$omp end target + + !$omp target update from (present: c) + do i = 1, N + if (c(i) /= 23 * i) stop 1 + end do print *, "CheCKpOInT" ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } ! This should fail as b has not been allocated. - ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } + ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } !$omp target update to (present: b) !$omp target exit data map (from: c) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c index 61c8109a7e0..02fbfdaf9ed 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c @@ -48,5 +48,5 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "present clause: !acc_is_present" } */ +/* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" } */ /* { dg-shouldfail "" } */