From: Tobias Burnus <tobias@codesourcery.com>
To: gcc-patches <gcc-patches@gcc.gnu.org>
Subject: [committed] OpenMP: Cleanups related to the 'present' modifier
Date: Mon, 12 Jun 2023 18:44:23 +0200 [thread overview]
Message-ID: <64017201-8206-fd22-70e4-897c858ae049@codesourcery.com> (raw)
[-- Attachment #1: Type: text/plain, Size: 1094 bytes --]
Cleanup follow up to
r14-1579-g4ede915d5dde93 "openmp: Add support for the 'present' modifier"
committed 6 days ago.
Namely:
* Replace for the program → libgomp ABI GOMP_MAP_PRESENT_[ALLOC,TO,FROM,TOFROM]
by the preexisting GOMP_MAP_FORCE_PRESENT but keep the other enum values
(and use them until gimplifcation).
* Improve wording if a non-existing/unsupported map-type modifier was used
by not referring to 'omp target' as it could be also target (enter/exit) data.
+ Add a testcase for enter/exit data + data.
* Unify + improve wording shown for 'present' when not present on the device.
* Extend in the testcases to check that data actually gets copied with
'target update' and 'map when the 'present' modifier is present.
Committed as Rev. r14-1736-g38944ec2a6fa10
Tobias
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Attachment #2: committed.diff --]
[-- Type: text/x-patch, Size: 40444 bytes --]
commit 38944ec2a6fa108d24e5cfbb24c52020f9aa3015
Author: Tobias Burnus <tobias@codesourcery.com>
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 %<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 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 %<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 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 "" } */
next prev reply other threads:[~2023-06-12 16:44 UTC|newest]
Thread overview: 13+ messages / expand[flat|nested] mbox.gz Atom feed top
2023-02-03 13:44 [PATCH] openmp: Add support for " Kwok Cheung Yeung
2023-02-09 21:17 ` [OG12][committed] openmp: Add support for the " Kwok Cheung Yeung
2023-02-14 22:44 ` [og12] Address cast to pointer from integer of different size in 'libgomp/target.c:gomp_target_rev' (was: [OG12][committed] openmp: Add support for the 'present' modifier) Thomas Schwinge
2023-02-15 0:00 ` [OG12][committed] openmp: Add support for the 'present' modifier Kwok Cheung Yeung
2023-02-15 19:02 ` [og12] Fix 'libgomp.{c-c++-common,fortran}/target-present-*' test cases (was: [OG12][committed] openmp: Add support for the 'present' modifier) Thomas Schwinge
2023-06-07 11:25 ` [committed] testsuite/libgomp.*/target-present-*.{c,f90}: Improve and fix (was: Re: [og12] Fix 'libgomp.{c-c++-common,fortran}/target-present-*' test cases) Tobias Burnus
2023-06-07 11:26 ` Tobias Burnus
2023-06-12 16:44 ` Tobias Burnus [this message]
2023-06-14 8:42 ` [committed] OpenMP: Cleanups related to the 'present' modifier Thomas Schwinge
2023-06-14 10:00 ` Tobias Burnus
2023-02-17 11:45 ` [PATCHv2] openmp: Add support for " Kwok Cheung Yeung
2023-04-28 17:26 ` Tobias Burnus
2023-06-06 14:55 ` [committed] " Tobias Burnus
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=64017201-8206-fd22-70e4-897c858ae049@codesourcery.com \
--to=tobias@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).