diff mbox series

[06/40] Add a "combined" flag for "acc kernels loop" etc directives.

Message ID 20211215155447.19379-7-frederik@codesourcery.com
State New
Headers
Series OpenACC "kernels" Improvements |

Commit Message

Frederik Harwath Dec. 15, 2021, 3:54 p.m. UTC 
  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
diff mbox series

Patch

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 80dd61d599ef..1258b48693de 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -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);

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 4c2075742d6a..c834d25b028f 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -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));

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index e81c5588c53c..618e106791e5 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -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 = &block;
   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:
diff --git a/gcc/tree.h b/gcc/tree.h
index 7542d97ce121..15e5147f40b0 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -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, \