Fix OpenMP's target update directive in templated code

[Thomas Schwinge <thomas@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 6347 bytes --]

Hi Jakub!

When adding support for C++ templates usage with OpenACC directives
(gcc/cp/pt.c:tsubst_expr), we found:

On Mon, 27 Apr 2015 10:57:14 -0700, Cesar Philippidis <cesar_philippidis@mentor.com> wrote:
> On 04/27/2015 10:40 AM, Thomas Schwinge wrote:
> > Cesar wrote:
> >>      case OMP_TARGET_UPDATE:
> >>        tmp = tsubst_omp_clauses (OMP_TARGET_UPDATE_CLAUSES (t), false,
> >>  				args, complain, in_decl);
> >> @@ -14253,6 +14292,16 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
> >>        add_stmt (t);
> >>        break;
> >>  
> >> +    case OACC_ENTER_DATA:
> >> +    case OACC_EXIT_DATA:
> >> +    case OACC_UPDATE:
> >> +      tmp = tsubst_omp_clauses (TREE_OPERAND (t, 0), false,
> >> +				args, complain, in_decl);
> >> +      t = copy_node (t);
> >> +      TREE_OPERAND (t, 0) = tmp;
> >> +      add_stmt (t);
> >> +      break;
> > 
> > Given the handling for similar codes, why not handle those together with
> > OMP_TARGET_UPDATE, replacing OMP_TARGET_UPDATE_CLAUSES with plain
> > OMP_CLAUSES?
> 
> I left it them separate because OMP_CLAUSES expects the clauses to
> appear in tree_operand 1, but oacc_update, oacc_enter_data and
> oacc_exit_data have the clauses at tree_operand 0. I think this may have
> been a bug with openmp, so I was planning on informing Jakub when I
> posted this patch upstream.
> 
> The attached patch cleans up [...]

> --- a/gcc/cp/pt.c
> +++ b/gcc/cp/pt.c
> @@ -14277,21 +14277,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
>        add_stmt (t);
>        break;
>  
>      case OMP_TARGET_UPDATE:
> -      tmp = tsubst_omp_clauses (OMP_TARGET_UPDATE_CLAUSES (t), false,
> -				args, complain, in_decl);
> -      t = copy_node (t);
> -      OMP_CLAUSES (t) = tmp;
> -      add_stmt (t);
> -      break;
> -
>      case OACC_ENTER_DATA:
>      case OACC_EXIT_DATA:
>      case OACC_UPDATE:

I guess nobody so far ;-) has been using OpenMP's target update directive
in templated code -- OK to commit the following, and to which branches
(4.9, 5, trunk)?

commit 5ea85b95c1d0ccb98d73345f105bf76839b16433
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Tue Apr 28 20:29:26 2015 +0200

    Fix OpenMP's target update directive in templated code.
    
        FAIL: g++.dg/gomp/tpl-target-update.C  -std=c++98 (internal compiler error)
        FAIL: g++.dg/gomp/tpl-target-update.C  -std=c++98 (test for excess errors)
        FAIL: g++.dg/gomp/tpl-target-update.C  -std=c++11 (internal compiler error)
        FAIL: g++.dg/gomp/tpl-target-update.C  -std=c++11 (test for excess errors)
        FAIL: g++.dg/gomp/tpl-target-update.C  -std=c++14 (internal compiler error)
        FAIL: g++.dg/gomp/tpl-target-update.C  -std=c++14 (test for excess errors)
    
        [...]/source-gcc/gcc/testsuite/g++.dg/gomp/tpl-target-update.C: In instantiation of 'void f(T, T) [with T = int]':
        [...]/source-gcc/gcc/testsuite/g++.dg/gomp/tpl-target-update.C:19:9:   required from here
        [...]/source-gcc/gcc/testsuite/g++.dg/gomp/tpl-target-update.C:10:9: internal compiler error: tree check: expected oacc_parallel or oacc_kernels or oacc_data or oacc_host_data or omp_parallel or omp_task or omp_for or omp_simd or cilk_simd or cilk_for or omp_distribute or oacc_loop or omp_teams or omp_target_data or omp_target or omp_sections or omp_single, have omp_target_update in tsubst_expr, at cp/pt.c:14209
        0xf5aae1 tree_range_check_failed(tree_node const*, char const*, int, char const*, tree_code, tree_code)
                [...]/source-gcc/gcc/tree.c:9384
        0x66e201 tree_range_check
                [...]/source-gcc/gcc/tree.h:2979
        0x66e201 tsubst_expr
                [...]/source-gcc/gcc/cp/pt.c:14209
        0x6695e3 tsubst_expr
                [...]/source-gcc/gcc/cp/pt.c:13752
        0x66ac07 tsubst_expr
                [...]/source-gcc/gcc/cp/pt.c:13938
        0x667c41 instantiate_decl(tree_node*, int, bool)
                [...]/source-gcc/gcc/cp/pt.c:20367
        0x6ae386 instantiate_pending_templates(int)
                [...]/source-gcc/gcc/cp/pt.c:20484
        0x6edc3d cp_write_global_declarations()
                [...]/source-gcc/gcc/cp/decl2.c:4456
    
    	gcc/cp/
    	* pt.c (tsubst_expr) <OMP_TARGET_UPDATE>: Use
    	OMP_TARGET_UPDATE_CLAUSES instead of OMP_CLAUSES.
    	gcc/testsuite/
    	* g++.dg/gomp/tpl-target-update.C: New file.
---
 gcc/cp/pt.c                                   |    2 +-
 gcc/testsuite/g++.dg/gomp/tpl-target-update.C |   20 ++++++++++++++++++++
 2 files changed, 21 insertions(+), 1 deletion(-)

diff --git gcc/cp/pt.c gcc/cp/pt.c
index f9a5c3b..129517a 100644
--- gcc/cp/pt.c
+++ gcc/cp/pt.c
@@ -14206,7 +14206,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
       tmp = tsubst_omp_clauses (OMP_TARGET_UPDATE_CLAUSES (t), false,
 				args, complain, in_decl);
       t = copy_node (t);
-      OMP_CLAUSES (t) = tmp;
+      OMP_TARGET_UPDATE_CLAUSES (t) = tmp;
       add_stmt (t);
       break;
 
diff --git gcc/testsuite/g++.dg/gomp/tpl-target-update.C gcc/testsuite/g++.dg/gomp/tpl-target-update.C
new file mode 100644
index 0000000..6226ebf
--- /dev/null
+++ gcc/testsuite/g++.dg/gomp/tpl-target-update.C
@@ -0,0 +1,20 @@
+// { dg-do compile }
+
+template <typename T>
+void f(T A, T B)
+{
+  extern int *v;
+  T a = 2;
+  T b = 4;
+
+#pragma omp target update to(v[a:b])
+  v[a] = 0;
+
+#pragma omp target update to(v[A:B])
+  v[a] = 0;
+}
+
+void g()
+{
+  f(1, 5);
+}


That said, what is the preferred approach to add support for
OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_UPDATE?  I'm not sure hard-coding
TREE_OPERAND (t, 0) in gcc/cp/pt.c:tsubst_expr is the way to go, and also
duplicating the OMP_TARGET_UPDATE code for each of the three OACC_*
doesn't sound appealing -- especially given that we just "switch"ed on
the respective tree code, so the O*_CHECK of the respective O*_CLAUSES
macro is completely redundant.  Is it OK to extend gcc/tree.h:OMP_CLAUSES
so that it can also be used for tree codes that keep clauses in operand
0, and then use that here (and also in
gcc/gimplify.c:gimplify_omp_target_update, for example)?


Grüße,
 Thomas

[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]