public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] Add a "combined" flag for "acc kernels loop" etc directives.
@ 2022-06-29 14:37 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2022-06-29 14:37 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:652fff3f4cb42d95e095991f1f4e173f4347d630
commit 652fff3f4cb42d95e095991f1f4e173f4347d630
Author: Sandra Loosemore <sandra@codesourcery.com>
Date: Wed Aug 19 19:13:55 2020 -0700
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:
---
gcc/ChangeLog.omp | 6 ++++++
gcc/c/ChangeLog.omp | 6 ++++++
gcc/c/c-parser.cc | 3 +++
gcc/cp/ChangeLog.omp | 6 ++++++
gcc/cp/parser.cc | 3 +++
gcc/fortran/ChangeLog.omp | 7 +++++++
gcc/fortran/trans-openmp.cc | 34 +++++++++++++++++++++-------------
gcc/tree.h | 5 +++++
8 files changed, 57 insertions(+), 13 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 750f7a2aad2..bc3360e053c 100644
--- 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 35c7a52c357..52af6cf6316 100644
--- 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 3d4f03c98a6..85bc30634e7 100644
--- 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 ead9de41a2c..d8156d528c6 100644
--- 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 f98854dc8ec..89d27520a27 100644
--- 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 a28bc3dc789..df352c58706 100644
--- 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 dd7f548ac24..ec9996af7de 100644
--- 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 9b9342011c4..4377fa88464 100644
--- 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, \
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-06-29 14:37 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-06-29 14:37 [gcc/devel/omp/gcc-12] Add a "combined" flag for "acc kernels loop" etc directives Kwok Yeung
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).