diff options
Diffstat (limited to 'libgomp')
| -rw-r--r-- | libgomp/ChangeLog | 15 | ||||
| -rw-r--r-- | libgomp/target.c | 23 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c | 49 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c | 30 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c | 25 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c | 82 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c | 49 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 | 39 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 | 90 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90 | 39 |
10 files changed, 441 insertions, 0 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 779f276..e1f6bb4 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,18 @@ +2019-12-19 Julian Brown <julian@codesourcery.com> + Maciej W. Rozycki <macro@codesourcery.com> + Tobias Burnus <tobias@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + + * target.c (gomp_map_vars_async): Support GOMP_MAP_NO_ALLOC. + * testsuite/libgomp.oacc-c-c++-common/no_create-1.c: New test. + * testsuite/libgomp.oacc-c-c++-common/no_create-2.c: New test. + * testsuite/libgomp.oacc-c-c++-common/no_create-3.c: New test. + * testsuite/libgomp.oacc-c-c++-common/no_create-4.c: New test. + * testsuite/libgomp.oacc-c-c++-common/no_create-5.c: New test. + * testsuite/libgomp.oacc-fortran/no_create-1.f90: New test. + * testsuite/libgomp.oacc-fortran/no_create-2.f90: New test. + * testsuite/libgomp.oacc-fortran/no_create-3.F90: New test. + 2019-12-18 Thomas Schwinge <thomas@codesourcery.com> * oacc-mem.c (goacc_enter_data): Refactor, so that it can be diff --git a/libgomp/target.c b/libgomp/target.c index 67cd80a..d83b353d 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -707,6 +707,21 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, { tgt->list[i].key = NULL; + if ((kind & typemask) == GOMP_MAP_IF_PRESENT) + { + /* Not present, hence, skip entry - including its MAP_POINTER, + when existing. */ + tgt->list[i].offset = 0; + if (i + 1 < mapnum + && ((typemask & get_kind (short_mapkind, kinds, i + 1)) + == GOMP_MAP_POINTER)) + { + ++i; + tgt->list[i].key = NULL; + tgt->list[i].offset = 0; + } + continue; + } size_t align = (size_t) 1 << (kind >> rshift); not_found_cnt++; if (tgt_align < align) @@ -893,6 +908,14 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start - n->host_start; continue; + case GOMP_MAP_IF_PRESENT: + /* Not present - otherwise handled above. Skip over its + MAP_POINTER as well. */ + if (i + 1 < mapnum + && ((typemask & get_kind (short_mapkind, kinds, i + 1)) + == GOMP_MAP_POINTER)) + ++i; + continue; default: break; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c new file mode 100644 index 0000000..22e0c20 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c @@ -0,0 +1,49 @@ +/* Test 'no_create' clause on compute construct, with data present on the + device. */ + +#include <stdlib.h> +#include <stdio.h> +#include <openacc.h> + +#define N 128 + +int +main (int argc, char *argv[]) +{ + int var; + int *arr = (int *) malloc (N * sizeof (*arr)); + int *devptr[2]; + + acc_copyin (&var, sizeof (var)); + acc_copyin (arr, N * sizeof (*arr)); + +#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr) + { + devptr[0] = &var; + devptr[1] = &arr[2]; + } + + if (acc_hostptr (devptr[0]) != (void *) &var) + __builtin_abort (); + if (acc_hostptr (devptr[1]) != (void *) &arr[2]) + __builtin_abort (); + + acc_delete (&var, sizeof (var)); + acc_delete (arr, N * sizeof (*arr)); + +#if ACC_MEM_SHARED + if (devptr[0] != &var) + __builtin_abort (); + if (devptr[1] != &arr[2]) + __builtin_abort (); +#else + if (devptr[0] == &var) + __builtin_abort (); + if (devptr[1] == &arr[2]) + __builtin_abort (); +#endif + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c new file mode 100644 index 0000000..fbd01a2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c @@ -0,0 +1,30 @@ +/* Test 'no_create' clause on compute construct, with data not present on the + device. */ + +#include <stdlib.h> +#include <stdio.h> + +#define N 128 + +int +main (int argc, char *argv[]) +{ + int var; + int *arr = (int *) malloc (N * sizeof (*arr)); + int *devptr[2]; + +#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr) + { + devptr[0] = &var; + devptr[1] = &arr[2]; + } + + if (devptr[0] != &var) + __builtin_abort (); + if (devptr[1] != &arr[2]) + __builtin_abort (); + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c new file mode 100644 index 0000000..18466b8 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c @@ -0,0 +1,25 @@ +#include <float.h> /* For FLT_EPSILON. */ +#include <math.h> /* For fabs. */ +#include <stdlib.h> /* For abort. */ + + +int main() +{ +#define N 100 + float b[N]; + float c[N]; + +#pragma acc enter data create(b) + +#pragma acc parallel loop no_create(b) no_create(c) + for (int i = 0; i < N; ++i) + b[i] = i; + +#pragma acc exit data copyout(b) + + for (int i = 0; i < N; ++i) + if (fabs (b[i] - i) > 10.0*FLT_EPSILON) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c new file mode 100644 index 0000000..963cb3a --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c @@ -0,0 +1,82 @@ +/* Test 'no_create' clause on 'data' construct and nested compute construct, + with data present on the device. */ + +#include <stdlib.h> +#include <stdio.h> +#include <openacc.h> + +#define N 128 + +int +main (int argc, char *argv[]) +{ + int var; + int *arr = (int *) malloc (N * sizeof (*arr)); + int *devptr[2]; + + acc_copyin (&var, sizeof (var)); + acc_copyin (arr, N * sizeof (*arr)); + +#pragma acc data no_create(var, arr[0:N]) + { + devptr[0] = (int *) acc_deviceptr (&var); + devptr[1] = (int *) acc_deviceptr (&arr[2]); + + if (devptr[0] == NULL) + __builtin_abort (); + if (devptr[1] == NULL) + __builtin_abort (); + + if (acc_hostptr (devptr[0]) != (void *) &var) + __builtin_abort (); + if (acc_hostptr (devptr[1]) != (void *) &arr[2]) + __builtin_abort (); + +#if ACC_MEM_SHARED + if (devptr[0] != &var) + __builtin_abort (); + if (devptr[1] != &arr[2]) + __builtin_abort (); +#else + if (devptr[0] == &var) + __builtin_abort (); + if (devptr[1] == &arr[2]) + __builtin_abort (); +#endif + +#pragma acc parallel copyout(devptr) + { + devptr[0] = &var; + devptr[1] = &arr[2]; + } + + if (devptr[0] == NULL) + __builtin_abort (); + if (devptr[1] == NULL) + __builtin_abort (); + + if (acc_hostptr (devptr[0]) != (void *) &var) + __builtin_abort (); + if (acc_hostptr (devptr[1]) != (void *) &arr[2]) + __builtin_abort (); + +#if ACC_MEM_SHARED + if (devptr[0] != &var) + __builtin_abort (); + if (devptr[1] != &arr[2]) + __builtin_abort (); +#else + if (devptr[0] == &var) + __builtin_abort (); + if (devptr[1] == &arr[2]) + __builtin_abort (); +#endif + } + + acc_delete (&var, sizeof (var)); + acc_delete (arr, N * sizeof (*arr)); + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c new file mode 100644 index 0000000..6f0ace5 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c @@ -0,0 +1,49 @@ +/* Test 'no_create' clause on 'data' construct and nested compute construct, + with data not present on the device. */ + +#include <stdlib.h> +#include <stdio.h> +#include <openacc.h> + +#define N 128 + +int +main (int argc, char *argv[]) +{ + int var; + int *arr = (int *) malloc (N * sizeof (*arr)); + int *devptr[2]; + +#pragma acc data no_create(var, arr[0:N]) + { + devptr[0] = (int *) acc_deviceptr (&var); + devptr[1] = (int *) acc_deviceptr (&arr[2]); + +#if ACC_MEM_SHARED + if (devptr[0] == NULL) + __builtin_abort (); + if (devptr[1] == NULL) + __builtin_abort (); +#else + if (devptr[0] != NULL) + __builtin_abort (); + if (devptr[1] != NULL) + __builtin_abort (); +#endif + +#pragma acc parallel copyout(devptr) // TODO implicit 'copy(var)' -- huh?! + { + devptr[0] = &var; + devptr[1] = &arr[2]; + } + + if (devptr[0] != &var) + __builtin_abort (); // { dg-xfail-run-if "TODO" { *-*-* } { "-DACC_MEM_SHARED=0" } } + if (devptr[1] != &arr[2]) + __builtin_abort (); + } + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 new file mode 100644 index 0000000..4a1d5da --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 @@ -0,0 +1,39 @@ +! { dg-do run } + +! Test no_create clause with data construct when data is present/not present. + +program no_create + use openacc + implicit none + logical :: shared_memory + integer, parameter :: n = 512 + integer :: myvar, myarr(n) + integer i + + shared_memory = .false. + !$acc kernels copyin (shared_memory) + shared_memory = .true. + !$acc end kernels + + myvar = 77 + do i = 1, n + myarr(i) = 0 + end do + + !$acc data no_create (myvar, myarr) + if (acc_is_present (myvar) .neqv. shared_memory) stop 10 + if (acc_is_present (myarr) .neqv. shared_memory) stop 11 + !$acc end data + + !$acc enter data copyin (myvar, myarr) + !$acc data no_create (myvar, myarr) + if (acc_is_present (myvar) .eqv. .false.) stop 20 + if (acc_is_present (myarr) .eqv. .false.) stop 21 + !$acc end data + !$acc exit data copyout (myvar, myarr) + + if (myvar .ne. 77) stop 30 + do i = 1, n + if (myarr(i) .ne. 0) stop 31 + end do +end program no_create diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 new file mode 100644 index 0000000..0b11f45 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 @@ -0,0 +1,90 @@ +! { dg-do run } + +! Test no_create clause with data/parallel constructs. + +program no_create + use openacc + implicit none + logical :: shared_memory + integer, parameter :: n = 512 + integer :: myvar, myarr(n) + integer i + + shared_memory = .false. + !$acc kernels copyin (shared_memory) + shared_memory = .true. + !$acc end kernels + + myvar = 55 + do i = 1, n + myarr(i) = 0 + end do + + call do_on_target(myvar, n, myarr) + + if (shared_memory) then + if (myvar .ne. 44) stop 10 + else + if (myvar .ne. 33) stop 11 + end if + do i = 1, n + if (shared_memory) then + if (myarr(i) .ne. i * 2) stop 20 + else + if (myarr(i) .ne. i) stop 21 + end if + end do + + myvar = 55 + do i = 1, n + myarr(i) = 0 + end do + + !$acc enter data copyin(myvar, myarr) + call do_on_target(myvar, n, myarr) + !$acc exit data copyout(myvar, myarr) + + if (myvar .ne. 44) stop 30 + do i = 1, n + if (myarr(i) .ne. i * 2) stop 31 + end do +end program no_create + +subroutine do_on_target (var, n, arr) + use openacc + implicit none + integer :: var, n, arr(n) + integer :: i + +!$acc data no_create (var, arr) + +if (acc_is_present(var)) then + ! The no_create clause is meant for partially shared-memory machines. This + ! test is written to work on non-shared-memory machines, though this is not + ! necessarily a useful way to use the no_create clause in practice. + + !$acc parallel !no_create (var) + var = 44 + !$acc end parallel +else + var = 33 +end if +if (acc_is_present(arr)) then + ! The no_create clause is meant for partially shared-memory machines. This + ! test is written to work on non-shared-memory machines, though this is not + ! necessarily a useful way to use the no_create clause in practice. + + !$acc parallel loop !no_create (arr) + do i = 1, n + arr(i) = i * 2 + end do + !$acc end parallel loop +else + do i = 1, n + arr(i) = i + end do +end if + +!$acc end data + +end subroutine do_on_target diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90 new file mode 100644 index 0000000..4362688 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90 @@ -0,0 +1,39 @@ +! { dg-do run } + +program main + use iso_c_binding, only: c_sizeof + use openacc, only: acc_is_present + implicit none + integer i + integer, parameter :: n = 100 + real*4 b(n), c(n) + real :: d(n), e(n) + common /BLOCK/ d, e + + !$acc enter data create(b) create(d) + + if (.not. acc_is_present(b, c_sizeof(b))) stop 1 + if (.not. acc_is_present(d, c_sizeof(d))) stop 2 +#if !ACC_MEM_SHARED + if (acc_is_present(c, 1) .or. acc_is_present(c, c_sizeof(c))) stop 3 + if (acc_is_present(e, 1) .or. acc_is_present(e, c_sizeof(d))) stop 4 +#endif + + !$acc parallel loop no_create(b) no_create(c) no_create(/BLOCK/) + do i = 1, n + b(i) = i + d(i) = -i + end do + !$acc end parallel loop + + if (.not. acc_is_present(b, c_sizeof(b))) stop 5 + if (.not. acc_is_present(d, c_sizeof(d))) stop 6 +#if !ACC_MEM_SHARED + if (acc_is_present(c, 1) .or. acc_is_present(c, c_sizeof(c))) stop 7 + if (acc_is_present(e, 1) .or. acc_is_present(e, c_sizeof(e))) stop 8 +#endif + + !$acc exit data copyout(b) copyout(d) + if (any(abs(b - [(real(i), i = 1, n)]) > 10*epsilon(b))) stop 9 + if (any(abs(d - [(real(-i), i = 1, n)]) > 10*epsilon(d))) stop 10 +end program main |