On 01/20/2016 07:46 PM, Cesar Philippidis wrote:
> I've applied this patch to gomp-4_0-branch which fixes of problems
> involving reference type variables in openacc data clauses. The first
> problem was, the c++ front end was incorrectly handling reference types
> in general in openacc. Instead of mapping the variable, it would map the
> pointer to the variable by itself. The second problem was, if the
> gimplifier saw a pointer mapping for a data clause, it would propagate
> it to omp-lower. That's bad because if you have something like this
>
> int &var = ...
>
> #pragma acc data copy (var)
> {
> ...var...
> }
>
> where the var inside the data region would have some uninitialized value
> because omplower installs a new variable for it. The gimpifier is
> already handling openmp target data regions properly, so this patch
> extends it to ignore pointer mappings in acc enter/exit and data constructs.
>
> Ultimately this patch will need to go in trunk, but the c++ changes
> don't apply cleanly. I'll need to work on that later.
And here's the patch.
Cesar
2016-01-20 Cesar Philippidis <ce...@codesourcery.com>
gcc/cp/
* parser.c (cp_parser_oacc_all_clauses): Call finish_omp_clauses
with allow_fields set to true.
(cp_parser_oacc_cache): Likewise.
(cp_parser_oacc_loop): Likewise.
* semantics.c (finish_omp_clauses): Ensure that is_oacc is properly
set when calling hanlde_omp_array_sections.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Consider OACC_{DATA,
PARALLEL, KERNELS} when processing firstprivate pointers and
references, and setting target_kind_firstprivatize_array_bases.
libgomp/
* testsuite/libgomp.oacc-c++/non-scalar-data.C: New test.
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d88877a..4882b19 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -32324,7 +32324,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
if (finish_p)
- return finish_omp_clauses (clauses, true, false);
+ return finish_omp_clauses (clauses, true, true);
return clauses;
}
@@ -35140,7 +35140,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
tree stmt, clauses;
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE);
- clauses = finish_omp_clauses (clauses, true, false);
+ clauses = finish_omp_clauses (clauses, true, true);
cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer));
@@ -35471,9 +35471,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
{
clauses = c_oacc_split_loop_clauses (clauses, cclauses);
if (*cclauses)
- finish_omp_clauses (*cclauses, true, false);
+ finish_omp_clauses (*cclauses, true, true);
if (clauses)
- finish_omp_clauses (clauses, true, false);
+ finish_omp_clauses (clauses, true, true);
}
tree block = begin_omp_structured_block ();
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 3ca6137..e161186 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5807,7 +5807,7 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c, allow_fields))
+ if (handle_omp_array_sections (c, allow_fields && !is_oacc))
{
remove = true;
break;
@@ -6567,7 +6567,7 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
}
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c, allow_fields))
+ if (handle_omp_array_sections (c, allow_fields && !is_oacc))
remove = true;
break;
}
@@ -6601,7 +6601,7 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c, allow_fields))
+ if (handle_omp_array_sections (c, allow_fields && !is_oacc))
remove = true;
else
{
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index cdb5b96..152942f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6092,7 +6092,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
{
unsigned nflags = flags;
if (ctx->target_map_pointers_as_0len_arrays
- || ctx->target_map_scalars_firstprivate)
+ || ctx->target_map_scalars_firstprivate)
{
bool is_declare_target = false;
bool is_scalar = false;
@@ -6456,7 +6456,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
+ case OACC_DATA:
case OACC_HOST_DATA:
+ case OACC_PARALLEL:
+ case OACC_KERNELS:
ctx->target_firstprivatize_array_bases = true;
default:
break;
@@ -6726,7 +6729,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
+ case OACC_DATA:
case OACC_HOST_DATA:
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_REFERENCE))
diff --git a/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
new file mode 100644
index 0000000..180e86f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
@@ -0,0 +1,109 @@
+// Ensure that a non-scalar dummy arguments which are implicitly used inside
+// offloaded regions are properly mapped using present_or_copy.
+
+// { dg-do run }
+
+#include <cassert>
+
+const int n = 100;
+
+struct data {
+ int v;
+};
+
+void
+kernels_present (data &d, int &x)
+{
+#pragma acc kernels present (d, x) default (none)
+ {
+ d.v = x;
+ }
+}
+
+void
+parallel_present (data &d, int &x)
+{
+#pragma acc parallel present (d, x) default (none)
+ {
+ d.v = x;
+ }
+}
+
+void
+kernels_implicit (data &d, int &x)
+{
+#pragma acc kernels
+ {
+ d.v = x;
+ }
+}
+
+void
+parallel_implicit (data &d, int &x)
+{
+#pragma acc parallel
+ {
+ d.v = x;
+ }
+}
+
+void
+reference_data (data &d, int &x)
+{
+#pragma acc data copy(d, x)
+ {
+ kernels_present (d, x);
+
+#pragma acc update host(d)
+ assert (d.v == x);
+
+ x = 200;
+#pragma acc update device(x)
+
+ parallel_present (d, x);
+ }
+
+ assert (d.v = x);
+
+ x = 300;
+ kernels_implicit (d, x);
+ assert (d.v = x);
+
+ x = 400;
+ parallel_implicit (d, x);
+ assert (d.v = x);
+}
+
+int
+main ()
+{
+ data d;
+ int x = 100;
+
+#pragma acc data copy(d, x)
+ {
+ kernels_present (d, x);
+
+#pragma acc update host(d)
+ assert (d.v == x);
+
+ x = 200;
+#pragma acc update device(x)
+
+ parallel_present (d, x);
+ }
+
+ assert (d.v = x);
+
+ x = 300;
+ kernels_implicit (d, x);
+ assert (d.v = x);
+
+ x = 400;
+ parallel_implicit (d, x);
+ assert (d.v = x);
+
+ reference_data (d, x);
+
+ return 0;
+}
