From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 39039 invoked by alias); 14 Jul 2015 05:46:10 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 39030 invoked by uid 89); 14 Jul 2015 05:46:09 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.5 required=5.0 tests=AWL,BAYES_00,KAM_ASCII_DIVIDERS,RCVD_IN_DNSWL_LOW,SPF_PASS autolearn=no version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 14 Jul 2015 05:46:08 +0000 Received: from svr-orw-fem-06.mgc.mentorg.com ([147.34.97.120]) by relay1.mentorg.com with esmtp id 1ZEt2Y-00023P-K2 from ChungLin_Tang@mentor.com ; Mon, 13 Jul 2015 22:46:02 -0700 Received: from [0.0.0.0] (147.34.91.1) by SVR-ORW-FEM-06.mgc.mentorg.com (147.34.97.120) with Microsoft SMTP Server id 14.3.224.2; Mon, 13 Jul 2015 22:46:02 -0700 Message-ID: <55A4A21C.1070004@codesourcery.com> Date: Tue, 14 Jul 2015 05:46:00 -0000 From: Chung-Lin Tang User-Agent: Mozilla/5.0 (Windows NT 6.1; WOW64; rv:31.0) Gecko/20100101 Thunderbird/31.7.0 MIME-Version: 1.0 To: gcc-patches , Tom de Vries CC: Thomas Schwinge , Jakub Jelinek Subject: [PATCH, gomp4] Propagate independent clause for OpenACC kernels pass Content-Type: multipart/mixed; boundary="------------040300030406060400000401" X-IsSubscribed: yes X-SW-Source: 2015-07/txt/msg01099.txt.bz2 --------------040300030406060400000401 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: 7bit Content-length: 1174 Hi Tom, this patch provides a 'bool independent' field in struct loop, which will be switched on by an "independent" clause in a #pragma acc loop directive. I assume you'll be wiring it to the kernels parloops pass in a followup patch. Note: there are already a few other similar fields in struct loop, namely 'safelen' and 'can_be_parallel', used by OMP simd safelen and GRAPHITE respectively. The intention and/or setting of these fields are all a bit different, so I've decided to add a new bool for OpenACC. Tested and committed to gomp-4_0-branch. Chung-Lin 2015-07-14 Chung-Lin Tang * cfgloop.h (struct loop): Add 'bool marked_independent' field. * gimplify.c (gimplify_scan_omp_clauses): Keep OMP_CLAUSE_INDEPENDENT. * omp-low.c (struct omp_region): Add 'int kind' and 'bool independent' fields. (expand_omp_for): Set 'marked_independent' field for loop corresponding to region. (find_omp_for_region_data): New function. (find_omp_target_region_data): Set kind field. (build_omp_regions_1): Call find_omp_for_region_data() for GIMPLE_OMP_FOR statements. --------------040300030406060400000401 Content-Type: text/x-patch; name="acc-loop-independent.patch" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="acc-loop-independent.patch" Content-length: 4392 Index: cfgloop.h =================================================================== --- cfgloop.h (revision 225758) +++ cfgloop.h (working copy) @@ -194,6 +194,10 @@ struct GTY ((chain_next ("%h.next"))) loop { /* True if the loop is part of an oacc kernels region. */ bool in_oacc_kernels_region; + /* True if loop is tagged as having independent iterations by user, + e.g. the OpenACC independent clause. */ + bool marked_independent; + /* For SIMD loops, this is a unique identifier of the loop, referenced by IFN_GOMP_SIMD_VF, IFN_GOMP_SIMD_LANE and IFN_GOMP_SIMD_LAST_LANE builtins. */ Index: gimplify.c =================================================================== --- gimplify.c (revision 225758) +++ gimplify.c (working copy) @@ -6602,7 +6602,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_se break; case OMP_CLAUSE_DEVICE_RESIDENT: - case OMP_CLAUSE_INDEPENDENT: remove = true; break; @@ -6612,6 +6611,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_se case OMP_CLAUSE_COLLAPSE: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_MERGEABLE: case OMP_CLAUSE_PROC_BIND: case OMP_CLAUSE_SAFELEN: Index: omp-low.c =================================================================== --- omp-low.c (revision 225758) +++ omp-low.c (working copy) @@ -136,8 +136,16 @@ struct omp_region /* True if this is nested inside an OpenACC kernels construct. */ bool inside_kernels_p; + /* Records a generic kind field. */ + int kind; + /* For an OpenACC loop, the level of parallelism requested. */ int gwv_this; + + /* For an OpenACC loop directive, true if has the 'independent' clause. */ + bool independent; + + tree broadcast_array; }; /* Context structure. Used to store information about each parallel @@ -8273,8 +8281,15 @@ expand_omp_for (struct omp_region *region, gimple loops_state_set (LOOPS_NEED_FIXUP); if (region->inside_kernels_p) - expand_omp_for_generic (region, &fd, BUILT_IN_NONE, BUILT_IN_NONE, - inner_stmt); + { + expand_omp_for_generic (region, &fd, BUILT_IN_NONE, BUILT_IN_NONE, + inner_stmt); + if (region->independent && region->cont->loop_father) + { + struct loop *loop = region->cont->loop_father; + loop->marked_independent = true; + } + } else if (gimple_omp_for_kind (fd.for_stmt) & GF_OMP_FOR_SIMD) expand_omp_simd (region, &fd); else if (gimple_omp_for_kind (fd.for_stmt) == GF_OMP_FOR_KIND_CILKFOR) @@ -9943,6 +9958,34 @@ find_omp_for_region_gwv (gimple stmt) return tmp; } +static void +find_omp_for_region_data (struct omp_region *region, gomp_for *stmt) +{ + region->gwv_this = find_omp_for_region_gwv (stmt); + region->kind = gimple_omp_for_kind (stmt); + + if (region->kind == GF_OMP_FOR_KIND_OACC_LOOP) + { + struct omp_region *target_region = region->outer; + while (target_region + && target_region->type != GIMPLE_OMP_TARGET) + target_region = target_region->outer; + if (!target_region) + return; + + tree clauses = gimple_omp_for_clauses (stmt); + + if (target_region->kind == GF_OMP_TARGET_KIND_OACC_PARALLEL + && !find_omp_clause (clauses, OMP_CLAUSE_SEQ)) + /* In OpenACC parallel constructs, 'independent' is implied on all + loop directives without a 'seq' clause. */ + region->independent = true; + else if (target_region->kind == GF_OMP_TARGET_KIND_OACC_KERNELS + && find_omp_clause (clauses, OMP_CLAUSE_INDEPENDENT)) + region->independent = true; + } +} + /* Fill in additional data for a region REGION associated with an OMP_TARGET STMT. */ @@ -9960,6 +10003,7 @@ find_omp_target_region_data (struct omp_region *re region->gwv_this |= OACC_LOOP_MASK (OACC_worker); if (find_omp_clause (clauses, OMP_CLAUSE_VECTOR_LENGTH)) region->gwv_this |= OACC_LOOP_MASK (OACC_vector); + region->kind = gimple_omp_target_kind (stmt); } /* Helper for build_omp_regions. Scan the dominator tree starting at @@ -10046,7 +10090,7 @@ build_omp_regions_1 (basic_block bb, struct omp_re } } else if (code == GIMPLE_OMP_FOR) - region->gwv_this = find_omp_for_region_gwv (stmt); + find_omp_for_region_data (region, as_a (stmt)); /* ..., this directive becomes the parent for a new region. */ if (region) parent = region; --------------040300030406060400000401--