Hi! On Thu, 30 Oct 2014 17:11:04 -0700, Cesar Philippidis <[email protected]> wrote: > gcc/fortran/ > * gfortran.h (enum OMP_LIST_HOST): Remove. > (enum OMP_LIST_DEVICE, OMP_LIST_DEVICE): Remove. > * dump-parse-tree.c (show_omp_clauses): Remove OMP_LIST_HOST and > OMP_LIST_DEVICE from here also. > * openmp.c (OMP_CLAUSE_SELF): New define. > (gfc_match_omp_clauses): Update handling of OMP_CLAUSE_HOST and > OMP_CLAUSE_DEVICE. Add support for OMP_CLAUSE_SELF. > * trans-openmp.c (gfc_trans_omp_clauses): Remove support for > OMP_LIST_HOST and OMP_LIST_DEVICE since they are treated as memory > maps now. > (gfc_trans_oacc_executable_directive): Remove stale EXEC_OACC_WAIT.
Applied to gomp-4_0-branch in r217148:
commit a7bba5ecc7c62a022616f55ff1d8fb48266fcb67
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Nov 5 16:54:07 2014 +0000
OpenACC update host/self maintenance.
gcc/c/
* c-parser.c (c_parser_omp_clause_name) <"host">: Return
PRAGMA_OMP_CLAUSE_HOST.
gcc/c/
(c_parser_oacc_data_clause): Group PRAGMA_OMP_CLAUSE_SELF next to
PRAGMA_OMP_CLAUSE_HOST.
gcc/cp/
* parser.c (cp_parser_oacc_data_clause): Group
PRAGMA_OMP_CLAUSE_SELF next to PRAGMA_OMP_CLAUSE_HOST.
gcc/fortran/
* openmp.c (OMP_CLAUSE_HOST, OMP_CLAUSE_SELF): Merge into the new
OMP_CLAUSE_HOST_SELF. Update all users.
gcc/
* tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE_HOST and
OMP_CLAUSE_OACC_DEVICE. Update all users.
gcc/testsuite/
* c-c++-common/goacc/update-1.c: Extend.
* gfortran.dg/goacc/assumed.f95: Likewise.
* gfortran.dg/goacc/coarray.f95: Likewise.
* gfortran.dg/goacc/cray.f95: Likewise.
* gfortran.dg/goacc/literal.f95: Likewise.
* gfortran.dg/goacc/parameter.f95: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/update-1-2.c: New file.
* testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/data-4.f90: In one instance, use
the self clause instead of host clause.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217148
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 3 +
gcc/c/ChangeLog.gomp | 6 +
gcc/c/c-parser.c | 10 +-
gcc/cp/ChangeLog.gomp | 3 +
gcc/cp/parser.c | 8 +-
gcc/fortran/ChangeLog.gomp | 3 +
gcc/fortran/openmp.c | 21 +-
gcc/gimplify.c | 4 -
gcc/omp-low.c | 4 -
gcc/testsuite/ChangeLog.gomp | 7 +
gcc/testsuite/c-c++-common/goacc/update-1.c | 5 +
gcc/testsuite/gfortran.dg/goacc/assumed.f95 | 8 +-
gcc/testsuite/gfortran.dg/goacc/coarray.f95 | 3 +-
gcc/testsuite/gfortran.dg/goacc/cray.f95 | 6 +-
gcc/testsuite/gfortran.dg/goacc/literal.f95 | 5 +-
gcc/testsuite/gfortran.dg/goacc/parameter.f95 | 3 +-
gcc/tree-core.h | 10 +-
gcc/tree-pretty-print.c | 6 -
gcc/tree.c | 6 -
libgomp/ChangeLog.gomp | 5 +
.../libgomp.oacc-c-c++-common/update-1-2.c | 282 +++++++++++++++++++++
.../{data-4.f90 => data-4-2.f90} | 8 +-
libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 | 2 +-
23 files changed, 353 insertions(+), 65 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 2c2b349..d140a35 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,8 @@
2014-11-05 Thomas Schwinge <[email protected]>
+ * tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE_HOST and
+ OMP_CLAUSE_OACC_DEVICE. Update all users.
+
* gimplify.c (gimplify_oacc_cache): New function.
(gimplify_expr): Use it for OACC_CACHE.
(gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index 70278b9..a223a17 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,5 +1,11 @@
2014-11-05 Thomas Schwinge <[email protected]>
+ * c-parser.c (c_parser_omp_clause_name) <"host">: Return
+ PRAGMA_OMP_CLAUSE_HOST.
+
+ * c-parser.c (c_parser_oacc_data_clause): Group
+ PRAGMA_OMP_CLAUSE_SELF next to PRAGMA_OMP_CLAUSE_HOST.
+
* c-parser.c (c_parser_oacc_cache): Generate OACC_CACHE.
* c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE__CACHE_.
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 40d4314..bd2864f 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -9832,7 +9832,7 @@ c_parser_omp_clause_name (c_parser *parser)
break;
case 'h':
if (!strcmp ("host", p))
- result = PRAGMA_OMP_CLAUSE_SELF;
+ result = PRAGMA_OMP_CLAUSE_HOST;
break;
case 'i':
if (!strcmp ("inbranch", p))
@@ -10187,8 +10187,6 @@ c_parser_oacc_data_clause (c_parser *parser,
pragma_omp_clause c_kind,
enum omp_clause_map_kind kind;
switch (c_kind)
{
- default:
- gcc_unreachable ();
case PRAGMA_OMP_CLAUSE_COPY:
kind = OMP_CLAUSE_MAP_FORCE_TOFROM;
break;
@@ -10208,6 +10206,7 @@ c_parser_oacc_data_clause (c_parser *parser,
pragma_omp_clause c_kind,
kind = OMP_CLAUSE_MAP_FORCE_TO;
break;
case PRAGMA_OMP_CLAUSE_HOST:
+ case PRAGMA_OMP_CLAUSE_SELF:
kind = OMP_CLAUSE_MAP_FORCE_FROM;
break;
case PRAGMA_OMP_CLAUSE_PRESENT:
@@ -10225,9 +10224,8 @@ c_parser_oacc_data_clause (c_parser *parser,
pragma_omp_clause c_kind,
case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE:
kind = OMP_CLAUSE_MAP_ALLOC;
break;
- case PRAGMA_OMP_CLAUSE_SELF:
- kind = OMP_CLAUSE_MAP_FORCE_FROM;
- break;
+ default:
+ gcc_unreachable ();
}
tree nl, c;
nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list);
diff --git gcc/cp/ChangeLog.gomp gcc/cp/ChangeLog.gomp
index 46d4912..f5d400f 100644
--- gcc/cp/ChangeLog.gomp
+++ gcc/cp/ChangeLog.gomp
@@ -1,5 +1,8 @@
2014-11-05 Thomas Schwinge <[email protected]>
+ * parser.c (cp_parser_oacc_data_clause): Group
+ PRAGMA_OMP_CLAUSE_SELF next to PRAGMA_OMP_CLAUSE_HOST.
+
* parser.c (cp_parser_oacc_cache): Generate OACC_CACHE.
* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE__CACHE_.
diff --git gcc/cp/parser.c gcc/cp/parser.c
index ea4ad2f..9c2c3ca 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -27812,8 +27812,6 @@ cp_parser_oacc_data_clause (cp_parser *parser,
pragma_omp_clause c_kind,
enum omp_clause_map_kind kind;
switch (c_kind)
{
- default:
- gcc_unreachable ();
case PRAGMA_OMP_CLAUSE_COPY:
kind = OMP_CLAUSE_MAP_FORCE_TOFROM;
break;
@@ -27833,6 +27831,7 @@ cp_parser_oacc_data_clause (cp_parser *parser,
pragma_omp_clause c_kind,
kind = OMP_CLAUSE_MAP_FORCE_TO;
break;
case PRAGMA_OMP_CLAUSE_HOST:
+ case PRAGMA_OMP_CLAUSE_SELF:
kind = OMP_CLAUSE_MAP_FORCE_FROM;
break;
case PRAGMA_OMP_CLAUSE_PRESENT:
@@ -27850,9 +27849,8 @@ cp_parser_oacc_data_clause (cp_parser *parser,
pragma_omp_clause c_kind,
case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE:
kind = OMP_CLAUSE_MAP_ALLOC;
break;
- case PRAGMA_OMP_CLAUSE_SELF:
- kind = OMP_CLAUSE_MAP_FORCE_FROM;
- break;
+ default:
+ gcc_unreachable ();
}
tree nl, c;
nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list);
diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp
index 98e3971..d10560e 100644
--- gcc/fortran/ChangeLog.gomp
+++ gcc/fortran/ChangeLog.gomp
@@ -1,5 +1,8 @@
2014-11-05 Thomas Schwinge <[email protected]>
+ * openmp.c (OMP_CLAUSE_HOST, OMP_CLAUSE_SELF): Merge into the new
+ OMP_CLAUSE_HOST_SELF. Update all users.
+
* gfortran.texi: Update for OpenACC.
* intrinsic.texi: Likewise.
* invoke.texi: Likewise.
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index c7af004..959798a 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -445,13 +445,12 @@ match_oacc_clause_gang (gfc_omp_clauses *cp)
#define OMP_CLAUSE_INDEPENDENT (1ULL << 49)
#define OMP_CLAUSE_USE_DEVICE (1ULL << 50)
#define OMP_CLAUSE_DEVICE_RESIDENT (1ULL << 51)
-#define OMP_CLAUSE_HOST (1ULL << 52)
+#define OMP_CLAUSE_HOST_SELF (1ULL << 52)
#define OMP_CLAUSE_OACC_DEVICE (1ULL << 53)
#define OMP_CLAUSE_WAIT (1ULL << 54)
#define OMP_CLAUSE_DELETE (1ULL << 55)
#define OMP_CLAUSE_AUTO (1ULL << 56)
#define OMP_CLAUSE_TILE (1ULL << 57)
-#define OMP_CLAUSE_SELF (1ULL << 58)
/* Helper function for OpenACC and OpenMP clauses involving memory
mapping. */
@@ -682,24 +681,20 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, unsigned
long long mask,
true)
== MATCH_YES)
continue;
- if ((mask & OMP_CLAUSE_HOST)
- && gfc_match ("host ( ") == MATCH_YES
- && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_FROM))
- continue;
if ((mask & OMP_CLAUSE_OACC_DEVICE)
&& gfc_match ("device ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_FORCE_TO))
continue;
+ if ((mask & OMP_CLAUSE_HOST_SELF)
+ && (gfc_match ("host ( ") == MATCH_YES
+ || gfc_match ("self ( ") == MATCH_YES)
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_FORCE_FROM))
+ continue;
if ((mask & OMP_CLAUSE_TILE)
&& match_oacc_expr_list ("tile (", &c->tile_list, true) == MATCH_YES)
continue;
- if ((mask & OMP_CLAUSE_SELF)
- && gfc_match ("self ( ") == MATCH_YES
- && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_FROM))
- continue;
if ((mask & OMP_CLAUSE_SEQ) && !c->seq
&& gfc_match ("seq") == MATCH_YES)
{
@@ -1170,7 +1165,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, unsigned
long long mask,
| OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
| OMP_CLAUSE_PRESENT_OR_CREATE)
#define OACC_UPDATE_CLAUSES \
- (OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST | OMP_CLAUSE_SELF \
+ (OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST_SELF \
| OMP_CLAUSE_OACC_DEVICE | OMP_CLAUSE_WAIT)
#define OACC_ENTER_DATA_CLAUSES \
(OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT | OMP_CLAUSE_COPYIN \
diff --git gcc/gimplify.c gcc/gimplify.c
index d58876f..5aeb726 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -6288,8 +6288,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
remove = true;
break;
- case OMP_CLAUSE_HOST:
- case OMP_CLAUSE_OACC_DEVICE:
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
@@ -6692,8 +6690,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree
*list_p)
case OMP_CLAUSE_VECTOR_LENGTH:
break;
- case OMP_CLAUSE_HOST:
- case OMP_CLAUSE_OACC_DEVICE:
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
diff --git gcc/omp-low.c gcc/omp-low.c
index 1c9d942..0b45e69 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1977,8 +1977,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_local (decl, ctx);
break;
- case OMP_CLAUSE_HOST:
- case OMP_CLAUSE_OACC_DEVICE:
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
@@ -2125,8 +2123,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_WAIT:
break;
- case OMP_CLAUSE_HOST:
- case OMP_CLAUSE_OACC_DEVICE:
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 1faf0fa..a02f58a 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,5 +1,12 @@
2014-11-05 Thomas Schwinge <[email protected]>
+ * c-c++-common/goacc/update-1.c: Extend.
+ * gfortran.dg/goacc/assumed.f95: Likewise.
+ * gfortran.dg/goacc/coarray.f95: Likewise.
+ * gfortran.dg/goacc/cray.f95: Likewise.
+ * gfortran.dg/goacc/literal.f95: Likewise.
+ * gfortran.dg/goacc/parameter.f95: Likewise.
+
* c-c++-common/goacc/cache-1.c: New file.
* gfortran.dg/goacc/data-tree.f95: Remove dg-prune-output directive.
diff --git gcc/testsuite/c-c++-common/goacc/update-1.c
gcc/testsuite/c-c++-common/goacc/update-1.c
index 2a3a910..97e9379 100644
--- gcc/testsuite/c-c++-common/goacc/update-1.c
+++ gcc/testsuite/c-c++-common/goacc/update-1.c
@@ -8,5 +8,10 @@ f (void)
#pragma acc update device(i)
#pragma acc update host(i)
#pragma acc update self(i)
+#pragma acc update device(a[1:3])
+#pragma acc update host(a[1:3])
+#pragma acc update self(a[1:3])
+#pragma acc update device(a(1:3)) /* { dg-error "expected '\\\)' before '\\\('
token" } */
#pragma acc update host(a(1:3)) /* { dg-error "expected '\\\)' before '\\\('
token" } */
+#pragma acc update self(a(1:3)) /* { dg-error "expected '\\\)' before '\\\('
token" } */
}
diff --git gcc/testsuite/gfortran.dg/goacc/assumed.f95
gcc/testsuite/gfortran.dg/goacc/assumed.f95
index 15bfa0c..3287241 100644
--- gcc/testsuite/gfortran.dg/goacc/assumed.f95
+++ gcc/testsuite/gfortran.dg/goacc/assumed.f95
@@ -19,8 +19,9 @@ contains
do i = 1,5
enddo
!$acc end parallel loop
- !$acc update host (a) ! { dg-error "Assumed size" }
!$acc update device (a) ! { dg-error "Assumed size" }
+ !$acc update host (a) ! { dg-error "Assumed size" }
+ !$acc update self (a) ! { dg-error "Assumed size" }
end subroutine assumed_size
subroutine assumed_rank(a)
implicit none
@@ -39,7 +40,8 @@ contains
do i = 1,5
enddo
!$acc end parallel loop
- !$acc update host (a) ! { dg-error "Assumed rank" }
!$acc update device (a) ! { dg-error "Assumed rank" }
+ !$acc update host (a) ! { dg-error "Assumed rank" }
+ !$acc update self (a) ! { dg-error "Assumed rank" }
end subroutine assumed_rank
-end module test
\ No newline at end of file
+end module test
diff --git gcc/testsuite/gfortran.dg/goacc/coarray.f95
gcc/testsuite/gfortran.dg/goacc/coarray.f95
index ab13157..4f1224e 100644
--- gcc/testsuite/gfortran.dg/goacc/coarray.f95
+++ gcc/testsuite/gfortran.dg/goacc/coarray.f95
@@ -27,8 +27,9 @@ contains
!$acc cache (a)
enddo
!$acc end parallel loop
- !$acc update host (a)
!$acc update device (a)
+ !$acc update host (a)
+ !$acc update self (a)
end subroutine oacc1
end module test
! { dg-prune-output "ACC cache unimplemented" }
diff --git gcc/testsuite/gfortran.dg/goacc/cray.f95
gcc/testsuite/gfortran.dg/goacc/cray.f95
index 3225b28..8f2c077 100644
--- gcc/testsuite/gfortran.dg/goacc/cray.f95
+++ gcc/testsuite/gfortran.dg/goacc/cray.f95
@@ -28,8 +28,9 @@ contains
!$acc cache (pointee) ! TODO: This must fail, as in openacc-1_0-branch
enddo
!$acc end parallel loop
- !$acc update host (pointee) ! { dg-error "Cray pointee" }
!$acc update device (pointee) ! { dg-error "Cray pointee" }
+ !$acc update host (pointee) ! { dg-error "Cray pointee" }
+ !$acc update self (pointee) ! { dg-error "Cray pointee" }
!$acc data copy (ptr)
!$acc end data
!$acc data deviceptr (ptr) ! { dg-error "Cray pointer" }
@@ -47,8 +48,9 @@ contains
!$acc cache (ptr) ! TODO: This must fail, as in openacc-1_0-branch
enddo
!$acc end parallel loop
- !$acc update host (ptr)
!$acc update device (ptr)
+ !$acc update host (ptr)
+ !$acc update self (ptr)
end subroutine oacc1
end module test
! { dg-prune-output "unimplemented" }
diff --git gcc/testsuite/gfortran.dg/goacc/literal.f95
gcc/testsuite/gfortran.dg/goacc/literal.f95
index bdbf66d..e6760d0 100644
--- gcc/testsuite/gfortran.dg/goacc/literal.f95
+++ gcc/testsuite/gfortran.dg/goacc/literal.f95
@@ -23,7 +23,8 @@ contains
!$acc cache (10) ! { dg-error "Syntax error" }
enddo
!$acc end parallel loop
- !$acc update host (10) ! { dg-error "Syntax error" }
!$acc update device (10) ! { dg-error "Syntax error" }
+ !$acc update host (10) ! { dg-error "Syntax error" }
+ !$acc update self (10) ! { dg-error "Syntax error" }
end subroutine oacc1
-end module test
\ No newline at end of file
+end module test
diff --git gcc/testsuite/gfortran.dg/goacc/parameter.f95
gcc/testsuite/gfortran.dg/goacc/parameter.f95
index 785d7f9..1364181 100644
--- gcc/testsuite/gfortran.dg/goacc/parameter.f95
+++ gcc/testsuite/gfortran.dg/goacc/parameter.f95
@@ -24,8 +24,9 @@ contains
!$acc cache (a) ! TODO: This must fail, as in openacc-1_0-branch
enddo
!$acc end parallel loop
- !$acc update host (a) ! { dg-error "not a variable" }
!$acc update device (a) ! { dg-error "not a variable" }
+ !$acc update host (a) ! { dg-error "not a variable" }
+ !$acc update self (a) ! { dg-error "not a variable" }
end subroutine oacc1
end module test
! { dg-prune-output "unimplemented" }
diff --git gcc/tree-core.h gcc/tree-core.h
index 42ad6a0..cd38861 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -259,8 +259,8 @@ enum omp_clause_code {
OMP_CLAUSE_TO,
/* OpenACC clauses: {copy, copyin, copyout, create, delete, deviceptr,
- present, present_or_copy (pcopy), present_or_copyin (pcopyin),
- present_or_copyout (pcopyout), present_or_create (pcreate)}
+ device, host (self), present, present_or_copy (pcopy), present_or_copyin
+ (pcopyin), present_or_copyout (pcopyout), present_or_create (pcreate)}
(variable-list).
OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
@@ -270,12 +270,6 @@ enum omp_clause_code {
#pragma acc cache (variable-list). */
OMP_CLAUSE__CACHE_,
- /* OpenACC clause: host (variable_list). */
- OMP_CLAUSE_HOST,
-
- /* OpenACC clause: device (variable_list). */
- OMP_CLAUSE_OACC_DEVICE,
-
/* OpenACC clause: device_resident (variable_list). */
OMP_CLAUSE_DEVICE_RESIDENT,
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index d678f36..1458913 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -335,12 +335,6 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int
spc, int flags)
case OMP_CLAUSE__LOOPTEMP_:
name = "_looptemp_";
goto print_remap;
- case OMP_CLAUSE_HOST:
- name = "host";
- goto print_remap;
- case OMP_CLAUSE_OACC_DEVICE:
- name = "device";
- goto print_remap;
case OMP_CLAUSE_DEVICE_RESIDENT:
name = "device_resident";
goto print_remap;
diff --git gcc/tree.c gcc/tree.c
index f39c63f..16d156b 100644
--- gcc/tree.c
+++ gcc/tree.c
@@ -271,8 +271,6 @@ unsigned const char omp_clause_num_ops[] =
2, /* OMP_CLAUSE_TO */
2, /* OMP_CLAUSE_MAP */
2, /* OMP_CLAUSE__CACHE_ */
- 1, /* OMP_CLAUSE_HOST */
- 1, /* OMP_CLAUSE_OACC_DEVICE */
1, /* OMP_CLAUSE_DEVICE_RESIDENT */
1, /* OMP_CLAUSE_USE_DEVICE */
1, /* OMP_CLAUSE_GANG */
@@ -330,8 +328,6 @@ const char * const omp_clause_code_name[] =
"to",
"map",
"_cache_",
- "host",
- "device",
"device_resident",
"use_device",
"gang",
@@ -11120,8 +11116,6 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE:
switch (OMP_CLAUSE_CODE (*tp))
{
- case OMP_CLAUSE_HOST:
- case OMP_CLAUSE_OACC_DEVICE:
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 096a2a9..d9b92fe 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,5 +1,10 @@
2014-11-05 Thomas Schwinge <[email protected]>
+ * testsuite/libgomp.oacc-c-c++-common/update-1-2.c: New file.
+ * testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-4.f90: In one instance, use
+ the self clause instead of host clause.
+
* testsuite/libgomp.oacc-c/cache-1.c: Remove directives that are
expected to fail, and rename the file to...
* testsuite/libgomp.oacc-c-c++-common/cache-1.c: ... this.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
new file mode 100644
index 0000000..c7e7257
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
@@ -0,0 +1,282 @@
+/* Copy of update-1.c with self exchanged with host for #pragma acc update. */
+
+/* { dg-do run } */
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdbool.h>
+
+int
+main (int argc, char **argv)
+{
+ int N = 8;
+ float *a, *b, *c;
+ float *d_a, *d_b, *d_c;
+ int i;
+
+ a = (float *) malloc (N * sizeof (float));
+ b = (float *) malloc (N * sizeof (float));
+ c = (float *) malloc (N * sizeof (float));
+
+ d_a = (float *) acc_malloc (N * sizeof (float));
+ d_b = (float *) acc_malloc (N * sizeof (float));
+ d_c = (float *) acc_malloc (N * sizeof (float));
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ }
+
+ acc_map_data (a, d_a, N * sizeof (float));
+ acc_map_data (b, d_b, N * sizeof (float));
+ acc_map_data (c, d_c, N * sizeof (float));
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 3.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ b[i] = 1.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ b[i] = 1.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update host (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 6.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 6.0)
+ abort ();
+
+ if (b[i] != 6.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 7.0;
+ b[i] = 2.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 7.0)
+ abort ();
+
+ if (b[i] != 7.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc update device (a[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 9.0)
+ abort ();
+
+ if (b[i] != 9.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ }
+
+#pragma acc update device (a[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 6.0;
+ }
+
+#pragma acc update device (a[0:N >> 1])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < (N >> 1); i++)
+ {
+ if (a[i] != 6.0)
+ abort ();
+
+ if (b[i] != 6.0)
+ abort ();
+ }
+
+ for (i = (N >> 1); i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
similarity index 91%
copy from libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
copy to libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
index 41c45fb..16a8598 100644
--- libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
@@ -1,3 +1,5 @@
+! Copy of data-4.f90 with self exchanged with host for !acc update.
+
! { dg-do run }
program asyncwait
@@ -24,7 +26,7 @@ program asyncwait
end do
!$acc end parallel
- !$acc update host (a(1:N), b(1:N)) async wait
+ !$acc update self (a(1:N), b(1:N)) async wait
!$acc wait
do i = 1, N
@@ -78,7 +80,7 @@ program asyncwait
end do
!$acc end parallel
- !$acc update host (a(1:N), b(1:N), c(1:N), d(1:N)) async (1) wait (1)
+ !$acc update self (a(1:N), b(1:N), c(1:N), d(1:N)) async (1) wait (1)
!$acc wait (1)
@@ -122,7 +124,7 @@ program asyncwait
end do
!$acc end parallel
- !$acc update host (a(1:N), b(1:N), c(1:N), d(1:N), e(1:N)) async (1) wait (1)
+ !$acc update self (a(1:N), b(1:N), c(1:N), d(1:N), e(1:N)) async (1) wait (1)
!$acc wait (1)
!$acc exit data delete (N, a(1:N), b(1:N), c(1:N), d(1:N), e(1:N))
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
index 41c45fb..f6886b0 100644
--- libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
@@ -44,7 +44,7 @@ program asyncwait
end do
!$acc end parallel
- !$acc update host (a(1:N), b(1:N)) async (1) wait (1)
+ !$acc update self (a(1:N), b(1:N)) async (1) wait (1)
!$acc wait (1)
do i = 1, N
Grüße,
Thomas
pgpHwaQbcYMJ4.pgp
Description: PGP signature
