From 398e3feb8a47aa00994903f0209c8f45c0b653c3 Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Wed, 7 Aug 2019 09:27:10 +0200 Subject: [PATCH] tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR OpenMP description. * tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR OpenMP description. Add OMP_CLAUSE_USE_DEVICE_ADDR clause. * tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries for OMP_CLAUSE_USE_DEVICE_ADDR clause. (walk_tree_1): Handle OMP_CLAUSE_USE_DEVICE_ADDR. * tree-pretty-print.c (dump_omp_clause): Likewise. * tree-nested.c (convert_nonlocal_omp_clauses, convert_local_omp_clauses): Likewise. * gimplify.c (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Likewise. * omp-low.c (scan_sharing_clauses, lower_omp_target): Likewise. Treat OMP_CLAUSE_USE_DEVICE_ADDR like OMP_CLAUSE_USE_DEVICE_PTR clause with array or reference to array types, no matter what type except for reference it has. gcc/c-family/ * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. Set PRAGMA_OACC_CLAUSE_USE_DEVICE equal to PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR instead of being a separate enumeration value. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause. (c_parser_omp_clause_use_device_addr): New function. (c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. (OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. (c_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no map or use_device_* clauses. * c-typeck.c (c_finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR in OpenMP, require pointer type rather than pointer or array type. Handle OMP_CLAUSE_USE_DEVICE_ADDR. gcc/cp/ * parser.c (cp_parser_omp_clause_name): Parse use_device_addr clause. (cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. (OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. (cp_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no map or use_device_* clauses. * semantics.c (finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR in OpenMP, require pointer or reference to pointer type rather than pointer or array or reference to pointer or array type. Handle OMP_CLAUSE_USE_DEVICE_ADDR. * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE_ADDR. gcc/testsuite/ * c-c++-common/gomp/target-data-1.c (foo): Use use_device_addr clause instead of use_device_ptr clause where required by OpenMP 5.0, add further tests for both use_device_ptr and use_device_addr clauses. libgomp/ * testsuite/libgomp.c/target-18.c (struct S): New type. (foo): Use use_device_addr clause instead of use_device_ptr clause where required by OpenMP 5.0, add further tests for both use_device_ptr and use_device_addr clauses. * testsuite/libgomp.c++/target-9.C (struct S): New type. (foo): Use use_device_addr clause instead of use_device_ptr clause where required by OpenMP 5.0, add further tests for both use_device_ptr and use_device_addr clauses. Add t and u arguments. (main): Adjust caller. From-SVN: r274159 --- gcc/ChangeLog | 17 ++++++++ gcc/c-family/ChangeLog | 7 +++ gcc/c-family/c-pragma.h | 5 ++- gcc/c/ChangeLog | 13 ++++++ gcc/c/c-parser.c | 25 +++++++++-- gcc/c/c-typeck.c | 28 +++++++++--- gcc/cp/ChangeLog | 14 ++++++ gcc/cp/parser.c | 16 +++++-- gcc/cp/pt.c | 2 + gcc/cp/semantics.c | 39 +++++++++++++---- gcc/gimplify.c | 4 +- gcc/omp-low.c | 31 +++++++++---- gcc/testsuite/ChangeLog | 6 +++ .../c-c++-common/gomp/target-data-1.c | 30 +++++++++++-- gcc/tree-core.h | 5 ++- gcc/tree-nested.c | 2 + gcc/tree-pretty-print.c | 3 ++ gcc/tree.c | 3 ++ libgomp/ChangeLog | 12 ++++++ libgomp/testsuite/libgomp.c++/target-9.C | 43 ++++++++++++++++--- libgomp/testsuite/libgomp.c/target-18.c | 30 +++++++++++-- 21 files changed, 290 insertions(+), 45 deletions(-) diff --git a/gcc/ChangeLog b/gcc/ChangeLog index b85a1dba105..23837ab3a9a 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,20 @@ +2019-08-07 Jakub Jelinek + + * tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR + OpenMP description. Add OMP_CLAUSE_USE_DEVICE_ADDR clause. + * tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries + for OMP_CLAUSE_USE_DEVICE_ADDR clause. + (walk_tree_1): Handle OMP_CLAUSE_USE_DEVICE_ADDR. + * tree-pretty-print.c (dump_omp_clause): Likewise. + * tree-nested.c (convert_nonlocal_omp_clauses, + convert_local_omp_clauses): Likewise. + * gimplify.c (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): + Likewise. + * omp-low.c (scan_sharing_clauses, lower_omp_target): Likewise. + Treat OMP_CLAUSE_USE_DEVICE_ADDR like OMP_CLAUSE_USE_DEVICE_PTR + clause with array or reference to array types, no matter what type + except for reference it has. + 2019-08-07 Kewen Lin * config/rs6000/vector.md (vrotr3): New define_expand. diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index 693a546347c..873b9ad5b7a 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,10 @@ +2019-08-07 Jakub Jelinek + + * c-pragma.h (enum pragma_omp_clause): Add + PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. Set PRAGMA_OACC_CLAUSE_USE_DEVICE + equal to PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR instead of being a separate + enumeration value. + 2019-08-05 Marek Polacek PR c++/91338 - Implement P1161R3: Deprecate a[b,c]. diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index e8a509f7073..803dc1b2ac6 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -137,6 +137,7 @@ enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_UNIFORM, PRAGMA_OMP_CLAUSE_UNTIED, PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, + PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR, /* Clauses for OpenACC. */ PRAGMA_OACC_CLAUSE_ASYNC, @@ -157,7 +158,6 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_SELF, PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_TILE, - PRAGMA_OACC_CLAUSE_USE_DEVICE, PRAGMA_OACC_CLAUSE_VECTOR, PRAGMA_OACC_CLAUSE_VECTOR_LENGTH, PRAGMA_OACC_CLAUSE_WAIT, @@ -171,7 +171,8 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF, PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE, PRAGMA_OACC_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION, - PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK + PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK, + PRAGMA_OACC_CLAUSE_USE_DEVICE = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR }; extern struct cpp_reader* parse_in; diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 18c4e413e2d..18c1ce0b6df 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,16 @@ +2019-08-07 Jakub Jelinek + + * c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause. + (c_parser_omp_clause_use_device_addr): New function. + (c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. + (OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. + (c_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR + like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no + map or use_device_* clauses. + * c-typeck.c (c_finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR + in OpenMP, require pointer type rather than pointer or array type. + Handle OMP_CLAUSE_USE_DEVICE_ADDR. + 2019-07-31 Jakub Jelinek PR c/91192 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index c7c13a6d9ae..9b5cfcf50d9 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -11866,6 +11866,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_UNTIED; else if (!strcmp ("use_device", p)) result = PRAGMA_OACC_CLAUSE_USE_DEVICE; + else if (!strcmp ("use_device_addr", p)) + result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR; else if (!strcmp ("use_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; break; @@ -13121,6 +13123,16 @@ c_parser_omp_clause_use_device_ptr (c_parser *parser, tree list) list); } +/* OpenMP 5.0: + use_device_addr ( variable-list ) */ + +static tree +c_parser_omp_clause_use_device_addr (c_parser *parser, tree list) +{ + return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE_ADDR, + list); +} + /* OpenMP 4.5: is_device_ptr ( variable-list ) */ @@ -15321,6 +15333,10 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_omp_clause_use_device_ptr (parser, clauses); c_name = "use_device_ptr"; break; + case PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR: + clauses = c_parser_omp_clause_use_device_addr (parser, clauses); + c_name = "use_device_addr"; + break; case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR: clauses = c_parser_omp_clause_is_device_ptr (parser, clauses); c_name = "is_device_ptr"; @@ -18288,7 +18304,8 @@ c_parser_omp_teams (location_t loc, c_parser *parser, ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR)) static tree c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) @@ -18323,7 +18340,8 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) *pc = OMP_CLAUSE_CHAIN (*pc); continue; } - else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR) + else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR + || OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_ADDR) map_seen = 3; pc = &OMP_CLAUSE_CHAIN (*pc); } @@ -18333,7 +18351,8 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) if (map_seen == 0) error_at (loc, "%<#pragma omp target data%> must contain at least " - "one % or % clause"); + "one %, % or % " + "clause"); return NULL_TREE; } diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 9a1a9106a40..bda90fdd273 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -14609,16 +14609,32 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_PTR: t = OMP_CLAUSE_DECL (c); - if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE - && TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE) + if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE) { - error_at (OMP_CLAUSE_LOCATION (c), - "%qs variable is neither a pointer nor an array", - omp_clause_code_name[OMP_CLAUSE_CODE (c)]); - remove = true; + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR + && ort == C_ORT_OMP) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qs variable is not a pointer", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + else if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qs variable is neither a pointer nor an array", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } } goto check_dup_generic; + case OMP_CLAUSE_USE_DEVICE_ADDR: + t = OMP_CLAUSE_DECL (c); + if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) + c_mark_addressable (t); + goto check_dup_generic; + case OMP_CLAUSE_NOWAIT: if (copyprivate_seen) { diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 9b0a34c98aa..c1e977b6926 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,17 @@ +2019-08-07 Jakub Jelinek + + * parser.c (cp_parser_omp_clause_name): Parse use_device_addr clause. + (cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. + (OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. + (cp_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR + like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no + map or use_device_* clauses. + * semantics.c (finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR + in OpenMP, require pointer or reference to pointer type rather than + pointer or array or reference to pointer or array type. Handle + OMP_CLAUSE_USE_DEVICE_ADDR. + * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE_ADDR. + 2019-08-06 Jason Merrill PR c++/91378 - ICE with noexcept and auto return type. diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 79da7b52eb9..4d07a6a3011 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -32648,6 +32648,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_UNTIED; else if (!strcmp ("use_device", p)) result = PRAGMA_OACC_CLAUSE_USE_DEVICE; + else if (!strcmp ("use_device_addr", p)) + result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR; else if (!strcmp ("use_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; break; @@ -35637,6 +35639,11 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses); c_name = "use_device_ptr"; break; + case PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE_ADDR, + clauses); + c_name = "use_device_addr"; + break; case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR: clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_IS_DEVICE_PTR, clauses); @@ -38715,7 +38722,8 @@ cp_parser_omp_teams (cp_parser *parser, cp_token *pragma_tok, ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR)) static tree cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) @@ -38751,7 +38759,8 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) *pc = OMP_CLAUSE_CHAIN (*pc); continue; } - else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR) + else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR + || OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_ADDR) map_seen = 3; pc = &OMP_CLAUSE_CHAIN (*pc); } @@ -38761,7 +38770,8 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) if (map_seen == 0) error_at (pragma_tok->location, "%<#pragma omp target data%> must contain at least " - "one % or % clause"); + "one %, % or % " + "clause"); return NULL_TREE; } diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index b71fbaad789..b1ad99d1481 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -16303,6 +16303,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, case OMP_CLAUSE_MAP: case OMP_CLAUSE_NONTEMPORAL: case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_INCLUSIVE: case OMP_CLAUSE_EXCLUSIVE: @@ -16427,6 +16428,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, case OMP_CLAUSE_IN_REDUCTION: case OMP_CLAUSE_TASK_REDUCTION: case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_INCLUSIVE: case OMP_CLAUSE_EXCLUSIVE: diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index fa6962454bf..77e7a6dced2 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -7524,20 +7524,41 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { tree type = TREE_TYPE (t); if (!TYPE_PTR_P (type) - && TREE_CODE (type) != ARRAY_TYPE - && (!TYPE_REF_P (type) - || (!TYPE_PTR_P (TREE_TYPE (type)) - && TREE_CODE (TREE_TYPE (type)) != ARRAY_TYPE))) + && (!TYPE_REF_P (type) || !TYPE_PTR_P (TREE_TYPE (type)))) { - error_at (OMP_CLAUSE_LOCATION (c), - "%qs variable is neither a pointer, nor an array " - "nor reference to pointer or array", - omp_clause_code_name[OMP_CLAUSE_CODE (c)]); - remove = true; + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR + && ort == C_ORT_OMP) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qs variable is neither a pointer " + "nor reference to pointer", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + else if (TREE_CODE (type) != ARRAY_TYPE + && (!TYPE_REF_P (type) + || TREE_CODE (TREE_TYPE (type)) != ARRAY_TYPE)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qs variable is neither a pointer, nor an " + "array nor reference to pointer or array", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } } } goto check_dup_generic; + case OMP_CLAUSE_USE_DEVICE_ADDR: + field_ok = true; + t = OMP_CLAUSE_DECL (c); + if (!processing_template_decl + && (VAR_P (t) || TREE_CODE (t) == PARM_DECL) + && !TYPE_REF_P (TREE_TYPE (t)) + && !cxx_mark_addressable (t)) + remove = true; + goto check_dup_generic; + case OMP_CLAUSE_NOWAIT: case OMP_CLAUSE_DEFAULT: case OMP_CLAUSE_UNTIED: diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 10b9b68d5ce..978df86f72a 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -9015,8 +9015,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, goto do_notice; case OMP_CLAUSE_USE_DEVICE_PTR: - flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT; - goto do_add; + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT; goto do_add; @@ -10264,6 +10263,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_CLAUSE_ORDER: case OMP_CLAUSE_BIND: case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 4a6ea0a1b71..ca8edde6cdb 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1238,8 +1238,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: decl = OMP_CLAUSE_DECL (c); - if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) + 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); else install_var_field (decl, false, 3, ctx); @@ -1635,6 +1638,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_ORDER: case OMP_CLAUSE_BIND: case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_NONTEMPORAL: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: @@ -11465,6 +11469,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: var = OMP_CLAUSE_DECL (c); map_cnt++; @@ -11481,7 +11486,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) SET_DECL_VALUE_EXPR (new_var, x); DECL_HAS_VALUE_EXPR_P (new_var) = 1; } - else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) + else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + && !omp_is_reference (var)) + || TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) { tree new_var = lookup_decl (var, ctx); tree type = build_pointer_type (TREE_TYPE (var)); @@ -11846,23 +11853,27 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: 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_USE_DEVICE_PTR) + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR) tkind = GOMP_MAP_USE_DEVICE_PTR; else tkind = GOMP_MAP_FIRSTPRIVATE_INT; type = TREE_TYPE (ovar); - if (TREE_CODE (type) == ARRAY_TYPE) + if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + && !omp_is_reference (ovar)) + || TREE_CODE (type) == ARRAY_TYPE) var = build_fold_addr_expr (var); else { if (omp_is_reference (ovar)) { type = TREE_TYPE (type); - if (TREE_CODE (type) != ARRAY_TYPE) + if (TREE_CODE (type) != ARRAY_TYPE + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR) var = build_simple_mem_ref (var); var = fold_convert (TREE_TYPE (x), var); } @@ -12017,9 +12028,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } break; case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: var = OMP_CLAUSE_DECL (c); - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR) + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR) x = build_sender_ref (var, ctx); else x = build_receiver_ref (var, false, ctx); @@ -12034,7 +12046,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq_add_stmt (&new_body, gimple_build_assign (new_var, x)); } - else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) + else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + && !omp_is_reference (var)) + || TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) { tree new_var = lookup_decl (var, ctx); new_var = DECL_VALUE_EXPR (new_var); @@ -12052,7 +12066,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (omp_is_reference (var)) { type = TREE_TYPE (type); - if (TREE_CODE (type) != ARRAY_TYPE) + if (TREE_CODE (type) != ARRAY_TYPE + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR) { tree v = create_tmp_var_raw (type, get_name (var)); gimple_add_tmp_var (v); diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index a372d7fbcf1..d524f72b4d2 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2019-08-07 Jakub Jelinek + + * c-c++-common/gomp/target-data-1.c (foo): Use use_device_addr clause + instead of use_device_ptr clause where required by OpenMP 5.0, add + further tests for both use_device_ptr and use_device_addr clauses. + 2019-08-07 Kewen Lin * gcc.target/powerpc/vec_rotate-1.c: New test. diff --git a/gcc/testsuite/c-c++-common/gomp/target-data-1.c b/gcc/testsuite/c-c++-common/gomp/target-data-1.c index 0d4975bcc03..7aa111a92a4 100644 --- a/gcc/testsuite/c-c++-common/gomp/target-data-1.c +++ b/gcc/testsuite/c-c++-common/gomp/target-data-1.c @@ -4,15 +4,39 @@ void foo (void) { int a[4] = { 1, 2, 3, 4 }; + int *p = &a[0]; + int x = 5; + #pragma omp target data map(to:p[:4]) + #pragma omp target data use_device_ptr(p) + #pragma omp target is_device_ptr(p) + { + p[0]++; + } #pragma omp target data map(to:a) - #pragma omp target data use_device_ptr(a) + #pragma omp target data use_device_addr(a) #pragma omp target is_device_ptr(a) { - a[0]++; + p[0]++; + } + #pragma omp target data map(to:x) + #pragma omp target data use_device_addr(x) + { + int *q = &x; + #pragma omp target is_device_ptr(q) + { + q[0]++; + } } #pragma omp target data /* { dg-error "must contain at least one" } */ a[0]++; + #pragma omp target data map(to:p) + #pragma omp target data use_device_ptr(p) use_device_ptr(p) /* { dg-error "appears more than once in data clauses" } */ + a[0]++; #pragma omp target data map(to:a) - #pragma omp target data use_device_ptr(a) use_device_ptr(a) /* { dg-error "appears more than once in data clauses" } */ + #pragma omp target data use_device_addr(a) use_device_addr(a) /* { dg-error "appears more than once in data clauses" } */ a[0]++; + #pragma omp target data map(to:a) + #pragma omp target data use_device_ptr(a) /* { dg-error "'use_device_ptr' variable is not a pointer" "" { target c } } */ + /* { dg-error "'use_device_ptr' variable is neither a pointer nor reference to pointer" "" { target c++ } .-1 } */ + a[0]++; /* { dg-error "must contain at least one" "" { target *-*-* } .-2 } */ } diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 47195626912..8216a606718 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -307,9 +307,12 @@ enum omp_clause_code { OMP_CLAUSE_MAP, /* OpenACC clause: use_device (variable-list). - OpenMP clause: use_device_ptr (variable-list). */ + OpenMP clause: use_device_ptr (ptr-list). */ OMP_CLAUSE_USE_DEVICE_PTR, + /* OpenMP clause: use_device_addr (variable-list). */ + OMP_CLAUSE_USE_DEVICE_ADDR, + /* OpenMP clause: is_device_ptr (variable-list). */ OMP_CLAUSE_IS_DEVICE_PTR, diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c index 5d189572e19..4df07edcf0a 100644 --- a/gcc/tree-nested.c +++ b/gcc/tree-nested.c @@ -1227,6 +1227,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_TO_DECLARE: case OMP_CLAUSE_LINK: case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: decl = OMP_CLAUSE_DECL (clause); @@ -1947,6 +1948,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_TO_DECLARE: case OMP_CLAUSE_LINK: case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: decl = OMP_CLAUSE_DECL (clause); diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 9bea132f7f0..8091a0af7f3 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -465,6 +465,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case OMP_CLAUSE_USE_DEVICE_PTR: name = "use_device_ptr"; goto print_remap; + case OMP_CLAUSE_USE_DEVICE_ADDR: + name = "use_device_addr"; + goto print_remap; case OMP_CLAUSE_IS_DEVICE_PTR: name = "is_device_ptr"; goto print_remap; diff --git a/gcc/tree.c b/gcc/tree.c index efa49e99d65..1f2aac4ffa5 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -299,6 +299,7 @@ unsigned const char omp_clause_num_ops[] = 2, /* OMP_CLAUSE_TO */ 2, /* OMP_CLAUSE_MAP */ 1, /* OMP_CLAUSE_USE_DEVICE_PTR */ + 1, /* OMP_CLAUSE_USE_DEVICE_ADDR */ 1, /* OMP_CLAUSE_IS_DEVICE_PTR */ 1, /* OMP_CLAUSE_INCLUSIVE */ 1, /* OMP_CLAUSE_EXCLUSIVE */ @@ -382,6 +383,7 @@ const char * const omp_clause_code_name[] = "to", "map", "use_device_ptr", + "use_device_addr", "is_device_ptr", "inclusive", "exclusive", @@ -12384,6 +12386,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE_TO_DECLARE: case OMP_CLAUSE_LINK: case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_INCLUSIVE: case OMP_CLAUSE_EXCLUSIVE: diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 0331c51c639..01440f7a252 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,15 @@ +2019-08-07 Jakub Jelinek + + * testsuite/libgomp.c/target-18.c (struct S): New type. + (foo): Use use_device_addr clause instead of use_device_ptr clause + where required by OpenMP 5.0, add further tests for both use_device_ptr + and use_device_addr clauses. + * testsuite/libgomp.c++/target-9.C (struct S): New type. + (foo): Use use_device_addr clause instead of use_device_ptr clause + where required by OpenMP 5.0, add further tests for both use_device_ptr + and use_device_addr clauses. Add t and u arguments. + (main): Adjust caller. + 2019-08-06 Jakub Jelinek * testsuite/libgomp.c++/loop-13.C: New test. diff --git a/libgomp/testsuite/libgomp.c++/target-9.C b/libgomp/testsuite/libgomp.c++/target-9.C index a5d171b0b3d..83a61cf599b 100644 --- a/libgomp/testsuite/libgomp.c++/target-9.C +++ b/libgomp/testsuite/libgomp.c++/target-9.C @@ -1,10 +1,13 @@ extern "C" void abort (void); +struct S { int e, f; }; void -foo (int *&p, int (&s)[5], int n) +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) @@ -30,7 +33,7 @@ foo (int *&p, int (&s)[5], int n) if (err) abort (); #pragma omp target data map(to:b) - #pragma omp target data use_device_ptr(b) map(from:err) + #pragma omp target data use_device_addr(b) map(from:err) #pragma omp target is_device_ptr(b) private(i) map(from:err) { err = 0; @@ -41,7 +44,7 @@ foo (int *&p, int (&s)[5], int n) if (err) abort (); #pragma omp target data map(to:c) - #pragma omp target data use_device_ptr(c) map(from:err) + #pragma omp target data use_device_addr(c) map(from:err) #pragma omp target is_device_ptr(c) private(i) map(from:err) { err = 0; @@ -52,7 +55,7 @@ foo (int *&p, int (&s)[5], int n) if (err) abort (); #pragma omp target data map(to:s[:5]) - #pragma omp target data use_device_ptr(s) map(from:err) + #pragma omp target data use_device_addr(s) map(from:err) #pragma omp target is_device_ptr(s) private(i) map(from:err) { err = 0; @@ -62,6 +65,34 @@ foo (int *&p, int (&s)[5], int n) } if (err) abort (); + #pragma omp target data map(to: v) map(to:u) + #pragma omp target data use_device_addr (v) 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, w) + #pragma omp target data use_device_addr (t, 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 @@ -69,5 +100,7 @@ main () { int a[4] = { 0, 1, 2, 3 }, b[5] = { 17, 18, 19, 20, 21 }; int *p = a + 1; - foo (p, b, 9); + int t = 24; + S u = { 25, 26 }; + foo (p, b, t, u, 9); } diff --git a/libgomp/testsuite/libgomp.c/target-18.c b/libgomp/testsuite/libgomp.c/target-18.c index cbacaf6a77a..dd511fb98cb 100644 --- a/libgomp/testsuite/libgomp.c/target-18.c +++ b/libgomp/testsuite/libgomp.c/target-18.c @@ -1,9 +1,11 @@ extern void abort (void); +struct S { int e, f; }; void foo (int n) { - int a[4] = { 0, 1, 2, 3 }, b[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; @@ -21,7 +23,7 @@ foo (int n) for (i = 0; i < 4; i++) a[i] = 23 + i; #pragma omp target data map(to:a) - #pragma omp target data use_device_ptr(a) map(from:err) + #pragma omp target data use_device_addr(a) map(from:err) #pragma omp target is_device_ptr(a) private(i) map(from:err) { err = 0; @@ -32,7 +34,7 @@ foo (int n) if (err) abort (); #pragma omp target data map(to:b) - #pragma omp target data use_device_ptr(b) map(from:err) + #pragma omp target data use_device_addr(b) map(from:err) #pragma omp target is_device_ptr(b) private(i) map(from:err) { err = 0; @@ -42,6 +44,28 @@ foo (int n) } if (err) abort (); + #pragma omp target data map(to:c) + #pragma omp target data 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 map(to:d) + #pragma omp target data use_device_addr(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 -- 2.30.2