+2019-08-08 Jakub Jelinek <jakub@redhat.com>
+
+ * gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT
+ for VLA helper variables on target data even if not GOVD_FIRSTPRIVATE.
+ (gimplify_scan_omp_clauses): For OMP_CLAUSE_USE_DEVICE_* use just
+ GOVD_EXPLICIT flags.
+ (gimplify_omp_workshare): For OMP_TARGET_DATA move all
+ OMP_CLAUSE_USE_DEVICE_* clauses to the end of clauses chain.
+ * omp-low.c (scan_sharing_clauses): For OMP_CLAUSE_USE_DEVICE_*
+ call install_var_field with mask 11 instead of 3.
+ (lower_omp_target): For OMP_CLAUSE_USE_DEVICE_* use pass
+ (splay_tree_key) &DECL_UID (var) to build_sender_ref instead of var.
+
2019-08-07 Richard Sandiford <richard.sandiford@arm.com>
* config/aarch64/constraints.md (Z): Handle floating-point zeros too.
+2019-08-08 Jakub Jelinek <jakub@redhat.com>
+
+ * c-typeck.c (c_finish_omp_clauses): For C_ORT_OMP
+ OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
+ instead of generic_head to track duplicates.
+
2019-08-07 Jakub Jelinek <jakub@redhat.com>
* c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause.
/* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead. */
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
- /* If ort == C_ORT_OMP used as nontemporal_head instead. */
+ /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
+ instead. */
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
if (ort & C_ORT_ACC)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
- else if (ort == C_ORT_ACC
- && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ else if ((ort == C_ORT_ACC
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ || (ort == C_ORT_OMP
+ && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
+ || (OMP_CLAUSE_CODE (c)
+ == OMP_CLAUSE_USE_DEVICE_ADDR))))
{
if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
- "%qD appears more than once in reduction clauses",
+ ort == C_ORT_ACC
+ ? "%qD appears more than once in reduction clauses"
+ : "%qD appears more than once in data clauses",
t);
remove = true;
}
+2019-08-08 Jakub Jelinek <jakub@redhat.com>
+
+ * semantics.c (finish_omp_clauses): For C_ORT_OMP
+ OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
+ instead of generic_head to track duplicates.
+
2019-08-07 Marek Polacek <polacek@redhat.com>
PR c++/81429 - wrong parsing of constructor with C++11 attribute.
/* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead. */
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
- /* If ort == C_ORT_OMP used as nontemporal_head instead. */
+ /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
+ instead. */
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
if (ort & C_ORT_ACC)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
- else if (ort == C_ORT_ACC
- && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ else if ((ort == C_ORT_ACC
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ || (ort == C_ORT_OMP
+ && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
+ || (OMP_CLAUSE_CODE (c)
+ == OMP_CLAUSE_USE_DEVICE_ADDR))))
{
if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
- "%qD appears more than once in reduction clauses",
+ ort == C_ORT_ACC
+ ? "%qD appears more than once in reduction clauses"
+ : "%qD appears more than once in data clauses",
t);
remove = true;
}
nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
else if (flags & GOVD_PRIVATE)
nflags = GOVD_PRIVATE;
- else if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
- && (flags & GOVD_FIRSTPRIVATE))
+ else if (((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
+ && (flags & GOVD_FIRSTPRIVATE))
+ || (ctx->region_type == ORT_TARGET_DATA
+ && (flags & GOVD_DATA_SHARE_CLASS) == 0))
nflags = GOVD_PRIVATE | GOVD_EXPLICIT;
else
nflags = GOVD_FIRSTPRIVATE;
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
+ flags = GOVD_EXPLICIT;
+ goto do_add;
+
case OMP_CLAUSE_IS_DEVICE_PTR:
flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
goto do_add;
OMP_CLAUSES (expr));
break;
case OMP_TARGET_DATA:
- stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_DATA,
- OMP_CLAUSES (expr));
+ /* Put use_device_{ptr,addr} clauses last, as map clauses are supposed
+ to be evaluated before the use_device_{ptr,addr} clauses if they
+ refer to the same variables. */
+ {
+ tree use_device_clauses;
+ tree *pc, *uc = &use_device_clauses;
+ for (pc = &OMP_CLAUSES (expr); *pc; )
+ if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR
+ || OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_ADDR)
+ {
+ *uc = *pc;
+ *pc = OMP_CLAUSE_CHAIN (*pc);
+ uc = &OMP_CLAUSE_CHAIN (*uc);
+ }
+ else
+ pc = &OMP_CLAUSE_CHAIN (*pc);
+ *uc = NULL_TREE;
+ *pc = use_device_clauses;
+ stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_DATA,
+ OMP_CLAUSES (expr));
+ }
break;
case OMP_TEAMS:
stmt = gimple_build_omp_teams (body, OMP_CLAUSES (expr));
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
&& !omp_is_reference (decl))
|| TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
- install_var_field (decl, true, 3, ctx);
+ install_var_field (decl, true, 11, ctx);
else
- install_var_field (decl, false, 3, ctx);
+ install_var_field (decl, false, 11, ctx);
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
case OMP_CLAUSE_IS_DEVICE_PTR:
ovar = OMP_CLAUSE_DECL (c);
var = lookup_decl_in_outer_ctx (ovar, ctx);
- x = build_sender_ref (ovar, ctx);
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
- tkind = GOMP_MAP_USE_DEVICE_PTR;
+ {
+ tkind = GOMP_MAP_USE_DEVICE_PTR;
+ x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), ctx);
+ }
else
- tkind = GOMP_MAP_FIRSTPRIVATE_INT;
+ {
+ tkind = GOMP_MAP_FIRSTPRIVATE_INT;
+ x = build_sender_ref (ovar, ctx);
+ }
type = TREE_TYPE (ovar);
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
&& !omp_is_reference (ovar))
case OMP_CLAUSE_IS_DEVICE_PTR:
var = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
- x = build_sender_ref (var, ctx);
+ x = build_sender_ref ((splay_tree_key) &DECL_UID (var), ctx);
else
x = build_receiver_ref (var, false, ctx);
if (is_variable_sized (var))
+2019-08-08 Jakub Jelinek <jakub@redhat.com>
+
+ * target.c (gomp_map_vars_internal): For GOMP_MAP_USE_DEVICE_PTR
+ perform the lookup in the first loop only if !not_found_cnt, otherwise
+ perform lookups for it in the second loop guarded with
+ if (not_found_cnt || has_firstprivate).
+ * testsuite/libgomp.c/target-37.c: New test.
+ * testsuite/libgomp.c++/target-22.C: New test.
+
2019-08-07 Jakub Jelinek <jakub@redhat.com>
* testsuite/libgomp.c/target-18.c (struct S): New type.
}
else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR)
{
- cur_node.host_start = (uintptr_t) hostaddrs[i];
- cur_node.host_end = cur_node.host_start;
- splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
- if (n == NULL)
+ tgt->list[i].key = NULL;
+ if (!not_found_cnt)
{
- gomp_mutex_unlock (&devicep->lock);
- gomp_fatal ("use_device_ptr pointer wasn't mapped");
}
- cur_node.host_start -= n->host_start;
- hostaddrs[i]
- = (void *) (n->tgt->tgt_start + n->tgt_offset
- + cur_node.host_start);
- tgt->list[i].key = NULL;
- tgt->list[i].offset = ~(uintptr_t) 0;
+ else
+ tgt->list[i].offset = 0;
continue;
}
else if ((kind & typemask) == GOMP_MAP_STRUCT)
tgt_size += len;
continue;
case GOMP_MAP_FIRSTPRIVATE_INT:
- case GOMP_MAP_USE_DEVICE_PTR:
case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
continue;
+ case GOMP_MAP_USE_DEVICE_PTR:
+ if (tgt->list[i].offset == 0)
+ {
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start;
+ n = gomp_map_lookup (mem_map, &cur_node);
+ if (n == NULL)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("use_device_ptr pointer wasn't mapped");
+ }
+ cur_node.host_start -= n->host_start;
+ hostaddrs[i]
+ = (void *) (n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start);
+ tgt->list[i].offset = ~(uintptr_t) 0;
+ }
+ continue;
case GOMP_MAP_STRUCT:
first = i + 1;
last = i + sizes[i];
--- /dev/null
+extern "C" void abort (void);
+struct S { int e, f; };
+
+void
+foo (int *&p, int (&s)[5], int &t, S &u, int n)
+{
+ int a[4] = { 7, 8, 9, 10 }, b[n], c[3] = { 20, 21, 22 };
+ int *r = a + 1, *q = p - 1, i, err;
+ int v = 27;
+ S w = { 28, 29 };
+ for (i = 0; i < n; i++)
+ b[i] = 9 + i;
+ #pragma omp target data map(to:a) use_device_ptr(r) map(from:err)
+ #pragma omp target is_device_ptr(r) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 4; i++)
+ if (r[i - 1] != 7 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data use_device_ptr(p) map(from:err) map(to:q[:4])
+ #pragma omp target is_device_ptr(p) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 4; i++)
+ if (p[i - 1] != i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data map(to:b) use_device_addr(b) map(from:err)
+ #pragma omp target is_device_ptr(b) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < n; i++)
+ if (b[i] != 9 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data use_device_addr(c) map(to:c) map(from:err)
+ #pragma omp target is_device_ptr(c) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 3; i++)
+ if (c[i] != 20 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data map(to:s[:5]) use_device_addr(s) map(from:err)
+ #pragma omp target is_device_ptr(s) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 5; i++)
+ if (s[i] != 17 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data use_device_addr (v) map(to: v) map(to:u) use_device_addr (u) map(from:err)
+ {
+ int *z = &v;
+ S *x = &u;
+ #pragma omp target is_device_ptr (z, x) map(from:err)
+ {
+ err = 0;
+ if (*z != 27 || x->e != 25 || x->f != 26)
+ err = 1;
+ }
+ }
+ if (err)
+ abort ();
+ #pragma omp target data map(to: t) use_device_addr (t, w) map (to: w) map(from:err)
+ {
+ int *z = &t;
+ S *x = &w;
+ #pragma omp target is_device_ptr (z) is_device_ptr (x) map(from:err)
+ {
+ err = 0;
+ if (*z != 24 || x->e != 28 || x->f != 29)
+ err = 1;
+ }
+ }
+ if (err)
+ abort ();
+}
+
+int
+main ()
+{
+ int a[4] = { 0, 1, 2, 3 }, b[5] = { 17, 18, 19, 20, 21 };
+ int *p = a + 1;
+ int t = 24;
+ S u = { 25, 26 };
+ foo (p, b, t, u, 9);
+}
--- /dev/null
+extern void abort (void);
+struct S { int e, f; };
+
+void
+foo (int n)
+{
+ int a[4] = { 0, 1, 2, 3 }, b[n], c = 4;
+ struct S d = { 5, 6 };
+ int *p = a + 1, i, err;
+ for (i = 0; i < n; i++)
+ b[i] = 9 + i;
+ #pragma omp target data use_device_ptr(p) map(from:err) map(to:a)
+ #pragma omp target is_device_ptr(p) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 4; i++)
+ if (p[i - 1] != i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < 4; i++)
+ a[i] = 23 + i;
+ #pragma omp target data map(to:a) use_device_addr(a) map(from:err)
+ #pragma omp target is_device_ptr(a) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 4; i++)
+ if (a[i] != 23 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data use_device_addr(b) map(from:err) map(to:b)
+ #pragma omp target is_device_ptr(b) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 4; i++)
+ if (b[i] != 9 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data map(to:c) use_device_addr(c) map(from:err)
+ {
+ int *q = &c;
+ #pragma omp target is_device_ptr(q) map(from:err)
+ {
+ err = *q != 4;
+ }
+ }
+ if (err)
+ abort ();
+ #pragma omp target data use_device_addr(d) map(to:d) map(from:err)
+ {
+ struct S *r = &d;
+ #pragma omp target is_device_ptr(r) map(from:err)
+ {
+ err = r->e != 5 || r->f != 6;
+ }
+ }
+ if (err)
+ abort ();
+}
+
+int
+main ()
+{
+ foo (9);
+ return 0;
+}
projects / gcc.git / commitdiff