public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] Add a "combined" flag for "acc kernels loop" etc directives.
@ 2021-05-13 16:17 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2021-05-13 16:17 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:506098b0fab2cfbd81f813de97101f2b242bc18e
commit 506098b0fab2cfbd81f813de97101f2b242bc18e
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.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:
---
gcc/ChangeLog.omp | 6 ++++++
gcc/c/ChangeLog.omp | 6 ++++++
gcc/c/c-parser.c | 3 +++
gcc/cp/ChangeLog.omp | 6 ++++++
gcc/cp/parser.c | 3 +++
gcc/fortran/ChangeLog.omp | 7 +++++++
gcc/fortran/trans-openmp.c | 30 +++++++++++++++++++-----------
gcc/tree.h | 5 +++++
8 files changed, 55 insertions(+), 11 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 29bcb59eff5..77d163266ee 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-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 e6fadbb1e89..0f6ac422f31 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.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 4225c730351..2679b3e6405 100644
--- 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 adc28b921ac..070f5633308 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.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 e42f32c1164..245a044f5b6 100644
--- 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 e619434815e..c608163f77f 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.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 66e47fdbc7e..51414a12910 100644
--- 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 0ae76b89356..edf4fd37979 100644
--- 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, \
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-13 16:17 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-13 16:17 [gcc/devel/omp/gcc-11] 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).