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
This commit is contained in:
parent
c3cbcd4517
commit
2260d19d4b
|
@ -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.
|
||||
|
|
|
@ -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
|
||||
|
|
|
@ -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
|
||||
|
|
|
@ -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 ();
|
||||
}
|
Loading…
Reference in New Issue