2016-03-14 Tom de Vries <tom@codesourcery.com>
* omp-low.c (install_var_field): Handle base_pointers_restrict for
pointers.
(omp_target_base_pointers_restrict_p): Handle GOMP_MAP_POINTER.
* tree.h (OMP_CLAUSE_MAP_POINTER_TO_FORCED): define.
* c-typeck.c (handle_omp_array_sections): Set
OMP_CLAUSE_MAP_POINTER_TO_FORCED on GOMP_MAP_POINTER clause.
* c-c++-common/goacc/kernels-alias-10.c: New test.
* c-c++-common/goacc/kernels-alias-9.c: New test.
Handle non-declared variables in kernels alias analysis
---
gcc/c/c-typeck.c | 15 ++++++-
gcc/omp-low.c | 48 ++++++++++++++++++++++
.../c-c++-common/goacc/kernels-alias-10.c | 29 +++++++++++++
gcc/testsuite/c-c++-common/goacc/kernels-alias-9.c | 29 +++++++++++++
gcc/tree.h | 3 ++
5 files changed, 123 insertions(+), 1 deletion(-)
@@ -12446,7 +12446,20 @@ handle_omp_array_sections (tree c, bool is_omp)
}
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
if (!is_omp)
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_FORCE_TOFROM:
+ OMP_CLAUSE_MAP_POINTER_TO_FORCED (c2) = 1;
+ break;
+ default:
+ break;
+ }
+ }
else if (TREE_CODE (t) == COMPONENT_REF)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
else
@@ -1429,6 +1429,9 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx,
}
else if (by_ref)
{
+ if (base_pointers_restrict
+ && POINTER_TYPE_P (type))
+ type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
type = build_pointer_type (type);
if (base_pointers_restrict)
type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
@@ -3132,6 +3135,47 @@ omp_target_base_pointers_restrict_p (tree clauses)
Because both mappings have the force prefix, we know that they will be
allocated when calling the corresponding offloaded function, which means we
can mark the base pointers for a and b in the offloaded function as
+ restrict.
+
+ II. GOMP_MAP_POINTER example:
+
+ void foo (unsigned int *a, unsigned int *b)
+ {
+ #pragma acc kernels copyout (a[0:2]) copyout (b[0:2])
+ {
+ a[0] = 0;
+ b[0] = 1;
+ }
+ }
+
+ After gimplification, we have:
+
+ foo (unsigned int * a, unsigned int * b)
+ {
+ unsigned int * b.0;
+ unsigned int * a.1;
+
+ b.0 = b;
+ a.1 = a;
+ #pragma omp target oacc_kernels \
+ map(force_from:*a.1 (*a) [len: 8]) \
+ map(alloc:a [pointer assign, bias: 0]) \
+ map(force_from:*b.0 (*b) [len: 8]) \
+ map(alloc:b [pointer assign, bias: 0])
+ {
+ unsigned int * a.2;
+ unsigned int * b.3;
+
+ a.2 = a;
+ *a.2 = 0;
+ b.3 = b;
+ *b.3 = 1;
+ }
+ }
+
+ By testing for OMP_CLAUSE_MAP_POINTER_TO_FORCED, we can known for both
+ pointer assign mappings that they point to a force-prefixed mapping, so
+ we can mark the base pointers for a and b in the offloaded function as
restrict. */
tree c;
@@ -3147,6 +3191,10 @@ omp_target_base_pointers_restrict_p (tree clauses)
case GOMP_MAP_FORCE_FROM:
case GOMP_MAP_FORCE_TOFROM:
break;
+ case GOMP_MAP_POINTER:
+ if (!OMP_CLAUSE_MAP_POINTER_TO_FORCED (c))
+ return false;
+ break;
default:
return false;
}
new file mode 100644
@@ -0,0 +1,29 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-ealias-all" } */
+
+#define N 2
+
+void
+foo (void)
+{
+ unsigned int a[N];
+ unsigned int b[N];
+ unsigned int c[N];
+ unsigned int d[N];
+
+#pragma acc kernels copyin (a[0:N]) create (b[0:N]) copyout (c[0:N]) copy (d[0:N])
+ {
+ a[0] = 0;
+ b[0] = 0;
+ c[0] = 0;
+ d[0] = 0;
+ }
+}
+
+/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */
+
new file mode 100644
@@ -0,0 +1,29 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-ealias-all" } */
+
+#define N 2
+
+void
+foo (unsigned int *a, unsigned int *b, unsigned int *c, unsigned int *d)
+{
+
+#pragma acc kernels copyin (a[0:N]) create (b[0:N]) copyout (c[0:N]) copy (d[0:N])
+ {
+ a[0] = 0;
+ b[0] = 0;
+ c[0] = 0;
+ d[0] = 0;
+ }
+}
+
+/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 6" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 7" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 8" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 9" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 12 "ealias" } } */
+
@@ -1533,6 +1533,9 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION(NODE) \
TREE_PROTECTED (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+#define OMP_CLAUSE_MAP_POINTER_TO_FORCED(NODE) \
+ TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+
#define OMP_CLAUSE_PROC_BIND_KIND(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_PROC_BIND)->omp_clause.subcode.proc_bind_kind)