From b605f6639c14383eb274e90c4fa6f995af514fbc Mon Sep 17 00:00:00 2001 From: Chung-Lin Tang Date: Fri, 3 Jun 2016 14:25:12 +0000 Subject: [PATCH] c-typeck.c (c_finish_omp_clauses): Mark OpenACC reduction arguments as addressable when async clause exists. 2016-06-03 Chung-Lin Tang c/ * c-typeck.c (c_finish_omp_clauses): Mark OpenACC reduction arguments as addressable when async clause exists. cp/ * semantics.c (finish_omp_clauses): Mark OpenACC reduction arguments as addressable when async clause exists. fortran/ * trans-openmp.c (gfc_trans_omp_reduction_list): Add mark_addressable bool parameter, set reduction clause DECLs as addressable when true. (gfc_trans_omp_clauses): Pass clauses->async to gfc_trans_omp_reduction_list, add comment describing OpenACC situation. libgomp/ * testsuite/libgomp.oacc-fortran/reduction-8.f90: New testcase. * testsuite/libgomp.oacc-c-c++-common/reduction-8.c: New testcase. From-SVN: r237070 --- gcc/c/ChangeLog | 5 +++ gcc/c/c-typeck.c | 11 +++++ gcc/cp/ChangeLog | 5 +++ gcc/cp/semantics.c | 11 +++++ gcc/fortran/ChangeLog | 7 ++++ gcc/fortran/trans-openmp.c | 9 +++- .../libgomp.oacc-c-c++-common/reduction-8.c | 30 ++++++++++++++ .../libgomp.oacc-fortran/reduction-8.f90 | 41 +++++++++++++++++++ 8 files changed, 117 insertions(+), 2 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-8.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/reduction-8.f90 diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index d2a8947d1b3..1262c82cedb 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,8 @@ +2016-06-03 Chung-Lin Tang + + * c-typeck.c (c_finish_omp_clauses): Mark OpenACC reduction + arguments as addressable when async clause exists. + 2016-05-30 Jakub Jelinek PR c++/71349 diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 1520c20c653..0ff28f2ec28 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12529,6 +12529,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) tree *nowait_clause = NULL; bool ordered_seen = false; tree schedule_clause = NULL_TREE; + bool oacc_async = false; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -12539,6 +12540,14 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_initialize (&map_field_head, &bitmap_default_obstack); bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); + if (ort & C_ORT_ACC) + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC) + { + oacc_async = true; + break; + } + for (pc = &clauses, c = clauses; c ; c = *pc) { bool remove = false; @@ -12575,6 +12584,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } + if (oacc_async) + c_mark_addressable (t); type = TREE_TYPE (t); if (TREE_CODE (t) == MEM_REF) type = TREE_TYPE (type); diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 69d3da17c9e..bec708c9f81 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,8 @@ +2016-06-03 Chung-Lin Tang + + * semantics.c (finish_omp_clauses): Mark OpenACC reduction + arguments as addressable when async clause exists. + 2016-06-02 Jan Hubicka * cp-gimplify.c (genericize_continue_stmt): Force addition of diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 8e682c53795..536509127e4 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5775,6 +5775,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bool branch_seen = false; bool copyprivate_seen = false; bool ordered_seen = false; + bool oacc_async = false; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -5785,6 +5786,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_initialize (&map_field_head, &bitmap_default_obstack); bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); + if (ort & C_ORT_ACC) + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC) + { + oacc_async = true; + break; + } + for (pc = &clauses, c = clauses; c ; c = *pc) { bool remove = false; @@ -5828,6 +5837,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = n; goto check_dup_generic_t; } + if (oacc_async) + cxx_mark_addressable (t); goto check_dup_generic; case OMP_CLAUSE_COPYPRIVATE: copyprivate_seen = true; diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index 2a9c9ceab93..8c04089f8d0 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,10 @@ +2016-06-03 Chung-Lin Tang + + * trans-openmp.c (gfc_trans_omp_reduction_list): Add mark_addressable + bool parameter, set reduction clause DECLs as addressable when true. + (gfc_trans_omp_clauses): Pass clauses->async to + gfc_trans_omp_reduction_list, add comment describing OpenACC situation. + 2016-06-01 Jerry DeLisle PR fortran/52393 diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index d3276f9331f..ab07fe45be9 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -1646,7 +1646,7 @@ gfc_trans_omp_array_reduction_or_udr (tree c, gfc_omp_namelist *n, locus where) static tree gfc_trans_omp_reduction_list (gfc_omp_namelist *namelist, tree list, - locus where) + locus where, bool mark_addressable) { for (; namelist != NULL; namelist = namelist->next) if (namelist->sym->attr.referenced) @@ -1657,6 +1657,8 @@ gfc_trans_omp_reduction_list (gfc_omp_namelist *namelist, tree list, tree node = build_omp_clause (where.lb->location, OMP_CLAUSE_REDUCTION); OMP_CLAUSE_DECL (node) = t; + if (mark_addressable) + TREE_ADDRESSABLE (t) = 1; switch (namelist->u.reduction_op) { case OMP_REDUCTION_PLUS: @@ -1747,7 +1749,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, switch (list) { case OMP_LIST_REDUCTION: - omp_clauses = gfc_trans_omp_reduction_list (n, omp_clauses, where); + /* An OpenACC async clause indicates the need to set reduction + arguments addressable, to allow asynchronous copy-out. */ + omp_clauses = gfc_trans_omp_reduction_list (n, omp_clauses, where, + clauses->async); break; case OMP_LIST_PRIVATE: clause_code = OMP_CLAUSE_PRIVATE; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-8.c new file mode 100644 index 00000000000..d25060572b5 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-8.c @@ -0,0 +1,30 @@ +const int n = 100; + +// Check async over parallel construct with reduction + +int +async_sum (int c) +{ + int s = 0; + +#pragma acc parallel loop num_gangs (10) gang reduction (+:s) async + for (int i = 0; i < n; i++) + s += i+c; + +#pragma acc wait + return s; +} + +int +main() +{ + int result = 0; + + for (int i = 0; i < n; i++) + result += i+1; + + if (async_sum (1) != result) + __builtin_abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-8.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-8.f90 new file mode 100644 index 00000000000..5cf4681191b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-8.f90 @@ -0,0 +1,41 @@ +! { dg-do run } + +program reduction + implicit none + integer, parameter :: n = 100 + integer :: i, h1, h2, s1, s2, a1, a2 + + h1 = 0 + h2 = 0 + do i = 1, n + h1 = h1 + 1 + h2 = h2 + 2 + end do + + s1 = 0 + s2 = 0 + !$acc parallel loop reduction(+:s1, s2) + do i = 1, n + s1 = s1 + 1 + s2 = s2 + 2 + end do + !$acc end parallel loop + + a1 = 0 + a2 = 0 + !$acc parallel loop reduction(+:a1, a2) async(1) + do i = 1, n + a1 = a1 + 1 + a2 = a2 + 2 + end do + !$acc end parallel loop + + if (h1 .ne. s1) call abort () + if (h2 .ne. s2) call abort () + + !$acc wait(1) + + if (h1 .ne. a1) call abort () + if (h2 .ne. a2) call abort () + +end program reduction -- 2.30.2