From patchwork Wed Dec 9 17:44:30 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 554733 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id D90E21402DE for ; Thu, 10 Dec 2015 04:44:54 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=WAn7PzBD; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:to:cc :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=A9luD8JoOxy8ei1Q5sF4/mxeceSVgDb4WkTDzqsx9EgRTrmeMH z9l5ANC3KoCK8ZachZsuGZouK+oh49qqc4/uszqzEs5IviLT9pZU2p2ZocRchZfG Xcjn0J9PHRGtkuxZQS1kcsfsjV0LKkv8qxc3sOKjQAGxcO62yZU5nieHM= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:to:cc :from:subject:message-id:date:mime-version:content-type; s= default; bh=4dKXhnp34ql/d6wXJQOTg7f7dsc=; b=WAn7PzBDy6h4InAlVVm9 1gHUDvfmWz/Uy+RSeNIIfQdcbXnfSIv/aJPlaxr+1ewzGUvdaz/s/OCLBDF3CtXO d3TkQA2A+JzyelkS02YvlQYakp9LH/h70I1nvsm96pKWn0egPtPDD5+SOAtGC8yV NWPpo6d1+TrZzeaJu2NopJQ= Received: (qmail 29158 invoked by alias); 9 Dec 2015 17:44:47 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 29149 invoked by uid 89); 9 Dec 2015 17:44:46 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL, BAYES_00, SPF_PASS, T_RP_MATCHES_RCVD autolearn=ham version=3.3.2 X-HELO: fencepost.gnu.org Received: from fencepost.gnu.org (HELO fencepost.gnu.org) (208.118.235.10) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-SHA encrypted) ESMTPS; Wed, 09 Dec 2015 17:44:45 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:41487) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1a6inD-0002B0-5I for gcc-patches@gnu.org; Wed, 09 Dec 2015 12:44:43 -0500 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1a6in8-0001Db-Qb for gcc-patches@gnu.org; Wed, 09 Dec 2015 12:44:42 -0500 Received: from relay1.mentorg.com ([192.94.38.131]:41553) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1a6in8-0001Bz-Ij for gcc-patches@gnu.org; Wed, 09 Dec 2015 12:44:38 -0500 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1a6in6-0007eY-Eq from Tom_deVries@mentor.com ; Wed, 09 Dec 2015 09:44:36 -0800 Received: from [127.0.0.1] (137.202.0.76) by SVR-IES-FEM-02.mgc.mentorg.com (137.202.0.106) with Microsoft SMTP Server id 14.3.224.2; Wed, 9 Dec 2015 17:44:35 +0000 To: "gcc-patches@gnu.org" CC: Jakub Jelinek , Thomas Schwinge From: Tom de Vries Subject: [gomp4, committed] backport: "Fix oacc kernels default mapping for scalars" Message-ID: <5668687E.4010908@mentor.com> Date: Wed, 9 Dec 2015 18:44:30 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.4.0 MIME-Version: 1.0 X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 Hi, I've backported the patch "Fix oacc kernels default mapping for scalars" from trunk to gomp-4_0-branch. Committed to gomp-4_0-branch. Thanks, - Tom backport: "Fix oacc kernels default mapping for scalars" 2015-12-02 Tom de Vries backport from trunk: * gimplify.c (enum gimplify_omp_var_data): Add enum value GOVD_MAP_FORCE. (oacc_default_clause): Fix default for scalars in oacc kernels. (gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_FORCE. * c-c++-common/goacc/kernels-default-2.c: New test. * c-c++-common/goacc/kernels-default.c: New test. --- gcc/gimplify.c | 21 +++++++++++++++------ .../c-c++-common/goacc/kernels-default-2.c | 17 +++++++++++++++++ gcc/testsuite/c-c++-common/goacc/kernels-default.c | 14 ++++++++++++++ gcc/testsuite/gfortran.dg/goacc/reduction-2.f95 | 2 +- 4 files changed, 47 insertions(+), 7 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index e8964c6..d305165 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -90,8 +90,11 @@ enum gimplify_omp_var_data /* Flag for shared vars that are or might be stored to in the region. */ GOVD_WRITTEN = 131072, + /* Flag for GOVD_MAP, if it is a forced mapping. */ + GOVD_MAP_FORCE = 262144, + /* OpenACC deviceptr clause. */ - GOVD_USE_DEVPTR = 1 << 18, + GOVD_USE_DEVPTR = 1 << 19, GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR @@ -5988,8 +5991,12 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) gcc_unreachable (); case ORT_ACC_KERNELS: - /* Everything under kernels are default 'present_or_copy'. */ + /* Scalars are default 'copy' under kernels, non-scalars are default + 'present_or_copy'. */ flags |= GOVD_MAP; + if (!AGGREGATE_TYPE_P (TREE_TYPE (decl))) + flags |= GOVD_MAP_FORCE; + rkind = "kernels"; break; @@ -7700,10 +7707,12 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) } else if (code == OMP_CLAUSE_MAP) { - OMP_CLAUSE_SET_MAP_KIND (clause, - flags & GOVD_MAP_TO_ONLY - ? GOMP_MAP_TO - : GOMP_MAP_TOFROM); + int kind = (flags & GOVD_MAP_TO_ONLY + ? GOMP_MAP_TO + : GOMP_MAP_TOFROM); + if (flags & GOVD_MAP_FORCE) + kind |= GOMP_MAP_FLAG_FORCE; + OMP_CLAUSE_SET_MAP_KIND (clause, kind); if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c new file mode 100644 index 0000000..232b123 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c @@ -0,0 +1,17 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 2 + +void +foo (void) +{ + unsigned int a[N]; + +#pragma acc kernels + { + a[0]++; + } +} + +/* { dg-final { scan-tree-dump-times "map\\(tofrom" 1 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-default.c b/gcc/testsuite/c-c++-common/goacc/kernels-default.c new file mode 100644 index 0000000..58cd5e1 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-default.c @@ -0,0 +1,14 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void +foo (void) +{ + unsigned int i; +#pragma acc kernels + { + i++; + } +} + +/* { dg-final { scan-tree-dump-times "map\\(force_tofrom" 1 "gimple" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/reduction-2.f95 b/gcc/testsuite/gfortran.dg/goacc/reduction-2.f95 index 3f46eac..9d34cbd 100644 --- a/gcc/testsuite/gfortran.dg/goacc/reduction-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/reduction-2.f95 @@ -17,6 +17,6 @@ end subroutine ! { dg-final { scan-tree-dump-times "target oacc_parallel firstprivate.a." 1 "gimple" } } ! { dg-final { scan-tree-dump-times "acc loop reduction..:a. private.p." 1 "gimple" } } -! { dg-final { scan-tree-dump-times "target oacc_kernels map.tofrom:a .len: 4.." 1 "gimple" } } +! { dg-final { scan-tree-dump-times "target oacc_kernels map.force_tofrom:a .len: 4.." 1 "gimple" } } ! { dg-final { scan-tree-dump-times "acc loop reduction..:a. private.k." 1 "gimple" } }