Hi! OpenMP 5.0 also newly says: "If one or more of the use_device_ptr or use_device_addr clauses and one or more map clauses are present on the same construct, the address conversions of use_device_addr and use_device_ptr clauses will occur as if performed after all variables are mapped according to those map clauses." Before this rule, one had to use a separate target data or target enter data to map the variables first before it was possible to use use_device_* on another target data.
The patch allows such cases in the FEs, sorts the use_device_* clauses last in the gimplifier, ensures omp lowering doesn't ICE on that and finally in libgomp arranges that if the mapping clauses really precede all use_device_* clauses and something has not been already found during the map clause processing, we defer the use_device_* lookups until after the new variables are mapped too. Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk. 2019-08-08 Jakub Jelinek <ja...@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. 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. --- gcc/gimplify.c.jj 2019-08-07 09:24:35.646096085 +0200 +++ gcc/gimplify.c 2019-08-07 12:32:49.927990656 +0200 @@ -6932,8 +6932,10 @@ omp_add_variable (struct gimplify_omp_ct 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, 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, gi 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)); --- gcc/omp-low.c.jj 2019-08-07 09:24:35.647096069 +0200 +++ gcc/omp-low.c 2019-08-07 11:12:44.137779196 +0200 @@ -1243,9 +1243,9 @@ scan_sharing_clauses (tree clauses, omp_ 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 * 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 * 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)) --- gcc/c/c-typeck.c.jj 2019-08-07 09:24:36.094089357 +0200 +++ gcc/c/c-typeck.c 2019-08-07 13:14:21.034893880 +0200 @@ -13680,7 +13680,8 @@ c_finish_omp_clauses (tree clauses, enum /* 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 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; } --- gcc/cp/semantics.c.jj 2019-08-07 09:24:36.124088907 +0200 +++ gcc/cp/semantics.c 2019-08-07 13:35:38.249901249 +0200 @@ -6146,7 +6146,8 @@ finish_omp_clauses (tree clauses, enum c /* 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) @@ -6404,13 +6405,19 @@ finish_omp_clauses (tree clauses, enum c 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; } --- libgomp/target.c.jj 2019-06-10 19:37:21.981343123 +0200 +++ libgomp/target.c 2019-08-07 17:09:50.495141003 +0200 @@ -580,20 +580,12 @@ gomp_map_vars_internal (struct gomp_devi } 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_devi 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]; --- libgomp/testsuite/libgomp.c/target-37.c.jj 2019-08-07 11:19:56.875299065 +0200 +++ libgomp/testsuite/libgomp.c/target-37.c 2019-08-06 15:47:47.283529318 +0200 @@ -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; +} --- libgomp/testsuite/libgomp.c++/target-22.C.jj 2019-08-07 11:20:22.865909863 +0200 +++ libgomp/testsuite/libgomp.c++/target-22.C 2019-08-07 11:22:05.407374333 +0200 @@ -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); +} Jakub