From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 38160 invoked by alias); 7 Apr 2017 07:58: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 37405 invoked by uid 89); 7 Apr 2017 07:58:00 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-24.5 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,RP_MATCHES_RCVD,SPF_PASS autolearn=ham version=3.3.2 spammy= X-HELO: smtp.ispras.ru Received: from bran.ispras.ru (HELO smtp.ispras.ru) (83.149.199.196) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 07 Apr 2017 07:57:58 +0000 Received: from monopod.intra.ispras.ru (monopod.intra.ispras.ru [10.10.3.121]) by smtp.ispras.ru (Postfix) with ESMTP id 7C30A611BA; Fri, 7 Apr 2017 10:57:57 +0300 (MSK) Date: Fri, 07 Apr 2017 07:58:00 -0000 From: Alexander Monakov To: Jakub Jelinek cc: gcc-patches@gcc.gnu.org Subject: [PATCH] omp-low: fix lastprivate/linear lowering for SIMT In-Reply-To: Message-ID: References: <1490197595-31938-1-git-send-email-amonakov@ispras.ru> <1490197595-31938-3-git-send-email-amonakov@ispras.ru> <20170323103159.GM11094@tucnak> User-Agent: Alpine 2.20.13 (LNX 116 2015-12-14) MIME-Version: 1.0 Content-Type: text/plain; charset=US-ASCII X-SW-Source: 2017-04/txt/msg00349.txt.bz2 Ping. > I've noticed while re-reading that this patch incorrectly handled SIMT case > in lower_lastprivate_clauses. The code was changed to look for variables > with "omp simt private" attribute, and was left under > 'simduid && DECL_HAS_VALUE_EXPR_P (new_var)' condition. This effectively > constrained processing to privatized (i.e. addressable) variables, as > non-addressable variables now have neither the value-expr nor the attribute. > > This wasn't caught in testing, as apparently all testcases that have target > simd loops with linear/lastprivate clauses also have the corresponding variables > mentioned in target map clause, which makes them addressable (is that necessary?), > and I didn't think to check something like that manually. > > The following patch fixes it by adjusting the if's in lower_lastprivate_clauses; > alternatively it may be possible to keep that code as-is, and instead set up > value-expr and "omp simt private" attributes for all formally private variables, > not just addressable ones, but to me that sounds less preferable. OK for trunk? > > I think it's possible to improve test coverage for NVPTX by running all OpenMP > testcases via nvptx-none-run (it'll need some changes, but shouldn't be hard). > > gcc/ > * omp-low.c (lower_lastprivate_clauses): Correct handling of linear and > lastprivate clauses in SIMT case. > > libgomp/ > * testsuite/libgomp.c/target-36.c: New testcase. > > Thanks. > Alexander > > diff --git a/gcc/omp-low.c b/gcc/omp-low.c > index 253dc85..02b681c 100644 > --- a/gcc/omp-low.c > +++ b/gcc/omp-low.c > @@ -4768,11 +4768,10 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list, > TREE_NO_WARNING (new_var) = 1; > } > > - if (simduid && DECL_HAS_VALUE_EXPR_P (new_var)) > + if (!maybe_simt && simduid && DECL_HAS_VALUE_EXPR_P (new_var)) > { > tree val = DECL_VALUE_EXPR (new_var); > - if (!maybe_simt > - && TREE_CODE (val) == ARRAY_REF > + if (TREE_CODE (val) == ARRAY_REF > && VAR_P (TREE_OPERAND (val, 0)) > && lookup_attribute ("omp simd array", > DECL_ATTRIBUTES (TREE_OPERAND (val, > @@ -4792,26 +4791,26 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list, > TREE_OPERAND (val, 0), lastlane, > NULL_TREE, NULL_TREE); > } > - else if (maybe_simt > - && VAR_P (val) > - && lookup_attribute ("omp simt private", > - DECL_ATTRIBUTES (val))) > + } > + else if (maybe_simt) > + { > + tree val = (DECL_HAS_VALUE_EXPR_P (new_var) > + ? DECL_VALUE_EXPR (new_var) > + : new_var); > + if (simtlast == NULL) > { > - if (simtlast == NULL) > - { > - simtlast = create_tmp_var (unsigned_type_node); > - gcall *g = gimple_build_call_internal > - (IFN_GOMP_SIMT_LAST_LANE, 1, simtcond); > - gimple_call_set_lhs (g, simtlast); > - gimple_seq_add_stmt (stmt_list, g); > - } > - x = build_call_expr_internal_loc > - (UNKNOWN_LOCATION, IFN_GOMP_SIMT_XCHG_IDX, > - TREE_TYPE (val), 2, val, simtlast); > - new_var = unshare_expr (new_var); > - gimplify_assign (new_var, x, stmt_list); > - new_var = unshare_expr (new_var); > + simtlast = create_tmp_var (unsigned_type_node); > + gcall *g = gimple_build_call_internal > + (IFN_GOMP_SIMT_LAST_LANE, 1, simtcond); > + gimple_call_set_lhs (g, simtlast); > + gimple_seq_add_stmt (stmt_list, g); > } > + x = build_call_expr_internal_loc > + (UNKNOWN_LOCATION, IFN_GOMP_SIMT_XCHG_IDX, > + TREE_TYPE (val), 2, val, simtlast); > + new_var = unshare_expr (new_var); > + gimplify_assign (new_var, x, stmt_list); > + new_var = unshare_expr (new_var); > } > > if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE > diff --git a/libgomp/testsuite/libgomp.c/target-36.c b/libgomp/testsuite/libgomp.c/target-36.c > new file mode 100644 > index 0000000..6925a2a > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c/target-36.c > @@ -0,0 +1,18 @@ > +int > +main () > +{ > + int ah, bh, n = 1024; > +#pragma omp target map(from: ah, bh) > + { > + int a, b; > +#pragma omp simd lastprivate(b) > + for (a = 0; a < n; a++) > + { > + b = a + n + 1; > + asm volatile ("" : "+r"(b)); > + } > + ah = a, bh = b; > + } > + if (ah != n || bh != 2 * n) > + __builtin_abort (); > +} >