From 2260d19d4be6218c0dfffe6f87fc7665a09d5ba5 Mon Sep 17 00:00:00 2001 From: Alexander Monakov Date: Thu, 20 Apr 2017 20:21:50 +0300 Subject: [PATCH] omp-low: fix lastprivate/linear lowering for SIMT 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. From-SVN: r247029 --- gcc/ChangeLog | 5 +++ gcc/omp-low.c | 41 ++++++++++++------------- libgomp/ChangeLog | 4 +++ libgomp/testsuite/libgomp.c/target-36.c | 18 +++++++++++ 4 files changed, 47 insertions(+), 21 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c/target-36.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index e1a889daea3..d96d987040f 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2017-04-20 Alexander Monakov + + * omp-low.c (lower_lastprivate_clauses): Correct handling of linear and + lastprivate clauses in SIMT case. + 2017-04-20 Volker Reichelt * doc/invoke.texi (-Wextra-semi): Document new warning option. diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 22772ba2891..9cc29964dfa 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -4770,11 +4770,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, @@ -4794,26 +4793,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/ChangeLog b/libgomp/ChangeLog index 5ab1b2a9066..633ba37a62b 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,7 @@ +2017-04-20 Alexander Monakov + + * testsuite/libgomp.c/target-36.c: New testcase. + 2017-04-13 Jakub Jelinek * plugin/plugin-nvptx.c (cuda_lib_inited): Use signed char type diff --git a/libgomp/testsuite/libgomp.c/target-36.c b/libgomp/testsuite/libgomp.c/target-36.c new file mode 100644 index 00000000000..6925a2a63de --- /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 (); +} -- 2.30.2