Add a "combined" flag for "acc kernels loop" etc directives.

2020-08-19  Sandra Loosemore  <sandra@codesourcery.com>

	gcc/
	* tree.h (OACC_LOOP_COMBINED): New.

	gcc/c/
	* c-parser.cc (c_parser_oacc_loop): Set OACC_LOOP_COMBINED.

	gcc/cp/
	* parser.cc (cp_parser_oacc_loop): Set OACC_LOOP_COMBINED.

	gcc/fortran/
	* trans-openmp.cc (gfc_trans_omp_do): Add combined parameter,
	use it to set OACC_LOOP_COMBINED.  Update all call sites.

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 750f7a2aad298ac7455176777b382f2ca7e0992b..bc3360e053ccca0704216026a53e668d984eac18 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,9 @@
+2020-08-19  Sandra Loosemore  <sandra@codesourcery.com>
+
+       Add a "combined" flag for "acc kernels loop" etc directives.
+
+       * tree.h (OACC_LOOP_COMBINED): New.
+
 2020-06-03  Tobias Burnus  <tobias@codesourcery.com>
 
        * gimplify.cc (localize_reductions): Do not create local

diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index 35c7a52c35751c917af627e41a26d16b361a281b..52af6cf63167f9b624798baab543e012049cde1e 100644 (file)
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,9 @@
+2020-08-19  Sandra Loosemore  <sandra@codesourcery.com>
+
+       Add a "combined" flag for "acc kernels loop" etc directives.
+
+       * c-parser.cc (c_parser_oacc_loop): Set OACC_LOOP_COMBINED.
+
 2020-03-27  Sandra Loosemore  <sandra@codesourcery.com>
 
        * c-decl.cc (c_unwrap_for_init): New.

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 3d4f03c98a6a2887c57c9db1df8a5c60f37163fb..85bc30634e74a0c2997826fc7b3c61491f5840fa 100644 (file)
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -17515,6 +17515,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;
@@ -17533,6 +17534,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 ead9de41a2c9c42281528a33eeaeb3cffbae1216..d8156d528c6b2e163e9ec0381576b9e7369c5bc9 100644 (file)
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,9 @@
+2020-08-19  Sandra Loosemore  <sandra@codesourcery.com>
+
+       Add a "combined" flag for "acc kernels loop" etc directives.
+
+       * parser.cc (cp_parser_oacc_loop): Set OACC_LOOP_COMBINED.
+
 2020-03-31  Sandra Loosemore  <sandra@codesourcery.com>
 
        * semantics.cc (handle_omp_array_sections_1): Call STRIP_NOPS

diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index f98854dc8ec107a7a646714846f91dc199d6713b..89d27520a27343375a67886fdfb2b728aa6c9174 100644 (file)
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -44814,6 +44814,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;
@@ -44832,6 +44833,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 a28bc3dc7892dc47703dcdbf00b1eb4d60ce35b9..df352c58706e55a445761b5b04637c01a2158969 100644 (file)
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,10 @@
+2020-08-19  Sandra Loosemore  <sandra@codesourcery.com>
+
+       Add a "combined" flag for "acc kernels loop" etc directives.
+
+       * trans-openmp.cc (gfc_trans_omp_do): Add combined parameter,
+       use it to set OACC_LOOP_COMBINED.  Update all call sites.
+
 2020-03-27  Sandra Loosemore  <sandra@codesourcery.com>
            Gergö Barany <gergo@codesourcery.com>
 

diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index dd7f548ac24f909b82ac630cc38099840a14646f..ec9996af7def313a421eb4a212fbc30ac5541b0b 100644 (file)
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -5071,7 +5071,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;
@@ -5435,7 +5436,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 ();
     }
 
@@ -5529,7 +5533,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));
@@ -6472,7 +6477,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)
@@ -6530,7 +6535,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)
@@ -6817,7 +6822,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
@@ -6876,13 +6882,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);
@@ -6962,7 +6968,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
@@ -7037,7 +7044,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
@@ -7081,7 +7089,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));
@@ -7479,7 +7487,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:
@@ -7519,7 +7527,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 9b9342011c4edce0a195888df9fcfe9fd3cbdd55..4377fa884641d3e692d6bad60d30f8695cc1ad70 100644 (file)
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1535,6 +1535,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, \