@@ -17371,6 +17371,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;
@@ -17389,6 +17390,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);
@@ -44580,6 +44580,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;
@@ -44598,6 +44599,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));
@@ -4855,7 +4855,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;
@@ -5219,7 +5220,10 @@ gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock,
case EXEC_OMP_DISTRIBUTE: stmt = make_node (OMP_DISTRIBUTE); break;
case EXEC_OMP_LOOP: stmt = make_node (OMP_LOOP); 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 ();
}
@@ -5313,7 +5317,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));
@@ -6151,7 +6156,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)
@@ -6209,7 +6214,7 @@ gfc_trans_omp_parallel_do (gfc_code *code, bool is_loop, stmtblock_t *pblock,
}
stmt = gfc_trans_omp_do (code, is_loop ? EXEC_OMP_LOOP : EXEC_OMP_DO,
new_pblock, &clausesa[GFC_OMP_SPLIT_DO],
- omp_clauses);
+ omp_clauses, false);
if (pblock == NULL)
{
if (TREE_CODE (stmt) != BIND_EXPR)
@@ -6496,7 +6501,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
@@ -6555,13 +6561,13 @@ 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;
case EXEC_OMP_TARGET_TEAMS_LOOP:
case EXEC_OMP_TEAMS_LOOP:
stmt = gfc_trans_omp_do (code, EXEC_OMP_LOOP, NULL,
&clausesa[GFC_OMP_SPLIT_DO],
- NULL);
+ NULL, false);
break;
default:
stmt = gfc_trans_omp_distribute (code, clausesa);
@@ -6641,7 +6647,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
@@ -6712,7 +6719,8 @@ gfc_trans_omp_taskloop (gfc_code *code, gfc_exec_op op)
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
@@ -6756,7 +6764,7 @@ gfc_trans_omp_master_masked_taskloop (gfc_code *code, gfc_exec_op op)
stmt = gfc_trans_omp_do (code, EXEC_OMP_TASKLOOP, NULL,
code->op != EXEC_OMP_MASTER_TASKLOOP
? &clausesa[GFC_OMP_SPLIT_TASKLOOP]
- : code->ext.omp_clauses, NULL);
+ : code->ext.omp_clauses, NULL, false);
}
if (TREE_CODE (stmt) != BIND_EXPR)
stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
@@ -7119,7 +7127,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:
@@ -7159,7 +7167,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:
@@ -1524,6 +1524,11 @@ class auto_suppress_location_wrappers
#define OMP_MASKED_COMBINED(NODE) \
(OMP_MASKED_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, \
From: Sandra Loosemore <sandra@codesourcery.com> 2020-08-19 Sandra Loosemore <sandra@codesourcery.com> 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. --- gcc/c/c-parser.c | 3 +++ gcc/cp/parser.c | 3 +++ gcc/fortran/trans-openmp.c | 34 +++++++++++++++++++++------------- gcc/tree.h | 5 +++++ 4 files changed, 32 insertions(+), 13 deletions(-) -- 2.33.0 ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955