From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 40774 invoked by alias); 22 Jul 2015 21:14:01 -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 40711 invoked by uid 89); 22 Jul 2015 21:14:00 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.3 required=5.0 tests=AWL,BAYES_05,KAM_LAZY_DOMAIN_SECURITY,RP_MATCHES_RCVD,SPF_HELO_PASS autolearn=ham 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; Wed, 22 Jul 2015 21:13:55 +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 3073C8E904; Wed, 22 Jul 2015 21:13:54 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-43.ams2.redhat.com [10.36.116.43]) by int-mx14.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t6MLDq8o008435 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Wed, 22 Jul 2015 17:13:53 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.9/8.14.9) with ESMTP id t6MLDoMY031974; Wed, 22 Jul 2015 23:13:50 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.9/8.14.9/Submit) id t6MLDm0t031973; Wed, 22 Jul 2015 23:13:48 +0200 Date: Thu, 23 Jul 2015 00:50:00 -0000 From: Jakub Jelinek To: Ilya Verbin Cc: Thomas Schwinge , gcc-patches@gcc.gnu.org, Kirill Yukhin Subject: Re: [gomp4.1] Initial support for some OpenMP 4.1 construct parsing Message-ID: <20150722211348.GA1750@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <874mnzrw1z.fsf@schwinge.name> <20150429120644.GG1751@tucnak.redhat.com> <20150609183608.GA47936@msticlxl57.ims.intel.com> <20150609202426.GG10247@tucnak.redhat.com> <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> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20150720181041.GE1780@tucnak.redhat.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-07/txt/msg01893.txt.bz2 On Mon, Jul 20, 2015 at 08:10:41PM +0200, Jakub Jelinek wrote: > And here is untested incremental libgomp side of the proposed > GOMP_MAP_FIRSTPRIVATE_POINTER. Actually, that seems unnecessary, for the array section maps we already have there a pointer, so we can easily implement that just on the compiler side. Here is a WIP patch. Unfortunately, in order not to break numerous examples-4/ testcases that were doing target data map of array sections with target region without any explicit maps, with the new way where target{, enter, exit} data no longer map the base pointer, I had to implement the new implicit pointer mapping semantics (map (alloc:ptr[0:0])) already in this patch. And, that patch really requires that if there is ptr[0:something] for something > 0 already mapped that we use the ptr[0:something] mapping rather than ptr[0:0]. See the libgomp changes for that. Unfortunately, that occassionally breaks the target8.f90 testcase at -O0, where we map zero-sized FRAME.6 object which happens to be adjacent to the array. And that reveals IMNSHO very serious flaw in the current standard draft, no idea what can be done about that... --- libgomp/testsuite/libgomp.c++/target-7.C.jj 2015-07-22 11:36:53.042867520 +0200 +++ libgomp/testsuite/libgomp.c++/target-7.C 2015-07-22 11:32:00.000000000 +0200 @@ -0,0 +1,90 @@ +extern "C" void abort (); + +void +foo (int *x, int *&y, int (&z)[15]) +{ + int a[10], b[15], err, i; + for (i = 0; i < 10; i++) + a[i] = 7 * i; + for (i = 0; i < 15; i++) + b[i] = 8 * i; + #pragma omp target map(to:x[5:10], y[5:10], z[5:10], a[0:10], b[5:10]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if (x[5 + i] != 20 + 4 * i + || y[5 + i] != 25 + 5 * i + || z[5 + i] != 30 + 6 * i + || a[i] != 7 * i + || b[5 + i] != 40 + 8 * i) + err = 1; + } + if (err) + abort (); +} + +void +bar (int n, int v) +{ + int a[n], b[n], c[n], d[n], e[n], err, i; + int (*x)[n] = &c; + int (*y2)[n] = &d; + int (*&y)[n] = y2; + int (&z)[n] = e; + for (i = 0; i < n; i++) + { + (*x)[i] = 4 * i; + (*y)[i] = 5 * i; + z[i] = 6 * i; + a[i] = 7 * i; + b[i] = 8 * i; + } + #pragma omp target map(to:x[0][5:10], y[0][5:10], z[5:10], a[0:10], b[5:10]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if ((*x)[5 + i] != 20 + 4 * i + || (*y)[5 + i] != 25 + 5 * i + || z[5 + i] != 30 + 6 * i + || a[i] != 7 * i + || b[5 + i] != 40 + 8 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + { + (*x)[i] = 9 * i; + (*y)[i] = 10 * i; + z[i] = 11 * i; + a[i] = 12 * i; + b[i] = 13 * i; + } + #pragma omp target map(to:x[0][v:v+5], y[0][v:v+5], z[v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if ((*x)[5 + i] != 45 + 9 * i + || (*y)[5 + i] != 50 + 10 * i + || z[5 + i] != 55 + 11 * i + || a[i] != 12 * i + || b[5 + i] != 65 + 13 * i) + err = 1; + } + if (err) + abort (); +} + +int +main () +{ + int x[15], y2[15], z[15], *y = y2, i; + for (i = 0; i < 15; i++) + { + x[i] = 4 * i; + y[i] = 5 * i; + z[i] = 6 * i; + } + foo (x, y, z); + bar (15, 5); +} --- libgomp/testsuite/libgomp.c/target-15.c.jj 2015-07-22 11:37:11.655612690 +0200 +++ libgomp/testsuite/libgomp.c/target-15.c 2015-07-22 11:38:54.590203394 +0200 @@ -0,0 +1,74 @@ +extern void abort (); + +void +foo (int *x) +{ + int a[10], b[15], err, i; + for (i = 0; i < 10; i++) + a[i] = 7 * i; + for (i = 0; i < 15; i++) + b[i] = 8 * i; + #pragma omp target map(to:x[5:10], a[0:10], b[5:10]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if (x[5 + i] != 20 + 4 * i + || a[i] != 7 * i + || b[5 + i] != 40 + 8 * i) + err = 1; + } + if (err) + abort (); +} + +void +bar (int n, int v) +{ + int a[n], b[n], c[n], d[n], e[n], err, i; + int (*x)[n] = &c; + for (i = 0; i < n; i++) + { + (*x)[i] = 4 * i; + a[i] = 7 * i; + b[i] = 8 * i; + } + #pragma omp target map(to:x[0][5:10], a[0:10], b[5:10]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if ((*x)[5 + i] != 20 + 4 * i + || a[i] != 7 * i + || b[5 + i] != 40 + 8 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + { + (*x)[i] = 9 * i; + a[i] = 12 * i; + b[i] = 13 * i; + } + #pragma omp target map(to:x[0][v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if ((*x)[5 + i] != 45 + 9 * i + || a[i] != 12 * i + || b[5 + i] != 65 + 13 * i) + err = 1; + } + if (err) + abort (); +} + +int +main () +{ + int x[15], i; + for (i = 0; i < 15; i++) + x[i] = 4 * i; + foo (x); + bar (15, 5); + return 0; +} --- libgomp/target.c.jj 2015-07-21 09:07:23.690851224 +0200 +++ libgomp/target.c 2015-07-22 21:12:22.438213557 +0200 @@ -142,7 +142,26 @@ resolve_device (int device_id) } -/* Handle the case where splay_tree_lookup found oldn for newn. +static inline splay_tree_key +gomp_map_lookup (splay_tree mem_map, splay_tree_key key) +{ + if (key->host_start != key->host_end) + return splay_tree_lookup (mem_map, key); + + key->host_end++; + splay_tree_key n = splay_tree_lookup (mem_map, key); + key->host_end--; + if (n) + return n; + key->host_start--; + n = splay_tree_lookup (mem_map, key); + key->host_start++; + if (n) + return n; + return splay_tree_lookup (mem_map, key); +} + +/* Handle the case where gmp_map_lookup found oldn for newn. Helper function of gomp_map_vars. */ static inline void @@ -204,20 +223,8 @@ gomp_map_pointer (struct target_mem_desc } /* Add bias to the pointer value. */ cur_node.host_start += bias; - cur_node.host_end = cur_node.host_start + 1; - splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); - if (n == NULL) - { - /* Could be possibly zero size array section. */ - cur_node.host_end--; - n = splay_tree_lookup (mem_map, &cur_node); - if (n == NULL) - { - cur_node.host_start--; - n = splay_tree_lookup (mem_map, &cur_node); - cur_node.host_start++; - } - } + cur_node.host_end = cur_node.host_start; + splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n == NULL) { gomp_mutex_unlock (&devicep->lock); @@ -293,7 +300,7 @@ gomp_map_vars (struct gomp_device_descr has_firstprivate = true; continue; } - splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); + splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n) gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i], kind & typemask); @@ -392,7 +399,7 @@ gomp_map_vars (struct gomp_device_descr k->host_end = k->host_start + sizes[i]; else k->host_end = k->host_start + sizeof (void *); - splay_tree_key n = splay_tree_lookup (mem_map, k); + splay_tree_key n = gomp_map_lookup (mem_map, k); if (n) gomp_map_vars_existing (devicep, n, k, &tgt->list[i], kind & typemask); @@ -526,7 +533,8 @@ gomp_map_vars (struct gomp_device_descr } else cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start - + tgt->list[i].key->tgt_offset; + + tgt->list[i].key->tgt_offset + + tgt->list[i].offset; /* FIXME: see above FIXME comment. */ devicep->host2dev_func (devicep->target_id, (void *) (tgt->tgt_start @@ -1289,20 +1297,8 @@ omp_target_is_present (void *ptr, size_t struct splay_tree_key_s cur_node; cur_node.host_start = (uintptr_t) ptr + offset; - cur_node.host_end = cur_node.host_start + 1; - splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); - if (n == NULL) - { - /* Could be possibly zero size array section. */ - cur_node.host_end--; - n = splay_tree_lookup (mem_map, &cur_node); - if (n == NULL) - { - cur_node.host_start--; - n = splay_tree_lookup (mem_map, &cur_node); - cur_node.host_start++; - } - } + cur_node.host_end = cur_node.host_start; + splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); int ret = n != NULL; gomp_mutex_unlock (&devicep->lock); return ret; @@ -1524,7 +1520,7 @@ omp_target_associate_ptr (void *host_ptr cur_node.host_start = (uintptr_t) host_ptr; cur_node.host_end = cur_node.host_start + size; - splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); + splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n) { if (n->tgt->tgt_start + n->tgt_offset @@ -1584,13 +1580,8 @@ omp_target_disassociate_ptr (void *ptr, int ret = EINVAL; cur_node.host_start = (uintptr_t) ptr; - cur_node.host_end = cur_node.host_start + 1; - splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); - if (n == NULL) - { - cur_node.host_end--; - n = splay_tree_lookup (mem_map, &cur_node); - } + cur_node.host_end = cur_node.host_start; + splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n && n->host_start == cur_node.host_start && n->refcount == REFCOUNT_INFINITY --- libgomp/libgomp.h.jj 2015-07-15 13:00:32.000000000 +0200 +++ libgomp/libgomp.h 2015-07-22 21:09:39.023307107 +0200 @@ -647,11 +647,9 @@ struct target_var_desc { bool copy_from; /* True if data always should be copied from device to host at the end. */ bool always_copy_from; - /* Used for unmapping of array sections, can be nonzero only when - always_copy_from is true. */ + /* Relative offset against key host_start. */ uintptr_t offset; - /* Used for unmapping of array sections, can be less than the size of the - whole object only when always_copy_from is true. */ + /* Actual length. */ uintptr_t length; }; --- include/gomp-constants.h.jj 2015-07-21 09:07:23.689851239 +0200 +++ include/gomp-constants.h 2015-07-21 15:01:05.384829637 +0200 @@ -95,7 +95,11 @@ enum gomp_map_kind GOMP_MAP_DELETE = GOMP_MAP_FORCE_DEALLOC, /* Decrement usage count and deallocate if zero. */ GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_ALWAYS - | GOMP_MAP_FORCE_DEALLOC) + | GOMP_MAP_FORCE_DEALLOC), + + /* Internal to GCC, not used in libgomp. */ + /* Do not map, but pointer assign a pointer instead. */ + GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1) }; #define GOMP_MAP_COPY_TO_P(X) \ --- gcc/cp/parser.c.jj 2015-07-21 09:06:42.000000000 +0200 +++ gcc/cp/parser.c 2015-07-21 12:55:03.966656803 +0200 @@ -32276,27 +32276,28 @@ cp_parser_omp_target_data (cp_parser *pa for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc)) - { - case GOMP_MAP_TO: - case GOMP_MAP_ALWAYS_TO: - case GOMP_MAP_FROM: - case GOMP_MAP_ALWAYS_FROM: - case GOMP_MAP_TOFROM: - case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: - map_seen = 3; - break; - default: - map_seen |= 1; - error_at (OMP_CLAUSE_LOCATION (*pc), - "%<#pragma omp target data%> with map-type other " - "than %, %, % or % " - "on % clause"); - *pc = OMP_CLAUSE_CHAIN (*pc); - continue; - } + switch (OMP_CLAUSE_MAP_KIND (*pc)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_ALLOC: + map_seen = 3; + break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; + default: + map_seen |= 1; + error_at (OMP_CLAUSE_LOCATION (*pc), + "%<#pragma omp target data%> with map-type other " + "than %, %, % or % " + "on % clause"); + *pc = OMP_CLAUSE_CHAIN (*pc); + continue; + } pc = &OMP_CLAUSE_CHAIN (*pc); } @@ -32370,22 +32371,23 @@ cp_parser_omp_target_enter_data (cp_pars for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc)) - { - case GOMP_MAP_TO: - case GOMP_MAP_ALWAYS_TO: - case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: - map_seen = 3; - break; - default: - map_seen |= 1; - error_at (OMP_CLAUSE_LOCATION (*pc), - "%<#pragma omp target enter data%> with map-type other " - "than % or % on % clause"); - *pc = OMP_CLAUSE_CHAIN (*pc); - continue; - } + switch (OMP_CLAUSE_MAP_KIND (*pc)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALLOC: + map_seen = 3; + break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; + default: + map_seen |= 1; + error_at (OMP_CLAUSE_LOCATION (*pc), + "%<#pragma omp target enter data%> with map-type other " + "than % or % on % clause"); + *pc = OMP_CLAUSE_CHAIN (*pc); + continue; + } pc = &OMP_CLAUSE_CHAIN (*pc); } @@ -32455,24 +32457,25 @@ cp_parser_omp_target_exit_data (cp_parse for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc)) - { - case GOMP_MAP_FROM: - case GOMP_MAP_ALWAYS_FROM: - case GOMP_MAP_RELEASE: - case GOMP_MAP_DELETE: - case GOMP_MAP_POINTER: - map_seen = 3; - break; - default: - map_seen |= 1; - error_at (OMP_CLAUSE_LOCATION (*pc), - "%<#pragma omp target exit data%> with map-type other " - "than %, % or % on %" - " clause"); - *pc = OMP_CLAUSE_CHAIN (*pc); - continue; - } + switch (OMP_CLAUSE_MAP_KIND (*pc)) + { + case GOMP_MAP_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + map_seen = 3; + break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; + default: + map_seen |= 1; + error_at (OMP_CLAUSE_LOCATION (*pc), + "%<#pragma omp target exit data%> with map-type other " + "than %, % or % on %" + " clause"); + *pc = OMP_CLAUSE_CHAIN (*pc); + continue; + } pc = &OMP_CLAUSE_CHAIN (*pc); } @@ -32697,7 +32700,7 @@ check_clauses: case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: + case GOMP_MAP_FIRSTPRIVATE_POINTER: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), --- gcc/cp/semantics.c.jj 2015-07-17 13:59:27.000000000 +0200 +++ gcc/cp/semantics.c 2015-07-22 13:01:26.296499686 +0200 @@ -4650,7 +4650,7 @@ handle_omp_array_sections_1 (tree c, tre /* Handle array sections for clause C. */ static bool -handle_omp_array_sections (tree c) +handle_omp_array_sections (tree c, bool is_omp) { bool maybe_zero_len = false; unsigned int first_non_one = 0; @@ -4828,8 +4828,9 @@ handle_omp_array_sections (tree c) return false; tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); - if (!cxx_mark_addressable (t)) + OMP_CLAUSE_SET_MAP_KIND (c2, is_omp ? GOMP_MAP_FIRSTPRIVATE_POINTER + : GOMP_MAP_POINTER); + if (!is_omp && !cxx_mark_addressable (t)) return false; OMP_CLAUSE_DECL (c2) = t; t = build_fold_addr_expr (first); @@ -4847,7 +4848,8 @@ handle_omp_array_sections (tree c) OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = c2; ptr = OMP_CLAUSE_DECL (c2); - if (TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE + if (!is_omp + && TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE && POINTER_TYPE_P (TREE_TYPE (TREE_TYPE (ptr)))) { tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), @@ -5569,7 +5571,7 @@ finish_omp_clauses (tree clauses, bool a t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, allow_fields)) { remove = true; break; @@ -6155,7 +6157,7 @@ finish_omp_clauses (tree clauses, bool a } if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, allow_fields)) remove = true; break; } @@ -6189,7 +6191,7 @@ finish_omp_clauses (tree clauses, bool a t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, allow_fields)) remove = true; else { @@ -6242,7 +6244,9 @@ finish_omp_clauses (tree clauses, bool a && !cxx_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_POINTER))) && !type_dependent_expression_p (t) && !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE) --- gcc/tree.h.jj 2015-07-16 17:56:41.000000000 +0200 +++ gcc/tree.h 2015-07-21 16:40:35.759401307 +0200 @@ -1445,7 +1445,7 @@ extern void protected_set_expr_location ((enum gomp_map_kind) OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind) #define OMP_CLAUSE_SET_MAP_KIND(NODE, MAP_KIND) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind \ - = (unsigned char) (MAP_KIND)) + = (unsigned int) (MAP_KIND)) /* Nonzero if this map clause is for array (rather than pointer) based array section with zero bias. Both the non-decl OMP_CLAUSE_MAP and corresponding --- gcc/gimplify.c.jj 2015-07-16 17:56:41.000000000 +0200 +++ gcc/gimplify.c 2015-07-22 18:27:58.545933111 +0200 @@ -90,6 +90,8 @@ enum gimplify_omp_var_data /* Flag for GOVD_LINEAR or GOVD_LASTPRIVATE: no outer reference. */ GOVD_LINEAR_LASTPRIVATE_NO_OUTER = 16384, + GOVD_MAP_0LEN_ARRAY = 32768, + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -156,6 +158,8 @@ struct gimplify_omp_ctx enum omp_region_type region_type; bool combined_loop; bool distribute; + bool target_map_scalars_firstprivate; + bool target_map_pointers_as_0len_arrays; }; static struct gimplify_ctx *gimplify_ctxp; @@ -5821,6 +5825,38 @@ omp_notice_variable (struct gimplify_omp "a mappable type", decl); omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags); } + else if (ctx->target_map_pointers_as_0len_arrays + || ctx->target_map_scalars_firstprivate) + { + bool is_declare_target = false; + if (is_global_var (decl) + && varpool_node::get_create (decl)->offloadable) + { + struct gimplify_omp_ctx *octx; + for (octx = ctx->outer_context; + octx; octx = octx->outer_context) + { + n = splay_tree_lookup (octx->variables, + (splay_tree_key)decl); + if (n + && (n->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED + && (n->value & GOVD_DATA_SHARE_CLASS) != 0) + break; + } + is_declare_target = octx == NULL; + } + if (is_declare_target) + omp_add_variable (ctx, decl, GOVD_MAP | flags); + else if (ctx->target_map_pointers_as_0len_arrays + && (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE + || (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE + && TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) + == POINTER_TYPE))) + omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_0LEN_ARRAY + | flags); + else + omp_add_variable (ctx, decl, GOVD_MAP | flags); + } else omp_add_variable (ctx, decl, GOVD_MAP | flags); } @@ -6144,6 +6180,13 @@ gimplify_scan_omp_clauses (tree *list_p, ctx = new_omp_context (region_type); outer_ctx = ctx->outer_context; + if (code == OMP_TARGET && !lang_GNU_Fortran ()) + { + ctx->target_map_pointers_as_0len_arrays = true; + /* FIXME: For Fortran we want to set this too, when + the Fortran FE is updated to OpenMP 4.1. */ + ctx->target_map_scalars_firstprivate = true; + } while ((c = *list_p) != NULL) { @@ -6319,10 +6362,24 @@ gimplify_scan_omp_clauses (tree *list_p, case OMP_CLAUSE_MAP: decl = OMP_CLAUSE_DECL (c); if (error_operand_p (decl)) + remove = true; + switch (code) { - remove = true; + case OMP_TARGET: + break; + case OMP_TARGET_DATA: + case OMP_TARGET_ENTER_DATA: + case OMP_TARGET_EXIT_DATA: + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + /* For target {,enter ,exit }data only the array slice is + mapped, but not the pointer to it. */ + remove = true; + break; + default: break; } + if (remove) + break; if (OMP_CLAUSE_SIZE (c) == NULL_TREE) OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl) : TYPE_SIZE_UNIT (TREE_TYPE (decl)); @@ -6332,6 +6389,14 @@ gimplify_scan_omp_clauses (tree *list_p, remove = true; break; } + else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST) + { + OMP_CLAUSE_SIZE (c) + = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL); + omp_add_variable (ctx, OMP_CLAUSE_SIZE (c), + GOVD_FIRSTPRIVATE | GOVD_SEEN); + } if (!DECL_P (decl)) { if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, @@ -6643,7 +6708,10 @@ gimplify_scan_omp_clauses (tree *list_p, case OMP_CLAUSE_NOGROUP: case OMP_CLAUSE_THREADS: case OMP_CLAUSE_SIMD: + break; + case OMP_CLAUSE_DEFAULTMAP: + ctx->target_map_scalars_firstprivate = false; break; case OMP_CLAUSE_ALIGNED: @@ -6759,6 +6827,29 @@ gimplify_adjust_omp_clauses_1 (splay_tre OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1; else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF)) OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1; + else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0) + { + tree nc = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (nc) = decl; + if (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE + && TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) == POINTER_TYPE) + OMP_CLAUSE_DECL (clause) + = build_simple_mem_ref_loc (input_location, decl); + OMP_CLAUSE_DECL (clause) + = build2 (MEM_REF, char_type_node, OMP_CLAUSE_DECL (clause), + build_int_cst (build_pointer_type (char_type_node), 0)); + OMP_CLAUSE_SIZE (clause) = size_zero_node; + OMP_CLAUSE_SIZE (nc) = size_zero_node; + OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_ALLOC); + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER); + OMP_CLAUSE_CHAIN (nc) = *list_p; + OMP_CLAUSE_CHAIN (clause) = nc; + struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; + gimplify_omp_ctxp = ctx->outer_context; + gimplify_expr (&TREE_OPERAND (OMP_CLAUSE_DECL (clause), 0), + pre_p, NULL, is_gimple_val, fb_rvalue); + gimplify_omp_ctxp = ctx; + } else if (code == OMP_CLAUSE_MAP) { OMP_CLAUSE_SET_MAP_KIND (clause, @@ -6915,7 +7006,8 @@ gimplify_adjust_omp_clauses (gimple_seq remove = true; else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST - && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER) + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER) { /* For GOMP_MAP_FORCE_DEVICEPTR, we'll never enter here, because for these, TREE_CODE (DECL_SIZE (decl)) will always be --- gcc/c/c-tree.h.jj 2015-07-01 12:50:49.000000000 +0200 +++ gcc/c/c-tree.h 2015-07-22 12:47:49.185826677 +0200 @@ -649,7 +649,7 @@ extern tree c_begin_omp_task (void); extern tree c_finish_omp_task (location_t, tree, tree); extern void c_finish_omp_cancel (location_t, tree); extern void c_finish_omp_cancellation_point (location_t, tree); -extern tree c_finish_omp_clauses (tree, bool = false); +extern tree c_finish_omp_clauses (tree, bool, bool = false); extern tree c_build_va_arg (location_t, tree, tree); extern tree c_finish_transaction (location_t, tree, int); extern bool c_tree_equal (tree, tree); --- gcc/c/c-typeck.c.jj 2015-07-17 13:06:58.000000000 +0200 +++ gcc/c/c-typeck.c 2015-07-22 13:00:21.130399057 +0200 @@ -11850,7 +11850,7 @@ handle_omp_array_sections_1 (tree c, tre /* Handle array sections for clause C. */ static bool -handle_omp_array_sections (tree c) +handle_omp_array_sections (tree c, bool is_omp) { bool maybe_zero_len = false; unsigned int first_non_one = 0; @@ -12031,8 +12031,10 @@ handle_omp_array_sections (tree c) return false; gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR); tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); - if (!c_mark_addressable (t)) + OMP_CLAUSE_SET_MAP_KIND (c2, is_omp + ? GOMP_MAP_FIRSTPRIVATE_POINTER + : GOMP_MAP_POINTER); + if (!is_omp && !c_mark_addressable (t)) return false; OMP_CLAUSE_DECL (c2) = t; t = build_fold_addr_expr (first); @@ -12097,7 +12099,7 @@ c_find_omp_placeholder_r (tree *tp, int Remove any elements from the list that are invalid. */ tree -c_finish_omp_clauses (tree clauses, bool declare_simd) +c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) { bitmap_head generic_head, firstprivate_head, lastprivate_head; bitmap_head aligned_head, map_head; @@ -12136,7 +12138,7 @@ c_finish_omp_clauses (tree clauses, bool t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, is_omp)) { remove = true; break; @@ -12496,7 +12498,7 @@ c_finish_omp_clauses (tree clauses, bool } if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, is_omp)) remove = true; break; } @@ -12519,7 +12521,7 @@ c_finish_omp_clauses (tree clauses, bool t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, is_omp)) remove = true; else { @@ -12556,6 +12558,8 @@ c_finish_omp_clauses (tree clauses, bool else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_POINTER) + || (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR))) && !lang_hooks.types.omp_mappable_type (TREE_TYPE (t))) { --- gcc/c/c-parser.c.jj 2015-07-21 09:06:42.000000000 +0200 +++ gcc/c/c-parser.c 2015-07-22 12:28:35.987814464 +0200 @@ -12435,7 +12435,7 @@ c_parser_oacc_all_clauses (c_parser *par c_parser_skip_to_pragma_eol (parser); if (finish_p) - return c_finish_omp_clauses (clauses); + return c_finish_omp_clauses (clauses, false); return clauses; } @@ -12720,8 +12720,8 @@ c_parser_omp_all_clauses (c_parser *pars if (finish_p) { if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0) - return c_finish_omp_clauses (clauses, true); - return c_finish_omp_clauses (clauses); + return c_finish_omp_clauses (clauses, true, true); + return c_finish_omp_clauses (clauses, true); } return clauses; @@ -12755,7 +12755,7 @@ c_parser_oacc_cache (location_t loc, c_p tree stmt, clauses; clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL); - clauses = c_finish_omp_clauses (clauses); + clauses = c_finish_omp_clauses (clauses, false); c_parser_skip_to_pragma_eol (parser); @@ -13902,7 +13902,7 @@ omp_split_clauses (location_t loc, enum c_omp_split_clauses (loc, code, mask, clauses, cclauses); for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++) if (cclauses[i]) - cclauses[i] = c_finish_omp_clauses (cclauses[i]); + cclauses[i] = c_finish_omp_clauses (cclauses[i], true); } /* OpenMP 4.0: @@ -14668,9 +14668,10 @@ c_parser_omp_target_data (location_t loc case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: map_seen = 3; break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), @@ -14800,9 +14801,10 @@ c_parser_omp_target_enter_data (location case GOMP_MAP_TO: case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: map_seen = 3; break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), @@ -14885,9 +14887,10 @@ c_parser_omp_target_exit_data (location_ case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: - case GOMP_MAP_POINTER: map_seen = 3; break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), @@ -15078,7 +15081,7 @@ check_clauses: case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: + case GOMP_MAP_FIRSTPRIVATE_POINTER: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), @@ -16379,7 +16382,7 @@ c_parser_cilk_for (c_parser *parser, tre tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE); OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR; OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain; - clauses = c_finish_omp_clauses (clauses); + clauses = c_finish_omp_clauses (clauses, false); tree block = c_begin_compound_stmt (true); tree sb = push_stmt_list (); @@ -16444,7 +16447,7 @@ c_parser_cilk_for (c_parser *parser, tre OMP_CLAUSE_OPERAND (c, 0) = cilk_for_number_of_iterations (omp_for); OMP_CLAUSE_CHAIN (c) = clauses; - OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c); + OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c, true); add_stmt (omp_par); } --- gcc/tree-core.h.jj 2015-07-17 09:30:44.000000000 +0200 +++ gcc/tree-core.h 2015-07-21 16:28:48.524156167 +0200 @@ -1354,7 +1354,7 @@ struct GTY(()) tree_omp_clause { enum omp_clause_schedule_kind schedule_kind; enum omp_clause_depend_kind depend_kind; /* See include/gomp-constants.h for enum gomp_map_kind's values. */ - unsigned char map_kind; + unsigned int map_kind; enum omp_clause_proc_bind_kind proc_bind_kind; enum tree_code reduction_code; enum omp_clause_linear_kind linear_kind; --- gcc/omp-low.c.jj 2015-07-21 09:07:23.000000000 +0200 +++ gcc/omp-low.c 2015-07-22 13:30:17.507589880 +0200 @@ -2025,6 +2025,21 @@ scan_sharing_clauses (tree clauses, omp_ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)) break; } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + { + if (DECL_SIZE (decl) + && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) + { + tree decl2 = DECL_VALUE_EXPR (decl); + gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); + decl2 = TREE_OPERAND (decl2, 0); + gcc_assert (DECL_P (decl2)); + install_var_local (decl2, ctx); + } + install_var_local (decl, ctx); + break; + } if (DECL_P (decl)) { if (DECL_SIZE (decl) @@ -2201,7 +2216,8 @@ scan_sharing_clauses (tree clauses, omp_ break; if (DECL_P (decl)) { - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER + if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE && !COMPLETE_TYPE_P (TREE_TYPE (decl))) { @@ -12883,6 +12899,7 @@ lower_omp_target (gimple_stmt_iterator * case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_FIRSTPRIVATE_POINTER: break; case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: @@ -12918,6 +12935,25 @@ lower_omp_target (gimple_stmt_iterator * var = var2; } + if (offloaded + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + { + if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) + { + tree type = build_pointer_type (TREE_TYPE (var)); + tree new_var = lookup_decl (var, ctx); + const char *name = NULL; + if (DECL_NAME (new_var)) + name = IDENTIFIER_POINTER (DECL_NAME (new_var)); + x = create_tmp_var_raw (type, name); + gimple_add_tmp_var (x); + x = build_simple_mem_ref (x); + SET_DECL_VALUE_EXPR (new_var, x); + DECL_HAS_VALUE_EXPR_P (new_var) = 1; + } + continue; + } + if (!maybe_lookup_field (var, ctx)) continue; @@ -12925,6 +12961,7 @@ lower_omp_target (gimple_stmt_iterator * { x = build_receiver_ref (var, true, ctx); tree new_var = lookup_decl (var, ctx); + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) && TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) @@ -13044,6 +13081,10 @@ lower_omp_target (gimple_stmt_iterator * } else { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_POINTER) + break; if (DECL_SIZE (ovar) && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST) { @@ -13239,6 +13280,7 @@ lower_omp_target (gimple_stmt_iterator * if (offloaded) { + tree prev = NULL_TREE; for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) switch (OMP_CLAUSE_CODE (c)) { @@ -13290,6 +13332,93 @@ lower_omp_target (gimple_stmt_iterator * } break; } + /* Handle GOMP_MAP_FIRSTPRIVATE_POINTER in second pass, + so that firstprivate vars holding OMP_CLAUSE_SIZE if needed + are already handled. */ + for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + tree var; + default: + break; + case OMP_CLAUSE_MAP: + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + { + location_t clause_loc = OMP_CLAUSE_LOCATION (c); + gcc_assert (prev); + var = OMP_CLAUSE_DECL (c); + if (DECL_SIZE (var) + && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST) + { + tree var2 = DECL_VALUE_EXPR (var); + gcc_assert (TREE_CODE (var2) == INDIRECT_REF); + var2 = TREE_OPERAND (var2, 0); + gcc_assert (DECL_P (var2)); + var = var2; + } + tree new_var = lookup_decl (var, ctx), x; + tree type = TREE_TYPE (new_var); + bool is_ref = is_reference (var); + bool ref_to_array = false; + if (is_ref) + { + type = TREE_TYPE (type); + if (TREE_CODE (type) == ARRAY_TYPE) + { + type = build_pointer_type (type); + ref_to_array = true; + } + } + else if (TREE_CODE (type) == ARRAY_TYPE) + { + tree decl2 = DECL_VALUE_EXPR (new_var); + gcc_assert (TREE_CODE (decl2) == MEM_REF); + decl2 = TREE_OPERAND (decl2, 0); + gcc_assert (DECL_P (decl2)); + new_var = decl2; + type = TREE_TYPE (new_var); + } + x = build_receiver_ref (OMP_CLAUSE_DECL (prev), false, ctx); + x = fold_convert_loc (clause_loc, type, x); + if (!integer_zerop (OMP_CLAUSE_SIZE (c))) + { + tree bias = OMP_CLAUSE_SIZE (c); + if (DECL_P (bias)) + bias = lookup_decl (bias, ctx); + bias = fold_convert_loc (clause_loc, sizetype, bias); + bias = fold_build1_loc (clause_loc, NEGATE_EXPR, sizetype, + bias); + x = fold_build2_loc (clause_loc, POINTER_PLUS_EXPR, + TREE_TYPE (x), x, bias); + } + if (ref_to_array) + x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x); + gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); + if (is_ref && !ref_to_array) + { + const char *name = NULL; + if (DECL_NAME (var)) + name = IDENTIFIER_POINTER (DECL_NAME (new_var)); + + tree t = create_tmp_var_raw (type, name); + gimple_add_tmp_var (t); + TREE_ADDRESSABLE (t) = 1; + gimple_seq_add_stmt (&new_body, + gimple_build_assign (t, x)); + x = build_fold_addr_expr_loc (clause_loc, t); + } + gimple_seq_add_stmt (&new_body, + gimple_build_assign (new_var, x)); + prev = NULL_TREE; + } + else if (OMP_CLAUSE_CHAIN (c) + && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) + == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_FIRSTPRIVATE_POINTER) + prev = c; + break; + } gimple_seq_add_seq (&new_body, tgt_body); new_body = maybe_catch_exception (new_body); } --- gcc/tree-pretty-print.c.jj 2015-07-21 09:06:42.000000000 +0200 +++ gcc/tree-pretty-print.c 2015-07-22 13:53:51.406065024 +0200 @@ -639,6 +639,9 @@ dump_omp_clause (pretty_printer *pp, tre case GOMP_MAP_RELEASE: pp_string (pp, "release"); break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + pp_string (pp, "firstprivate"); + break; default: gcc_unreachable (); } @@ -649,7 +652,9 @@ dump_omp_clause (pretty_printer *pp, tre if (OMP_CLAUSE_SIZE (clause)) { if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER) + && (OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER + || OMP_CLAUSE_MAP_KIND (clause) + == GOMP_MAP_FIRSTPRIVATE_POINTER)) pp_string (pp, " [pointer assign, bias: "); else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_TO_PSET) Jakub