In OpenACC, if an offloaded region is lexically nested inside an acc
data region, then those variables should technically be marked as
present. For the most part we can get away with a simpler analysis and
making those variables present_or_copy, as that is the default for array
variables. However, that scheme falls apart when the data region
contains a data clause for a subarray of a local array. The problem
there is, the data region transfers the subarray to the accelerator, but
the offloaded region expects the entire local array to be present on the
device. And since only part of the array resides on the device, the
runtime ultimately generates an error for unmapped data.
There are two separate cases that are addressed in this patch. 1) The
creation of present and firstprivate_pointer data clauses for C and C++
subarrays and 2) the creation of a pointer data clause for fortran
subarrays. I suspect that fortran will eventually utilize firstprivate
pointers for subarrays too, but then this patch would stop working
because the fortran FE uses an separate internal pointer for the array
data, that doesn't correspond to the actual array decl like it does it
does in C/C++.
In the meantime, I'll apply this WIP patch to gomp-4_0-branch.
Cesar
2016-08-17 Cesar Philippidis <ce...@codesourcery.com>
PR middle-end/70828
gcc/
* gimplify.c (struct gimplify_omp_ctx): Add tree clauses member.
(new_omp_context): Initialize clauses to NULL_TREE.
(gimplify_scan_omp_clauses): Set clauses in the gimplify_omp_ctx.
(omp_clause_matching_array_ref): New function.
(gomp_needs_data_present): New function.
(gimplify_adjust_omp_clauses_1): Use preset or pointer omp clause map
kinds when creating implicit data clauses for OpenACC offloaded
variables defined used an acc data region as necessary.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test.
* testsuite/libgomp.oacc-fortran/pr70828.f90: New test.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 717b25f..9efb907 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -180,6 +180,7 @@ struct gimplify_omp_ctx
bool target_map_scalars_firstprivate;
bool target_map_pointers_as_0len_arrays;
bool target_firstprivatize_array_bases;
+ tree clauses;
};
static struct gimplify_ctx *gimplify_ctxp;
@@ -396,6 +397,7 @@ new_omp_context (enum omp_region_type region_type)
c->privatized_types = new hash_set<tree>;
c->location = input_location;
c->region_type = region_type;
+ c->clauses = NULL_TREE;
if ((region_type & ORT_TASK) == 0)
c->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
else
@@ -6546,6 +6548,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
tree *prev_list_p = NULL;
ctx = new_omp_context (region_type);
+ ctx->clauses = *list_p;
outer_ctx = ctx->outer_context;
if (code == OMP_TARGET && !lang_GNU_Fortran ())
{
@@ -7694,6 +7697,58 @@ struct gimplify_adjust_omp_clauses_data
gimple_seq *pre_p;
};
+/* Return true if clause contains an array_ref of DECL. */
+
+static bool
+omp_clause_matching_array_ref (tree clause, tree decl)
+{
+ tree cdecl = OMP_CLAUSE_DECL (clause);
+
+ if (TREE_CODE (cdecl) != ARRAY_REF)
+ return false;
+
+ return TREE_OPERAND (cdecl, 0) == decl;
+}
+
+/* Inside OpenACC parallel and kernels regions, the implicit data
+ clauses for arrays must respect the explicit data clauses set by a
+ containing acc data region. Specifically, care must be taken
+ pointers or if an subarray of a local array is specified in an acc
+ data region, so that the referenced array inside the offloaded
+ region has a present data clasue for that array with an
+ approporiate subarray argument. This function returns the tree
+ node of the acc data clause that utilizes DECL as an argument. */
+
+static tree
+gomp_needs_data_present (tree decl)
+{
+ gimplify_omp_ctx *ctx = NULL;
+ bool found_match = false;
+ tree c = NULL_TREE;
+
+ if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
+ return NULL_TREE;
+
+ if (gimplify_omp_ctxp->region_type != ORT_ACC_PARALLEL
+ && gimplify_omp_ctxp->region_type != ORT_ACC_KERNELS)
+ return NULL_TREE;
+
+ for (ctx = gimplify_omp_ctxp->outer_context; !found_match && ctx;
+ ctx = ctx->outer_context)
+ {
+ if (ctx->region_type != ORT_ACC_DATA)
+ break;
+
+ for (c = ctx->clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (omp_clause_matching_array_ref (c, decl)
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER))
+ return c;
+ }
+
+ return NULL_TREE;
+}
+
/* For all variables that were not actually used within the context,
remove PRIVATE, SHARED, and FIRSTPRIVATE clauses. */
@@ -7806,10 +7861,54 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
int kind = (flags & GOVD_MAP_TO_ONLY
? GOMP_MAP_TO
: GOMP_MAP_TOFROM);
+ tree c2 = NULL_TREE;
if (flags & GOVD_MAP_FORCE)
kind |= GOMP_MAP_FLAG_FORCE;
OMP_CLAUSE_SET_MAP_KIND (clause, kind);
- if (DECL_SIZE (decl)
+ c2 = gomp_needs_data_present (decl);
+ /* Handle OpenACC pointers that were declared inside acc data
+ regions. */
+ if (c2 != NULL && OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_POINTER)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SIZE (clause) = unshare_expr (OMP_CLAUSE_SIZE (c2));
+ }
+ /* Handle OpenACC subarrays that were declared inside acc data
+ regions. */
+ else if (c2 != NULL)
+ {
+ tree first = OMP_CLAUSE_DECL (c2);
+
+ /* Adjust the existing clause to make it a present data
+ clause with the proper subarray attributes. */
+ OMP_CLAUSE_DECL (clause) = unshare_expr (first);
+ OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_PRESENT);
+ OMP_CLAUSE_SIZE (clause) = unshare_expr (OMP_CLAUSE_SIZE (c2));
+
+ /* Create a new data clause for the firstprivate pointer. */
+ tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (nc) = decl;
+ OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
+
+ tree t = build_fold_addr_expr (first);
+ t = fold_convert_loc (OMP_CLAUSE_LOCATION (clause),
+ ptrdiff_type_node, t);
+ tree ptr = build_fold_addr_expr (decl);
+ t = fold_build2_loc (OMP_CLAUSE_LOCATION (clause), MINUS_EXPR,
+ ptrdiff_type_node, t,
+ fold_convert_loc (OMP_CLAUSE_LOCATION (clause),
+ ptrdiff_type_node, ptr));
+ OMP_CLAUSE_SIZE (nc) = t;
+
+ struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = ctx->outer_context;
+ gimplify_expr (&OMP_CLAUSE_SIZE (nc),
+ pre_p, NULL, is_gimple_val, fb_rvalue);
+ gimplify_omp_ctxp = ctx;
+ OMP_CLAUSE_CHAIN (clause) = nc;
+ }
+ else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
tree decl2 = DECL_VALUE_EXPR (decl);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c
new file mode 100644
index 0000000..c7dce2f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c
@@ -0,0 +1,25 @@
+#include <assert.h>
+
+int
+main ()
+{
+ int a[100], i;
+
+ for (i = 0; i < 100; i++)
+ a[i] = 0;
+
+#pragma acc data copy(a[10:80])
+ {
+ #pragma acc parallel loop
+ for (i = 10; i < 90; i++)
+ a[i] = i;
+ }
+
+ for (i = 0; i < 100; i++)
+ if (i >= 10 && i < 90)
+ assert (a[i] == i);
+ else
+ assert (a[i] == 0);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90
new file mode 100644
index 0000000..d1eba16
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90
@@ -0,0 +1,24 @@
+! Runtime data mapping error.
+
+program test
+ integer, parameter :: n = 100
+ integer i, data(n)
+
+ data(:) = 0
+
+ !$acc data copy(data(5:n-10))
+ !$acc parallel loop
+ do i = 10, n - 10
+ data(i) = i
+ end do
+ !$acc end parallel loop
+ !$acc end data
+
+ do i = 1, n
+ if ((i < 10 .or. i > n-10)) then
+ if ((data(i) .ne. 0)) call abort
+ else if (data(i) .ne. i) then
+ call abort
+ end if
+ end do
+end program test
