@@ -1574,10 +1574,24 @@ scan_sharing_clauses (tree clauses, omp_
}
if (DECL_P (decl))
{
- install_var_field (decl, true, 3, ctx);
- if (gimple_omp_target_kind (ctx->stmt)
- == GF_OMP_TARGET_KIND_REGION)
- install_var_local (decl, ctx);
+ if (DECL_SIZE (decl)
+ && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+ {
+ tree decl2 = DECL_VALUE_EXPR (decl);
+ gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+ decl2 = TREE_OPERAND (decl2, 0);
+ gcc_assert (DECL_P (decl2));
+ install_var_field (decl2, true, 3, ctx);
+ install_var_local (decl2, ctx);
+ install_var_local (decl, ctx);
+ }
+ else
+ {
+ install_var_field (decl, true, 3, ctx);
+ if (gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_REGION)
+ install_var_local (decl, ctx);
+ }
}
else
{
@@ -1600,6 +1614,7 @@ scan_sharing_clauses (tree clauses, omp_
tree field
= build_decl (OMP_CLAUSE_LOCATION (c),
FIELD_DECL, NULL_TREE, ptr_type_node);
+ DECL_ALIGN (field) = TYPE_ALIGN (ptr_type_node);
insert_field_into_struct (ctx->record_type, field);
splay_tree_insert (ctx->field_map, (splay_tree_key) decl,
(splay_tree_value) field);
@@ -1684,6 +1699,16 @@ scan_sharing_clauses (tree clauses, omp_
TREE_TYPE (new_decl)
= remap_type (TREE_TYPE (decl), &ctx->cb);
}
+ else if (DECL_SIZE (decl)
+ && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+ {
+ tree decl2 = DECL_VALUE_EXPR (decl);
+ gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+ decl2 = TREE_OPERAND (decl2, 0);
+ gcc_assert (DECL_P (decl2));
+ fixup_remapped_decl (decl2, ctx, false);
+ fixup_remapped_decl (decl, ctx, true);
+ }
else
fixup_remapped_decl (decl, ctx, false);
}
@@ -2126,6 +2151,16 @@ scan_omp_target (gimple stmt, omp_contex
ctx->record_type = ctx->receiver_decl = NULL;
else
{
+ TYPE_FIELDS (ctx->record_type)
+ = nreverse (TYPE_FIELDS (ctx->record_type));
+#ifdef ENABLE_CHECKING
+ tree field;
+ unsigned int align = DECL_ALIGN (TYPE_FIELDS (ctx->record_type));
+ for (field = TYPE_FIELDS (ctx->record_type);
+ field;
+ field = DECL_CHAIN (field))
+ gcc_assert (DECL_ALIGN (field) == align);
+#endif
layout_type (ctx->record_type);
if (kind == GF_OMP_TARGET_KIND_REGION)
fixup_child_record_type (ctx);
@@ -9201,7 +9236,18 @@ lower_omp_target (gimple_stmt_iterator *
map_cnt++;
continue;
}
- if (!lookup_sfield (var, ctx))
+
+ if (DECL_SIZE (var)
+ && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+ {
+ tree var2 = DECL_VALUE_EXPR (var);
+ gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
+ var2 = TREE_OPERAND (var2, 0);
+ gcc_assert (DECL_P (var2));
+ var = var2;
+ }
+
+ if (!maybe_lookup_field (var, ctx))
continue;
if (kind == GF_OMP_TARGET_KIND_REGION)
@@ -9293,8 +9339,20 @@ lower_omp_target (gimple_stmt_iterator *
nc = NULL_TREE;
}
}
- else if (!lookup_sfield (ovar, ctx))
- continue;
+ else
+ {
+ if (DECL_SIZE (ovar)
+ && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
+ {
+ tree ovar2 = DECL_VALUE_EXPR (ovar);
+ gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF);
+ ovar2 = TREE_OPERAND (ovar2, 0);
+ gcc_assert (DECL_P (ovar2));
+ ovar = ovar2;
+ }
+ if (!maybe_lookup_field (ovar, ctx))
+ continue;
+ }
if (nc)
{
@@ -508,7 +508,8 @@ dump_omp_clause (pretty_printer *buffer,
print_clause_size:
if (OMP_CLAUSE_SIZE (clause))
{
- if (OMP_CLAUSE_MAP_KIND (clause) == OMP_CLAUSE_MAP_POINTER)
+ if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (clause) == OMP_CLAUSE_MAP_POINTER)
pp_string (buffer, " [pointer assign, bias: ");
else
pp_string (buffer, " [len: ");
@@ -61,6 +61,7 @@ enum gimplify_omp_var_data
GOVD_PRIVATE_OUTER_REF = 1024,
GOVD_LINEAR = 2048,
GOVD_ALIGNED = 4096,
+ GOVD_MAP_TO_ONLY = 8192,
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -5740,11 +5741,16 @@ omp_firstprivatize_variable (struct gimp
{
if (n->value & GOVD_SHARED)
n->value = GOVD_FIRSTPRIVATE | (n->value & GOVD_SEEN);
+ else if (n->value & GOVD_MAP)
+ n->value |= GOVD_MAP_TO_ONLY;
else
return;
}
+ else if (ctx->region_type == ORT_TARGET)
+ omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
else if (ctx->region_type != ORT_WORKSHARE
- && ctx->region_type != ORT_SIMD)
+ && ctx->region_type != ORT_SIMD
+ && ctx->region_type != ORT_TARGET_DATA)
omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
ctx = ctx->outer_context;
@@ -5847,16 +5853,15 @@ omp_add_variable (struct gimplify_omp_ct
the parameters of the type. */
if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
- /* To be handled later. */
- gcc_assert ((flags & GOVD_MAP) == 0);
-
/* Add the pointer replacement variable as PRIVATE if the variable
replacement is private, else FIRSTPRIVATE since we'll need the
address of the original variable either for SHARED, or for the
copy into or out of the context. */
if (!(flags & GOVD_LOCAL))
{
- nflags = flags & GOVD_PRIVATE ? GOVD_PRIVATE : GOVD_FIRSTPRIVATE;
+ nflags = flags & GOVD_MAP
+ ? GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT
+ : flags & GOVD_PRIVATE ? GOVD_PRIVATE : GOVD_FIRSTPRIVATE;
nflags |= flags & GOVD_SEEN;
t = DECL_VALUE_EXPR (decl);
gcc_assert (TREE_CODE (t) == INDIRECT_REF);
@@ -5885,15 +5890,13 @@ omp_add_variable (struct gimplify_omp_ct
For local variables TYPE_SIZE_UNIT might not be gimplified yet,
in this case omp_notice_variable will be called later
on when it is gimplified. */
- else if (! (flags & GOVD_LOCAL)
+ else if (! (flags & (GOVD_LOCAL | GOVD_MAP))
&& DECL_P (TYPE_SIZE_UNIT (TREE_TYPE (decl))))
omp_notice_variable (ctx, TYPE_SIZE_UNIT (TREE_TYPE (decl)), true);
}
- else if (lang_hooks.decls.omp_privatize_by_reference (decl))
+ else if ((flags & GOVD_MAP) == 0
+ && lang_hooks.decls.omp_privatize_by_reference (decl))
{
- /* To be handled later. */
- gcc_assert ((flags & GOVD_MAP) == 0);
-
gcc_assert ((flags & GOVD_LOCAL) == 0);
omp_firstprivatize_type_sizes (ctx, TREE_TYPE (decl));
@@ -6562,7 +6565,35 @@ gimplify_adjust_omp_clauses_1 (splay_tre
else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
else if (code == OMP_CLAUSE_MAP)
- OMP_CLAUSE_MAP_KIND (clause) = OMP_CLAUSE_MAP_TOFROM;
+ {
+ OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY
+ ? OMP_CLAUSE_MAP_TO
+ : OMP_CLAUSE_MAP_TOFROM;
+ if (DECL_SIZE (decl)
+ && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+ {
+ tree decl2 = DECL_VALUE_EXPR (decl);
+ gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+ decl2 = TREE_OPERAND (decl2, 0);
+ gcc_assert (DECL_P (decl2));
+ tree mem = build_simple_mem_ref (decl2);
+ OMP_CLAUSE_DECL (clause) = mem;
+ OMP_CLAUSE_SIZE (clause) = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+ if (gimplify_omp_ctxp->outer_context)
+ {
+ struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp->outer_context;
+ omp_notice_variable (ctx, decl2, true);
+ omp_notice_variable (ctx, OMP_CLAUSE_SIZE (clause), true);
+ }
+ tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (nc) = decl;
+ OMP_CLAUSE_SIZE (nc) = size_zero_node;
+ OMP_CLAUSE_MAP_KIND (nc) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause);
+ OMP_CLAUSE_CHAIN (clause) = nc;
+ }
+ }
*list_p = clause;
lang_hooks.decls.omp_finish_clause (clause);
@@ -6687,6 +6718,56 @@ gimplify_adjust_omp_clauses (tree *list_
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
remove = true;
+ else if (DECL_SIZE (decl)
+ && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
+ && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_POINTER)
+ {
+ tree decl2 = DECL_VALUE_EXPR (decl);
+ gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+ decl2 = TREE_OPERAND (decl2, 0);
+ gcc_assert (DECL_P (decl2));
+ tree mem = build_simple_mem_ref (decl2);
+ OMP_CLAUSE_DECL (c) = mem;
+ OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+ if (ctx->outer_context)
+ {
+ omp_notice_variable (ctx->outer_context, decl2, true);
+ omp_notice_variable (ctx->outer_context,
+ OMP_CLAUSE_SIZE (c), true);
+ }
+ tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (nc) = decl;
+ OMP_CLAUSE_SIZE (nc) = size_zero_node;
+ OMP_CLAUSE_MAP_KIND (nc) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = nc;
+ c = nc;
+ }
+ break;
+
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ decl = OMP_CLAUSE_DECL (c);
+ if (!DECL_P (decl))
+ break;
+ if (DECL_SIZE (decl)
+ && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+ {
+ tree decl2 = DECL_VALUE_EXPR (decl);
+ gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+ decl2 = TREE_OPERAND (decl2, 0);
+ gcc_assert (DECL_P (decl2));
+ tree mem = build_simple_mem_ref (decl2);
+ OMP_CLAUSE_DECL (c) = mem;
+ OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+ if (ctx->outer_context)
+ {
+ omp_notice_variable (ctx->outer_context, decl2, true);
+ omp_notice_variable (ctx->outer_context,
+ OMP_CLAUSE_SIZE (c), true);
+ }
+ }
break;
case OMP_CLAUSE_REDUCTION:
@@ -6708,8 +6789,6 @@ gimplify_adjust_omp_clauses (tree *list_
case OMP_CLAUSE_MERGEABLE:
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE_SAFELEN:
- case OMP_CLAUSE_TO:
- case OMP_CLAUSE_FROM:
case OMP_CLAUSE_DEPEND:
break;
@@ -0,0 +1,88 @@
+extern
+#ifdef __cplusplus
+"C"
+#endif
+void abort (void);
+
+void
+fn1 (double *x, double *y, int z)
+{
+ int i;
+ for (i = 0; i < z; i++)
+ {
+ x[i] = i & 31;
+ y[i] = (i & 63) - 30;
+ }
+}
+
+double
+fn2 (int x)
+{
+ double s = 0;
+ double b[3 * x], c[3 * x], d[3 * x], e[3 * x];
+ int i;
+ fn1 (b, c, x);
+ fn1 (e, d + x, x);
+ #pragma omp target map(to: b, c[:x], d[x:x], e)
+ #pragma omp parallel for reduction(+:s)
+ for (i = 0; i < x; i++)
+ s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c);
+ return s;
+}
+
+double
+fn3 (int x)
+{
+ double s = 0;
+ double b[3 * x], c[3 * x], d[3 * x], e[3 * x];
+ int i;
+ fn1 (b, c, x);
+ fn1 (e, d, x);
+ #pragma omp target
+ #pragma omp parallel for reduction(+:s)
+ for (i = 0; i < x; i++)
+ s += b[i] * c[i] + d[i];
+ return s;
+}
+
+double
+fn4 (int x)
+{
+ double s = 0;
+ double b[3 * x], c[3 * x], d[3 * x], e[3 * x];
+ int i;
+ fn1 (b, c, x);
+ fn1 (e, d + x, x);
+ #pragma omp target data map(from: b, c[:x], d[x:x], e)
+ {
+ #pragma omp target update to(b, c[:x], d[x:x], e)
+ #pragma omp target map(c[:x], d[x:x])
+ #pragma omp parallel for reduction(+:s)
+ for (i = 0; i < x; i++)
+ {
+ s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c);
+ b[i] = i + 0.5;
+ c[i] = 0.5 - i;
+ d[x + i] = 0.5 * i;
+ }
+ }
+ for (i = 0; i < x; i++)
+ if (b[i] != i + 0.5 || c[i] != 0.5 - i || d[x + i] != 0.5 * i)
+ abort ();
+ return s;
+}
+
+int
+main ()
+{
+ double a = fn2 (128);
+ if (a != 14080.0)
+ abort ();
+ double b = fn3 (128);
+ if (a != b)
+ abort ();
+ double c = fn4 (256);
+ if (c != 28160.0)
+ abort ();
+ return 0;
+}
@@ -0,0 +1 @@
+#include "../libgomp.c/target-2.c"
Hi! This fixes mainly VLA handling in target{, data, update} constructs, but also deals with field alignments in the structure and field order. Committed to gomp-4_0-branch. 2013-09-06 Jakub Jelinek <jakub@redhat.com> * omp-low.c (scan_sharing_clauses): Handle VLAs in OMP_CLAUSE_{MAP,TO,FROM}. Set DECL_ALIGN (field) before calling insert_field_into_struct. (scan_omp_target): Reverse TYPE_FIELDS, verify that all field alignments are the same. (lower_omp_target): Use maybe_lookup_field instead of lookup_sfield to check if field is present. Handle VLAs. * tree-pretty-print.c (dump_omp_clause): Only check OMP_CLAUSE_MAP_KIND on OMP_CLAUSE_MAP clauses. * gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_TO_ONLY. (omp_firstprivatize_variable, omp_add_variable, gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Handle VLAs in OMP_CLAUSE_{MAP,TO,FROM}. libgomp/ * testsuite/libgomp.c/target-2.c: New test. * testsuite/libgomp.c++/target-3.C: New test. Jakub