Jakub,
On 11/06/2015 01:03 PM, Jakub Jelinek wrote:
On Fri, Nov 06, 2015 at 10:08:26AM -0600, James Norris wrote:
2015-10-27 James Norris <jnor...@codesourcery.com>
Joseph Myers <jos...@codesourcery.com>
gcc/
* c-family/c-pragma.c (oacc_pragmas): Add entry for declare directive.
* c-family/c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DECLARE.
c-family/, c/, cp/ have all their own ChangeLog files, and the entries
that go in there should not have the path prefixes.
Will fix.
+ for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+ {
+ location_t loc = OMP_CLAUSE_LOCATION (t);
+ tree decl = OMP_CLAUSE_DECL (t);
+ tree devres = NULL_TREE;
+ if (!DECL_P (decl))
+ {
+ error_at (loc, "subarray in %<#pragma acc declare%>");
Is that the OpenACC term that you use everywhere? In OpenMP
those are array sections.
I can change the term so that both OpenACC and OpenMP
use the same terminology.
+ case GOMP_MAP_LINK:
+ if (!global_bindings_p () && !DECL_EXTERNAL (decl))
+ {
+ error_at (loc,
+ "%qD must be a global variable in"
+ "%<#pragma acc declare link%>",
+ decl);
+ error = true;
+ continue;
What about TREE_STATIC?
Will fix.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index fa34858..fbd4a69 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -157,6 +157,7 @@ struct gimplify_omp_ctx
bool target_map_scalars_firstprivate;
bool target_map_pointers_as_0len_arrays;
bool target_firstprivatize_array_bases;
+ gomp_target *declare_returns;
};
This declare_returns stuff looks broken.
What exactly do you want to achieve? Why do you record it
in gimplify_omp_ctx, but then only look at it in gimplify_body?
The way you abuse gimplify_omp_ctx for that is just too ugly.
All the gimplification code expects that when you enter some OpenMP/OpenACC
construct, you create new context if it needs one, then process the body
of the construct, then pop it up. The declare_returns stuff
violates all of this. Do you want to enqueue all the statements
at the end of the body? Then just stick it into some gimple_seq
outside of the gimplify_omp_ctx, and process in there. Or if you
want to process it when leaving some construct, arrange for that.
For example:
void func (int *a)
{
int b[10];
#pragma acc declare copyout (b)
.
.
.
}
What I was trying to do was for a copyout to be done at the
time the function return'ed. It also may be the case where
a mapping operation must be done at entry and exit to a
function. I think the approach you outline below (target data/
acc data) is the approach to use.
@@ -5841,6 +5863,8 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree
decl,
flags |= GOVD_FIRSTPRIVATE;
break;
case OMP_CLAUSE_DEFAULT_UNSPECIFIED:
+ if (is_global_var (decl) && device_resident_p (decl))
+ flags |= GOVD_MAP_TO_ONLY | GOVD_MAP;
I don't think you want to do this except for (selected or all?)
OpenACC contexts. Say, I don't see why one couldn't e.g. try to mix
OpenMP host parallelization or tasking with OpenACC offloading,
and that affecting in weird way OpenMP semantics.
Will fix.
+ /* Any clauses that affect the state of a variable prior
+ to return are saved and dealt with after the body has
+ been gimplified. */
+
+ gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA,
+ OACC_DECLARE);
+
+ c = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = c->outer_context;
+ delete_omp_context (c);
Why don't you call gimplify_adjust_omp_clauses instead?
What I was attempting to do was to create to sets of clauses.
One used for mapping operations at entry to a function and
the other at exit. Obviously, I've made a mess of things with
this approach.
+ stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
+ clauses);
+ gimplify_omp_ctxp->declare_returns = stmt;
But the above is the main thing I have trouble with.
What happens if you have multiple #pragma acc declare directives?
Is the stmt from the other ones lost?
I'd expect something like gimple_seq in there instead and pushing
stmts into the sequence.
I don't know what is the semantics of the construct, if it is something
close to say target data in OpenMP or acc data in OpenACC, the addition
of code to unmap the variables is performed using a cleanup, otherwise
how do you handle exceptions, or goto and all other kinds of abnormal
returns from the function.
I think the approach you've outlined is the way to go.
@@ -160,6 +151,25 @@ varpool_node::get_create (tree decl)
node->force_output = 1;
#endif
}
+}
+
+/* Return varpool node assigned to DECL. Create new one when needed. */
+varpool_node *
+varpool_node::get_create (tree decl)
+{
+ varpool_node *node = varpool_node::get (decl);
+ gcc_checking_assert (TREE_CODE (decl) == VAR_DECL);
+ if (node)
+ {
+ if (!node->offloadable)
+ make_offloadable (node, decl);
I don't like this change at all, that is significant extra work
every time somebody calls this function.
Just do what we do on the OpenMP side, if you add some "omp declare target"
or similar attribute and the varpool node for it has been already created,
set offloadable in the FE.
Yes setting node->offloadable = 1 is much easier.
Jakub
Thank you for taking time for the review.
Will resubmit patch for review after corrections.
Thanks!
Jim
