@@ -1,5 +1,8 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * 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
@@ -1,5 +1,11 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * 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_.
@@ -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);
@@ -1,5 +1,8 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * 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_.
@@ -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);
@@ -1,5 +1,8 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * 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.
@@ -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 \
@@ -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:
@@ -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:
@@ -1,5 +1,12 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * 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.
@@ -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" } */
}
@@ -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
@@ -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" }
@@ -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" }
@@ -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
@@ -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" }
@@ -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,
@@ -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;
@@ -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:
@@ -1,5 +1,10 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * 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.
new file mode 100644
@@ -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;
+}
similarity index 91%
copy from libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
copy to 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))
@@ -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