From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
gcc/c/
* c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses):
Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
gcc/
* gimplify.c (gimplify_scan_omp_clauses)
(gimplify_adjust_omp_clauses): Handle
OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
* omp-low.c (scan_sharing_clauses, lower_oacc_offload)
(lower_omp_target): Likewise.
* tree-core.h (enum omp_clause_map_kind)
<OMP_CLAUSE_MAP_FORCE_DEVICEPTR>: Update comment.
gcc/testsuite/
* c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC
deviceptr clause is now supported.
* c-c++-common/goacc/deviceptr-1.c: Extend.
* c-c++-common/goacc/deviceptr-2.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@211278
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 8 +++
gcc/c/ChangeLog.gomp | 5 ++
gcc/c/c-typeck.c | 5 +-
gcc/gimplify.c | 7 ++-
gcc/omp-low.c | 60 +++++++++++++++++++---
gcc/testsuite/ChangeLog.gomp | 5 ++
.../c-c++-common/goacc/data-clause-duplicate-1.c | 4 +-
gcc/testsuite/c-c++-common/goacc/deviceptr-1.c | 22 +++++++-
gcc/testsuite/c-c++-common/goacc/deviceptr-2.c | 23 +++++++++
gcc/tree-core.h | 3 +-
10 files changed, 127 insertions(+), 15 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/deviceptr-2.c
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 7371aa5..88f09b3 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,13 @@
2014-06-05 Thomas Schwinge <tho...@codesourcery.com>
+ * gimplify.c (gimplify_scan_omp_clauses)
+ (gimplify_adjust_omp_clauses): Handle
+ OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
+ * omp-low.c (scan_sharing_clauses, lower_oacc_offload)
+ (lower_omp_target): Likewise.
+ * tree-core.h (enum omp_clause_map_kind)
+ <OMP_CLAUSE_MAP_FORCE_DEVICEPTR>: Update comment.
+
* gimplify.c (gimplify_scan_omp_clauses) <case OMP_CLAUSE_MAP>:
Don't block OMP_CLAUSE_MAP_FORCE_PRESENT.
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index 91978db..1e80031 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-06-05 Thomas Schwinge <tho...@codesourcery.com>
+
+ * c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses):
+ Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
+
2014-03-20 Thomas Schwinge <tho...@codesourcery.com>
* c-parser.c: Update comments.
diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c
index c4ba531..839cdf7 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -11747,6 +11747,7 @@ handle_omp_array_sections (tree c)
OMP_CLAUSE_SIZE (c) = size;
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
return false;
+ gcc_assert (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
if (!c_mark_addressable (t))
@@ -12168,7 +12169,9 @@ c_finish_omp_clauses (tree clauses)
else if (!c_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)))
&& !lang_hooks.types.omp_mappable_type (TREE_TYPE (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
diff --git gcc/gimplify.c gcc/gimplify.c
index 6eaf6fd..a1b6be6 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -6015,7 +6015,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
switch (OMP_CLAUSE_MAP_KIND (c))
{
case OMP_CLAUSE_MAP_FORCE_DEALLOC:
- case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
input_location = OMP_CLAUSE_LOCATION (c);
/* TODO. */
sorry ("data clause not yet implemented");
@@ -6533,6 +6532,12 @@ gimplify_adjust_omp_clauses (tree *list_p)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
&& OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_POINTER)
{
+ /* For OMP_CLAUSE_MAP_FORCE_DEVICEPTR, we'll never enter here,
+ because for these, TREE_CODE (DECL_SIZE (decl)) will always be
+ INTEGER_CST. */
+ gcc_assert (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
+
tree decl2 = DECL_VALUE_EXPR (decl);
gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
decl2 = TREE_OPERAND (decl2, 0);
diff --git gcc/omp-low.c gcc/omp-low.c
index 3e282c0..39f0598 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1708,6 +1708,18 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& !POINTER_TYPE_P (TREE_TYPE (decl)))
break;
}
+#if 0
+ /* In target regions that are not offloaded, libgomp won't pay
+ attention to OMP_CLAUSE_MAP_FORCE_DEVICEPTR -- but I think we need
+ to handle it here anyway, in order to create a visible copy of the
+ variable. */
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ {
+ if (!is_gimple_omp_offloaded (ctx->stmt))
+ break;
+ }
+#endif
if (DECL_P (decl))
{
if (DECL_SIZE (decl)
@@ -1723,6 +1735,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
}
else
{
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -1738,6 +1754,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
{
tree base = get_base_address (decl);
tree nc = OMP_CLAUSE_CHAIN (c);
+ gcc_assert (nc == NULL_TREE
+ || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (nc)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR));
if (DECL_P (base)
&& nc != NULL_TREE
&& OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
@@ -1867,6 +1887,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
}
if (DECL_P (decl))
{
+ gcc_assert ((OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE);
if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
&& !COMPLETE_TYPE_P (TREE_TYPE (decl)))
@@ -1878,6 +1901,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
+ gcc_assert (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
+
tree decl2 = DECL_VALUE_EXPR (decl);
gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
decl2 = TREE_OPERAND (decl2, 0);
@@ -9100,6 +9126,10 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
{
x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx);
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (var)) != ARRAY_TYPE);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -9199,6 +9229,10 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
{
tree var = lookup_decl_in_outer_ctx (ovar, ctx);
tree x = build_sender_ref (ovar, ctx);
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -9219,12 +9253,14 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
= OMP_CLAUSE_MAP_KIND (c);
if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
&& (map_kind & OMP_CLAUSE_MAP_TO))
- || map_kind == OMP_CLAUSE_MAP_POINTER)
+ || map_kind == OMP_CLAUSE_MAP_POINTER
+ || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
gimplify_assign (avar, var, &ilist);
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
- if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
- && (map_kind & OMP_CLAUSE_MAP_FROM))
+ if (((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+ && (map_kind & OMP_CLAUSE_MAP_FROM))
+ || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
&& !TYPE_READONLY (TREE_TYPE (var)))
{
x = build_sender_ref (ovar, ctx);
@@ -10606,6 +10642,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
{
x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx);
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (var)) != ARRAY_TYPE);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -10732,12 +10772,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
{
tree var = lookup_decl_in_outer_ctx (ovar, ctx);
tree x = build_sender_ref (ovar, ctx);
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
{
- gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
tree avar
= create_tmp_var (TREE_TYPE (TREE_TYPE (x)), NULL);
mark_addressable (avar);
@@ -10747,19 +10790,20 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
}
else if (is_gimple_reg (var))
{
- gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
tree avar = create_tmp_var (TREE_TYPE (var), NULL);
mark_addressable (avar);
enum omp_clause_map_kind map_kind
= OMP_CLAUSE_MAP_KIND (c);
if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
&& (map_kind & OMP_CLAUSE_MAP_TO))
- || map_kind == OMP_CLAUSE_MAP_POINTER)
+ || map_kind == OMP_CLAUSE_MAP_POINTER
+ || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
gimplify_assign (avar, var, &ilist);
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
- if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
- && (map_kind & OMP_CLAUSE_MAP_FROM))
+ if (((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+ && (map_kind & OMP_CLAUSE_MAP_FROM))
+ || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
&& !TYPE_READONLY (TREE_TYPE (var)))
{
x = build_sender_ref (ovar, ctx);
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 4e0ee28..08ec907 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,5 +1,10 @@
2014-06-05 Thomas Schwinge <tho...@codesourcery.com>
+ * c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC
+ deviceptr clause is now supported.
+ * c-c++-common/goacc/deviceptr-1.c: Extend.
+ * c-c++-common/goacc/deviceptr-2.c: New file.
+
* c-c++-common/goacc/data-clause-duplicate-1.c: Extend.
* c-c++-common/goacc/present-1.c: New file.
diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
index 5c5ab02..7a1cf68 100644
--- gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
+++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
@@ -6,9 +6,7 @@ fun (void)
;
#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* {
dg-error "'fp' appears more than once in map clauses" } */
;
-#pragma acc data create(fp[:10]) deviceptr(fp)
- /* { dg-error "'fp' appears more than once in map clauses" "" { target *-*-*
} 9 } */
- /* { dg-message "sorry, unimplemented: data clause not yet implemented" "" {
target *-*-* } 9 } */
+#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears
more than once in map clauses" } */
;
#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than
once in map clauses" } */
;
diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
index 1ac63bd..cf2d809 100644
--- gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
+++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
@@ -61,4 +61,24 @@ fun3 (void)
;
}
-/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" }
*/
+extern struct s s1;
+extern struct s s2[1]; /* { dg-error "array type has incomplete element type"
"" { target c } } */
+
+void
+fun4 (void)
+{
+ struct s *s1_p = &s1;
+ struct s *s2_p = &s2;
+
+#pragma acc parallel deviceptr(s1) /* { dg-error "'s1' is not a pointer
variable" } */
+ ;
+
+#pragma acc parallel deviceptr(s2)
+ ;
+
+#pragma acc parallel deviceptr(s1_p)
+ s1_p = 0;
+
+#pragma acc parallel deviceptr(s2_p)
+ s2_p = 0;
+}
diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-2.c
gcc/testsuite/c-c++-common/goacc/deviceptr-2.c
new file mode 100644
index 0000000..ac162b4
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/deviceptr-2.c
@@ -0,0 +1,23 @@
+void
+fun1 (void)
+{
+ char *a = 0;
+
+#pragma acc data deviceptr(a)
+ ++a;
+
+#pragma acc data deviceptr(a)
+#pragma acc parallel
+ ++a;
+
+#pragma acc data deviceptr(a)
+#pragma acc parallel deviceptr(a)
+ ++a;
+
+#pragma acc data
+#pragma acc parallel deviceptr(a)
+ ++a;
+
+#pragma acc parallel deviceptr(a)
+ ++a;
+}
diff --git gcc/tree-core.h gcc/tree-core.h
index 8603553..8b70c5b 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -1225,7 +1225,8 @@ enum omp_clause_map_kind
OMP_CLAUSE_MAP_FORCE_PRESENT = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_SPECIAL,
/* Deallocate a mapping, without copying from device. */
OMP_CLAUSE_MAP_FORCE_DEALLOC,
- /* Is a device pointer. */
+ /* Is a device pointer. OMP_CLAUSE_SIZE for these is unused; is implicitly
+ POINTER_SIZE / BITS_PER_UNIT. */
OMP_CLAUSE_MAP_FORCE_DEVICEPTR,
/* End marker. */
--
1.9.1
