[OpenACC 'exit data'] Evaluate 'copyfrom' individually for 'GOMP_MAP_STRUCT' entries
Currently, we don't at all evaluate 'copyfrom' for 'GOMP_MAP_STRUCT' entries. Fix this by copying/adapting the corresponding non-'GOMP_MAP_STRUCT' code. libgomp/ * oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>: Evaluate 'copyfrom' individually for each entry. * testsuite/libgomp.oacc-c-c++-common/struct-1.c: Update.
This commit is contained in:
parent
a02f1adbfe
commit
2c838a3e4e
|
@ -1194,6 +1194,12 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
|
|||
|| kind == GOMP_MAP_FORCE_DETACH)
|
||||
finalize = true;
|
||||
|
||||
copyfrom = false;
|
||||
if (kind == GOMP_MAP_FROM
|
||||
|| kind == GOMP_MAP_FORCE_FROM
|
||||
|| kind == GOMP_MAP_ALWAYS_FROM)
|
||||
copyfrom = true;
|
||||
|
||||
struct splay_tree_key_s k;
|
||||
k.host_start = (uintptr_t) hostaddrs[i + j];
|
||||
k.host_end = k.host_start + sizes[i + j];
|
||||
|
@ -1216,6 +1222,16 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
|
|||
else if (str->refcount > 0
|
||||
&& str->refcount != REFCOUNT_INFINITY)
|
||||
str->refcount--;
|
||||
|
||||
if (copyfrom
|
||||
&& (kind != GOMP_MAP_FROM || str->refcount == 0))
|
||||
gomp_copy_dev2host (acc_dev, aq, (void *) k.host_start,
|
||||
(void *) (str->tgt->tgt_start
|
||||
+ str->tgt_offset
|
||||
+ k.host_start
|
||||
- str->host_start),
|
||||
k.host_end - k.host_start);
|
||||
|
||||
if (str->refcount == 0)
|
||||
{
|
||||
if (aq)
|
||||
|
|
|
@ -1,4 +1,4 @@
|
|||
/* Test dynamic refcount of separate structure members. */
|
||||
/* Test dynamic refcount and copy behavior of separate structure members. */
|
||||
|
||||
#include <assert.h>
|
||||
#include <stdbool.h>
|
||||
|
@ -12,41 +12,45 @@ struct s
|
|||
|
||||
static void test(unsigned variant)
|
||||
{
|
||||
struct s s;
|
||||
struct s s = { .a = 73, .b = -22 };
|
||||
|
||||
#pragma acc enter data create(s.a, s.b)
|
||||
#pragma acc enter data copyin(s.a, s.b)
|
||||
assert(acc_is_present(&s.a, sizeof s.a));
|
||||
assert(acc_is_present(&s.b, sizeof s.b));
|
||||
|
||||
/* To verify that any following 'copyin' doesn't 'copyin' again. */
|
||||
s.a = -s.a;
|
||||
s.b = -s.b;
|
||||
|
||||
if (variant & 4)
|
||||
{
|
||||
if (variant & 8)
|
||||
{
|
||||
#pragma acc enter data create(s.b)
|
||||
#pragma acc enter data copyin(s.b)
|
||||
}
|
||||
else
|
||||
acc_create(&s.b, sizeof s.b);
|
||||
acc_copyin(&s.b, sizeof s.b);
|
||||
assert(acc_is_present(&s.a, sizeof s.a));
|
||||
assert(acc_is_present(&s.b, sizeof s.b));
|
||||
|
||||
if (variant & 16)
|
||||
{
|
||||
#pragma acc enter data create(s.a)
|
||||
#pragma acc enter data copyin(s.a)
|
||||
}
|
||||
else
|
||||
acc_create(&s.a, sizeof s.a);
|
||||
acc_copyin(&s.a, sizeof s.a);
|
||||
assert(acc_is_present(&s.a, sizeof s.a));
|
||||
assert(acc_is_present(&s.b, sizeof s.b));
|
||||
|
||||
if (variant & 32)
|
||||
{
|
||||
#pragma acc enter data create(s.a)
|
||||
acc_create(&s.b, sizeof s.b);
|
||||
#pragma acc enter data create(s.b)
|
||||
#pragma acc enter data create(s.b)
|
||||
acc_create(&s.a, sizeof s.a);
|
||||
acc_create(&s.a, sizeof s.a);
|
||||
acc_create(&s.a, sizeof s.a);
|
||||
#pragma acc enter data copyin(s.a)
|
||||
acc_copyin(&s.b, sizeof s.b);
|
||||
#pragma acc enter data copyin(s.b)
|
||||
#pragma acc enter data copyin(s.b)
|
||||
acc_copyin(&s.a, sizeof s.a);
|
||||
acc_copyin(&s.a, sizeof s.a);
|
||||
acc_copyin(&s.a, sizeof s.a);
|
||||
}
|
||||
assert(acc_is_present(&s.a, sizeof s.a));
|
||||
assert(acc_is_present(&s.b, sizeof s.b));
|
||||
|
@ -55,85 +59,122 @@ static void test(unsigned variant)
|
|||
#pragma acc parallel \
|
||||
copy(s.a, s.b)
|
||||
{
|
||||
#if ACC_MEM_SHARED
|
||||
if (s.a++ != -73)
|
||||
__builtin_abort();
|
||||
if (s.b-- != 22)
|
||||
__builtin_abort();
|
||||
#else
|
||||
if (s.a++ != 73)
|
||||
__builtin_abort();
|
||||
if (s.b-- != -22)
|
||||
__builtin_abort();
|
||||
#endif
|
||||
}
|
||||
#if ACC_MEM_SHARED
|
||||
assert(s.a == -72);
|
||||
assert(s.b == 21);
|
||||
#else
|
||||
assert(s.a == -73);
|
||||
assert(s.b == 22);
|
||||
#endif
|
||||
|
||||
if (variant & 32)
|
||||
{
|
||||
if (variant & 1)
|
||||
{
|
||||
#pragma acc exit data delete(s.a) finalize
|
||||
#pragma acc exit data copyout(s.a) finalize
|
||||
}
|
||||
else
|
||||
acc_delete_finalize(&s.a, sizeof s.a);
|
||||
acc_copyout_finalize(&s.a, sizeof s.a);
|
||||
}
|
||||
else
|
||||
{
|
||||
if (variant & 1)
|
||||
{
|
||||
#pragma acc exit data delete(s.a)
|
||||
#pragma acc exit data copyout(s.a)
|
||||
}
|
||||
else
|
||||
acc_delete(&s.a, sizeof s.a);
|
||||
acc_copyout(&s.a, sizeof s.a);
|
||||
if (variant & 4)
|
||||
{
|
||||
assert(acc_is_present(&s.a, sizeof s.a));
|
||||
assert(acc_is_present(&s.b, sizeof s.b));
|
||||
#if ACC_MEM_SHARED
|
||||
assert(s.a == -72);
|
||||
assert(s.b == 21);
|
||||
#else
|
||||
assert(s.a == -73);
|
||||
assert(s.b == 22);
|
||||
#endif
|
||||
if (variant & 1)
|
||||
{
|
||||
#pragma acc exit data delete(s.a)
|
||||
#pragma acc exit data copyout(s.a)
|
||||
}
|
||||
else
|
||||
acc_delete(&s.a, sizeof s.a);
|
||||
acc_copyout(&s.a, sizeof s.a);
|
||||
}
|
||||
}
|
||||
#if ACC_MEM_SHARED
|
||||
assert(acc_is_present(&s.a, sizeof s.a));
|
||||
assert(acc_is_present(&s.b, sizeof s.b));
|
||||
assert(s.a == -72);
|
||||
assert(s.b == 21);
|
||||
#else
|
||||
assert(!acc_is_present(&s.a, sizeof s.a));
|
||||
assert(acc_is_present(&s.b, sizeof s.b));
|
||||
assert(s.a == 74);
|
||||
assert(s.b == 22);
|
||||
#endif
|
||||
|
||||
if (variant & 32)
|
||||
{
|
||||
if (variant & 2)
|
||||
{
|
||||
#pragma acc exit data delete(s.b) finalize
|
||||
#pragma acc exit data copyout(s.b) finalize
|
||||
}
|
||||
else
|
||||
acc_delete_finalize(&s.b, sizeof s.b);
|
||||
acc_copyout_finalize(&s.b, sizeof s.b);
|
||||
}
|
||||
else
|
||||
{
|
||||
if (variant & 2)
|
||||
{
|
||||
#pragma acc exit data delete(s.b)
|
||||
#pragma acc exit data copyout(s.b)
|
||||
}
|
||||
else
|
||||
acc_delete(&s.b, sizeof s.b);
|
||||
acc_copyout(&s.b, sizeof s.b);
|
||||
if (variant & 4)
|
||||
{
|
||||
#if ACC_MEM_SHARED
|
||||
assert(acc_is_present(&s.a, sizeof s.a));
|
||||
assert(acc_is_present(&s.b, sizeof s.b));
|
||||
assert(s.a == -72);
|
||||
assert(s.b == 21);
|
||||
#else
|
||||
assert(!acc_is_present(&s.a, sizeof s.a));
|
||||
assert(acc_is_present(&s.b, sizeof s.b));
|
||||
assert(s.a == 74);
|
||||
assert(s.b == 22);
|
||||
#endif
|
||||
if (variant & 2)
|
||||
{
|
||||
#pragma acc exit data delete(s.b)
|
||||
#pragma acc exit data copyout(s.b)
|
||||
}
|
||||
else
|
||||
acc_delete(&s.b, sizeof s.b);
|
||||
acc_copyout(&s.b, sizeof s.b);
|
||||
}
|
||||
}
|
||||
#if ACC_MEM_SHARED
|
||||
assert(acc_is_present(&s.a, sizeof s.a));
|
||||
assert(acc_is_present(&s.b, sizeof s.b));
|
||||
assert(s.a == -72);
|
||||
assert(s.b == 21);
|
||||
#else
|
||||
assert(!acc_is_present(&s.a, sizeof s.a));
|
||||
assert(!acc_is_present(&s.b, sizeof s.b));
|
||||
assert(s.a == 74);
|
||||
assert(s.b == -23);
|
||||
#endif
|
||||
}
|
||||
|
||||
|
|
Loading…
Reference in New Issue