@@ -2704,7 +2704,14 @@ gimplify_call_expr (tree *expr_p, gimple
notice_special_calls (call);
gimplify_seq_add_stmt (pre_p, call);
gsi = gsi_last (*pre_p);
- fold_stmt (&gsi);
+ /* Don't fold stmts inside of target construct. We'll do it
+ during omplower pass instead. */
+ struct gimplify_omp_ctx *ctx;
+ for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
+ if (ctx->region_type == ORT_TARGET)
+ break;
+ if (ctx == NULL)
+ fold_stmt (&gsi);
*expr_p = NULL_TREE;
}
else
@@ -4961,7 +4968,14 @@ gimplify_modify_expr (tree *expr_p, gimp
gimplify_seq_add_stmt (pre_p, assign);
gsi = gsi_last (*pre_p);
- fold_stmt (&gsi);
+ /* Don't fold stmts inside of target construct. We'll do it
+ during omplower pass instead. */
+ struct gimplify_omp_ctx *ctx;
+ for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
+ if (ctx->region_type == ORT_TARGET)
+ break;
+ if (ctx == NULL)
+ fold_stmt (&gsi);
if (want_value)
{
@@ -134,6 +134,7 @@ struct omp_for_data
static splay_tree all_contexts;
static int taskreg_nesting_level;
+static int target_nesting_level;
struct omp_region *root_omp_region;
static bitmap task_shared_vars;
@@ -9213,7 +9214,13 @@ lower_omp_target (gimple_stmt_iterator *
map_cnt++;
}
- if (kind != GF_OMP_TARGET_KIND_UPDATE)
+ if (kind == GF_OMP_TARGET_KIND_REGION)
+ {
+ target_nesting_level++;
+ lower_omp (&tgt_body, ctx);
+ target_nesting_level--;
+ }
+ else if (kind == GF_OMP_TARGET_KIND_DATA)
lower_omp (&tgt_body, ctx);
if (kind == GF_OMP_TARGET_KIND_REGION)
@@ -9320,7 +9327,7 @@ lower_omp_target (gimple_stmt_iterator *
}
tree s = OMP_CLAUSE_SIZE (c);
if (s == NULL_TREE)
- s = TYPE_SIZE (TREE_TYPE (ovar));
+ s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
s = fold_convert (size_type_node, s);
tree purpose = size_int (map_idx++);
CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
@@ -9342,6 +9349,11 @@ lower_omp_target (gimple_stmt_iterator *
default:
gcc_unreachable ();
}
+ unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
+ if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
+ talign = DECL_ALIGN_UNIT (ovar);
+ talign = ceil_log2 (talign);
+ tkind |= talign << 3;
CONSTRUCTOR_APPEND_ELT (vkind, purpose,
build_int_cst (unsigned_char_type_node,
tkind));
@@ -9673,6 +9685,12 @@ lower_omp (gimple_seq *body, omp_context
gimple_stmt_iterator gsi;
for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi))
lower_omp_1 (&gsi, ctx);
+ /* Inside target region we haven't called fold_stmt during gimplification,
+ because it can break code by adding decl references that weren't in the
+ source. Call fold_stmt now. */
+ if (target_nesting_level)
+ for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi))
+ fold_stmt (&gsi);
input_location = saved_location;
}
@@ -4506,6 +4506,7 @@ handle_omp_array_sections (tree c)
t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
ptrdiff_type_node, t);
tree ptr = OMP_CLAUSE_DECL (c2);
+ ptr = convert_from_reference (ptr);
if (!POINTER_TYPE_P (TREE_TYPE (ptr)))
ptr = build_fold_addr_expr (ptr);
t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
@@ -4515,6 +4516,19 @@ handle_omp_array_sections (tree c)
OMP_CLAUSE_SIZE (c2) = t;
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
+ ptr = OMP_CLAUSE_DECL (c2);
+ if (TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE
+ && POINTER_TYPE_P (TREE_TYPE (TREE_TYPE (ptr))))
+ {
+ tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (c3) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_DECL (c3) = ptr;
+ OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr);
+ OMP_CLAUSE_SIZE (c3) = size_zero_node;
+ OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (c2);
+ OMP_CLAUSE_CHAIN (c2) = c3;
+ }
}
}
return false;
@@ -4943,6 +4957,9 @@ finish_omp_clauses (tree clauses)
{
if (processing_template_decl)
break;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
+ break;
if (DECL_P (t))
error ("%qD is not a variable in %qs clause", t,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
@@ -0,0 +1,58 @@
+// { dg-options "-O2 -fopenmp" }
+// { dg-additional-sources "target-2-aux.cc" }
+
+extern "C" 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 b[1024];
+double (&br) [1024] = b;
+double cbuf[1024];
+double *c = cbuf;
+double *&cr = c;
+extern double (&fr) [1024];
+extern double *&gr;
+
+double
+fn2 (int x, double (&dr) [1024], double *&er)
+{
+ double s = 0;
+ double h[1024];
+ double (&hr) [1024] = h;
+ double ibuf[1024];
+ double *i = ibuf;
+ double *&ir = i;
+ int j;
+ fn1 (hr + 2 * x, ir + 2 * x, x);
+ #pragma omp target map(to: br[:x], cr[0:x], dr[x:x], er[x:x]) \
+ map(to: fr[0:x], gr[0:x], hr[2 * x:x], ir[2 * x:x])
+ #pragma omp parallel for reduction(+:s)
+ for (j = 0; j < x; j++)
+ s += br[j] * cr[j] + dr[x + j] + er[x + j]
+ + fr[j] + gr[j] + hr[2 * x + j] + ir[2 * x + j];
+ return s;
+}
+
+int
+main ()
+{
+ double d[1024];
+ double ebuf[1024];
+ double *e = ebuf;
+ fn1 (br, cr, 128);
+ fn1 (d + 128, e + 128, 128);
+ fn1 (fr, gr, 128);
+ double h = fn2 (128, d, e);
+ if (h != 20416.0)
+ abort ();
+ return 0;
+}
@@ -0,0 +1,5 @@
+double f[1024];
+double (&fr) [1024] = f;
+double gbuf[1024];
+double *g = gbuf;
+double *&gr = g;