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.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 --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 29bcb59eff5f4cc8b477a12ec237ca08b5de9746..77d163266ee08c16898c59f6f7e794cda32bc270 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-07-30  Julian Brown  <julian@codesourcery.com>
 
        * 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 e6fadbb1e890922deba3937c708956e84982eb63..0f6ac422f3121fc7944b2378c68f405ebc174a42 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.c (c_parser_oacc_loop): Set OACC_LOOP_COMBINED.
+
 2020-03-27  Sandra Loosemore  <sandra@codesourcery.com>
 
        * c-decl.c (c_unwrap_for_init): New.

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 4225c730351810961314bbff56f3718cdeb472cb..2679b3e6405e75e74f5c02a66a1e48a68b43098d 100644 (file)
--- 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 adc28b921ac0db76c4837d8396c780c19ad31018..070f5633308c4a4090cff6e9f2da36bdb9f901f7 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.c (cp_parser_oacc_loop): Set OACC_LOOP_COMBINED.
+
 2020-03-31  Sandra Loosemore  <sandra@codesourcery.com>
 
        * semantics.c (handle_omp_array_sections_1): Call STRIP_NOPS

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index e42f32c116474110b0fdeb133c1ea2b6e9331438..245a044f5b68f0e640c0237de7d4fab194f4c946 100644 (file)
--- 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 e619434815e5ca89b08cb16d7be0bfcdb998b872..c608163f77fe6bfa7940eb5e0d3fb28fda77e6b1 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.c (gfc_trans_omp_do): Add combined parameter,
+       use it to set OACC_LOOP_COMBINED.  Update all call sites.
+
 2020-07-20  Frederik Harwath  <frederik@codesourcery.com>
 
        * openmp.c (oacc_is_parallel_or_serial): Removed function.

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 66e47fdbc7ef0cde0202dd5245ce83ad569b7a25..51414a12910b0ace867c11408d3df2b19c8734d7 100644 (file)
--- 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 0ae76b893569c33b5f760e10bf83d73a5a8b894c..edf4fd3797923dd9ac4e7f2d5fad158aafaaff8f 100644 (file)
--- 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, \