From patchwork Fri Sep 6 14:57:31 2013 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Jakub Jelinek X-Patchwork-Id: 273215 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 with cipher DHE-RSA-AES256-SHA (256/256 bits)) (Client CN "www.sourceware.org", Issuer "StartCom Class 1 Primary Intermediate Server CA" (not verified)) by ozlabs.org (Postfix) with ESMTPS id B7F752C0084 for ; Sat, 7 Sep 2013 00:57:44 +1000 (EST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:reply-to:mime-version :content-type; q=dns; s=default; b=ZCMlrjGxapqvJcfoPx0yo+RTpJIKe Sy8KAMLp5zuSRU6VDgx/j7Z29UQugi32r9rOPQ9jADYIKW/35skiMjQhWxrSLqnS dnfPa6Bs+dQd/JhVloUMdGKaKRw39EHnXH1q/mOcTUJl3rGP5Ek1SVIIykRanN8Y myJg1Se6eP5KsA= 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:date :from:to:cc:subject:message-id:reply-to:mime-version :content-type; s=default; bh=qv3TqvWrgV8V7pmp7Inb8WNIxf4=; b=RK1 pZcKjsffZW5QmQvbAgGyZhbKLKUA2vhTPHYV8MuhzxgBf4LqkdIEviwaDw/NWyWw oO/CzxHj/yEGTvB5uURvx/m5QOp6wBF6OIAV2iAmsRbjY20oKnsKsWs6U6c9rUJb BVVxYhAYnWUWo1r1HCyBRr/Gji1KZCeW6g92Tje4= Received: (qmail 24069 invoked by alias); 6 Sep 2013 14:57:37 -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 24059 invoked by uid 89); 6 Sep 2013 14:57:37 -0000 Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 06 Sep 2013 14:57:37 +0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-4.3 required=5.0 tests=AWL, BAYES_00, RP_MATCHES_RCVD autolearn=ham version=3.3.2 X-HELO: mx1.redhat.com Received: from int-mx09.intmail.prod.int.phx2.redhat.com (int-mx09.intmail.prod.int.phx2.redhat.com [10.5.11.22]) by mx1.redhat.com (8.14.4/8.14.4) with ESMTP id r86EvXra019319 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-SHA bits=256 verify=OK); Fri, 6 Sep 2013 10:57:34 -0400 Received: from tucnak.zalov.cz (vpn1-7-233.ams2.redhat.com [10.36.7.233]) by int-mx09.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id r86EvWCV031093 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-SHA bits=256 verify=NO); Fri, 6 Sep 2013 10:57:33 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.7/8.14.7) with ESMTP id r86EvVHA030290; Fri, 6 Sep 2013 16:57:31 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.7/8.14.7/Submit) id r86EvVZ5030289; Fri, 6 Sep 2013 16:57:31 +0200 Date: Fri, 6 Sep 2013 16:57:31 +0200 From: Jakub Jelinek To: Richard Henderson , "Michael V. Zolotukhin" Cc: gcc-patches@gcc.gnu.org Subject: [gomp4] Further accel fixes Message-ID: <20130906145731.GS23437@tucnak.redhat.com> Reply-To: Jakub Jelinek MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.5.21 (2010-09-15) X-IsSubscribed: yes Hi! This fixes mainly VLA handling in target{, data, update} constructs, but also deals with field alignments in the structure and field order. Committed to gomp-4_0-branch. 2013-09-06 Jakub Jelinek * omp-low.c (scan_sharing_clauses): Handle VLAs in OMP_CLAUSE_{MAP,TO,FROM}. Set DECL_ALIGN (field) before calling insert_field_into_struct. (scan_omp_target): Reverse TYPE_FIELDS, verify that all field alignments are the same. (lower_omp_target): Use maybe_lookup_field instead of lookup_sfield to check if field is present. Handle VLAs. * tree-pretty-print.c (dump_omp_clause): Only check OMP_CLAUSE_MAP_KIND on OMP_CLAUSE_MAP clauses. * gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_TO_ONLY. (omp_firstprivatize_variable, omp_add_variable, gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Handle VLAs in OMP_CLAUSE_{MAP,TO,FROM}. libgomp/ * testsuite/libgomp.c/target-2.c: New test. * testsuite/libgomp.c++/target-3.C: New test. Jakub --- gcc/omp-low.c.jj 2013-09-05 17:11:14.000000000 +0200 +++ gcc/omp-low.c 2013-09-06 16:15:16.367638718 +0200 @@ -1574,10 +1574,24 @@ scan_sharing_clauses (tree clauses, omp_ } if (DECL_P (decl)) { - install_var_field (decl, true, 3, ctx); - if (gimple_omp_target_kind (ctx->stmt) - == GF_OMP_TARGET_KIND_REGION) - install_var_local (decl, ctx); + 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_field (decl2, true, 3, ctx); + install_var_local (decl2, ctx); + install_var_local (decl, ctx); + } + else + { + install_var_field (decl, true, 3, ctx); + if (gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_REGION) + install_var_local (decl, ctx); + } } else { @@ -1600,6 +1614,7 @@ scan_sharing_clauses (tree clauses, omp_ tree field = build_decl (OMP_CLAUSE_LOCATION (c), FIELD_DECL, NULL_TREE, ptr_type_node); + DECL_ALIGN (field) = TYPE_ALIGN (ptr_type_node); insert_field_into_struct (ctx->record_type, field); splay_tree_insert (ctx->field_map, (splay_tree_key) decl, (splay_tree_value) field); @@ -1684,6 +1699,16 @@ scan_sharing_clauses (tree clauses, omp_ TREE_TYPE (new_decl) = remap_type (TREE_TYPE (decl), &ctx->cb); } + else 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)); + fixup_remapped_decl (decl2, ctx, false); + fixup_remapped_decl (decl, ctx, true); + } else fixup_remapped_decl (decl, ctx, false); } @@ -2126,6 +2151,16 @@ scan_omp_target (gimple stmt, omp_contex ctx->record_type = ctx->receiver_decl = NULL; else { + TYPE_FIELDS (ctx->record_type) + = nreverse (TYPE_FIELDS (ctx->record_type)); +#ifdef ENABLE_CHECKING + tree field; + unsigned int align = DECL_ALIGN (TYPE_FIELDS (ctx->record_type)); + for (field = TYPE_FIELDS (ctx->record_type); + field; + field = DECL_CHAIN (field)) + gcc_assert (DECL_ALIGN (field) == align); +#endif layout_type (ctx->record_type); if (kind == GF_OMP_TARGET_KIND_REGION) fixup_child_record_type (ctx); @@ -9201,7 +9236,18 @@ lower_omp_target (gimple_stmt_iterator * map_cnt++; continue; } - if (!lookup_sfield (var, ctx)) + + 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; + } + + if (!maybe_lookup_field (var, ctx)) continue; if (kind == GF_OMP_TARGET_KIND_REGION) @@ -9293,8 +9339,20 @@ lower_omp_target (gimple_stmt_iterator * nc = NULL_TREE; } } - else if (!lookup_sfield (ovar, ctx)) - continue; + else + { + if (DECL_SIZE (ovar) + && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST) + { + tree ovar2 = DECL_VALUE_EXPR (ovar); + gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF); + ovar2 = TREE_OPERAND (ovar2, 0); + gcc_assert (DECL_P (ovar2)); + ovar = ovar2; + } + if (!maybe_lookup_field (ovar, ctx)) + continue; + } if (nc) { --- gcc/tree-pretty-print.c.jj 2013-08-27 22:18:05.000000000 +0200 +++ gcc/tree-pretty-print.c 2013-09-06 15:36:48.481578926 +0200 @@ -508,7 +508,8 @@ dump_omp_clause (pretty_printer *buffer, print_clause_size: if (OMP_CLAUSE_SIZE (clause)) { - if (OMP_CLAUSE_MAP_KIND (clause) == OMP_CLAUSE_MAP_POINTER) + if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (clause) == OMP_CLAUSE_MAP_POINTER) pp_string (buffer, " [pointer assign, bias: "); else pp_string (buffer, " [len: "); --- gcc/gimplify.c.jj 2013-09-05 14:45:48.000000000 +0200 +++ gcc/gimplify.c 2013-09-06 14:47:13.153081283 +0200 @@ -61,6 +61,7 @@ enum gimplify_omp_var_data GOVD_PRIVATE_OUTER_REF = 1024, GOVD_LINEAR = 2048, GOVD_ALIGNED = 4096, + GOVD_MAP_TO_ONLY = 8192, GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -5740,11 +5741,16 @@ omp_firstprivatize_variable (struct gimp { if (n->value & GOVD_SHARED) n->value = GOVD_FIRSTPRIVATE | (n->value & GOVD_SEEN); + else if (n->value & GOVD_MAP) + n->value |= GOVD_MAP_TO_ONLY; else return; } + else if (ctx->region_type == ORT_TARGET) + omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY); else if (ctx->region_type != ORT_WORKSHARE - && ctx->region_type != ORT_SIMD) + && ctx->region_type != ORT_SIMD + && ctx->region_type != ORT_TARGET_DATA) omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE); ctx = ctx->outer_context; @@ -5847,16 +5853,15 @@ omp_add_variable (struct gimplify_omp_ct the parameters of the type. */ if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { - /* To be handled later. */ - gcc_assert ((flags & GOVD_MAP) == 0); - /* Add the pointer replacement variable as PRIVATE if the variable replacement is private, else FIRSTPRIVATE since we'll need the address of the original variable either for SHARED, or for the copy into or out of the context. */ if (!(flags & GOVD_LOCAL)) { - nflags = flags & GOVD_PRIVATE ? GOVD_PRIVATE : GOVD_FIRSTPRIVATE; + nflags = flags & GOVD_MAP + ? GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT + : flags & GOVD_PRIVATE ? GOVD_PRIVATE : GOVD_FIRSTPRIVATE; nflags |= flags & GOVD_SEEN; t = DECL_VALUE_EXPR (decl); gcc_assert (TREE_CODE (t) == INDIRECT_REF); @@ -5885,15 +5890,13 @@ omp_add_variable (struct gimplify_omp_ct For local variables TYPE_SIZE_UNIT might not be gimplified yet, in this case omp_notice_variable will be called later on when it is gimplified. */ - else if (! (flags & GOVD_LOCAL) + else if (! (flags & (GOVD_LOCAL | GOVD_MAP)) && DECL_P (TYPE_SIZE_UNIT (TREE_TYPE (decl)))) omp_notice_variable (ctx, TYPE_SIZE_UNIT (TREE_TYPE (decl)), true); } - else if (lang_hooks.decls.omp_privatize_by_reference (decl)) + else if ((flags & GOVD_MAP) == 0 + && lang_hooks.decls.omp_privatize_by_reference (decl)) { - /* To be handled later. */ - gcc_assert ((flags & GOVD_MAP) == 0); - gcc_assert ((flags & GOVD_LOCAL) == 0); omp_firstprivatize_type_sizes (ctx, TREE_TYPE (decl)); @@ -6562,7 +6565,35 @@ gimplify_adjust_omp_clauses_1 (splay_tre else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF)) OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1; else if (code == OMP_CLAUSE_MAP) - OMP_CLAUSE_MAP_KIND (clause) = OMP_CLAUSE_MAP_TOFROM; + { + OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY + ? OMP_CLAUSE_MAP_TO + : OMP_CLAUSE_MAP_TOFROM; + 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)); + tree mem = build_simple_mem_ref (decl2); + OMP_CLAUSE_DECL (clause) = mem; + OMP_CLAUSE_SIZE (clause) = TYPE_SIZE_UNIT (TREE_TYPE (decl)); + if (gimplify_omp_ctxp->outer_context) + { + struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp->outer_context; + omp_notice_variable (ctx, decl2, true); + omp_notice_variable (ctx, OMP_CLAUSE_SIZE (clause), true); + } + tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause), + OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (nc) = decl; + OMP_CLAUSE_SIZE (nc) = size_zero_node; + OMP_CLAUSE_MAP_KIND (nc) = OMP_CLAUSE_MAP_POINTER; + OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause); + OMP_CLAUSE_CHAIN (clause) = nc; + } + } *list_p = clause; lang_hooks.decls.omp_finish_clause (clause); @@ -6687,6 +6718,56 @@ gimplify_adjust_omp_clauses (tree *list_ n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN)) remove = true; + else if (DECL_SIZE (decl) + && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST + && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_POINTER) + { + tree decl2 = DECL_VALUE_EXPR (decl); + gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); + decl2 = TREE_OPERAND (decl2, 0); + gcc_assert (DECL_P (decl2)); + tree mem = build_simple_mem_ref (decl2); + OMP_CLAUSE_DECL (c) = mem; + OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (decl)); + if (ctx->outer_context) + { + omp_notice_variable (ctx->outer_context, decl2, true); + omp_notice_variable (ctx->outer_context, + OMP_CLAUSE_SIZE (c), true); + } + tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (nc) = decl; + OMP_CLAUSE_SIZE (nc) = size_zero_node; + OMP_CLAUSE_MAP_KIND (nc) = OMP_CLAUSE_MAP_POINTER; + OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = nc; + c = nc; + } + break; + + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: + decl = OMP_CLAUSE_DECL (c); + if (!DECL_P (decl)) + break; + 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)); + tree mem = build_simple_mem_ref (decl2); + OMP_CLAUSE_DECL (c) = mem; + OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (decl)); + if (ctx->outer_context) + { + omp_notice_variable (ctx->outer_context, decl2, true); + omp_notice_variable (ctx->outer_context, + OMP_CLAUSE_SIZE (c), true); + } + } break; case OMP_CLAUSE_REDUCTION: @@ -6708,8 +6789,6 @@ gimplify_adjust_omp_clauses (tree *list_ case OMP_CLAUSE_MERGEABLE: case OMP_CLAUSE_PROC_BIND: case OMP_CLAUSE_SAFELEN: - case OMP_CLAUSE_TO: - case OMP_CLAUSE_FROM: case OMP_CLAUSE_DEPEND: break; --- libgomp/testsuite/libgomp.c/target-2.c.jj 2013-09-06 16:24:29.213769868 +0200 +++ libgomp/testsuite/libgomp.c/target-2.c 2013-09-06 16:24:05.000000000 +0200 @@ -0,0 +1,88 @@ +extern +#ifdef __cplusplus +"C" +#endif +void abort (void); + +void +fn1 (double *x, double *y, int z) +{ + int i; + for (i = 0; i < z; i++) + { + x[i] = i & 31; + y[i] = (i & 63) - 30; + } +} + +double +fn2 (int x) +{ + double s = 0; + double b[3 * x], c[3 * x], d[3 * x], e[3 * x]; + int i; + fn1 (b, c, x); + fn1 (e, d + x, x); + #pragma omp target map(to: b, c[:x], d[x:x], e) + #pragma omp parallel for reduction(+:s) + for (i = 0; i < x; i++) + s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c); + return s; +} + +double +fn3 (int x) +{ + double s = 0; + double b[3 * x], c[3 * x], d[3 * x], e[3 * x]; + int i; + fn1 (b, c, x); + fn1 (e, d, x); + #pragma omp target + #pragma omp parallel for reduction(+:s) + for (i = 0; i < x; i++) + s += b[i] * c[i] + d[i]; + return s; +} + +double +fn4 (int x) +{ + double s = 0; + double b[3 * x], c[3 * x], d[3 * x], e[3 * x]; + int i; + fn1 (b, c, x); + fn1 (e, d + x, x); + #pragma omp target data map(from: b, c[:x], d[x:x], e) + { + #pragma omp target update to(b, c[:x], d[x:x], e) + #pragma omp target map(c[:x], d[x:x]) + #pragma omp parallel for reduction(+:s) + for (i = 0; i < x; i++) + { + s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c); + b[i] = i + 0.5; + c[i] = 0.5 - i; + d[x + i] = 0.5 * i; + } + } + for (i = 0; i < x; i++) + if (b[i] != i + 0.5 || c[i] != 0.5 - i || d[x + i] != 0.5 * i) + abort (); + return s; +} + +int +main () +{ + double a = fn2 (128); + if (a != 14080.0) + abort (); + double b = fn3 (128); + if (a != b) + abort (); + double c = fn4 (256); + if (c != 28160.0) + abort (); + return 0; +} --- libgomp/testsuite/libgomp.c++/target-3.C.jj 2013-09-06 16:24:40.674710426 +0200 +++ libgomp/testsuite/libgomp.c++/target-3.C 2013-09-06 16:24:44.594690089 +0200 @@ -0,0 +1 @@ +#include "../libgomp.c/target-2.c"