[PATCH 2/5] omp-low: implement SIMT privatization, part 1
Alexander Monakov
amonakov@ispras.ru
Fri Mar 31 16:05:00 GMT 2017
Hello Jakub,
I've noticed while re-reading that this patch incorrectly handled SIMT case
in lower_lastprivate_clauses. The code was changed to look for variables
with "omp simt private" attribute, and was left under
'simduid && DECL_HAS_VALUE_EXPR_P (new_var)' condition. This effectively
constrained processing to privatized (i.e. addressable) variables, as
non-addressable variables now have neither the value-expr nor the attribute.
This wasn't caught in testing, as apparently all testcases that have target
simd loops with linear/lastprivate clauses also have the corresponding variables
mentioned in target map clause, which makes them addressable (is that necessary?),
and I didn't think to check something like that manually.
The following patch fixes it by adjusting the if's in lower_lastprivate_clauses;
alternatively it may be possible to keep that code as-is, and instead set up
value-expr and "omp simt private" attributes for all formally private variables,
not just addressable ones, but to me that sounds less preferable. OK for trunk?
I think it's possible to improve test coverage for NVPTX by running all OpenMP
testcases via nvptx-none-run (it'll need some changes, but shouldn't be hard).
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.
Thanks.
Alexander
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 253dc85..02b681c 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -4768,11 +4768,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,
@@ -4792,26 +4791,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/testsuite/libgomp.c/target-36.c b/libgomp/testsuite/libgomp.c/target-36.c
new file mode 100644
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 ();
+}
More information about the Gcc-patches
mailing list