diff options
21 files changed, 153 insertions, 102 deletions
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 72f6fba..61487cc 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 %<always%>, %<close%> " - "or %<present%> on %<map%> clause"); + c_parser_error (parser, "%<map%> clause with map-type modifier other " + "than %<always%>, %<close%> or %<present%>"); parens.skip_until_found_close (parser); return list; } diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 09cba71..f5f9f5a 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 %<always%>, %<close%> " - "or %<present%> on %<map%> clause"); + cp_parser_error (parser, "%<map%> clause with map-type modifier other" + " than %<always%>, %<close%> or %<present%>"); 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 a2e74a5..8fb47fd 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 91640de..0e24b91 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 1857b5b..b882df0 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 1afff7e..b84f89b 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 8c5b7f7..5152d9d 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 4b4bd6d..fcf3125 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 669a623..e26d1e0 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 9eb956f..7ef9d46 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 74bd01f..ac9a0f8 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 f99bffe..a9db2f1 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 49b7dd8..8d4e8e8 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 a9e8005..e3c4121 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 aa34319..5eaa9cd 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 ad11023..07ae90b 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 455519a..582247d 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 768166f..fc13609 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 8f2c24e..524d01d 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 eb29c90..dd4af4c 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 61c8109..02fbfda 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 "" } */ |