From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1729) id 461DB3AA9813; Thu, 13 May 2021 16:17:09 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 461DB3AA9813 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Kwok Yeung To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] Add a "combined" flag for "acc kernels loop" etc directives. X-Act-Checkin: gcc X-Git-Author: Sandra Loosemore X-Git-Refname: refs/heads/devel/omp/gcc-11 X-Git-Oldrev: 72417bb208f13ce35ee135cc8ded43a8e519593f X-Git-Newrev: 506098b0fab2cfbd81f813de97101f2b242bc18e Message-Id: <20210513161709.461DB3AA9813@sourceware.org> Date: Thu, 13 May 2021 16:17:09 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Thu, 13 May 2021 16:17:09 -0000 https://gcc.gnu.org/g:506098b0fab2cfbd81f813de97101f2b242bc18e commit 506098b0fab2cfbd81f813de97101f2b242bc18e Author: Sandra Loosemore Date: Wed Aug 19 19:13:55 2020 -0700 Add a "combined" flag for "acc kernels loop" etc directives. 2020-08-19 Sandra Loosemore gcc/ * tree.h (OACC_LOOP_COMBINED): New. gcc/c/ * c-parser.c (c_parser_oacc_loop): Set OACC_LOOP_COMBINED. gcc/cp/ * parser.c (cp_parser_oacc_loop): Set OACC_LOOP_COMBINED. gcc/fortran/ * trans-openmp.c (gfc_trans_omp_do): Add combined parameter, use it to set OACC_LOOP_COMBINED. Update all call sites. Diff: --- gcc/ChangeLog.omp | 6 ++++++ gcc/c/ChangeLog.omp | 6 ++++++ gcc/c/c-parser.c | 3 +++ gcc/cp/ChangeLog.omp | 6 ++++++ gcc/cp/parser.c | 3 +++ gcc/fortran/ChangeLog.omp | 7 +++++++ gcc/fortran/trans-openmp.c | 30 +++++++++++++++++++----------- gcc/tree.h | 5 +++++ 8 files changed, 55 insertions(+), 11 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 29bcb59eff5..77d163266ee 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,9 @@ +2020-08-19 Sandra Loosemore + + Add a "combined" flag for "acc kernels loop" etc directives. + + * tree.h (OACC_LOOP_COMBINED): New. + 2020-07-30 Julian Brown * config/gcn/gcn-tree.c (gcn_goacc_get_worker_red_decl): Do not diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp index e6fadbb1e89..0f6ac422f31 100644 --- a/gcc/c/ChangeLog.omp +++ b/gcc/c/ChangeLog.omp @@ -1,3 +1,9 @@ +2020-08-19 Sandra Loosemore + + Add a "combined" flag for "acc kernels loop" etc directives. + + * c-parser.c (c_parser_oacc_loop): Set OACC_LOOP_COMBINED. + 2020-03-27 Sandra Loosemore * c-decl.c (c_unwrap_for_init): New. diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 4225c730351..2679b3e6405 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -16932,6 +16932,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, omp_clause_mask mask, tree *cclauses, bool *if_p) { bool is_parallel = ((mask >> PRAGMA_OACC_CLAUSE_REDUCTION) & 1) == 1; + bool is_combined = (cclauses != NULL); strcat (p_name, " loop"); mask |= OACC_LOOP_CLAUSE_MASK; @@ -16950,6 +16951,8 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, tree block = c_begin_compound_stmt (true); tree stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL, if_p); + if (stmt && stmt != error_mark_node) + OACC_LOOP_COMBINED (stmt) = is_combined; block = c_end_compound_stmt (loc, block, true); add_stmt (block); diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp index adc28b921ac..070f5633308 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,9 @@ +2020-08-19 Sandra Loosemore + + Add a "combined" flag for "acc kernels loop" etc directives. + + * parser.c (cp_parser_oacc_loop): Set OACC_LOOP_COMBINED. + 2020-03-31 Sandra Loosemore * semantics.c (handle_omp_array_sections_1): Call STRIP_NOPS diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index e42f32c1164..245a044f5b6 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -42475,6 +42475,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, omp_clause_mask mask, tree *cclauses, bool *if_p) { bool is_parallel = ((mask >> PRAGMA_OACC_CLAUSE_REDUCTION) & 1) == 1; + bool is_combined = (cclauses != NULL); strcat (p_name, " loop"); mask |= OACC_LOOP_CLAUSE_MASK; @@ -42493,6 +42494,8 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, tree block = begin_omp_structured_block (); int save = cp_parser_begin_omp_structured_block (parser); tree stmt = cp_parser_omp_for_loop (parser, OACC_LOOP, clauses, NULL, if_p); + if (stmt && stmt != error_mark_node) + OACC_LOOP_COMBINED (stmt) = is_combined; cp_parser_end_omp_structured_block (parser, save); add_stmt (finish_omp_structured_block (block)); diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index e619434815e..c608163f77f 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,10 @@ +2020-08-19 Sandra Loosemore + + Add a "combined" flag for "acc kernels loop" etc directives. + + * trans-openmp.c (gfc_trans_omp_do): Add combined parameter, + use it to set OACC_LOOP_COMBINED. Update all call sites. + 2020-07-20 Frederik Harwath * openmp.c (oacc_is_parallel_or_serial): Removed function. diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 66e47fdbc7e..51414a12910 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -4489,7 +4489,8 @@ typedef struct dovar_init_d { static tree gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock, - gfc_omp_clauses *do_clauses, tree par_clauses) + gfc_omp_clauses *do_clauses, tree par_clauses, + bool combined) { gfc_se se; tree dovar, stmt, from, to, step, type, init, cond, incr, orig_decls; @@ -4852,7 +4853,10 @@ gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock, case EXEC_OMP_DO: stmt = make_node (OMP_FOR); break; case EXEC_OMP_DISTRIBUTE: stmt = make_node (OMP_DISTRIBUTE); break; case EXEC_OMP_TASKLOOP: stmt = make_node (OMP_TASKLOOP); break; - case EXEC_OACC_LOOP: stmt = make_node (OACC_LOOP); break; + case EXEC_OACC_LOOP: + stmt = make_node (OACC_LOOP); + OACC_LOOP_COMBINED (stmt) = combined; + break; default: gcc_unreachable (); } @@ -4946,7 +4950,8 @@ gfc_trans_oacc_combined_directive (gfc_code *code) pblock = █ else pushlevel (); - stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, &loop_clauses, NULL); + stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, &loop_clauses, NULL, + true); protected_set_expr_location (stmt, loc); if (TREE_CODE (stmt) != BIND_EXPR) stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0)); @@ -5418,7 +5423,7 @@ gfc_trans_omp_do_simd (gfc_code *code, stmtblock_t *pblock, omp_do_clauses = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_DO], code->loc); body = gfc_trans_omp_do (code, EXEC_OMP_SIMD, pblock ? pblock : &block, - &clausesa[GFC_OMP_SPLIT_SIMD], omp_clauses); + &clausesa[GFC_OMP_SPLIT_SIMD], omp_clauses, false); if (pblock == NULL) { if (TREE_CODE (body) != BIND_EXPR) @@ -5471,7 +5476,7 @@ gfc_trans_omp_parallel_do (gfc_code *code, stmtblock_t *pblock, pushlevel (); } stmt = gfc_trans_omp_do (code, EXEC_OMP_DO, new_pblock, - &clausesa[GFC_OMP_SPLIT_DO], omp_clauses); + &clausesa[GFC_OMP_SPLIT_DO], omp_clauses, false); if (pblock == NULL) { if (TREE_CODE (stmt) != BIND_EXPR) @@ -5714,7 +5719,8 @@ gfc_trans_omp_distribute (gfc_code *code, gfc_omp_clauses *clausesa) case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD: stmt = gfc_trans_omp_do (code, EXEC_OMP_SIMD, &block, - &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE); + &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE, + false); if (TREE_CODE (stmt) != BIND_EXPR) stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0)); else @@ -5770,7 +5776,7 @@ gfc_trans_omp_teams (gfc_code *code, gfc_omp_clauses *clausesa, case EXEC_OMP_TEAMS_DISTRIBUTE: stmt = gfc_trans_omp_do (code, EXEC_OMP_DISTRIBUTE, NULL, &clausesa[GFC_OMP_SPLIT_DISTRIBUTE], - NULL); + NULL, false); break; default: stmt = gfc_trans_omp_distribute (code, clausesa); @@ -5844,7 +5850,8 @@ gfc_trans_omp_target (gfc_code *code) break; case EXEC_OMP_TARGET_SIMD: stmt = gfc_trans_omp_do (code, EXEC_OMP_SIMD, &block, - &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE); + &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE, + false); if (TREE_CODE (stmt) != BIND_EXPR) stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0)); else @@ -5914,7 +5921,8 @@ gfc_trans_omp_taskloop (gfc_code *code) break; case EXEC_OMP_TASKLOOP_SIMD: stmt = gfc_trans_omp_do (code, EXEC_OMP_SIMD, &block, - &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE); + &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE, + false); if (TREE_CODE (stmt) != BIND_EXPR) stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0)); else @@ -6229,7 +6237,7 @@ gfc_trans_oacc_directive (gfc_code *code) return gfc_trans_oacc_construct (code); case EXEC_OACC_LOOP: return gfc_trans_omp_do (code, code->op, NULL, code->ext.omp_clauses, - NULL); + NULL, false); case EXEC_OACC_UPDATE: case EXEC_OACC_CACHE: case EXEC_OACC_ENTER_DATA: @@ -6266,7 +6274,7 @@ gfc_trans_omp_directive (gfc_code *code) case EXEC_OMP_SIMD: case EXEC_OMP_TASKLOOP: return gfc_trans_omp_do (code, code->op, NULL, code->ext.omp_clauses, - NULL); + NULL, false); case EXEC_OMP_DISTRIBUTE_PARALLEL_DO: case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: case EXEC_OMP_DISTRIBUTE_SIMD: diff --git a/gcc/tree.h b/gcc/tree.h index 0ae76b89356..edf4fd37979 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1502,6 +1502,11 @@ class auto_suppress_location_wrappers #define OMP_TARGET_COMBINED(NODE) \ (OMP_TARGET_CHECK (NODE)->base.private_flag) +/* True on an OACC_LOOP statement if it is part of a combined construct, + for example "#pragma acc kernels loop". */ +#define OACC_LOOP_COMBINED(NODE) \ + (OACC_LOOP_CHECK (NODE)->base.private_flag) + /* Memory order for OMP_ATOMIC*. */ #define OMP_ATOMIC_MEMORY_ORDER(NODE) \ (TREE_RANGE_CHECK (NODE, OMP_ATOMIC, \