From: Frederik Harwath <frederik@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>
Cc: Sandra Loosemore <sandra@codesourcery.com>,
<thomas@codesourcery.com>, <joseph@codesourcery.com>,
<jason@redhat.com>, <nathan@acm.org>, <tobias@codesourcery.com>,
<fortran@gcc.gnu.org>
Subject: [PATCH 06/40] Add a "combined" flag for "acc kernels loop" etc directives.
Date: Wed, 15 Dec 2021 16:54:13 +0100 [thread overview]
Message-ID: <20211215155447.19379-7-frederik@codesourcery.com> (raw)
In-Reply-To: <20211215155447.19379-1-frederik@codesourcery.com>
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(-)
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 = █
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, \
--
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
next prev parent reply other threads:[~2021-12-15 15:55 UTC|newest]
Thread overview: 9+ messages / expand[flat|nested] mbox.gz Atom feed top
2021-12-15 15:54 [PATCH 00/40] OpenACC "kernels" Improvements Frederik Harwath
2021-12-15 15:54 ` [PATCH 03/40] Kernels loops annotation: Fortran Frederik Harwath
2021-12-15 15:54 ` [PATCH 04/40] Additional Fortran testsuite fixes for kernels loops annotation pass Frederik Harwath
2021-12-15 15:54 ` Frederik Harwath [this message]
2021-12-15 15:54 ` [PATCH 08/40] Annotate inner loops in "acc kernels loop" directives (Fortran) Frederik Harwath
2021-12-15 15:54 ` [PATCH 09/40] Permit calls to builtins and intrinsics in kernels loops Frederik Harwath
2021-12-15 15:54 ` [PATCH 10/40] Fix patterns in Fortran tests for kernels loop annotation Frederik Harwath
2021-12-15 15:54 ` [PATCH 13/40] Fortran: Delinearize array accesses Frederik Harwath
2021-12-16 12:00 ` [PATCH 40/40] openacc: Adjust testsuite to new "kernels" handling Frederik Harwath
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=20211215155447.19379-7-frederik@codesourcery.com \
--to=frederik@codesourcery.com \
--cc=fortran@gcc.gnu.org \
--cc=gcc-patches@gcc.gnu.org \
--cc=jason@redhat.com \
--cc=joseph@codesourcery.com \
--cc=nathan@acm.org \
--cc=sandra@codesourcery.com \
--cc=thomas@codesourcery.com \
--cc=tobias@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
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).