@@ -1,5 +1,19 @@
2017-05-24 Thomas Schwinge <thomas@codesourcery.com>
+ * openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
+ of preprocessor definitions.
+ * libgomp.h (strong_alias): Guard by "#ifdef
+ HAVE_ATTRIBUTE_ALIAS".
+ * oacc-mem.c: Provide "acc_pcreate" as alias for
+ "acc_present_or_create", and "acc_pcopyin" as alias for
+ "acc_present_or_copyin".
+ * libgomp.map: New version "OACC_2.0.1".
+ (OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate".
+ * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
+ its content into...
+ * testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
+ Extend testing.
+
* plugin/plugin-nvptx.c (nvptx_get_num_devices): Debugging output
when disabling nvptx offloading.
@@ -1060,8 +1060,6 @@ extern void gomp_set_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern void gomp_unset_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
-# define strong_alias(fn, al) \
- extern __typeof (fn) al __attribute__ ((alias (#fn)));
# define omp_lock_symver(fn) \
__asm (".symver g" #fn "_30, " #fn "@@OMP_3.0"); \
__asm (".symver g" #fn "_25, " #fn "@OMP_1.0");
@@ -1085,6 +1083,9 @@ extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
#endif
#ifdef HAVE_ATTRIBUTE_ALIAS
+# define strong_alias(fn, al) \
+ extern __typeof (fn) al __attribute__ ((alias (#fn)));
+
# define ialias_ulp ialias_str1(__USER_LABEL_PREFIX__)
# define ialias_str1(x) ialias_str2(x)
# define ialias_str2(x) #x
@@ -378,6 +378,12 @@ OACC_2.0 {
acc_set_cuda_stream;
};
+OACC_2.0.1 {
+ global:
+ acc_pcopyin;
+ acc_pcreate;
+} OACC_2.0;
+
GOACC_2.0 {
global:
GOACC_data_end;
@@ -514,12 +514,34 @@ acc_present_or_create (void *h, size_t s)
return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s);
}
+/* acc_pcreate is acc_present_or_create by a different name. */
+#ifdef HAVE_ATTRIBUTE_ALIAS
+strong_alias (acc_present_or_create, acc_pcreate)
+#else
+void *
+acc_pcreate (void *h, size_t s)
+{
+ return acc_present_or_create (h, s);
+}
+#endif
+
void *
acc_present_or_copyin (void *h, size_t s)
{
return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s);
}
+/* acc_pcopyin is acc_present_or_copyin by a different name. */
+#ifdef HAVE_ATTRIBUTE_ALIAS
+strong_alias (acc_present_or_copyin, acc_pcopyin)
+#else
+void *
+acc_pcopyin (void *h, size_t s)
+{
+ return acc_present_or_copyin (h, s);
+}
+#endif
+
#define FLAG_COPYOUT (1 << 0)
static void
@@ -91,8 +91,10 @@ void acc_free (void *) __GOACC_NOTHROW;
the standard specifies otherwise. */
void *acc_copyin (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_copyin (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcopyin (void *, size_t) __GOACC_NOTHROW;
void *acc_create (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_create (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcreate (void *, size_t) __GOACC_NOTHROW;
void acc_copyout (void *, size_t) __GOACC_NOTHROW;
void acc_delete (void *, size_t) __GOACC_NOTHROW;
void acc_update_device (void *, size_t) __GOACC_NOTHROW;
@@ -105,11 +107,6 @@ int acc_is_present (void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
-/* Old names. OpenACC does not specify whether these can or must
- not be macros, inlines or aliases for the new names. */
-#define acc_pcreate acc_present_or_create
-#define acc_pcopyin acc_present_or_copyin
-
/* CUDA-specific routines. */
void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
@@ -1,36 +1,245 @@
-/* { dg-do run } */
+/* acc_present_or_create, acc_present_or_copyin, etc. */
+#include <stdbool.h>
#include <stdlib.h>
#include <openacc.h>
int
main (int argc, char **argv)
{
- const int N = 256;
- unsigned char *h;
- void *d1, *d2;
+ int *h, *d;
+ const int N = 10000;
+ const int S = N * sizeof *h;
+ bool shared_mem;
- h = (unsigned char *) malloc (N);
-
- d1 = acc_present_or_create (h, N);
- if (!d1)
+ h = (int *) malloc (S);
+ if (!h)
abort ();
+ for (int i = 0; i < N; ++i)
+ h[i] = i + 0;
- d2 = acc_present_or_create (h, N);
- if (!d2)
- abort ();
+ shared_mem = acc_is_present (h, S);
- if (d1 != d2)
+ d = (int *) acc_present_or_create (h, S);
+ if (!d)
abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
+ abort ();
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ d[i] = i + 1;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 1 : 0))
+ abort ();
+ h[i] = i + 2;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_create (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 2 : 1))
+ abort ();
+ d[i] = i + 3;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 3 : 2))
+ abort ();
+ h[i] = i + 4;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 4 : 3))
+ abort ();
+ d[i] = i + 5;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 5 : 4))
+ abort ();
+ h[i] = i + 6;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_copyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 6 : 5))
+ abort ();
+ d[i] = i + 7;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 7 : 6))
+ abort ();
+ h[i] = i + 8;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 8 : 7))
+ abort ();
+ d[i] = i + 9;
+ }
- d2 = acc_pcreate (h, N);
- if (!d2)
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 9 : 8))
+ abort ();
+ h[i] = i + 10;
+ }
+
+ acc_copyout (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort ();
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ }
+
+ d = (int *) acc_pcopyin (h, S);
+ if (!d)
+ abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
abort ();
- if (d1 != d2)
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ d[i] = i + 11;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 11 : 9))
+ abort ();
+ h[i] = i + 12;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 12 : 11))
+ abort ();
+ d[i] = i + 13;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 13 : 12))
+ abort ();
+ h[i] = i + 14;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 14 : 13))
+ abort ();
+ d[i] = i + 15;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 15 : 14))
+ abort ();
+ h[i] = i + 16;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 16 : 15))
+ abort ();
+ d[i] = i + 17;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 17 : 16))
+ abort ();
+ h[i] = i + 18;
+ }
+
+ acc_update_self (h, S);
+ if (!acc_is_present (h, S))
abort ();
- acc_delete (h, N);
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 18 : 17))
+ abort ();
+ }
+
+ acc_delete (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort();
free (h);
deleted file mode 100644
@@ -1,64 +0,0 @@
-/* { dg-do run } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <string.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- int i;
- unsigned char *h;
- void *d1, *d2;
-
- h = (unsigned char *) malloc (N);
-
- for (i = 0; i < N; i++)
- {
- h[i] = i;
- }
-
- d1 = acc_present_or_copyin (h, N);
- if (!d1)
- abort ();
-
- for (i = 0; i < N; i++)
- {
- h[i] = 0xab;
- }
-
- d2 = acc_present_or_copyin (h, N);
- if (!d2)
- abort ();
-
- if (d1 != d2)
- abort ();
-
- memset (&h[0], 0, N);
-
- acc_copyout (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- d2 = acc_pcopyin (h, N);
- if (!d2)
- abort ();
-
- acc_copyout (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- free (h);
-
- return 0;
-}
Committed to gomp-4_0-branch in r248414:
commit 705bc833826559e64a7ccd821fcf419bfaf2b727
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed May 24 13:25:42 2017 +0000
C/C++ OpenACC: acc_pcopyin, acc_pcreate
libgomp/
* openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
of preprocessor definitions.
* libgomp.h (strong_alias): Guard by "#ifdef
HAVE_ATTRIBUTE_ALIAS".
* oacc-mem.c: Provide "acc_pcreate" as alias for
"acc_present_or_create", and "acc_pcopyin" as alias for
"acc_present_or_copyin".
* libgomp.map: New version "OACC_2.0.1".
(OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate".
* testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
its content into...
* testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
Extend testing.
trunk r248410
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@248414 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog.gomp | 14 ++
libgomp/libgomp.h | 5 +-
libgomp/libgomp.map | 8 +-
libgomp/oacc-mem.c | 24 +-
libgomp/openacc.h | 7 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-32.c | 241 +++++++++++++++++++--
.../testsuite/libgomp.oacc-c-c++-common/lib-38.c | 64 ------
7 files changed, 271 insertions(+), 92 deletions(-)
@@ -1,5 +1,19 @@
2017-05-24 Thomas Schwinge <thomas@codesourcery.com>
+ * openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
+ of preprocessor definitions.
+ * libgomp.h (strong_alias): Guard by "#ifdef
+ HAVE_ATTRIBUTE_ALIAS".
+ * oacc-mem.c: Provide "acc_pcreate" as alias for
+ "acc_present_or_create", and "acc_pcopyin" as alias for
+ "acc_present_or_copyin".
+ * libgomp.map: New version "OACC_2.0.1".
+ (OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate".
+ * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
+ its content into...
+ * testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
+ Extend testing.
+
* plugin/plugin-nvptx.c (nvptx_get_num_devices): Debugging output
when disabling nvptx offloading.
@@ -1049,8 +1049,6 @@ extern void gomp_set_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern void gomp_unset_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
-# define strong_alias(fn, al) \
- extern __typeof (fn) al __attribute__ ((alias (#fn)));
# define omp_lock_symver(fn) \
__asm (".symver g" #fn "_30, " #fn "@@OMP_3.0"); \
__asm (".symver g" #fn "_25, " #fn "@OMP_1.0");
@@ -1074,6 +1072,9 @@ extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
#endif
#ifdef HAVE_ATTRIBUTE_ALIAS
+# define strong_alias(fn, al) \
+ extern __typeof (fn) al __attribute__ ((alias (#fn)));
+
# define ialias_ulp ialias_str1(__USER_LABEL_PREFIX__)
# define ialias_str1(x) ialias_str2(x)
# define ialias_str2(x) #x
@@ -378,6 +378,12 @@ OACC_2.0 {
acc_set_cuda_stream;
};
+OACC_2.0.1 {
+ global:
+ acc_pcopyin;
+ acc_pcreate;
+} OACC_2.0;
+
OACC_2.5 {
global:
acc_copyin_async;
@@ -430,7 +436,7 @@ OACC_2.5 {
acc_update_self_async_32_h_;
acc_update_self_async_64_h_;
acc_update_self_async_array_h_;
-} OACC_2.0;
+} OACC_2.0.1;
GOACC_2.0 {
global:
@@ -654,15 +654,23 @@ acc_create_async (void *h, size_t s, int async)
present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s, async);
}
+/* acc_present_or_create used to be what acc_create is now. */
+/* acc_pcreate is acc_present_or_create by a different name. */
#ifdef HAVE_ATTRIBUTE_ALIAS
-extern void *acc_present_or_create (void *, size_t)
- __attribute__((alias ("acc_create")));
+strong_alias (acc_create, acc_present_or_create)
+strong_alias (acc_create, acc_pcreate)
#else
void *
acc_present_or_create (void *h, size_t s)
{
return acc_create (h, s);
}
+
+void *
+acc_pcreate (void *h, size_t s)
+{
+ return acc_create (h, s);
+}
#endif
void *
@@ -678,15 +686,23 @@ acc_copyin_async (void *h, size_t s, int async)
present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s, async);
}
+/* acc_present_or_copyin used to be what acc_copyin is now. */
+/* acc_pcopyin is acc_present_or_copyin by a different name. */
#ifdef HAVE_ATTRIBUTE_ALIAS
-extern void *acc_present_or_copyin (void *, size_t)
- __attribute__((alias ("acc_copyin")));
+strong_alias (acc_copyin, acc_present_or_copyin)
+strong_alias (acc_copyin, acc_pcopyin)
#else
void *
acc_present_or_copyin (void *h, size_t s)
{
return acc_copyin (h, s);
}
+
+void *
+acc_pcopyin (void *h, size_t s)
+{
+ return acc_copyin (h, s);
+}
#endif
#define FLAG_COPYOUT (1 << 0)
@@ -94,8 +94,10 @@ void acc_free (void *) __GOACC_NOTHROW;
the standard specifies otherwise. */
void *acc_copyin (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_copyin (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcopyin (void *, size_t) __GOACC_NOTHROW;
void *acc_create (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_create (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcreate (void *, size_t) __GOACC_NOTHROW;
void acc_copyout (void *, size_t) __GOACC_NOTHROW;
void acc_delete (void *, size_t) __GOACC_NOTHROW;
void acc_update_device (void *, size_t) __GOACC_NOTHROW;
@@ -124,11 +126,6 @@ void acc_copyout_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
void acc_delete_finalize (void *, size_t) __GOACC_NOTHROW;
void acc_delete_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
-/* Old names. OpenACC does not specify whether these can or must
- not be macros, inlines or aliases for the new names. */
-#define acc_pcreate acc_present_or_create
-#define acc_pcopyin acc_present_or_copyin
-
/* CUDA-specific routines. */
void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
@@ -1,36 +1,245 @@
-/* { dg-do run } */
+/* acc_present_or_create, acc_present_or_copyin, etc. */
+#include <stdbool.h>
#include <stdlib.h>
#include <openacc.h>
int
main (int argc, char **argv)
{
- const int N = 256;
- unsigned char *h;
- void *d1, *d2;
+ int *h, *d;
+ const int N = 10000;
+ const int S = N * sizeof *h;
+ bool shared_mem;
- h = (unsigned char *) malloc (N);
-
- d1 = acc_present_or_create (h, N);
- if (!d1)
+ h = (int *) malloc (S);
+ if (!h)
abort ();
+ for (int i = 0; i < N; ++i)
+ h[i] = i + 0;
- d2 = acc_present_or_create (h, N);
- if (!d2)
- abort ();
+ shared_mem = acc_is_present (h, S);
- if (d1 != d2)
+ d = (int *) acc_present_or_create (h, S);
+ if (!d)
abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
+ abort ();
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ d[i] = i + 1;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 1 : 0))
+ abort ();
+ h[i] = i + 2;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_create (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 2 : 1))
+ abort ();
+ d[i] = i + 3;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 3 : 2))
+ abort ();
+ h[i] = i + 4;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 4 : 3))
+ abort ();
+ d[i] = i + 5;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 5 : 4))
+ abort ();
+ h[i] = i + 6;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_copyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 6 : 5))
+ abort ();
+ d[i] = i + 7;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 7 : 6))
+ abort ();
+ h[i] = i + 8;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 8 : 7))
+ abort ();
+ d[i] = i + 9;
+ }
- d2 = acc_pcreate (h, N);
- if (!d2)
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 9 : 8))
+ abort ();
+ h[i] = i + 10;
+ }
+
+ acc_copyout_finalize (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort ();
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ }
+
+ d = (int *) acc_pcopyin (h, S);
+ if (!d)
+ abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
abort ();
- if (d1 != d2)
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ d[i] = i + 11;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 11 : 9))
+ abort ();
+ h[i] = i + 12;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 12 : 11))
+ abort ();
+ d[i] = i + 13;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 13 : 12))
+ abort ();
+ h[i] = i + 14;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 14 : 13))
+ abort ();
+ d[i] = i + 15;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 15 : 14))
+ abort ();
+ h[i] = i + 16;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 16 : 15))
+ abort ();
+ d[i] = i + 17;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 17 : 16))
+ abort ();
+ h[i] = i + 18;
+ }
+
+ acc_update_self (h, S);
+ if (!acc_is_present (h, S))
abort ();
- acc_delete (h, N);
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 18 : 17))
+ abort ();
+ }
+
+ acc_delete_finalize (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort();
free (h);
deleted file mode 100644
@@ -1,64 +0,0 @@
-/* { dg-do run } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <string.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- int i;
- unsigned char *h;
- void *d1, *d2;
-
- h = (unsigned char *) malloc (N);
-
- for (i = 0; i < N; i++)
- {
- h[i] = i;
- }
-
- d1 = acc_present_or_copyin (h, N);
- if (!d1)
- abort ();
-
- for (i = 0; i < N; i++)
- {
- h[i] = 0xab;
- }
-
- d2 = acc_present_or_copyin (h, N);
- if (!d2)
- abort ();
-
- if (d1 != d2)
- abort ();
-
- memset (&h[0], 0, N);
-
- acc_copyout_finalize (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- d2 = acc_pcopyin (h, N);
- if (!d2)
- abort ();
-
- acc_copyout (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- free (h);
-
- return 0;
-}