@@ -1,5 +1,14 @@
2014-11-05 Thomas Schwinge <thomas@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.
@@ -1,3 +1,11 @@
+2014-11-05 Thomas Schwinge <thomas@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 <jnorris@codesourcery.com>
* c-parser.c (c_parser_omp_variable_list): Handle
@@ -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:
@@ -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)
{
@@ -1,5 +1,11 @@
2014-11-05 Thomas Schwinge <thomas@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".
@@ -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:
@@ -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)
{
@@ -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:
@@ -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:
@@ -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:
@@ -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_,
@@ -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:
@@ -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));
@@ -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]
@@ -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) \