diff options
author | Thomas Schwinge <thomas@codesourcery.com> | 2017-05-19 15:32:48 +0200 |
---|---|---|
committer | Thomas Schwinge <tschwinge@gcc.gnu.org> | 2017-05-19 15:32:48 +0200 |
commit | 7fd549d24fda05c859fb17697c51c16886902dad (patch) | |
tree | a1ddeff9f474c8969d68b1b0b726f4ca763411f3 /libgomp | |
parent | 0d0afa9fafec1711b52259b9b8caebf51380216d (diff) | |
download | gcc-7fd549d24fda05c859fb17697c51c16886902dad.zip gcc-7fd549d24fda05c859fb17697c51c16886902dad.tar.gz gcc-7fd549d24fda05c859fb17697c51c16886902dad.tar.bz2 |
OpenACC 2.5 default (present) clause
gcc/c/
* c-parser.c (c_parser_omp_clause_default): Handle
"OMP_CLAUSE_DEFAULT_PRESENT".
gcc/cp/
* parser.c (cp_parser_omp_clause_default): Handle
"OMP_CLAUSE_DEFAULT_PRESENT".
gcc/fortran/
* gfortran.h (enum gfc_omp_default_sharing): Add
"OMP_DEFAULT_PRESENT".
* dump-parse-tree.c (show_omp_clauses): Handle it.
* openmp.c (gfc_match_omp_clauses): Likewise.
* trans-openmp.c (gfc_trans_omp_clauses): Likewise.
gcc/
* tree-core.h (enum omp_clause_default_kind): Add
"OMP_CLAUSE_DEFAULT_PRESENT".
* tree-pretty-print.c (dump_omp_clause): Handle it.
* gimplify.c (enum gimplify_omp_var_data): Add
"GOVD_MAP_FORCE_PRESENT".
(gimplify_adjust_omp_clauses_1): Map it to
"GOMP_MAP_FORCE_PRESENT".
(oacc_default_clause): Handle "OMP_CLAUSE_DEFAULT_PRESENT".
gcc/testsuite/
* c-c++-common/goacc/default-1.c: Update.
* c-c++-common/goacc/default-2.c: Likewise.
* c-c++-common/goacc/default-4.c: Likewise.
* gfortran.dg/goacc/default-1.f95: Likewise.
* gfortran.dg/goacc/default-4.f: Likewise.
* c-c++-common/goacc/default-5.c: New file.
* gfortran.dg/goacc/default-5.f: Likewise.
libgomp/
* testsuite/libgomp.oacc-c++/template-reduction.C: Update.
* testsuite/libgomp.oacc-c-c++-common/nested-2.c: Update.
* testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/default-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/non-scalar-data.f90: Likewise.
From-SVN: r248280
Diffstat (limited to 'libgomp')
-rw-r--r-- | libgomp/ChangeLog | 6 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c++/template-reduction.C | 25 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c | 31 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90 | 21 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-fortran/default-1.f90 | 10 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90 | 44 |
6 files changed, 125 insertions, 12 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 84d1c839..460605b 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,5 +1,11 @@ 2017-05-19 Thomas Schwinge <thomas@codesourcery.com> + * testsuite/libgomp.oacc-c++/template-reduction.C: Update. + * testsuite/libgomp.oacc-c-c++-common/nested-2.c: Update. + * testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise. + * testsuite/libgomp.oacc-fortran/default-1.f90: Likewise. + * testsuite/libgomp.oacc-fortran/non-scalar-data.f90: Likewise. + * plugin/plugin-hsa.c (DLSYM_FN, init_hsa_runtime_functions): Debug output for failure. diff --git a/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C b/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C index 6c85fba..ede0aae 100644 --- a/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C +++ b/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C @@ -32,6 +32,28 @@ sum () return s; } +// Check template with default (present) + +template<typename T> T +sum_default_present () +{ + T s = 0; + T array[n]; + + for (int i = 0; i < n; i++) + array[i] = i+1; + +#pragma acc enter data copyin (array) + +#pragma acc parallel loop num_gangs (10) gang reduction (+:s) default (present) + for (int i = 0; i < n; i++) + s += array[i]; + +#pragma acc exit data delete (array) + + return s; +} + // Check present and async template<typename T> T @@ -86,6 +108,9 @@ main() if (sum<int> () != result) __builtin_abort (); + if (sum_default_present<int> () != result) + __builtin_abort (); + #pragma acc enter data copyin (a) if (async_sum (a) != result) __builtin_abort (); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c index c164598..51b3b18 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c @@ -137,5 +137,36 @@ main (int argc, char *argv[]) abort (); } + +#pragma acc enter data create (a) + +#pragma acc parallel default (present) + { + for (int j = 0; j < N; ++j) + a[j] = j + 1; + } + +#pragma acc update host (a) + + for (i = 0; i < N; ++i) + { + if (a[i] != i + 1) + abort (); + } + +#pragma acc kernels default (present) + { + for (int j = 0; j < N; ++j) + a[j] = j + 2; + } + +#pragma acc exit data copyout (a) + + for (i = 0; i < N; ++i) + { + if (a[i] != i + 2) + abort (); + } + return 0; } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90 index 16a8598..df4aca0 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90 @@ -1,4 +1,5 @@ -! Copy of data-4.f90 with self exchanged with host for !acc update. +! Copy of data-4.f90 with self exchanged with host for !acc update, and with +! default (present) clauses added. ! { dg-do run } @@ -19,7 +20,7 @@ program asyncwait !$acc enter data copyin (a(1:N)) copyin (b(1:N)) copyin (N) async - !$acc parallel async wait + !$acc parallel default (present) async wait !$acc loop do i = 1, N b(i) = a(i) @@ -39,7 +40,7 @@ program asyncwait !$acc update device (a(1:N), b(1:N)) async (1) - !$acc parallel async (1) wait (1) + !$acc parallel default (present) async (1) wait (1) !$acc loop do i = 1, N b(i) = a(i) @@ -62,19 +63,19 @@ program asyncwait !$acc enter data copyin (c(1:N), d(1:N)) async (1) !$acc update device (a(1:N), b(1:N)) async (1) - !$acc parallel async (1) + !$acc parallel default (present) async (1) do i = 1, N b(i) = (a(i) * a(i) * a(i)) / a(i) end do !$acc end parallel - !$acc parallel async (1) + !$acc parallel default (present) async (1) do i = 1, N c(i) = (a(i) * 4) / a(i) end do !$acc end parallel - !$acc parallel async (1) + !$acc parallel default (present) async (1) do i = 1, N d(i) = ((a(i) * a(i) + a(i)) / a(i)) - a(i) end do @@ -100,25 +101,25 @@ program asyncwait !$acc enter data copyin (e(1:N)) async (1) !$acc update device (a(1:N), b(1:N), c(1:N), d(1:N)) async (1) - !$acc parallel async (1) + !$acc parallel default (present) async (1) do i = 1, N b(i) = (a(i) * a(i) * a(i)) / a(i) end do !$acc end parallel - !$acc parallel async (1) + !$acc parallel default (present) async (1) do i = 1, N c(i) = (a(i) * 4) / a(i) end do !$acc end parallel - !$acc parallel async (1) + !$acc parallel default (present) async (1) do i = 1, N d(i) = ((a(i) * a(i) + a(i)) / a(i)) - a(i) end do !$acc end parallel - !$acc parallel wait (1) async (1) + !$acc parallel default (present) wait (1) async (1) do i = 1, N e(i) = a(i) + b(i) + c(i) + d(i) end do diff --git a/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90 index 1059089..d8ceb60 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90 @@ -51,4 +51,14 @@ program main if (a .ne. 7.0) call abort + ! The default (present) clause doesn't affect scalar variables; these will + ! still get an implicit copy clause added. + !$acc kernels default (present) + c = a + a = 1.0 + a = a + c + !$acc end kernels + + if (a .ne. 8.0) call abort + end program main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90 b/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90 index 94e4228..53b2cd0 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90 @@ -1,5 +1,6 @@ ! Ensure that a non-scalar dummy arguments which are implicitly used inside -! offloaded regions are properly mapped using present_or_copy. +! offloaded regions are properly mapped using present_or_copy, or (default) +! present. ! { dg-do run } @@ -12,6 +13,7 @@ program main n = size !$acc data copy(array) + call kernels(array, n) !$acc update host(array) @@ -20,12 +22,29 @@ program main if (array(i) .ne. i) call abort end do + call kernels_default_present(array, n) + + !$acc update host(array) + + do i = 1, n + if (array(i) .ne. i+1) call abort + end do + call parallel(array, n) - !$acc end data + + !$acc update host(array) do i = 1, n if (array(i) .ne. i+i) call abort end do + + call parallel_default_present(array, n) + + !$acc end data + + do i = 1, n + if (array(i) .ne. i+i+1) call abort + end do end program main subroutine kernels (array, n) @@ -39,6 +58,16 @@ subroutine kernels (array, n) !$acc end kernels end subroutine kernels +subroutine kernels_default_present (array, n) + integer, dimension (n) :: array + integer :: n, i + + !$acc kernels default(present) + do i = 1, n + array(i) = i+1 + end do + !$acc end kernels +end subroutine kernels_default_present subroutine parallel (array, n) integer, dimension (n) :: array @@ -50,3 +79,14 @@ subroutine parallel (array, n) end do !$acc end parallel end subroutine parallel + +subroutine parallel_default_present (array, n) + integer, dimension (n) :: array + integer :: n, i + + !$acc parallel default(present) + do i = 1, n + array(i) = i+i+1 + end do + !$acc end parallel +end subroutine parallel_default_present |