From: Jakub Jelinek Date: Thu, 8 Aug 2019 06:39:02 +0000 (+0200) Subject: gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT for VLA helper variab... X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=8860d2706d9bd21da3885f980f313733fdc525d1;p=gcc.git gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT for VLA helper variables on target data even if... * 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. gcc/c/ * 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. gcc/cp/ * 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. libgomp/ * 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. From-SVN: r274206 --- diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 45050ed2f7b..00d8030d043 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,16 @@ +2019-08-08 Jakub Jelinek + + * 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 * config/aarch64/constraints.md (Z): Handle floating-point zeros too. diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 18c1ce0b6df..4944696a1d6 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,9 @@ +2019-08-08 Jakub Jelinek + + * 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 * c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause. diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index bda90fdd273..1d865aa7719 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13680,7 +13680,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) /* 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) @@ -14072,13 +14073,19 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) 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; } diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index bae5147fca9..f937bf5eb32 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,9 @@ +2019-08-08 Jakub Jelinek + + * 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 PR c++/81429 - wrong parsing of constructor with C++11 attribute. diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 8fe632f2239..9fd759c22c9 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -6148,7 +6148,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) /* 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) @@ -6406,13 +6407,19 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) 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; } diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 978df86f72a..daa0b71c191 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -6932,8 +6932,10 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags) 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; @@ -9016,6 +9018,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, 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; @@ -12404,8 +12409,27 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) 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)); diff --git a/gcc/omp-low.c b/gcc/omp-low.c index ca8edde6cdb..444610b0096 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1243,9 +1243,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) 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) { @@ -11857,11 +11857,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) 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)) @@ -12032,7 +12037,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) 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)) diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 01440f7a252..5356d005ff7 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,12 @@ +2019-08-08 Jakub Jelinek + + * 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 * testsuite/libgomp.c/target-18.c (struct S): New type. diff --git a/libgomp/target.c b/libgomp/target.c index 9416401806f..1c9ca68ba10 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -580,20 +580,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, } 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) @@ -791,9 +783,26 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, 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]; diff --git a/libgomp/testsuite/libgomp.c++/target-22.C b/libgomp/testsuite/libgomp.c++/target-22.C new file mode 100644 index 00000000000..9d9dea01c75 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-22.C @@ -0,0 +1,99 @@ +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); +} diff --git a/libgomp/testsuite/libgomp.c/target-37.c b/libgomp/testsuite/libgomp.c/target-37.c new file mode 100644 index 00000000000..b3cc6a25215 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-37.c @@ -0,0 +1,71 @@ +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; +}