From patchwork Wed Dec 15 15:54:13 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 48948 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 00D743857C77 for ; Wed, 15 Dec 2021 15:59:25 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 908E1385842F; Wed, 15 Dec 2021 15:55:19 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 908E1385842F Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: aLzVlocf82sWQygCJ+bv+FlK66JvYsc4v22g4vSHgwMhnrz6PjnbSzJz7RBjPHvcqg7rZ44NAe RXTpeCoN2a/Ymv/I1JKtQOj7GZw0+3kMsrVODY1s8R23jtRLNQje4SC/NdumfPfkYqYpdr1Knd PqpkyOO8AybwP9nPRcCwwvUB/tcK9hz4uFpP0iAqyVKgEhlCykZs5WHNqNk7xicWQk0GFUHzn+ /j9M+iYFIoaTxfTIUiEz+Z6h9YNEV4uGMusTrRcTEbUITv0krjrDNHBk6XeRODaNL6tkRoH03e /X6Pt5ZzqMdsrfUwqapihaEf X-IronPort-AV: E=Sophos;i="5.88,207,1635235200"; d="scan'208";a="69736538" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 15 Dec 2021 07:55:19 -0800 IronPort-SDR: ALLs/6P9Sk/TXEZe7ZJG0VQk0GyPxrjHfBB8ho6YGtKXsG+jIL4I/QtZbDtcQ6m+vFHB7KgslS BFWtoS2CXiSwzZjjnhKYTxhjrCAivOyFoSpr0vnGyD0vQ0p9dA9kOi3DSgqHss7RyB2SrEkUPj n5ek/c3VrCOLQnO5kbdOSxYqzpwRfMLv/yfHBG5xXJCPpu+2+LI2fU614RsL/4P1h6Tzkvwr/F 77dM/EDTCOegRf8X0iwzbj1e730Lgse77nPsM/AWgtzxZguVhfAm8xFcKzJnsXs+8lhgxk4rB2 TZg= From: Frederik Harwath To: Subject: [PATCH 06/40] Add a "combined" flag for "acc kernels loop" etc directives. Date: Wed, 15 Dec 2021 16:54:13 +0100 Message-ID: <20211215155447.19379-7-frederik@codesourcery.com> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211215155447.19379-1-frederik@codesourcery.com> References: <20211215155447.19379-1-frederik@codesourcery.com> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-2.6 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, UNWANTED_LANGUAGE_BODY autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: fortran@gcc.gnu.org, nathan@acm.org, Sandra Loosemore , tobias@codesourcery.com, thomas@codesourcery.com, joseph@codesourcery.com Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" From: Sandra Loosemore 2020-08-19 Sandra Loosemore 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 --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, \