From e8e44b733808997d06c0cdf9bf5756ce03530f42 Mon Sep 17 00:00:00 2001 From: tschwinge Date: Wed, 5 Nov 2014 16:35:30 +0000 Subject: [PATCH] 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) : Remove explicit mark_exp_read invocations. gcc/cp/ * parser.c (cp_parser_omp_var_list_no_open) : 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 a/gcc/ChangeLog.gomp b/gcc/ChangeLog.gomp index fc624c80ad3..2c2b3497c13 100644 --- a/gcc/ChangeLog.gomp +++ b/gcc/ChangeLog.gomp @@ -1,5 +1,14 @@ 2014-11-05 Thomas Schwinge + * 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 a/gcc/c/ChangeLog.gomp b/gcc/c/ChangeLog.gomp index 7acd7b343ae..70278b9ec43 100644 --- a/gcc/c/ChangeLog.gomp +++ b/gcc/c/ChangeLog.gomp @@ -1,3 +1,11 @@ +2014-11-05 Thomas Schwinge + + * 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) : + Remove explicit mark_exp_read invocations. + 2014-11-05 James Norris * c-parser.c (c_parser_omp_variable_list): Handle diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 410b19fe038..40d4314d842 100644 --- a/gcc/c/c-parser.c +++ b/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 a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 76503e4dca7..e315690d86f 100644 --- a/gcc/c/c-typeck.c +++ b/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 a/gcc/cp/ChangeLog.gomp b/gcc/cp/ChangeLog.gomp index 024e6a594ba..46d4912b138 100644 --- a/gcc/cp/ChangeLog.gomp +++ b/gcc/cp/ChangeLog.gomp @@ -1,5 +1,11 @@ 2014-11-05 Thomas Schwinge + * 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) : + 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 a/gcc/cp/parser.c b/gcc/cp/parser.c index 3ef2de7ab6c..ea4ad2f274e 100644 --- a/gcc/cp/parser.c +++ b/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 a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 2457a6fb955..6e35eef80f3 100644 --- a/gcc/cp/semantics.c +++ b/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 a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 97613ae024f..7dd4498cc1d 100644 --- a/gcc/fortran/trans-openmp.c +++ b/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 a/gcc/gimplify.c b/gcc/gimplify.c index bfd7f66bf65..d58876f2aca 100644 --- a/gcc/gimplify.c +++ b/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 a/gcc/omp-low.c b/gcc/omp-low.c index 49cf1ab5bc6..1c9d94257c6 100644 --- a/gcc/omp-low.c +++ b/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 a/gcc/tree-core.h b/gcc/tree-core.h index abdc2c979a9..42ad6a0e214 100644 --- a/gcc/tree-core.h +++ b/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 a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index f311ed9be17..d678f367dea 100644 --- a/gcc/tree-pretty-print.c +++ b/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 a/gcc/tree.c b/gcc/tree.c index 0475622a506..f39c63f6eff 100644 --- a/gcc/tree.c +++ b/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 a/gcc/tree.def b/gcc/tree.def index 871a7fb20fc..f44853a6bff 100644 --- a/gcc/tree.def +++ b/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 a/gcc/tree.h b/gcc/tree.h index c91e7165bb7..e1adbabbee5 100644 --- a/gcc/tree.h +++ b/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) \ -- 2.11.4.GIT