omp-low: fix lastprivate/linear lowering for SIMT

Backport from mainline r247029
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.



git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gcc-7-branch@247032 138bc75d-0d04-0410-961f-82ee72b054a4

4 files changed, 53 insertions, 21 deletions

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index f00ff88b65c..627ee35faef 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,11 @@
+2017-04-20  Alexander Monakov  <amonakov@ispras.ru>
+
+	Backport from mainline
+	2017-04-20  Alexander Monakov  <amonakov@ispras.ru>
+
+	* omp-low.c (lower_lastprivate_clauses): Correct handling of linear
+	and lastprivate clauses in SIMT case.
+
 2017-04-20  Matthew Fortune  <matthew.fortune@imgtec.com>
 
 	Backport from mainline
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..806986654a9 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,10 @@
+2017-04-20  Alexander Monakov  <amonakov@ispras.ru>
+
+	Backport from mainline
+	2017-04-20  Alexander Monakov  <amonakov@ispras.ru>
+
+	* testsuite/libgomp.c/target-36.c: New testcase.
+
 2017-04-13  Jakub Jelinek  <jakub@redhat.com>
 
 	* 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 ();
+}