@@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
break;
case GOMP_MAP_STRUCT:
- {
- int elems = sizes[i];
- for (int j = 1; j <= elems; j++)
- {
- struct splay_tree_key_s k;
- k.host_start = (uintptr_t) hostaddrs[i + j];
- k.host_end = k.host_start + sizes[i + j];
- splay_tree_key str;
- str = splay_tree_lookup (&acc_dev->mem_map, &k);
- if (str)
- {
- if (finalize)
- {
- if (str->refcount != REFCOUNT_INFINITY)
- str->refcount -= str->virtual_refcount;
- str->virtual_refcount = 0;
- }
- if (str->virtual_refcount > 0)
- {
- if (str->refcount != REFCOUNT_INFINITY)
- str->refcount--;
- str->virtual_refcount--;
- }
- else if (str->refcount > 0
- && str->refcount != REFCOUNT_INFINITY)
- str->refcount--;
- if (str->refcount == 0)
- gomp_remove_var_async (acc_dev, str, aq);
- }
- }
- i += elems;
- }
break;
default:
new file mode 100644
@@ -0,0 +1,38 @@
+#include <assert.h>
+
+struct str1 {
+ int a;
+ int b;
+};
+
+struct str2 {
+ int c;
+ int d;
+ struct str1 s;
+};
+
+int
+main (int argc, char *argv[])
+{
+ struct str2 t;
+
+ t.c = 1;
+ t.d = 2;
+ t.s.a = 3;
+ t.s.b = 4;
+
+ #pragma acc enter data copyin(t.s)
+
+ #pragma acc serial present(t.s) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+ {
+ t.s.a = 5;
+ t.s.b = 6;
+ }
+
+ #pragma acc exit data copyout(t.s)
+
+ assert (t.s.a == 5);
+ assert (t.s.b == 6);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,44 @@
+#include <assert.h>
+#include <stdlib.h>
+
+struct str1 {
+ int a;
+ int b;
+ int *c;
+};
+
+#define N 1024
+
+int
+main (int argc, char *argv[])
+{
+ struct str1 s;
+
+ s.a = 1;
+ s.b = 2;
+ s.c = (int *) malloc (sizeof (int) * N);
+
+ for (int i = 0; i < N; i++)
+ s.c[i] = i + 10;
+
+ #pragma acc enter data copyin(s.a, s.b, s.c[0:N])
+
+ #pragma acc serial present(s.a, s.b, s.c[0:N]) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+ {
+ s.a = 3;
+ s.b = 4;
+ for (int i = 0; i < N; i++)
+ s.c[i] = i + 20;
+ }
+
+ #pragma acc exit data copyout(s.a, s.b, s.c[0:N])
+
+ assert (s.a == 3);
+ assert (s.b == 4);
+ for (int i = 0; i < N; i++)
+ assert (s.c[i] == i + 20);
+
+ free (s.c);
+
+ return 0;
+}