From af8fd1a99d9a21f8088ebb11250cd06a3f275052 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Fri, 29 May 2020 14:11:27 +0200 Subject: [PATCH] Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854] libgomp/ PR libgomp/92854 * testsuite/libgomp.oacc-c-c++-common/pr92854-1.c: Extend some more. --- .../libgomp.oacc-c-c++-common/pr92854-1.c | 64 ++++++++++++++----- 1 file changed, 47 insertions(+), 17 deletions(-) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c index 6ba96b6bf8f..79cebf65c34 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c @@ -1,31 +1,61 @@ -/* Verify that 'acc_unmap_data' unmaps even in presence of dynamic reference - counts. */ +/* Verify that 'acc_unmap_data' unmaps even in presence of structured and + dynamic reference counts, but the device memory remains allocated. */ /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ #include #include +#include #include int main () { const int N = 180; - - char *h = (char *) malloc (N); - char *d = (char *) acc_malloc (N); - if (!d) - abort (); - acc_map_data (h, d, N); - - char *d_ = (char *) acc_create (h + 3, N - 77); - assert (d_ == d + 3); - - d_ = (char *) acc_create (h, N); - assert (d_ == d); - - acc_unmap_data (h); - assert (!acc_is_present (h, N)); + const int N_i = 537; + const int C = 37; + + unsigned char *h = (unsigned char *) malloc (N); + assert (h); + unsigned char *d = (unsigned char *) acc_malloc (N); + assert (d); + + for (int i = 0; i < N_i; ++i) + { + acc_map_data (h, d, N); + assert (acc_is_present (h, N)); +#pragma acc parallel present(h[0:N]) + { + if (i == 0) + memset (h, C, N); + } + + unsigned char *d_ = (unsigned char *) acc_create (h + 3, N - 77); + assert (d_ == d + 3); + +#pragma acc data create(h[6:N - 44]) + { + d_ = (unsigned char *) acc_create (h, N); + assert (d_ == d); + +#pragma acc enter data create(h[0:N]) + + assert (acc_is_present (h, N)); + acc_unmap_data (h); + assert (!acc_is_present (h, N)); + } + + /* We can however still access the device memory. */ +#pragma acc parallel loop deviceptr(d) + for (int j = 0; j < N; ++j) + d[j] += i * j; + } + + acc_memcpy_from_device(h, d, N); + for (int j = 0; j < N; ++j) + assert (h[j] == ((C + N_i * (N_i - 1) / 2 * j) % 256)); + + acc_free (d); return 0; } -- 2.30.2