| Message ID | 639a56ef-eeed-eb38-8a19-f5cf8d082973@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | [1/3,OpenMP] Target mapping changes for OpenMP 5.0, front-end parts | expand |
Ping this patch set. Thanks, Chung-Lin On 2020/9/1 9:16 PM, Chung-Lin Tang wrote: > Hi Jakub, > this patch set implements parts of the target mapping changes introduced > in OpenMP 5.0, mainly the attachment requirements for pointer-based > list items, and the clause ordering. > > The first patch here are the C/C++ front-end changes. > > The entire set of changes has been tested for without regressions for > the compiler and libgomp. Hope this is ready to commit to master. > > Thanks, > Chung-Lin > > gcc/c-family/ > * c-common.h (c_omp_adjust_clauses): New declaration. > * c-omp.c (c_omp_adjust_clauses): New function. > > gcc/c/ > * c-parser.c (c_parser_omp_target_data): Add use of > new c_omp_adjust_clauses function. Add GOMP_MAP_ATTACH_DETACH as > handled map clause kind. > (c_parser_omp_target_enter_data): Likewise. > (c_parser_omp_target_exit_data): Likewise. > (c_parser_omp_target): Likewise. > * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to > use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. > (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and > same struct field access to co-exist on OpenMP construct. > > gcc/cp/ > * parser.c (cp_parser_omp_target_data): Add use of > new c_omp_adjust_clauses function. Add GOMP_MAP_ATTACH_DETACH as > handled map clause kind. > (cp_parser_omp_target_enter_data): Likewise. > (cp_parser_omp_target_exit_data): Likewise. > (cp_parser_omp_target): Likewise. > * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to > use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix > interaction between reference case and attach/detach. > (finish_omp_clauses): Adjust bitmap checks to allow struct decl and > same struct field access to co-exist on OpenMP construct.
On Tue, Sep 01, 2020 at 09:16:23PM +0800, Chung-Lin Tang wrote: > this patch set implements parts of the target mapping changes introduced > in OpenMP 5.0, mainly the attachment requirements for pointer-based > list items, and the clause ordering. > > The first patch here are the C/C++ front-end changes. Do you think you could mention in detail which exact target mapping changes in the spec is the patchset attempting to implement? 5.0 unfortunately contains many target mapping changes and this patchset can't implement them all and it would be easier to see the list of rules (e.g. from openmp-diff-full-4.5-5.0.pdf, if you don't have that one, I can send it to you), rather than trying to guess them from the patchset. Thanks. > gcc/c-family/ > * c-common.h (c_omp_adjust_clauses): New declaration. > * c-omp.c (c_omp_adjust_clauses): New function. This function name is too broad, it should have target in it as it is for processing target* construct clauses only. Jakub
On 2020/9/29 6:16 PM, Jakub Jelinek wrote: > On Tue, Sep 01, 2020 at 09:16:23PM +0800, Chung-Lin Tang wrote: >> this patch set implements parts of the target mapping changes introduced >> in OpenMP 5.0, mainly the attachment requirements for pointer-based >> list items, and the clause ordering. >> >> The first patch here are the C/C++ front-end changes. > > Do you think you could mention in detail which exact target mapping changes > in the spec is the patchset attempting to implement? > 5.0 unfortunately contains many target mapping changes and this patchset > can't implement them all and it would be easier to see the list of rules > (e.g. from openmp-diff-full-4.5-5.0.pdf, if you don't have that one, I can > send it to you), rather than trying to guess them from the patchset. > > Thanks. Hi Jakub, the main implemented features are the clause ordering rules: "For a given construct, the effect of a map clause with the to, from, or tofrom map-type is ordered before the effect of a map clause with the alloc, release, or delete map-type." "If item1 is a list item in a map clause, and item2 is another list item in a map clause on the same construct that has a base pointer that is, or is part of, item1, then: * If the map clause(s) appear on a target, target data, or target enter data construct, then on entry to the corresponding region the effect of the map clause on item1 is ordered to occur before the effect of the map clause on item2. * If the map clause(s) appear on a target, target data, or target exit data construct then on exit from the corresponding region the effect of the map clause on item2 is ordered to occur before the effect of the map clause on item1." and the base-pointer attachment behavior: "If a list item in a map clause has a base pointer, and a pointer variable is present in the device data environment that corresponds to the base pointer when the effect of the map clause occurs, then if the corresponding pointer or the corresponding list item is created in the device data environment on entry to the construct, then: ... 2. The corresponding pointer variable becomes an attached pointer for the corresponding list item." (these passages are all in the "2.19.7.1 map Clause" section of the 5.0 spec, all are new as also verified from the diff PDFs you sent us) Also, because of the these new features, having multiple maps of the same variables now have meaning in OpenMP, so changes in the C/C++ frontends to relax the no-duplicate rules are also included. >> gcc/c-family/ >> * c-common.h (c_omp_adjust_clauses): New declaration. >> * c-omp.c (c_omp_adjust_clauses): New function. > > This function name is too broad, it should have target in it as it is > for processing target* construct clauses only. > > Jakub Sure, I'll update this naming in a later version. Thanks, Chung-Lin
On Tue, Sep 01, 2020 at 09:16:23PM +0800, Chung-Lin Tang wrote: > this patch set implements parts of the target mapping changes introduced > in OpenMP 5.0, mainly the attachment requirements for pointer-based > list items, and the clause ordering. > > The first patch here are the C/C++ front-end changes. > > The entire set of changes has been tested for without regressions for > the compiler and libgomp. Hope this is ready to commit to master. Sorry for the delay in patch review and thanks for the standard citations, that really helps. > gcc/c-family/ > * c-common.h (c_omp_adjust_clauses): New declaration. > * c-omp.c (c_omp_adjust_clauses): New function. Besides the naming, I wonder why is it done in a separate function and so early, can't what the function does be done either in {,c_}finish_omp_clauses (provided we'd pass separate ORT_OMP vs. ORT_OMP_TARGET to it to determine if it is target region vs. anything else), or perhaps even better during gimplification (gimplify_scan_omp_clauses)? > gcc/c/ > * c-parser.c (c_parser_omp_target_data): Add use of > new c_omp_adjust_clauses function. Add GOMP_MAP_ATTACH_DETACH as > handled map clause kind. > (c_parser_omp_target_enter_data): Likewise. > (c_parser_omp_target_exit_data): Likewise. > (c_parser_omp_target): Likewise. > * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to > use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. > (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and > same struct field access to co-exist on OpenMP construct. > > gcc/cp/ > * parser.c (cp_parser_omp_target_data): Add use of > new c_omp_adjust_clauses function. Add GOMP_MAP_ATTACH_DETACH as > handled map clause kind. > (cp_parser_omp_target_enter_data): Likewise. > (cp_parser_omp_target_exit_data): Likewise. > (cp_parser_omp_target): Likewise. > * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to > use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix > interaction between reference case and attach/detach. > (finish_omp_clauses): Adjust bitmap checks to allow struct decl and > same struct field access to co-exist on OpenMP construct. The changelog has some 8 space indented lines. > + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) > + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP > + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER > + && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) != ARRAY_TYPE) > + { > + tree ptr = OMP_CLAUSE_DECL (c); > + bool ptr_mapped = false; > + if (is_target) > + { > + for (tree m = clauses; m; m = OMP_CLAUSE_CHAIN (m)) Isn't this O(n^2) in number of clauses? I mean, e.g. for the equality comparisons (but see below) it could be dealt with e.g. using some bitmap with DECL_UIDs. > + if (OMP_CLAUSE_CODE (m) == OMP_CLAUSE_MAP > + && OMP_CLAUSE_DECL (m) == ptr Does it really need to be equality? I mean it will be for map(tofrom:ptr) map(tofrom:ptr[:32]) but what about e.g. map(tofrom:structx) map(tofrom:structx.ptr[:32]) ? It is true that likely we don't parse this yet though. > + && (OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALLOC > + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TO > + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_FROM > + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TOFROM)) What about the always modified mapping kinds? > + { > + ptr_mapped = true; > + break; > + } > + > + if (!ptr_mapped > + && DECL_P (ptr) > + && is_global_var (ptr) > + && lookup_attribute ("omp declare target", > + DECL_ATTRIBUTES (ptr))) > + ptr_mapped = true; > + } > + > + /* If the pointer variable was mapped, or if this is not an offloaded > + target region, adjust the map kind to attach/detach. */ > + if (ptr_mapped || !is_target) > + { > + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH); > + c_common_mark_addressable_vec (ptr); Though perhaps this is argument why it needs to be done in the FEs and not during gimplification, because it is hard to mark something addressable at that point. > --- a/gcc/c/c-typeck.c > +++ b/gcc/c/c-typeck.c > @@ -13580,16 +13580,17 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) > break; > } > tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); > if (ort != C_ORT_OMP && ort != C_ORT_ACC) > OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); > else if (TREE_CODE (t) == COMPONENT_REF) > { > - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH > - : GOMP_MAP_ALWAYS_POINTER; > + gomp_map_kind k > + = ((ort == C_ORT_ACC || ort == C_ORT_OMP) > + ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER); So what kind of C_ORT_* would be left after this change? C_ORT_*DECLARE_SIMD shouldn't have any kind of array sections in it. So maybe just > OMP_CLAUSE_SET_MAP_KIND (c2, k); OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); ? > if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) > { > - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) > + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) > + || bitmap_bit_p (&map_head, DECL_UID (t))) > break; Shall this change apply to OpenACC too? > } > } > if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) > { > error_at (OMP_CLAUSE_LOCATION (c), > "%qE is not a variable in %qs clause", t, > @@ -14751,29 +14753,36 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) > error_at (OMP_CLAUSE_LOCATION (c), > "%qD appears both in data and map clauses", t); > remove = true; > } > else > bitmap_set_bit (&generic_head, DECL_UID (t)); > } > - else if (bitmap_bit_p (&map_head, DECL_UID (t))) > + else if (bitmap_bit_p (&map_head, DECL_UID (t)) > + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) Ditto. Otherwise, what shall this diagnose now that the restriction that the same list item may not appear in multiple clauses is gone. > { > if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) > error_at (OMP_CLAUSE_LOCATION (c), > "%qD appears more than once in motion clauses", t); > else if (ort == C_ORT_ACC) > error_at (OMP_CLAUSE_LOCATION (c), > "%qD appears more than once in data clauses", t); > else > error_at (OMP_CLAUSE_LOCATION (c), > "%qD appears more than once in map clauses", t); > remove = true; So what is this supposed to diagnose now that the restriction that C++ ditto. Jakub
Hi Jakub, Thank you for the review. On 2020/10/13 9:01 PM, Jakub Jelinek wrote: >> gcc/c-family/ >> * c-common.h (c_omp_adjust_clauses): New declaration. >> * c-omp.c (c_omp_adjust_clauses): New function. > Besides the naming, I wonder why is it done in a separate function and so > early, can't what the function does be done either in > {,c_}finish_omp_clauses (provided we'd pass separate ORT_OMP vs. > ORT_OMP_TARGET to it to determine if it is target region vs. anything else), > or perhaps even better during gimplification (gimplify_scan_omp_clauses)? I figured that differentiating with something like "C_ORT_OMP_TARGET" could be more error prone to adjust changes related to C_ORT_OMP across the code, plus this has the added benefit of sharing a single place of handling logic across C/C++. You're right about the need for early addressable-marking. Learned that the hard way, one of my prior attempts tried to place this code somewhere in gimplify, didn't work. >> gcc/cp/ >> * parser.c (cp_parser_omp_target_data): Add use of >> new c_omp_adjust_clauses function. Add GOMP_MAP_ATTACH_DETACH as >> handled map clause kind. >> (cp_parser_omp_target_enter_data): Likewise. >> (cp_parser_omp_target_exit_data): Likewise. >> (cp_parser_omp_target): Likewise. >> * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to >> use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix >> interaction between reference case and attach/detach. >> (finish_omp_clauses): Adjust bitmap checks to allow struct decl and >> same struct field access to co-exist on OpenMP construct. > The changelog has some 8 space indented lines. I'll take care of that in the final git push. >> + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) >> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP >> + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER >> + && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) != ARRAY_TYPE) >> + { >> + tree ptr = OMP_CLAUSE_DECL (c); >> + bool ptr_mapped = false; >> + if (is_target) >> + { >> + for (tree m = clauses; m; m = OMP_CLAUSE_CHAIN (m)) > Isn't this O(n^2) in number of clauses? I mean, e.g. for the equality > comparisons (but see below) it could be dealt with e.g. using some bitmap > with DECL_UIDs. At this stage, we really don't assume any ordering of the clauses, nor try to modify its ordering yet, so the base-pointer map (if it exists) could be any where in the list (building some "visited set" isn't really suitable here). I don't think this is really that much an issue of concern though. >> + if (OMP_CLAUSE_CODE (m) == OMP_CLAUSE_MAP >> + && OMP_CLAUSE_DECL (m) == ptr > Does it really need to be equality? I mean it will be for > map(tofrom:ptr) map(tofrom:ptr[:32]) > but what about e.g. > map(tofrom:structx) map(tofrom:structx.ptr[:32]) > ? It is true that likely we don't parse this yet though. The code for COMPONENT_REF based expressions are actually handled quite differently in gimplify_scan_omp_clauses. Not completely sure there's nothing to handle for the code in this patch set, but will have to discover such testcases later. >> + && (OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALLOC >> + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TO >> + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_FROM >> + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TOFROM)) > What about the always modified mapping kinds? Took care of that. >> + { >> + ptr_mapped = true; >> + break; >> + } >> + >> + if (!ptr_mapped >> + && DECL_P (ptr) >> + && is_global_var (ptr) >> + && lookup_attribute ("omp declare target", >> + DECL_ATTRIBUTES (ptr))) >> + ptr_mapped = true; >> + } >> + >> + /* If the pointer variable was mapped, or if this is not an offloaded >> + target region, adjust the map kind to attach/detach. */ >> + if (ptr_mapped || !is_target) >> + { >> + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH); >> + c_common_mark_addressable_vec (ptr); > Though perhaps this is argument why it needs to be done in the FEs and not > during gimplification, because it is hard to mark something addressable at > that point. Discussed above. >> --- a/gcc/c/c-typeck.c >> +++ b/gcc/c/c-typeck.c >> @@ -13580,16 +13580,17 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) >> break; >> } >> tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); >> if (ort != C_ORT_OMP && ort != C_ORT_ACC) >> OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); >> else if (TREE_CODE (t) == COMPONENT_REF) >> { >> - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH >> - : GOMP_MAP_ALWAYS_POINTER; >> + gomp_map_kind k >> + = ((ort == C_ORT_ACC || ort == C_ORT_OMP) >> + ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER); > So what kind of C_ORT_* would be left after this change? > C_ORT_*DECLARE_SIMD shouldn't have any kind of array sections in it. > So maybe just > >> OMP_CLAUSE_SET_MAP_KIND (c2, k); > OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); > ? I have changed this code to just "OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);" >> if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) >> { >> - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) >> + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) >> + || bitmap_bit_p (&map_head, DECL_UID (t))) >> break; > Shall this change apply to OpenACC too? > >> } >> } >> if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) >> { >> error_at (OMP_CLAUSE_LOCATION (c), >> "%qE is not a variable in %qs clause", t, >> @@ -14751,29 +14753,36 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) >> error_at (OMP_CLAUSE_LOCATION (c), >> "%qD appears both in data and map clauses", t); >> remove = true; >> } >> else >> bitmap_set_bit (&generic_head, DECL_UID (t)); >> } >> - else if (bitmap_bit_p (&map_head, DECL_UID (t))) >> + else if (bitmap_bit_p (&map_head, DECL_UID (t)) >> + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) > Ditto. Otherwise, what shall this diagnose now that the restriction that > the same list item may not appear in multiple clauses is gone. Thanks for catching this, I've added "C_ORT_OMP" tests to these parts. Attached is the revised patch, dubbed "v2". Entire patch set re-tested with no regressions for gcc, g++, gfortran, and libgomp on x86_64-linux with nvptx offloading. Thanks, Chung-Lin diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index bb38e6c76a4..35ad417f9cc 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1221,6 +1221,7 @@ extern enum omp_clause_defaultmap_kind c_omp_predetermined_mapping (tree); extern tree c_omp_check_context_selector (location_t, tree); extern void c_omp_mark_declare_variant (location_t, tree, tree); extern const char *c_omp_map_clause_name (tree, bool); +extern void c_omp_adjust_clauses (tree, bool); /* Return next tree in the chain for chain_next walking of tree nodes. */ static inline tree diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index d7cff0f4cca..fd4995d67fb 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -2579,3 +2579,50 @@ c_omp_map_clause_name (tree clause, bool oacc) } return omp_clause_code_name[OMP_CLAUSE_CODE (clause)]; } + +/* Adjust map clauses after normal clause parsing, mainly to turn specific + base-pointer map cases into attach/detach and mark them addressable. */ +void +c_omp_adjust_clauses (tree clauses, bool is_target) +{ + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) != ARRAY_TYPE) + { + tree ptr = OMP_CLAUSE_DECL (c); + bool ptr_mapped = false; + if (is_target) + { + for (tree m = clauses; m; m = OMP_CLAUSE_CHAIN (m)) + if (OMP_CLAUSE_CODE (m) == OMP_CLAUSE_MAP + && OMP_CLAUSE_DECL (m) == ptr + && (OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALLOC + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TO + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_FROM + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TOFROM + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_TO + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_FROM + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_TOFROM)) + { + ptr_mapped = true; + break; + } + + if (!ptr_mapped + && DECL_P (ptr) + && is_global_var (ptr) + && lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (ptr))) + ptr_mapped = true; + } + + /* If the pointer variable was mapped, or if this is not an offloaded + target region, adjust the map kind to attach/detach. */ + if (ptr_mapped || !is_target) + { + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH); + c_common_mark_addressable_vec (ptr); + } + } +} diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index b6a7ef4c92b..24fd6afb9a3 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -19470,6 +19470,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data"); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -19487,6 +19488,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -19610,6 +19612,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser, tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, "#pragma omp target enter data"); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -19623,6 +19626,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser, break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -19694,7 +19698,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data"); - + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -19709,6 +19713,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -19918,6 +19923,8 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p) OMP_TARGET_CLAUSES (stmt) = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, "#pragma omp target"); + c_omp_adjust_clauses (OMP_TARGET_CLAUSES (stmt), true); + pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (); block = c_begin_compound_stmt (true); @@ -19942,6 +19949,7 @@ check_clauses: case GOMP_MAP_ALLOC: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 459090e227d..30c48e3a205 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13501,11 +13501,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if (ort != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) - { - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, k); - } + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER @@ -14604,7 +14600,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) + || (ort == C_ORT_OMP + && bitmap_bit_p (&map_head, DECL_UID (t)))) break; } } @@ -14673,7 +14671,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else bitmap_set_bit (&generic_head, DECL_UID (t)); } - else if (bitmap_bit_p (&map_head, DECL_UID (t))) + else if (bitmap_bit_p (&map_head, DECL_UID (t)) + && (ort != C_ORT_OMP + || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), @@ -14687,7 +14687,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + && ort == C_ORT_ACC) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data clauses", t); + remove = true; + } + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 7ec7d42773c..33cc5069ba0 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40510,6 +40510,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data", pragma_tok); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -40528,6 +40529,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -40611,6 +40613,7 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok, tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, "#pragma omp target enter data", pragma_tok); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -40625,6 +40628,7 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok, case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -40699,6 +40703,7 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data", pragma_tok); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -40714,6 +40719,7 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -40962,6 +40968,8 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, OMP_TARGET_CLAUSES (stmt) = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, "#pragma omp target", pragma_tok); + c_omp_adjust_clauses (OMP_TARGET_CLAUSES (stmt), true); + pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (true); OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p); @@ -40985,6 +40993,7 @@ check_clauses: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 1e42cd799c2..a8dc769db34 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5382,11 +5382,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) - { - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, k); - } + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { @@ -5424,8 +5420,12 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2)); OMP_CLAUSE_DECL (c3) = ptr; - if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER) - OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr); + if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ATTACH_DETACH) + { + OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + } else OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr); OMP_CLAUSE_SIZE (c3) = size_zero_node; @@ -7411,7 +7411,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); OMP_CLAUSE_DECL (c) = t; } - if (ort == C_ORT_ACC + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) && TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF) t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); @@ -7457,7 +7457,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) + || (ort == C_ORT_OMP + && bitmap_bit_p (&map_head, DECL_UID (t)))) goto handle_map_references; } } @@ -7551,13 +7553,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_set_bit (&generic_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t)) - && (ort != C_ORT_ACC - || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in motion clauses", t); - if (ort == C_ORT_ACC) + else if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else @@ -7566,7 +7567,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + && ort == C_ORT_ACC) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data clauses", t); + remove = true; + } + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), @@ -7602,17 +7609,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE) && (OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_ALWAYS_POINTER)) + != GOMP_MAP_ALWAYS_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_ATTACH_DETACH)) { tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if (TREE_CODE (t) == COMPONENT_REF) - { - gomp_map_kind k - = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, k); - } + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE);
On Wed, Oct 28, 2020 at 06:31:22PM +0800, Chung-Lin Tang wrote: > > > + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) > > > + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP > > > + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER > > > + && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) != ARRAY_TYPE) > > > + { > > > + tree ptr = OMP_CLAUSE_DECL (c); > > > + bool ptr_mapped = false; > > > + if (is_target) > > > + { > > > + for (tree m = clauses; m; m = OMP_CLAUSE_CHAIN (m)) > > Isn't this O(n^2) in number of clauses? I mean, e.g. for the equality > > comparisons (but see below) it could be dealt with e.g. using some bitmap > > with DECL_UIDs. > > At this stage, we really don't assume any ordering of the clauses, nor try to > modify its ordering yet, so the base-pointer map (if it exists) could be any > where in the list (building some "visited set" isn't really suitable here). > I don't think this is really that much an issue of concern though. Many functions try hard to avoid O(n^2) issues, see e.g. all the bitmap handling in *finish_omp_clauses etc. One can have tens of thousands of clauses and then the quadraticness will hit hard. This does a mere OMP_CLAUSE_DECL (c) == ptr comparison, so it is only about the decls and decls can be very easily handled through DECL_UID (bitmaps, hash sets/maps/tables). > +extern void c_omp_adjust_clauses (tree, bool); So, can you please rename the function to either c_omp_adjust_target_clauses or c_omp_adjust_mapping_clauses or c_omp_adjust_map_clauses? > --- a/gcc/c-family/c-omp.c > +++ b/gcc/c-family/c-omp.c > @@ -2579,3 +2579,50 @@ c_omp_map_clause_name (tree clause, bool oacc) > } > return omp_clause_code_name[OMP_CLAUSE_CODE (clause)]; > } > + > +/* Adjust map clauses after normal clause parsing, mainly to turn specific > + base-pointer map cases into attach/detach and mark them addressable. */ > +void > +c_omp_adjust_clauses (tree clauses, bool is_target) > +{ > + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) > + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP > + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER If this is only meant to handle decls, perhaps there should be && DECL_P (OMP_CLAUSE_DECL (c)) ? > + && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) != ARRAY_TYPE) > + { > + tree ptr = OMP_CLAUSE_DECL (c); > + bool ptr_mapped = false; > + if (is_target) > + { > + for (tree m = clauses; m; m = OMP_CLAUSE_CHAIN (m)) > + if (OMP_CLAUSE_CODE (m) == OMP_CLAUSE_MAP > + && OMP_CLAUSE_DECL (m) == ptr > + && (OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALLOC > + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TO > + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_FROM > + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TOFROM > + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_TO > + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_FROM > + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_TOFROM)) > + { > + ptr_mapped = true; > + break; > + } What you could e.g. do is have this loop at the start of function, with && DECL_P (OMP_CLAUSE_DECL (m)) instead of the == ptr check, and perhaps && POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (m))) check and set a bit in a bitmap for each such decl, then in the GOMP_MAP_FIRSTPRIVATE_POINTER loop just check the bitmap. Or, keep it in the loop like it is above, but populate the bitmap lazily (upon seeing the first GOMP_MAP_FIRSTPRIVATE_POINTER) and for further ones just use it. Jakub
Hi Jakub, here is v3 of this patch set. On 2020/10/29 7:44 PM, Jakub Jelinek wrote: >> +extern void c_omp_adjust_clauses (tree, bool); > So, can you please rename the function to either > c_omp_adjust_target_clauses or c_omp_adjust_mapping_clauses or > c_omp_adjust_map_clauses? I've renamed it to 'c_omp_adjust_map_clauses'. >> --- a/gcc/c-family/c-omp.c >> +++ b/gcc/c-family/c-omp.c >> @@ -2579,3 +2579,50 @@ c_omp_map_clause_name (tree clause, bool oacc) >> } >> return omp_clause_code_name[OMP_CLAUSE_CODE (clause)]; >> } >> + >> +/* Adjust map clauses after normal clause parsing, mainly to turn specific >> + base-pointer map cases into attach/detach and mark them addressable. */ >> +void >> +c_omp_adjust_clauses (tree clauses, bool is_target) >> +{ >> + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) >> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP >> + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER > If this is only meant to handle decls, perhaps there should be > && DECL_P (OMP_CLAUSE_DECL (c)) > ? > >> + && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) != ARRAY_TYPE) >> + { >> + tree ptr = OMP_CLAUSE_DECL (c); >> + bool ptr_mapped = false; >> + if (is_target) >> + { >> + for (tree m = clauses; m; m = OMP_CLAUSE_CHAIN (m)) >> + if (OMP_CLAUSE_CODE (m) == OMP_CLAUSE_MAP >> + && OMP_CLAUSE_DECL (m) == ptr >> + && (OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALLOC >> + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TO >> + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_FROM >> + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TOFROM >> + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_TO >> + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_FROM >> + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_TOFROM)) >> + { >> + ptr_mapped = true; >> + break; >> + } > What you could e.g. do is have this loop at the start of function, with > && DECL_P (OMP_CLAUSE_DECL (m)) > instead of the == ptr check, and perhaps && POINTER_TYPE_P (TREE_TYPE > (OMP_CLAUSE_DECL (m))) check and set a bit in a bitmap for each such decl, > then in the GOMP_MAP_FIRSTPRIVATE_POINTER loop just check the bitmap. > Or, keep it in the loop like it is above, but populate the bitmap > lazily (upon seeing the first GOMP_MAP_FIRSTPRIVATE_POINTER) and for further > ones just use it. I re-wrote c_omp_adjust_map_clauses to address the complexity issues you mentioned, now it should be limited by a linear pass to collect and merge the firstprivate base pointer + existence of a mapping of it, using a hash_map. Patch set has been re-tested with no regressions for gcc, g++, gfortran, and libgomp. Thanks, Chung-Lin gcc/c-family/ * c-common.h (c_omp_adjust_map_clauses): New declaration. * c-omp.c (c_omp_adjust_map_clauses): New function. gcc/c/ * c-parser.c (c_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (c_parser_omp_target_enter_data): Likewise. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_target): Likewise. * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/cp/ * parser.c (cp_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (cp_parser_omp_target_enter_data): Likewise. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_target): Likewise. * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix interaction between reference case and attach/detach. (finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index bb38e6c76a4..3eb909a2946 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1221,6 +1221,7 @@ extern enum omp_clause_defaultmap_kind c_omp_predetermined_mapping (tree); extern tree c_omp_check_context_selector (location_t, tree); extern void c_omp_mark_declare_variant (location_t, tree, tree); extern const char *c_omp_map_clause_name (tree, bool); +extern void c_omp_adjust_map_clauses (tree, bool); /* Return next tree in the chain for chain_next walking of tree nodes. */ static inline tree diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index d7cff0f4cca..275c6afabe1 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -2579,3 +2579,92 @@ c_omp_map_clause_name (tree clause, bool oacc) } return omp_clause_code_name[OMP_CLAUSE_CODE (clause)]; } + +/* Adjust map clauses after normal clause parsing, mainly to turn specific + base-pointer map cases into attach/detach and mark them addressable. */ +void +c_omp_adjust_map_clauses (tree clauses, bool is_target) +{ + if (!is_target) + { + /* If this is not a target construct, just turn firstprivate pointers + into attach/detach, the runtime will check and do the rest. */ + + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + && DECL_P (OMP_CLAUSE_DECL (c)) + && POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (c)))) + { + tree ptr = OMP_CLAUSE_DECL (c); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH); + c_common_mark_addressable_vec (ptr); + } + return; + } + + struct map_clause + { + tree clause; + bool firstprivate_ptr_p; + bool decl_mapped; + bool omp_declare_target; + map_clause (void) : clause (NULL_TREE), firstprivate_ptr_p (false), + decl_mapped (false), omp_declare_target (false) { } + }; + + hash_map<tree, map_clause> maps; + + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && DECL_P (OMP_CLAUSE_DECL (c))) + { + /* If this is for a target construct, the firstprivate pointer + is changed to attach/detach if either is true: + (1) the base-pointer is mapped in this same construct, or + (2) the base-pointer is a variable place on the device by + "declare target" directives. + + Here we iterate through all map clauses collecting these cases, + and merge them with a hash_map to process below. */ + + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + && POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (c)))) + { + tree ptr = OMP_CLAUSE_DECL (c); + map_clause &mc = maps.get_or_insert (ptr); + if (mc.clause == NULL_TREE) + mc.clause = c; + mc.firstprivate_ptr_p = true; + + if (is_global_var (ptr) + && lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (ptr))) + mc.omp_declare_target = true; + } + else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TOFROM + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_FROM + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM) + { + map_clause &mc = maps.get_or_insert (OMP_CLAUSE_DECL (c)); + mc.decl_mapped = true; + } + } + + for (hash_map<tree, map_clause>::iterator i = maps.begin (); + i != maps.end (); ++i) + { + map_clause &mc = (*i).second; + + if (mc.firstprivate_ptr_p + && (mc.decl_mapped || mc.omp_declare_target)) + { + OMP_CLAUSE_SET_MAP_KIND (mc.clause, GOMP_MAP_ATTACH_DETACH); + c_common_mark_addressable_vec (OMP_CLAUSE_DECL (mc.clause)); + } + } +} diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index b6a7ef4c92b..ab7b0bbc29f 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -19470,6 +19470,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data"); + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -19487,6 +19488,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -19610,6 +19612,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser, tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, "#pragma omp target enter data"); + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -19623,6 +19626,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser, break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -19694,7 +19698,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data"); - + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -19709,6 +19713,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -19918,6 +19923,8 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p) OMP_TARGET_CLAUSES (stmt) = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, "#pragma omp target"); + c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); + pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (); block = c_begin_compound_stmt (true); @@ -19942,6 +19949,7 @@ check_clauses: case GOMP_MAP_ALLOC: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 459090e227d..ada6662fca7 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13501,11 +13501,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if (ort != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) - { - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, k); - } + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER @@ -14604,7 +14600,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) + || (ort == C_ORT_OMP + && bitmap_bit_p (&map_head, DECL_UID (t)))) break; } } @@ -14673,7 +14671,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else bitmap_set_bit (&generic_head, DECL_UID (t)); } - else if (bitmap_bit_p (&map_head, DECL_UID (t))) + else if (bitmap_bit_p (&map_head, DECL_UID (t)) + && (ort != C_ORT_OMP + || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), @@ -14687,7 +14687,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + && ort == C_ORT_ACC) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data clauses", t); + remove = true; + } + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 7ec7d42773c..8527a7d0478 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40510,6 +40510,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data", pragma_tok); + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -40528,6 +40529,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -40611,6 +40613,7 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok, tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, "#pragma omp target enter data", pragma_tok); + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -40625,6 +40628,7 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok, case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -40699,6 +40703,7 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data", pragma_tok); + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -40714,6 +40719,7 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -40962,6 +40968,8 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, OMP_TARGET_CLAUSES (stmt) = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, "#pragma omp target", pragma_tok); + c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); + pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (true); OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p); @@ -40985,6 +40993,7 @@ check_clauses: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 1e42cd799c2..a8dc769db34 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5382,11 +5382,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) - { - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, k); - } + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { @@ -5424,8 +5420,12 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2)); OMP_CLAUSE_DECL (c3) = ptr; - if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER) - OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr); + if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ATTACH_DETACH) + { + OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + } else OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr); OMP_CLAUSE_SIZE (c3) = size_zero_node; @@ -7411,7 +7411,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); OMP_CLAUSE_DECL (c) = t; } - if (ort == C_ORT_ACC + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) && TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF) t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); @@ -7457,7 +7457,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) + || (ort == C_ORT_OMP + && bitmap_bit_p (&map_head, DECL_UID (t)))) goto handle_map_references; } } @@ -7551,13 +7553,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_set_bit (&generic_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t)) - && (ort != C_ORT_ACC - || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in motion clauses", t); - if (ort == C_ORT_ACC) + else if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else @@ -7566,7 +7567,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + && ort == C_ORT_ACC) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data clauses", t); + remove = true; + } + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), @@ -7602,17 +7609,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE) && (OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_ALWAYS_POINTER)) + != GOMP_MAP_ALWAYS_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_ATTACH_DETACH)) { tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if (TREE_CODE (t) == COMPONENT_REF) - { - gomp_map_kind k - = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, k); - } + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE);
On Wed, Nov 04, 2020 at 02:02:25AM +0800, Chung-Lin Tang wrote: > gcc/c-family/ > * c-common.h (c_omp_adjust_map_clauses): New declaration. > * c-omp.c (c_omp_adjust_map_clauses): New function. > > gcc/c/ > * c-parser.c (c_parser_omp_target_data): Add use of > new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as > handled map clause kind. > (c_parser_omp_target_enter_data): Likewise. > (c_parser_omp_target_exit_data): Likewise. > (c_parser_omp_target): Likewise. > * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to > use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. > (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and > same struct field access to co-exist on OpenMP construct. > > gcc/cp/ > * parser.c (cp_parser_omp_target_data): Add use of > new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as > handled map clause kind. > (cp_parser_omp_target_enter_data): Likewise. > (cp_parser_omp_target_exit_data): Likewise. > (cp_parser_omp_target): Likewise. > * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to > use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix > interaction between reference case and attach/detach. > (finish_omp_clauses): Adjust bitmap checks to allow struct decl and > same struct field access to co-exist on OpenMP construct. Ok, thanks. Jakub
diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index 4fc64bc4aa6..9ef85b401f0 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1208,14 +1208,15 @@ extern tree c_omp_declare_simd_clauses_to_numbers (tree, tree); extern void c_omp_declare_simd_clauses_to_decls (tree, tree); extern bool c_omp_predefined_variable (tree); extern enum omp_clause_default_kind c_omp_predetermined_sharing (tree); extern enum omp_clause_defaultmap_kind c_omp_predetermined_mapping (tree); extern tree c_omp_check_context_selector (location_t, tree); extern void c_omp_mark_declare_variant (location_t, tree, tree); extern const char *c_omp_map_clause_name (tree, bool); +extern void c_omp_adjust_clauses (tree, bool); /* Return next tree in the chain for chain_next walking of tree nodes. */ static inline tree c_tree_chain_next (tree t) { /* TREE_CHAIN of a type is TYPE_STUB_DECL, which is different kind of object, never a long chain of nodes. Prefer diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index d7cff0f4cca..596f33cebfb 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -2575,7 +2575,51 @@ c_omp_map_clause_name (tree clause, bool oacc) case GOMP_MAP_DEVICE_RESIDENT: return "device_resident"; case GOMP_MAP_LINK: return "link"; case GOMP_MAP_FORCE_DEVICEPTR: return "deviceptr"; default: break; } return omp_clause_code_name[OMP_CLAUSE_CODE (clause)]; } + +/* Adjust map clauses after normal clause parsing, mainly to turn specific + base-pointer map cases into attach/detach and mark them addressable. */ +void +c_omp_adjust_clauses (tree clauses, bool is_target) +{ + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) != ARRAY_TYPE) + { + tree ptr = OMP_CLAUSE_DECL (c); + bool ptr_mapped = false; + if (is_target) + { + for (tree m = clauses; m; m = OMP_CLAUSE_CHAIN (m)) + if (OMP_CLAUSE_CODE (m) == OMP_CLAUSE_MAP + && OMP_CLAUSE_DECL (m) == ptr + && (OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALLOC + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TO + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_FROM + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TOFROM)) + { + ptr_mapped = true; + break; + } + + if (!ptr_mapped + && DECL_P (ptr) + && is_global_var (ptr) + && lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (ptr))) + ptr_mapped = true; + } + + /* If the pointer variable was mapped, or if this is not an offloaded + target region, adjust the map kind to attach/detach. */ + if (ptr_mapped || !is_target) + { + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH); + c_common_mark_addressable_vec (ptr); + } + } +} diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index a8bc301ffad..92dfe3b6a4a 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -19452,14 +19452,15 @@ c_parser_omp_teams (location_t loc, c_parser *parser, static tree c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) { tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data"); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_TO: @@ -19469,14 +19470,15 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target data%> with map-type other " "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> " "on %<map%> clause"); @@ -19592,27 +19594,29 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser, c_parser_skip_to_pragma_eol (parser, false); return NULL_TREE; } tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, "#pragma omp target enter data"); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_TO: case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALLOC: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target enter data%> with map-type other " "than %<to%> or %<alloc%> on %<map%> clause"); *pc = OMP_CLAUSE_CHAIN (*pc); @@ -19676,29 +19680,30 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, c_parser_skip_to_pragma_eol (parser, false); return NULL_TREE; } tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data"); - + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_FROM: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target exit data%> with map-type other " "than %<from%>, %<release%> or %<delete%> on %<map%>" " clause"); @@ -19900,14 +19905,16 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p) stmt = make_node (OMP_TARGET); TREE_TYPE (stmt) = void_type_node; OMP_TARGET_CLAUSES (stmt) = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, "#pragma omp target"); + c_omp_adjust_clauses (OMP_TARGET_CLAUSES (stmt), true); + pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (); block = c_begin_compound_stmt (true); add_stmt (c_parser_omp_structured_block (parser, if_p)); OMP_TARGET_BODY (stmt) = c_end_compound_stmt (loc, block, true); SET_EXPR_LOCATION (stmt, loc); @@ -19924,14 +19931,15 @@ check_clauses: case GOMP_MAP_FROM: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target%> with map-type other " "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> " "on %<map%> clause"); *pc = OMP_CLAUSE_CHAIN (*pc); diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 0d639b60ea3..17ac2f566da 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13580,16 +13580,17 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) break; } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if (ort != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) { - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; + gomp_map_kind k + = ((ort == C_ORT_ACC || ort == C_ORT_OMP) + ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER); OMP_CLAUSE_SET_MAP_KIND (c2, k); } else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER && !c_mark_addressable (t)) return false; @@ -14682,15 +14683,16 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); } } if (remove) break; if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) + || bitmap_bit_p (&map_head, DECL_UID (t))) break; } } if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { error_at (OMP_CLAUSE_LOCATION (c), "%qE is not a variable in %qs clause", t, @@ -14751,29 +14753,36 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears both in data and map clauses", t); remove = true; } else bitmap_set_bit (&generic_head, DECL_UID (t)); } - else if (bitmap_bit_p (&map_head, DECL_UID (t))) + else if (bitmap_bit_p (&map_head, DECL_UID (t)) + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in motion clauses", t); else if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in map clauses", t); remove = true; } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + && ort == C_ORT_ACC) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data clauses", t); + remove = true; + } + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else error_at (OMP_CLAUSE_LOCATION (c), "%qD appears both in data and map clauses", t); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 7cc2dbed5fe..7773f9d4f79 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40449,14 +40449,15 @@ cp_parser_omp_teams (cp_parser *parser, cp_token *pragma_tok, static tree cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) { tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data", pragma_tok); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_TO: @@ -40467,14 +40468,15 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target data%> with map-type other " "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> " "on %<map%> clause"); @@ -40550,28 +40552,30 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok, cp_parser_skip_to_pragma_eol (parser, pragma_tok); return NULL_TREE; } tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, "#pragma omp target enter data", pragma_tok); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_TO: case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALLOC: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target enter data%> with map-type other " "than %<to%> or %<alloc%> on %<map%> clause"); *pc = OMP_CLAUSE_CHAIN (*pc); @@ -40638,14 +40642,15 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, cp_parser_skip_to_pragma_eol (parser, pragma_tok); return NULL_TREE; } tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data", pragma_tok); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_FROM: @@ -40653,14 +40658,15 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target exit data%> with map-type other " "than %<from%>, %<release%> or %<delete%> on %<map%>" " clause"); @@ -40901,14 +40907,16 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, stmt = make_node (OMP_TARGET); TREE_TYPE (stmt) = void_type_node; OMP_TARGET_CLAUSES (stmt) = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, "#pragma omp target", pragma_tok); + c_omp_adjust_clauses (OMP_TARGET_CLAUSES (stmt), true); + pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (true); OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p); SET_EXPR_LOCATION (stmt, pragma_tok->location); add_stmt (stmt); @@ -40924,14 +40932,15 @@ check_clauses: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target%> with map-type other " "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> " "on %<map%> clause"); *pc = OMP_CLAUSE_CHAIN (*pc); diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index b71ca0729a8..0f6b36f2dab 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5373,16 +5373,17 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) { - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; + gomp_map_kind k + = ((ort == C_ORT_ACC || ort == C_ORT_OMP) + ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER); OMP_CLAUSE_SET_MAP_KIND (c2, k); } else if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { t = TREE_OPERAND (t, 0); gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH @@ -5414,16 +5415,20 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) && TYPE_REF_P (TREE_TYPE (ptr)) && INDIRECT_TYPE_P (TREE_TYPE (TREE_TYPE (ptr)))) { tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2)); OMP_CLAUSE_DECL (c3) = ptr; - if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER) - OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr); + if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ATTACH_DETACH) + { + OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + } else 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; } } @@ -7400,15 +7405,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) OMP_CLAUSE_SIZE (c) = size_zero_node; if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { t = TREE_OPERAND (t, 0); OMP_CLAUSE_DECL (c) = t; } - if (ort == C_ORT_ACC + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) && TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF) t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); if (TREE_CODE (t) == COMPONENT_REF && ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP || ort == C_ORT_ACC) && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) @@ -7446,15 +7451,16 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } if (remove) break; if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) + || bitmap_bit_p (&map_head, DECL_UID (t))) goto handle_map_references; } } if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { if (processing_template_decl && TREE_CODE (t) != OVERLOAD) break; @@ -7540,30 +7546,35 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) "%qD appears both in data and map clauses", t); remove = true; } else bitmap_set_bit (&generic_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t)) - && (ort != C_ORT_ACC - || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in motion clauses", t); - if (ort == C_ORT_ACC) + else if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in map clauses", t); remove = true; } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + && ort == C_ORT_ACC) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data clauses", t); + remove = true; + } + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else error_at (OMP_CLAUSE_LOCATION (c), "%qD appears both in data and map clauses", t); @@ -7591,23 +7602,25 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (t))); } else if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE) && (OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_ALWAYS_POINTER)) + != GOMP_MAP_ALWAYS_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_ATTACH_DETACH)) { tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if (TREE_CODE (t) == COMPONENT_REF) { gomp_map_kind k - = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; + = ((ort == C_ORT_ACC || ort == C_ORT_OMP) + ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER); OMP_CLAUSE_SET_MAP_KIND (c2, k); } else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE); OMP_CLAUSE_DECL (c2) = t; OMP_CLAUSE_SIZE (c2) = size_zero_node;
Hi Jakub, this patch set implements parts of the target mapping changes introduced in OpenMP 5.0, mainly the attachment requirements for pointer-based list items, and the clause ordering. The first patch here are the C/C++ front-end changes. The entire set of changes has been tested for without regressions for the compiler and libgomp. Hope this is ready to commit to master. Thanks, Chung-Lin gcc/c-family/ * c-common.h (c_omp_adjust_clauses): New declaration. * c-omp.c (c_omp_adjust_clauses): New function. gcc/c/ * c-parser.c (c_parser_omp_target_data): Add use of new c_omp_adjust_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (c_parser_omp_target_enter_data): Likewise. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_target): Likewise. * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/cp/ * parser.c (cp_parser_omp_target_data): Add use of new c_omp_adjust_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (cp_parser_omp_target_enter_data): Likewise. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_target): Likewise. * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix interaction between reference case and attach/detach. (finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct.