| Message ID | 87383kt7kx.fsf@schwinge.name |
|---|---|
| State | New |
| Headers | show |
On Tue, Apr 28, 2015 at 08:45:50PM +0200, Thomas Schwinge wrote: > 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)? Seems I've missed testcases for target {,update,data} in templates indeed, generally for C++ I'm trying to add testcases for templates, both when relevant types are type dependent and whey aren't. > 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. This is ok for trunk, 5.2 and 4.9.3, thanks for fixing that. > 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)? How could it work when it is operand 1 on some and operand 0 in others. IMHO, if you want to reuse the same code for OMP_TARGET_UPDATE, various OACC_* standalone directives (and OMP_TARGET_ENTER_DATA/OMP_TARGET_EXIT_DATA in OpenMP 4.1), then you should make sure they are consecutive in target.def and define OMP_STANDALONE_CLAUSES or OMP_TARGET_STANDALONE_CLAUSES as a range check between OMP_TARGET_UPDATE and the last OpenACC directive without body, just clauses. Jakub
Hi Jakub! On Wed, 29 Apr 2015 10:53:32 +0200, Jakub Jelinek <jakub@redhat.com> wrote: > On Tue, Apr 28, 2015 at 08:45:50PM +0200, Thomas Schwinge wrote: > > 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)? > > Seems I've missed testcases for target {,update,data} in templates indeed, > generally for C++ I'm trying to add testcases for templates, both when > relevant types are type dependent and whey aren't. > > > 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. > > This is ok for trunk, 5.2 and 4.9.3, thanks for fixing that. Committed in r222564, r222565, r222566, respectively. Grüße, Thomas
Hi Jakub! On Wed, 29 Apr 2015 10:53:32 +0200, Jakub Jelinek <jakub@redhat.com> wrote: > On Tue, Apr 28, 2015 at 08:45:50PM +0200, Thomas Schwinge wrote: > > 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)? > > How could it work when it is operand 1 on some and operand 0 in others. Something like (untested): #define OMP_CLAUSES(NODE) \ - TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_SINGLE), 1) \ + (tree_code (NODE) <= OMP_SINGLE) \ + ? TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_SINGLE), 1) \ + : TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_CACHE, OMP_TARGET_UPDATE), 0) Rationale: I'm not expecting another variant to be added later on (clauses are either in operand 0 or 1). Encoding explicit tree code (ordering) assuptions is already done with the usage of TREE_RANGE_CHECK macros, so the additional tree code check doesn't make this much worse. It has the benefit of offering the same known OMP_CLAUSES interface. Yet, if that's a non-starter, I'll pursue this one: > IMHO, if you want to reuse the same code for OMP_TARGET_UPDATE, > various OACC_* standalone directives (and > OMP_TARGET_ENTER_DATA/OMP_TARGET_EXIT_DATA in OpenMP 4.1), then > you should make sure they are consecutive in target.def (tree.def; we've already made sure they're grouped consecutively.) > and > define OMP_STANDALONE_CLAUSES or OMP_TARGET_STANDALONE_CLAUSES as > a range check between OMP_TARGET_UPDATE and the last OpenACC directive > without body, just clauses. Grüße, Thomas
On Wed, Apr 29, 2015 at 11:28:55AM +0200, Thomas Schwinge wrote:
> Yet, if that's a non-starter, I'll pursue this one:
Yeah, it is a non-starter, it has unnecessary runtime overhead everywhere
where it is used.
Jakub
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); +}