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

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index e1a889daea35852e7c4b0a3c3732cbfb077ef45c..d96d987040fca9e0dbf92d9d8f8da63e60936b92 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,8 @@
+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  Volker Reichelt  <v.reichelt@netcologne.de>
 
        * doc/invoke.texi (-Wextra-semi): Document new warning option.

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 22772ba28918eeeab039c8b7db0562334e1de7d0..9cc29964dfa03c7f2eb4ba4043d11bc2216f1b15 100644 (file)
--- 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 5ab1b2a906652e14de0e6114fe2b05aa1710f35f..633ba37a62b3f0bcafd5c52fc3c938b4f976a932 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,7 @@
+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 (file)
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 ();
+}