From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 36161 invoked by alias); 31 Mar 2017 15:38:13 -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 36145 invoked by uid 89); 31 Mar 2017 15:38:12 -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=formally 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, 31 Mar 2017 15:38:10 +0000 Received: from monopod.intra.ispras.ru (monopod.intra.ispras.ru [10.10.3.121]) by smtp.ispras.ru (Postfix) with ESMTP id 49B6F60912; Fri, 31 Mar 2017 18:38:09 +0300 (MSK) Date: Fri, 31 Mar 2017 16:05:00 -0000 From: Alexander Monakov To: Jakub Jelinek cc: gcc-patches@gcc.gnu.org Subject: Re: [PATCH 2/5] omp-low: implement SIMT privatization, part 1 In-Reply-To: <20170323103159.GM11094@tucnak> 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-03/txt/msg01579.txt.bz2 Hello Jakub, 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 (); +}