Hi! On Fri, 24 Jan 2014 20:33:35 +0100, I wrote: > On Thu, 23 Jan 2014 22:04:45 +0400, Ilmir Usmanov <i.usma...@samsung.com> > wrote: > > Subject: [PATCH 4/6] OpenACC GENERIC nodes
> > --- a/gcc/tree-core.h
> > +++ b/gcc/tree-core.h
> > @@ -216,12 +216,18 @@ enum omp_clause_code {
> > + /* Internal structure to hold OpenACC cache directive's variable-list.
> > + #pragma acc cache (variable-_ist). */
> > + OACC_NO_CLAUSE_CACHE,
>
> Hmm, yeah, while *_NO_CLAUSE_* perhaps isn't the most beautiful approach,
> I think it's fine at least for now.
In r217146, I applied the following to gomp-4_0-branch:
commit e8e44b733808997d06c0cdf9bf5756ce03530f42
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Nov 5 16:35:30 2014 +0000
OpenACC cache directive maintenance.
gcc/c/
* c-parser.c (c_parser_oacc_cache): Generate OACC_CACHE.
* c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE__CACHE_.
gcc/cp/
* parser.c (cp_parser_oacc_cache): Generate OACC_CACHE.
* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE__CACHE_.
gcc/
* gimplify.c (gimplify_oacc_cache): New function.
(gimplify_expr): Use it for OACC_CACHE.
(gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle
OMP_CLAUSE__CACHE_.
gcc/c/
* c-parser.c (c_parser_omp_variable_list) <OMP_CLAUSE__CACHE_>:
Remove explicit mark_exp_read invocations.
gcc/cp/
* parser.c (cp_parser_omp_var_list_no_open) <OMP_CLAUSE__CACHE_>:
Remove explicit mark_exp_read invocations.
gcc/
* tree-core.h (enum omp_clause_code): Move OMP_NO_CLAUSE_CACHE
next to, and handle it like a data clause. Rename it to
OMP_CLAUSE__CACHE_. Update all users.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217146
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 9 +++++++++
gcc/c/ChangeLog.gomp | 8 ++++++++
gcc/c/c-parser.c | 23 +++++++++++++++--------
gcc/c/c-typeck.c | 1 +
gcc/cp/ChangeLog.gomp | 6 ++++++
gcc/cp/parser.c | 24 +++++++++++++++---------
gcc/cp/semantics.c | 1 +
gcc/fortran/trans-openmp.c | 2 +-
gcc/gimplify.c | 25 ++++++++++++++++++++++---
gcc/omp-low.c | 4 ++--
gcc/tree-core.h | 8 ++++----
gcc/tree-pretty-print.c | 11 +++++++----
gcc/tree.c | 6 +++---
gcc/tree.def | 5 +++--
gcc/tree.h | 2 +-
15 files changed, 98 insertions(+), 37 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index fc624c8..2c2b349 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,14 @@
2014-11-05 Thomas Schwinge <tho...@codesourcery.com>
+ * gimplify.c (gimplify_oacc_cache): New function.
+ (gimplify_expr): Use it for OACC_CACHE.
+ (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle
+ OMP_CLAUSE__CACHE_.
+
+ * tree-core.h (enum omp_clause_code): Move OMP_NO_CLAUSE_CACHE
+ next to, and handle it like a data clause. Rename it to
+ OMP_CLAUSE__CACHE_. Update all users.
+
* invoke.texi: Update for OpenACC.
* sourcebuild.texi: Likewise.
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index 7acd7b3..70278b9 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,11 @@
+2014-11-05 Thomas Schwinge <tho...@codesourcery.com>
+
+ * c-parser.c (c_parser_oacc_cache): Generate OACC_CACHE.
+ * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE__CACHE_.
+
+ * c-parser.c (c_parser_omp_variable_list) <OMP_CLAUSE__CACHE_>:
+ Remove explicit mark_exp_read invocations.
+
2014-11-05 James Norris <jnor...@codesourcery.com>
* c-parser.c (c_parser_omp_variable_list): Handle
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 410b19f..40d4314 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -10053,7 +10053,7 @@ c_parser_omp_variable_list (c_parser *parser,
{
switch (kind)
{
- case OMP_NO_CLAUSE_CACHE:
+ case OMP_CLAUSE__CACHE_:
if (c_parser_peek_token (parser)->type != CPP_OPEN_SQUARE)
{
c_parser_error (parser, "expected %<[%>");
@@ -10100,11 +10100,8 @@ c_parser_omp_variable_list (c_parser *parser,
break;
}
- if (kind == OMP_NO_CLAUSE_CACHE)
+ if (kind == OMP_CLAUSE__CACHE_)
{
- mark_exp_read (low_bound);
- mark_exp_read (length);
-
if (TREE_CODE (low_bound) != INTEGER_CST
&& !TREE_READONLY (low_bound))
{
@@ -11901,12 +11898,22 @@ c_parser_omp_structured_block (c_parser *parser)
*/
static tree
-c_parser_oacc_cache (location_t loc __attribute__((unused)), c_parser *parser)
+c_parser_oacc_cache (location_t loc, c_parser *parser)
{
- c_parser_omp_var_list_parens (parser, OMP_NO_CLAUSE_CACHE, NULL);
+ tree stmt, clauses;
+
+ clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL);
+ clauses = c_finish_omp_clauses (clauses);
+
c_parser_skip_to_pragma_eol (parser);
- return NULL_TREE;
+ stmt = make_node (OACC_CACHE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_CACHE_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, loc);
+ add_stmt (stmt);
+
+ return stmt;
}
/* OpenACC 2.0:
diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c
index 76503e4..e315690d 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -12204,6 +12204,7 @@ c_finish_omp_clauses (tree clauses)
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ case OMP_CLAUSE__CACHE_:
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
diff --git gcc/cp/ChangeLog.gomp gcc/cp/ChangeLog.gomp
index 024e6a5..46d4912 100644
--- gcc/cp/ChangeLog.gomp
+++ gcc/cp/ChangeLog.gomp
@@ -1,5 +1,11 @@
2014-11-05 Thomas Schwinge <tho...@codesourcery.com>
+ * parser.c (cp_parser_oacc_cache): Generate OACC_CACHE.
+ * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE__CACHE_.
+
+ * parser.c (cp_parser_omp_var_list_no_open) <OMP_CLAUSE__CACHE_>:
+ Remove explicit mark_exp_read invocations.
+
* parser.c (cp_parser_omp_clause_name): Also look for "pcopy",
"pcopyin", "pcopyout", "pcreate". Look for "wait" instead of
"WAIT".
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 3ef2de7..ea4ad2f 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -27669,7 +27669,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum
omp_clause_code kind,
{
switch (kind)
{
- case OMP_NO_CLAUSE_CACHE:
+ case OMP_CLAUSE__CACHE_:
if (cp_lexer_peek_token (parser->lexer)->type != CPP_OPEN_SQUARE)
{
error_at (token->location, "expected %<[%>");
@@ -27708,11 +27708,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser,
enum omp_clause_code kind,
RT_CLOSE_SQUARE))
goto skip_comma;
- if (kind == OMP_NO_CLAUSE_CACHE)
+ if (kind == OMP_CLAUSE__CACHE_)
{
- mark_exp_read (low_bound);
- mark_exp_read (length);
-
if (TREE_CODE (low_bound) != INTEGER_CST
&& !TREE_READONLY (low_bound))
{
@@ -31410,13 +31407,22 @@ cp_parser_omp_target (cp_parser *parser, cp_token
*pragma_tok,
*/
static tree
-cp_parser_oacc_cache (cp_parser *parser,
- cp_token *pragma_tok __attribute__((unused)))
+cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
{
- cp_parser_omp_var_list (parser, OMP_NO_CLAUSE_CACHE, NULL_TREE);
+ tree stmt, clauses;
+
+ clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE);
+ clauses = finish_omp_clauses (clauses);
+
cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer));
- return NULL_TREE;
+ stmt = make_node (OACC_CACHE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_CACHE_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+ add_stmt (stmt);
+
+ return stmt;
}
/* OpenACC 2.0:
diff --git gcc/cp/semantics.c gcc/cp/semantics.c
index 2457a6f..6e35eef 100644
--- gcc/cp/semantics.c
+++ gcc/cp/semantics.c
@@ -5704,6 +5704,7 @@ finish_omp_clauses (tree clauses)
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ case OMP_CLAUSE__CACHE_:
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c
index 97613ae..7dd4498 100644
--- gcc/fortran/trans-openmp.c
+++ gcc/fortran/trans-openmp.c
@@ -1807,7 +1807,7 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
clause_code = OMP_CLAUSE_DEVICE_RESIDENT;
goto add_clause;
case OMP_LIST_CACHE:
- clause_code = OMP_NO_CLAUSE_CACHE;
+ clause_code = OMP_CLAUSE__CACHE_;
goto add_clause;
add_clause:
diff --git gcc/gimplify.c gcc/gimplify.c
index bfd7f66..d58876f 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -6114,6 +6114,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ case OMP_CLAUSE__CACHE_:
decl = OMP_CLAUSE_DECL (c);
if (error_operand_p (decl))
{
@@ -6294,7 +6295,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
case OMP_CLAUSE_GANG:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
- case OMP_NO_CLAUSE_CACHE:
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
@@ -6641,6 +6641,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree
*list_p)
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ case OMP_CLAUSE__CACHE_:
decl = OMP_CLAUSE_DECL (c);
if (!DECL_P (decl))
break;
@@ -6698,7 +6699,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree
*list_p)
case OMP_CLAUSE_GANG:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
- case OMP_NO_CLAUSE_CACHE:
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
@@ -6722,6 +6722,21 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree
*list_p)
delete_omp_context (ctx);
}
+/* Gimplify OACC_CACHE. */
+
+static void
+gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p)
+{
+ tree expr = *expr_p;
+
+ gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE);
+ gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr));
+
+ /* TODO: Do something sensible with this information. */
+
+ *expr_p = NULL_TREE;
+}
+
/* Gimplify the contents of an OMP_PARALLEL statement. This involves
gimplification of the body, as well as scanning the body for used
variables. We need to do this scan now, because variable-sized
@@ -8312,7 +8327,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p,
gimple_seq *post_p,
case OACC_HOST_DATA:
case OACC_DECLARE:
- case OACC_CACHE:
sorry ("directive not yet implemented");
ret = GS_ALL_DONE;
break;
@@ -8352,6 +8366,11 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p,
gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
+ case OACC_CACHE:
+ gimplify_oacc_cache (expr_p, pre_p);
+ ret = GS_ALL_DONE;
+ break;
+
case OACC_DATA:
case OMP_SECTIONS:
case OMP_SINGLE:
diff --git gcc/omp-low.c gcc/omp-low.c
index 49cf1ab..1c9d942 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1982,7 +1982,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
- case OMP_NO_CLAUSE_CACHE:
+ case OMP_CLAUSE__CACHE_:
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
@@ -2130,7 +2130,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
- case OMP_NO_CLAUSE_CACHE:
+ case OMP_CLAUSE__CACHE_:
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
diff --git gcc/tree-core.h gcc/tree-core.h
index abdc2c9..42ad6a0 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -266,6 +266,10 @@ enum omp_clause_code {
OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
OMP_CLAUSE_MAP,
+ /* Internal structure to hold OpenACC cache directive's variable-list.
+ #pragma acc cache (variable-list). */
+ OMP_CLAUSE__CACHE_,
+
/* OpenACC clause: host (variable_list). */
OMP_CLAUSE_HOST,
@@ -292,10 +296,6 @@ enum omp_clause_code {
/* OpenACC clause/directive: wait [(integer-expression-list)]. */
OMP_CLAUSE_WAIT,
- /* Internal structure to hold OpenACC cache directive's variable-list.
- #pragma acc cache (variable-list). */
- OMP_NO_CLAUSE_CACHE,
-
/* Internal clause: temporary for combined loops expansion. */
OMP_CLAUSE__LOOPTEMP_,
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index f311ed9..d678f36 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -347,9 +347,6 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int
spc, int flags)
case OMP_CLAUSE_USE_DEVICE:
name = "use_device";
goto print_remap;
- case OMP_NO_CLAUSE_CACHE:
- name = "_cache_";
- goto print_remap;
print_remap:
pp_string (buffer, name);
pp_left_paren (buffer);
@@ -599,6 +596,12 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int
spc, int flags)
spc, flags, false);
goto print_clause_size;
+ case OMP_CLAUSE__CACHE_:
+ pp_string (buffer, "(");
+ dump_generic_node (buffer, OMP_CLAUSE_DECL (clause),
+ spc, flags, false);
+ goto print_clause_size;
+
case OMP_CLAUSE_NUM_TEAMS:
pp_string (buffer, "num_teams(");
dump_generic_node (buffer, OMP_CLAUSE_NUM_TEAMS_EXPR (clause),
@@ -2548,7 +2551,7 @@ dump_generic_node (pretty_printer *buffer, tree node, int
spc, int flags,
case OACC_CACHE:
pp_string (buffer, "#pragma acc cache");
- dump_omp_clauses (buffer, OACC_CACHE_CLAUSES(node), spc, flags);
+ dump_omp_clauses (buffer, OACC_CACHE_CLAUSES (node), spc, flags);
break;
case OMP_PARALLEL:
diff --git gcc/tree.c gcc/tree.c
index 0475622..f39c63f 100644
--- gcc/tree.c
+++ gcc/tree.c
@@ -270,6 +270,7 @@ unsigned const char omp_clause_num_ops[] =
2, /* OMP_CLAUSE_FROM */
2, /* OMP_CLAUSE_TO */
2, /* OMP_CLAUSE_MAP */
+ 2, /* OMP_CLAUSE__CACHE_ */
1, /* OMP_CLAUSE_HOST */
1, /* OMP_CLAUSE_OACC_DEVICE */
1, /* OMP_CLAUSE_DEVICE_RESIDENT */
@@ -277,7 +278,6 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_GANG */
1, /* OMP_CLAUSE_ASYNC */
1, /* OMP_CLAUSE_WAIT */
- 1, /* OMP_NO_CLAUSE_CACHE */
1, /* OMP_CLAUSE__LOOPTEMP_ */
1, /* OMP_CLAUSE_IF */
1, /* OMP_CLAUSE_NUM_THREADS */
@@ -329,6 +329,7 @@ const char * const omp_clause_code_name[] =
"from",
"to",
"map",
+ "_cache_",
"host",
"device",
"device_resident",
@@ -336,7 +337,6 @@ const char * const omp_clause_code_name[] =
"gang",
"async",
"wait",
- "_cache_",
"_looptemp_",
"if",
"num_threads",
@@ -11127,7 +11127,6 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_GANG:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
- case OMP_NO_CLAUSE_CACHE:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
case OMP_CLAUSE_NUM_GANGS:
@@ -11194,6 +11193,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_MAP:
+ case OMP_CLAUSE__CACHE_:
WALK_SUBTREE (OMP_CLAUSE_DECL (*tp));
WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 1));
WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
diff --git gcc/tree.def gcc/tree.def
index 871a7fb..f44853a 100644
--- gcc/tree.def
+++ gcc/tree.def
@@ -1163,8 +1163,9 @@ DEFTREECODE (OACC_ENTER_DATA, "oacc_enter_data",
tcc_statement, 1)
Operand 0: OACC_EXIT_DATA_CLAUSES: List of clauses. */
DEFTREECODE (OACC_EXIT_DATA, "oacc_exit_data", tcc_statement, 1)
-/* OpenACC - #pragma acc cache [clause1 ... clauseN]
- Operand 0: OACC_CACHE_CLAUSES: List of clauses. */
+/* OpenACC - #pragma acc cache (variable1 ... variableN)
+ Operand 0: OACC_CACHE_CLAUSES: List of variables (transformed into
+ OMP_CLAUSE__CACHE_ clauses). */
DEFTREECODE (OACC_CACHE, "oacc_cache", tcc_statement, 1)
/* OpenMP - #pragma omp target update [clause1 ... clauseN]
diff --git gcc/tree.h gcc/tree.h
index c91e716..e1adbab 100644
--- gcc/tree.h
+++ gcc/tree.h
@@ -1254,7 +1254,7 @@ extern void protected_set_expr_location (tree,
location_t);
#define OMP_CLAUSE_SIZE(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \
OMP_CLAUSE_FROM, \
- OMP_CLAUSE_MAP), 1)
+ OMP_CLAUSE__CACHE_), 1)
#define OMP_CLAUSE_CHAIN(NODE) TREE_CHAIN (OMP_CLAUSE_CHECK (NODE))
#define OMP_CLAUSE_DECL(NODE) \
Grüße,
Thomas
pgp43sCISoujB.pgp
Description: PGP signature
