From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 92902 invoked by alias); 7 Sep 2015 12:48:23 -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 92885 invoked by uid 89); 7 Sep 2015 12:48:21 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.4 required=5.0 tests=AWL,BAYES_00,KAM_LAZY_DOMAIN_SECURITY,SPF_HELO_PASS,T_RP_MATCHES_RCVD autolearn=no version=3.3.2 X-HELO: mx1.redhat.com Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Mon, 07 Sep 2015 12:48:20 +0000 Received: from int-mx14.intmail.prod.int.phx2.redhat.com (int-mx14.intmail.prod.int.phx2.redhat.com [10.5.11.27]) by mx1.redhat.com (Postfix) with ESMTPS id 6B90E8AE6C; Mon, 7 Sep 2015 12:48:19 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-44.ams2.redhat.com [10.36.116.44]) by int-mx14.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t87CmH30014911 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Mon, 7 Sep 2015 08:48:18 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id t87CmG4g021273; Mon, 7 Sep 2015 14:48:16 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id t87CmF2O021272; Mon, 7 Sep 2015 14:48:15 +0200 Date: Mon, 07 Sep 2015 12:48:00 -0000 From: Jakub Jelinek To: Ilya Verbin Cc: gcc-patches@gcc.gnu.org, Kirill Yukhin Subject: Re: [gomp4.1] Various accelerator updates from OpenMP 4.1 Message-ID: <20150907124815.GR1847@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <20150625194529.GB33078@msticlxl57.ims.intel.com> <20150625201058.GK10247@tucnak.redhat.com> <20150717163136.GB15252@msticlxl57.ims.intel.com> <20150717164306.GT1780@tucnak.redhat.com> <20150720161422.GC1780@tucnak.redhat.com> <20150720181041.GE1780@tucnak.redhat.com> <20150722211348.GA1750@tucnak.redhat.com> <20150724200457.GB1750@tucnak.redhat.com> <20150729171907.GI1780@tucnak.redhat.com> <20150904180702.GB10772@msticlxl57.ims.intel.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20150904180702.GB10772@msticlxl57.ims.intel.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-09/txt/msg00432.txt.bz2 On Fri, Sep 04, 2015 at 09:07:02PM +0300, Ilya Verbin wrote: > It seems that there is a bug some here: > > Here is the reproducer: > > #pragma omp declare target > int a[1]; > #pragma omp end declare target > > void foo () > { > #pragma omp target map(to: a[0:1]) > a; > } > > > lookup_decl (var, ctx) tries to lookup for 'a', but ctx->cb.decl_map->get () > returns null-pointer. Fixed thusly, tested on x86_64-linux, committed to gomp4.1 branch. 2015-09-07 Jakub Jelinek * omp-low.c (scan_sharing_clauses): Don't ignore map with declare target vars for GOMP_MAP_FIRSTPRIVATE_POINTER, unless the decl is an array. (lower_omp_target): Ignore GOMP_MAP_FIRSTPRIVATE_POINTER map with declare target var if it is an array. * testsuite/libgomp.c/target-26.c: New test. --- gcc/omp-low.c.jj 2015-09-04 11:34:45.000000000 +0200 +++ gcc/omp-low.c 2015-09-07 14:10:36.198517500 +0200 @@ -2060,6 +2060,8 @@ scan_sharing_clauses (tree clauses, omp_ directly. */ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && DECL_P (decl) + && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER + || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable) break; @@ -2284,6 +2286,8 @@ scan_sharing_clauses (tree clauses, omp_ break; decl = OMP_CLAUSE_DECL (c); if (DECL_P (decl) + && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER + || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable) break; @@ -13358,6 +13362,10 @@ lower_omp_target (gimple_stmt_iterator * { if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) { + if (is_global_var (maybe_lookup_decl_in_outer_ctx (var, ctx)) + && varpool_node::get_create (var)->offloadable) + continue; + tree type = build_pointer_type (TREE_TYPE (var)); tree new_var = lookup_decl (var, ctx); x = create_tmp_var_raw (type, get_name (new_var)); @@ -14081,6 +14089,12 @@ lower_omp_target (gimple_stmt_iterator * HOST_WIDE_INT offset = 0; gcc_assert (prev); var = OMP_CLAUSE_DECL (c); + if (DECL_P (var) + && TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE + && is_global_var (maybe_lookup_decl_in_outer_ctx (var, + ctx)) + && varpool_node::get_create (var)->offloadable) + break; if (TREE_CODE (var) == INDIRECT_REF && TREE_CODE (TREE_OPERAND (var, 0)) == COMPONENT_REF) var = TREE_OPERAND (var, 0); --- libgomp/testsuite/libgomp.c/target-26.c.jj 2015-09-07 11:59:08.665425993 +0200 +++ libgomp/testsuite/libgomp.c/target-26.c 2015-09-07 12:38:56.000000000 +0200 @@ -0,0 +1,36 @@ +extern void abort (void); +#pragma omp declare target +int a[4] = { 2, 3, 4, 5 }, *b; +#pragma omp end declare target + +int +main () +{ + int err; + int c[3] = { 6, 7, 8 }; + b = c; + #pragma omp target map(to: a[0:2], b[0:2]) map(from: err) + err = a[0] != 2 || a[1] != 3 || a[2] != 4 || a[3] != 5 || b[0] != 6 || b[1] != 7; + if (err) + abort (); + a[1] = 9; + a[2] = 10; + #pragma omp target map(always,to:a[1:2]) map(from: err) + err = a[0] != 2 || a[1] != 9 || a[2] != 10 || a[3] != 5; + if (err) + abort (); + #pragma omp parallel firstprivate(a, b, c, err) num_threads (2) + #pragma omp single + { + b = c + 1; + a[0] = 11; + a[2] = 13; + c[1] = 14; + int d = 0; + #pragma omp target map(to: a[0:3], b[d:2]) map (from: err) + err = a[0] != 11 || a[1] != 9 || a[2] != 13 || b[0] != 14 || b[1] != 8; + if (err) + abort (); + } + return 0; +} Jakub