OpenE2K/gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

omp-low: fix lastprivate/linear lowering for SIMT

Browse Source 
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.

From-SVN: r247032
This commit is contained in:
Alexander Monakov 2017-04-20 20:59:25 +03:00 committed by Alexander Monakov
parent 120fbfd0ee
commit 70cae6302a
4 changed files with 53 additions and 21 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
gcc/ChangeLog
View File

@ -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

41
gcc/omp-low.c
View File

@ -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

7
libgomp/ChangeLog
View File

@ -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

18
libgomp/testsuite/libgomp.c/target-36.c Normal file
View File

@ -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 ();
}