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.
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");
#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
acc_set_cuda_stream;
};
+OACC_2.0.1 {
+ global:
+ acc_pcopyin;
+ acc_pcreate;
+} OACC_2.0;
+
GOACC_2.0 {
global:
GOACC_data_end;
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
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;
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;
-/* { 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);
+++ /dev/null
-/* { 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;
-}
projects / gcc.git / commitdiff