diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index 0251aec5348..f5314d66f27 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -1862,6 +1862,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code, case OMP_CLAUSE_DEVICE: case OMP_CLAUSE_MAP: case OMP_CLAUSE_IS_DEVICE_PTR: + case OMP_CLAUSE_HAS_DEVICE_ADDR: case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE_DEPEND: s = C_OMP_CLAUSE_SPLIT_TARGET; diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 6e6c806ccf7..54864c2ec41 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -89,8 +89,8 @@ enum pragma_kind { }; -/* All clauses defined by OpenACC 2.0, and OpenMP 2.5, 3.0, 3.1, 4.0, 4.5 - and 5.0. Used internally by both C and C++ parsers. */ +/* All clauses defined by OpenACC 2.0, and OpenMP 2.5, 3.0, 3.1, 4.0, 4.5, 5.0, + and 5.1. Used internally by both C and C++ parsers. */ enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_NONE = 0, @@ -114,6 +114,7 @@ enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_FOR, PRAGMA_OMP_CLAUSE_FROM, PRAGMA_OMP_CLAUSE_GRAINSIZE, + PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR, PRAGMA_OMP_CLAUSE_HINT, PRAGMA_OMP_CLAUSE_IF, PRAGMA_OMP_CLAUSE_IN_REDUCTION, diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 7e81c33de11..3b1d2d45add 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -12771,7 +12771,9 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_GRAINSIZE; break; case 'h': - if (!strcmp ("hint", p)) + if (!strcmp ("has_device_addr", p)) + result = PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR; + else if (!strcmp ("hint", p)) result = PRAGMA_OMP_CLAUSE_HINT; else if (!strcmp ("host", p)) result = PRAGMA_OACC_CLAUSE_HOST; @@ -13164,6 +13166,7 @@ c_parser_omp_variable_list (c_parser *parser, case OMP_CLAUSE_REDUCTION: case OMP_CLAUSE_IN_REDUCTION: case OMP_CLAUSE_TASK_REDUCTION: + case OMP_CLAUSE_HAS_DEVICE_ADDR: array_section_p = false; dims.truncate (0); while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE)) @@ -14324,6 +14327,16 @@ c_parser_omp_clause_use_device_addr (c_parser *parser, tree list) list); } +/* OpenMP 5.1: + has_device_addr ( variable-list ) */ + +static tree +c_parser_omp_clause_has_device_addr (c_parser *parser, tree list) +{ + return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_HAS_DEVICE_ADDR, + list); +} + /* OpenMP 4.5: is_device_ptr ( variable-list ) */ @@ -17052,6 +17065,10 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_omp_clause_use_device_addr (parser, clauses); c_name = "use_device_addr"; break; + case PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR: + clauses = c_parser_omp_clause_has_device_addr (parser, clauses); + c_name = "has_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"; @@ -21034,7 +21051,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR)) static bool c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p) diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 39094cc6f50..3075c883548 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -13804,6 +13804,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) } first = c_fully_fold (first, false, NULL); OMP_CLAUSE_DECL (c) = first; + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) + return false; if (size) size = c_fully_fold (size, false, NULL); OMP_CLAUSE_SIZE (c) = size; @@ -14109,7 +14111,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { bitmap_head generic_head, firstprivate_head, lastprivate_head; bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head; - bitmap_head oacc_reduction_head; + bitmap_head oacc_reduction_head, is_on_device_head; tree c, t, type, *pc; tree simdlen = NULL_TREE, safelen = NULL_TREE; bool branch_seen = false; @@ -14145,6 +14147,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head instead and for ort == C_ORT_OMP_TARGET used as in_reduction_head. */ bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); + bitmap_initialize (&is_on_device_head, &bitmap_default_obstack); if (ort & C_ORT_ACC) for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) @@ -14573,7 +14576,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) "%qE appears more than once in data clauses", t); remove = true; } - else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE + else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) && bitmap_bit_p (&map_head, DECL_UID (t))) { if (ort == C_ORT_ACC) @@ -15187,7 +15192,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) "%qD appears more than once in data clauses", t); remove = true; } - else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)) + || bitmap_bit_p (&is_on_device_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), @@ -15272,6 +15278,8 @@ 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 (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) + bitmap_set_bit (&is_on_device_head, DECL_UID (t)); if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE) { if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR @@ -15292,6 +15300,24 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } goto check_dup_generic; + case OMP_CLAUSE_HAS_DEVICE_ADDR: + t = OMP_CLAUSE_DECL (c); + if (TREE_CODE (t) == TREE_LIST) + { + if (handle_omp_array_sections (c, ort)) + remove = true; + else + { + t = OMP_CLAUSE_DECL (c); + while (TREE_CODE (t) == ARRAY_REF) + t = TREE_OPERAND (t, 0); + } + } + bitmap_set_bit (&is_on_device_head, DECL_UID (t)); + if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) + c_mark_addressable (t); + goto check_dup_generic_t; + case OMP_CLAUSE_USE_DEVICE_ADDR: t = OMP_CLAUSE_DECL (c); if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 94a5c64be4c..03d99aba13e 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -36341,7 +36341,9 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_GRAINSIZE; break; case 'h': - if (!strcmp ("hint", p)) + if (!strcmp ("has_device_addr", p)) + result = PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR; + else if (!strcmp ("hint", p)) result = PRAGMA_OMP_CLAUSE_HINT; else if (!strcmp ("host", p)) result = PRAGMA_OACC_CLAUSE_HOST; @@ -36644,6 +36646,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, case OMP_CLAUSE_REDUCTION: case OMP_CLAUSE_IN_REDUCTION: case OMP_CLAUSE_TASK_REDUCTION: + case OMP_CLAUSE_HAS_DEVICE_ADDR: array_section_p = false; dims.truncate (0); while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE)) @@ -40085,6 +40088,11 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses); c_name = "is_device_ptr"; break; + case PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_HAS_DEVICE_ADDR, + clauses); + c_name = "has_device_addr"; + break; case PRAGMA_OMP_CLAUSE_IF: clauses = cp_parser_omp_clause_if (parser, clauses, token->location, true); @@ -44265,7 +44273,8 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok, | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ALLOCATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR)) static bool cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 466d6b56871..0cb17a6a8ab 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -5648,6 +5648,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) return false; } OMP_CLAUSE_DECL (c) = first; + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) + return false; OMP_CLAUSE_SIZE (c) = size; if (TREE_CODE (t) == FIELD_DECL) t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); @@ -6677,7 +6679,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { bitmap_head generic_head, firstprivate_head, lastprivate_head; bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head; - bitmap_head oacc_reduction_head; + bitmap_head oacc_reduction_head, is_on_device_head; tree c, t, *pc; tree safelen = NULL_TREE; bool branch_seen = false; @@ -6710,6 +6712,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head instead and for ort == C_ORT_OMP_TARGET used as in_reduction_head. */ bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); + bitmap_initialize (&is_on_device_head, &bitmap_default_obstack); if (ort & C_ORT_ACC) for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) @@ -7008,7 +7011,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) "%qD appears more than once in data clauses", t); remove = true; } - else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE + else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) && bitmap_bit_p (&map_head, DECL_UID (t))) { if (ort == C_ORT_ACC) @@ -8232,7 +8237,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) "%qD appears more than once in data clauses", t); remove = true; } - else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)) + || bitmap_bit_p (&is_on_device_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), @@ -8491,6 +8497,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_USE_DEVICE_PTR: field_ok = (ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP; t = OMP_CLAUSE_DECL (c); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) + bitmap_set_bit (&is_on_device_head, DECL_UID (t)); if (!type_dependent_expression_p (t)) { tree type = TREE_TYPE (t); @@ -8520,6 +8528,25 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } goto check_dup_generic; + case OMP_CLAUSE_HAS_DEVICE_ADDR: + t = OMP_CLAUSE_DECL (c); + if (TREE_CODE (t) == TREE_LIST) + { + if (handle_omp_array_sections (c, ort)) + remove = true; + else + { + t = OMP_CLAUSE_DECL (c); + while (TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == ARRAY_REF) + t = TREE_OPERAND (t, 0); + } + } + bitmap_set_bit (&is_on_device_head, DECL_UID (t)); + if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) + cxx_mark_addressable (t); + goto check_dup_generic_t; + case OMP_CLAUSE_USE_DEVICE_ADDR: field_ok = true; t = OMP_CLAUSE_DECL (c); diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 3112caec053..2a2f9901b08 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1683,6 +1683,7 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses) case OMP_LIST_CACHE: type = "CACHE"; break; case OMP_LIST_IS_DEVICE_PTR: type = "IS_DEVICE_PTR"; break; case OMP_LIST_USE_DEVICE_PTR: type = "USE_DEVICE_PTR"; break; + case OMP_LIST_HAS_DEVICE_ADDR: type = "HAS_DEVICE_ADDR"; break; case OMP_LIST_USE_DEVICE_ADDR: type = "USE_DEVICE_ADDR"; break; case OMP_LIST_NONTEMPORAL: type = "NONTEMPORAL"; break; case OMP_LIST_ALLOCATE: type = "ALLOCATE"; break; diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 993879feda4..cb136f875f4 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1393,7 +1393,8 @@ enum OMP_LIST_USE_DEVICE_ADDR, OMP_LIST_NONTEMPORAL, OMP_LIST_ALLOCATE, - OMP_LIST_NUM + OMP_LIST_HAS_DEVICE_ADDR, + OMP_LIST_NUM /* Must be the last. */ }; /* Because a symbol can belong to multiple namelists, they must be diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 38c67e1f640..33b372f3430 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -926,7 +926,7 @@ enum omp_mask1 OMP_MASK1_LAST }; -/* OpenACC 2.0+ specific clauses. */ +/* More OpenMP clauses and OpenACC 2.0+ specific clauses. */ enum omp_mask2 { OMP_CLAUSE_ASYNC, @@ -955,6 +955,7 @@ enum omp_mask2 OMP_CLAUSE_FINALIZE, OMP_CLAUSE_ATTACH, OMP_CLAUSE_NOHOST, + OMP_CLAUSE_HAS_DEVICE_ADDR, /* OpenMP 5.1 */ /* This must come last. */ OMP_MASK2_LAST }; @@ -2151,6 +2152,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, } break; case 'h': + if ((mask & OMP_CLAUSE_HAS_DEVICE_ADDR) + && gfc_match_omp_variable_list + ("has_device_addr (", &c->lists[OMP_LIST_HAS_DEVICE_ADDR], + false, NULL, NULL, true) == MATCH_YES) + continue; if ((mask & OMP_CLAUSE_HINT) && (m = gfc_match_dupl_check (!c->hint, "hint", true, &c->hint)) != MATCH_NO) @@ -2923,8 +2929,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, continue; if ((mask & OMP_CLAUSE_USE_DEVICE_ADDR) && gfc_match_omp_variable_list - ("use_device_addr (", - &c->lists[OMP_LIST_USE_DEVICE_ADDR], false) == MATCH_YES) + ("use_device_addr (", &c->lists[OMP_LIST_USE_DEVICE_ADDR], + false, NULL, NULL, true) == MATCH_YES) continue; break; case 'v': @@ -3651,7 +3657,8 @@ cleanup: | OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOWAIT | OMP_CLAUSE_PRIVATE \ | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULTMAP \ | OMP_CLAUSE_IS_DEVICE_PTR | OMP_CLAUSE_IN_REDUCTION \ - | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE) + | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE \ + | OMP_CLAUSE_HAS_DEVICE_ADDR) #define OMP_TARGET_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_DEVICE) | OMP_CLAUSE_MAP | OMP_CLAUSE_IF \ | OMP_CLAUSE_USE_DEVICE_PTR | OMP_CLAUSE_USE_DEVICE_ADDR) @@ -6283,7 +6290,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, "IN_REDUCTION", "TASK_REDUCTION", "DEVICE_RESIDENT", "LINK", "USE_DEVICE", "CACHE", "IS_DEVICE_PTR", "USE_DEVICE_PTR", "USE_DEVICE_ADDR", - "NONTEMPORAL", "ALLOCATE" }; + "NONTEMPORAL", "ALLOCATE", "HAS_DEVICE_ADDR" }; STATIC_ASSERT (ARRAY_SIZE (clause_names) == OMP_LIST_NUM); if (omp_clauses == NULL) @@ -7132,6 +7139,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, n->sym->name, name, &n->where); } break; + case OMP_LIST_HAS_DEVICE_ADDR: case OMP_LIST_USE_DEVICE_PTR: case OMP_LIST_USE_DEVICE_ADDR: /* FIXME: Handle OMP_LIST_USE_DEVICE_PTR. */ diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index d5a6b2d6ee3..0eba0b3c3e1 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -1910,7 +1910,17 @@ gfc_trans_omp_variable_list (enum omp_clause_code code, tree t = gfc_trans_omp_variable (namelist->sym, declare_simd); if (t != error_mark_node) { - tree node = build_omp_clause (input_location, code); + tree node; + /* For HAS_DEVICE_ADDR of an array descriptor, firstprivatize the + descriptor such that the bounds are available; its data component + is unmodified; it is handled as device address inside target. */ + if (code == OMP_CLAUSE_HAS_DEVICE_ADDR + && (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (t)) + || (POINTER_TYPE_P (TREE_TYPE (t)) + && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (t)))))) + node = build_omp_clause (input_location, OMP_CLAUSE_FIRSTPRIVATE); + else + node = build_omp_clause (input_location, code); OMP_CLAUSE_DECL (node) = t; list = gfc_trans_add_clause (node, list); @@ -2604,6 +2614,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, case OMP_LIST_IS_DEVICE_PTR: clause_code = OMP_CLAUSE_IS_DEVICE_PTR; goto add_clause; + case OMP_LIST_HAS_DEVICE_ADDR: + clause_code = OMP_CLAUSE_HAS_DEVICE_ADDR; + goto add_clause; case OMP_LIST_NONTEMPORAL: clause_code = OMP_CLAUSE_NONTEMPORAL; goto add_clause; diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 875b115d02d..8d676fb96c8 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -10278,6 +10278,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, flags = GOVD_EXPLICIT; goto do_add; + case OMP_CLAUSE_HAS_DEVICE_ADDR: + decl = OMP_CLAUSE_DECL (c); + while (TREE_CODE (decl) == INDIRECT_REF + || TREE_CODE (decl) == ARRAY_REF) + decl = TREE_OPERAND (decl, 0); + flags = GOVD_EXPLICIT; + goto do_add_decl; + case OMP_CLAUSE_IS_DEVICE_PTR: flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT; goto do_add; @@ -11428,6 +11436,16 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } break; + case OMP_CLAUSE_HAS_DEVICE_ADDR: + decl = OMP_CLAUSE_DECL (c); + while (TREE_CODE (decl) == INDIRECT_REF + || TREE_CODE (decl) == ARRAY_REF) + decl = TREE_OPERAND (decl, 0); + n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); + remove = n == NULL || !(n->value & GOVD_SEEN); + break; + + case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_NONTEMPORAL: decl = OMP_CLAUSE_DECL (c); n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); @@ -11729,7 +11747,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_CLAUSE_DETACH: 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: case OMP_CLAUSE_INDEPENDENT: diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index c33b3daa439..065208464b2 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1375,7 +1375,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) decl = OMP_CLAUSE_DECL (c); do_private: if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE - || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) && is_gimple_omp_offloaded (ctx->stmt)) { if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) @@ -1383,8 +1384,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) by_ref = !omp_privatize_by_reference (decl); install_var_field (decl, by_ref, 3, ctx); } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) + { + if (TREE_CODE (decl) == INDIRECT_REF) + decl = TREE_OPERAND (decl, 0); + install_var_field (decl, true, 3, ctx); + } else if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) - install_var_field (decl, true, 3, ctx); + install_var_field (decl, true, 3, ctx); else install_var_field (decl, false, 3, ctx); } @@ -1452,6 +1459,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) install_var_local (decl, ctx); break; + case OMP_CLAUSE_HAS_DEVICE_ADDR: + decl = OMP_CLAUSE_DECL (c); + while (TREE_CODE (decl) == INDIRECT_REF + || TREE_CODE (decl) == ARRAY_REF) + decl = TREE_OPERAND (decl, 0); + goto do_private; + case OMP_CLAUSE_IS_DEVICE_PTR: decl = OMP_CLAUSE_DECL (c); goto do_private; @@ -1729,12 +1743,21 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_PRIVATE: case OMP_CLAUSE_LINEAR: + case OMP_CLAUSE_HAS_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: decl = OMP_CLAUSE_DECL (c); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) + { + while (TREE_CODE (decl) == INDIRECT_REF + || TREE_CODE (decl) == ARRAY_REF) + decl = TREE_OPERAND (decl, 0); + } + if (is_variable_sized (decl)) { if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE - || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) && is_gimple_omp_offloaded (ctx->stmt)) { tree decl2 = DECL_VALUE_EXPR (decl); @@ -12819,8 +12842,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_ADDR: + case OMP_CLAUSE_HAS_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: var = OMP_CLAUSE_DECL (c); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) + { + while (TREE_CODE (var) == INDIRECT_REF + || TREE_CODE (var) == ARRAY_REF) + var = TREE_OPERAND (var, 0); + } map_cnt++; if (is_variable_sized (var)) { @@ -12835,7 +12865,8 @@ 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 ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + else if (((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) && !omp_privatize_by_reference (var) && !omp_is_allocatable_or_ptr (var) && !lang_hooks.decls.omp_array_data (var, true)) @@ -13301,17 +13332,26 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_ADDR: + case OMP_CLAUSE_HAS_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: ovar = OMP_CLAUSE_DECL (c); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) + { + while (TREE_CODE (ovar) == INDIRECT_REF + || TREE_CODE (ovar) == ARRAY_REF) + ovar = TREE_OPERAND (ovar, 0); + } var = lookup_decl_in_outer_ctx (ovar, ctx); if (lang_hooks.decls.omp_array_data (ovar, true)) { - tkind = (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR + tkind = ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR) ? GOMP_MAP_USE_DEVICE_PTR : GOMP_MAP_FIRSTPRIVATE_INT); x = build_sender_ref ((splay_tree_key) &DECL_NAME (ovar), ctx); } - else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR) + else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR) { tkind = GOMP_MAP_USE_DEVICE_PTR; x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), ctx); @@ -13333,7 +13373,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) type = TREE_TYPE (ovar); if (lang_hooks.decls.omp_array_data (ovar, true)) var = lang_hooks.decls.omp_array_data (ovar, false); - else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + else if (((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) && !omp_privatize_by_reference (ovar) && !omp_is_allocatable_or_ptr (ovar)) || TREE_CODE (type) == ARRAY_TYPE) @@ -13348,6 +13389,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (POINTER_TYPE_P (type) && TREE_CODE (type) != ARRAY_TYPE && ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR && !omp_is_allocatable_or_ptr (ovar)) || (omp_privatize_by_reference (ovar) && omp_is_allocatable_or_ptr (ovar)))) @@ -13545,6 +13587,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_HAS_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: tree new_var; gimple_seq assign_body; @@ -13555,12 +13598,21 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) var = OMP_CLAUSE_DECL (c); is_array_data = lang_hooks.decls.omp_array_data (var, true) != NULL; - if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR) + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR) x = build_sender_ref (is_array_data ? (splay_tree_key) &DECL_NAME (var) : (splay_tree_key) &DECL_UID (var), ctx); else - x = build_receiver_ref (var, false, ctx); + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) + { + while (TREE_CODE (var) == INDIRECT_REF + || TREE_CODE (var) == ARRAY_REF) + var = TREE_OPERAND (var, 0); + } + x = build_receiver_ref (var, false, ctx); + } if (is_array_data) { @@ -13607,7 +13659,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq_add_stmt (&assign_body, gimple_build_assign (new_var, x)); } - else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + else if (((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) && !omp_privatize_by_reference (var) && !omp_is_allocatable_or_ptr (var)) || TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) @@ -13630,7 +13683,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) type = TREE_TYPE (type); if (POINTER_TYPE_P (type) && TREE_CODE (type) != ARRAY_TYPE - && (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR + && ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR) || (omp_privatize_by_reference (var) && omp_is_allocatable_or_ptr (var)))) { @@ -13653,7 +13707,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_build_assign (new_var, x)); } tree present; - present = (do_optional_check + present = ((do_optional_check + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR) ? omp_check_optional_argument (OMP_CLAUSE_DECL (c), true) : NULL_TREE); if (present) diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-1.c b/gcc/testsuite/c-c++-common/gomp/clauses-1.c index 3ff49e0a298..71ca41c5804 100644 --- a/gcc/testsuite/c-c++-common/gomp/clauses-1.c +++ b/gcc/testsuite/c-c++-common/gomp/clauses-1.c @@ -102,7 +102,7 @@ baz (int d, int m, int i1, int i2, int p, int *idp, int s, } void -bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, +bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int hda, int s, int nte, int tl, int nth, int g, int nta, int fi, int pp, int *q, int *dd, int ntm) { #pragma omp for simd \ @@ -138,20 +138,20 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, #pragma omp target parallel \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ - nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) + nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda) ; #pragma omp target parallel for \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ lastprivate (l) linear (ll:1) ordered schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) \ - allocate (omp_default_mem_alloc:f) in_reduction(+:r2) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda) for (int i = 0; i < 64; i++) ll++; #pragma omp target parallel for \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) order(concurrent) \ - allocate (omp_default_mem_alloc:f) in_reduction(+:r2) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda) for (int i = 0; i < 64; i++) ll++; #pragma omp target parallel for simd \ @@ -159,18 +159,19 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) \ safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) order(concurrent) \ - allocate (omp_default_mem_alloc:f) in_reduction(+:r2) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda) for (int i = 0; i < 64; i++) ll++; #pragma omp target teams \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ shared(s) default(shared) reduction(+:r) num_teams(nte - 1:nte) thread_limit(tl) nowait depend(inout: dd[0]) \ - allocate (omp_default_mem_alloc:f) in_reduction(+:r2) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda) ; #pragma omp target teams distribute \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) order(concurrent) \ - collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) + collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) \ + has_device_addr(hda) for (int i = 0; i < 64; i++) ; #pragma omp target teams distribute parallel for \ @@ -179,7 +180,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, collapse(1) dist_schedule(static, 16) \ if (parallel: i2) num_threads (nth) proc_bind(spread) \ lastprivate (l) schedule(static, 4) nowait depend(inout: dd[0]) order(concurrent) \ - allocate (omp_default_mem_alloc:f) in_reduction(+:r2) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda) for (int i = 0; i < 64; i++) ll++; #pragma omp target teams distribute parallel for simd \ @@ -189,7 +190,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, if (parallel: i2) num_threads (nth) proc_bind(spread) \ lastprivate (l) schedule(static, 4) order(concurrent) \ safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) \ - allocate (omp_default_mem_alloc:f) in_reduction(+:r2) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda) for (int i = 0; i < 64; i++) ll++; #pragma omp target teams distribute simd \ @@ -197,14 +198,14 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) \ collapse(1) dist_schedule(static, 16) order(concurrent) \ safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) \ - allocate (omp_default_mem_alloc:f) in_reduction(+:r2) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda) for (int i = 0; i < 64; i++) ll++; #pragma omp target simd \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ safelen(8) simdlen(4) lastprivate (l) linear(ll: 1) aligned(q: 32) reduction(+:r) \ nowait depend(inout: dd[0]) nontemporal(ntm) if(simd:i3) order(concurrent) \ - allocate (omp_default_mem_alloc:f) in_reduction(+:r2) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda) for (int i = 0; i < 64; i++) ll++; #pragma omp taskgroup task_reduction(+:r2) allocate (r2) @@ -430,28 +431,28 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ nowait depend(inout: dd[0]) lastprivate (l) bind(parallel) order(concurrent) collapse(1) \ - allocate (omp_default_mem_alloc: f) in_reduction(+:r2) + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda) for (l = 0; l < 64; ++l) ; #pragma omp target parallel loop \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ nowait depend(inout: dd[0]) lastprivate (l) order(concurrent) collapse(1) \ - allocate (omp_default_mem_alloc: f) in_reduction(+:r2) + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda) for (l = 0; l < 64; ++l) ; #pragma omp target teams loop \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) nowait depend(inout: dd[0]) \ lastprivate (l) bind(teams) collapse(1) \ - allocate (omp_default_mem_alloc: f) in_reduction(+:r2) + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda) for (l = 0; l < 64; ++l) ; #pragma omp target teams loop \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0]) \ lastprivate (l) order(concurrent) collapse(1) \ - allocate (omp_default_mem_alloc: f) in_reduction(+:r2) + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda) for (l = 0; l < 64; ++l) ; } diff --git a/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-1.c b/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-1.c new file mode 100644 index 00000000000..ebf55eea70b --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-1.c @@ -0,0 +1,65 @@ +/* { dg-do compile } */ + +void +foo () +{ + int * x; + #pragma omp target is_device_ptr(x) has_device_addr(x) /*{ dg-error "'x' appears more than once in data clauses" } */ + ; + #pragma omp target has_device_addr(x) is_device_ptr(x) /* { dg-error "'x' appears more than once in data clauses" } */ + ; + + int y = 42; + #pragma omp target has_device_addr(y) has_device_addr(y) /* { dg-error "'y' appears more than once in data clauses" } */ + ; + + #pragma omp target private(y) has_device_addr(y) /*{ dg-error "'y' appears more than once in data clauses" } */ + ; + #pragma omp target has_device_addr(y) private(y) /*{ dg-error "'y' appears more than once in data clauses" } */ + ; + #pragma omp target firstprivate(y) has_device_addr(y) /*{ dg-error "'y' appears more than once in data clauses" } */ + ; + + #pragma omp target has_device_addr(y) map(y) /* { dg-error "'y' appears both in data and map clauses" } */ + ; + #pragma omp target map(y) has_device_addr(y) /* { dg-error "'y' appears both in data and map clauses" } */ + ; + + int z[3] = { 2, 5, 7 }; + #pragma omp target data map(z[:3]) use_device_addr(z) + #pragma omp target has_device_addr(z[1:]) + ; + + #pragma omp target data map(z[:3]) use_device_addr(z) + #pragma omp target has_device_addr(z[1]) + ; + + #pragma omp target data map(z[:3]) use_device_addr(z) + #pragma omp target has_device_addr(z[1:2]) + ; + + #pragma omp target data map(z[:3]) use_device_addr(z) + #pragma omp target has_device_addr(z[:2]) + ; + + int w[3][4]; + #pragma omp target data map(w) use_device_addr(w) + #pragma omp target has_device_addr(w[1][2]) + ; + + #pragma omp target data map(w) use_device_addr(w) + #pragma omp target has_device_addr(w[:1][2:]) + ; + + int u[0]; + #pragma omp target data map(u) use_device_addr(u) + #pragma omp target has_device_addr(u) + ; + + struct S { int m; } s; + s.m = 42; + #pragma omp target data map (s) use_device_addr (s) + #pragma omp target has_device_addr (s) + ++s.m; + +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-2.c b/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-2.c new file mode 100644 index 00000000000..7378416964c --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-2.c @@ -0,0 +1,17 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ + +void +foo () +{ + int x, y; + + #pragma omp target data map(x, y) use_device_addr(x, y) + #pragma omp target has_device_addr(x, y) + { + x = 42; + } +} + +/* { dg-final { scan-tree-dump "has_device_addr\\(x\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump-not "has_device_addr\\(y\\)" "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr-1.c b/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr-1.c new file mode 100644 index 00000000000..ecf30ca65f7 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr-1.c @@ -0,0 +1,22 @@ +/* { dg-do compile } */ + +void +foo () +{ + int *x; + + #pragma omp target is_device_ptr(x) is_device_ptr(x) /* { dg-error "'x' appears more than once in data clauses" } */ + ; + + #pragma omp target private(x) is_device_ptr(x) /*{ dg-error "'x' appears more than once in data clauses" } */ + ; + #pragma omp target is_device_ptr(x) private(x) /*{ dg-error "'x' appears more than once in data clauses" } */ + ; + #pragma omp target firstprivate(x) is_device_ptr(x) /*{ dg-error "'x' appears more than once in data clauses" } */ + ; + + #pragma omp target is_device_ptr(x) map(x) /* { dg-error "'x' appears both in data and map clauses" } */ + ; + #pragma omp target map(x) is_device_ptr(x) /* { dg-error "'x' appears both in data and map clauses" } */ + ; +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr-2.c b/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr-2.c new file mode 100644 index 00000000000..df743dda166 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr-2.c @@ -0,0 +1,17 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ + +void +foo () +{ + int *x, *y; + + #pragma omp target data map(x, y) use_device_ptr(x, y) + #pragma omp target is_device_ptr(x, y) + { + *x = 42; + } +} + +/* { dg-final { scan-tree-dump "is_device_ptr\\(x\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump-not "is_device_ptr\\(y\\)" "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/attrs-1.C b/gcc/testsuite/g++.dg/gomp/attrs-1.C index 319ad3241de..f64b078db40 100644 --- a/gcc/testsuite/g++.dg/gomp/attrs-1.C +++ b/gcc/testsuite/g++.dg/gomp/attrs-1.C @@ -121,7 +121,7 @@ baz (int d, int m, int i1, int i2, int p, int *idp, int s, } void -bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, +bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int hda, int s, int nte, int tl, int nth, int g, int nta, int fi, int pp, int *q, int *dd, int ntm, const char *msg) { @@ -185,20 +185,20 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, [[omp::directive (target parallel device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) - nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]] + nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]] ; [[omp::directive (target parallel for device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) lastprivate (l) linear (ll:1) ordered schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) - allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[omp::directive (target parallel for device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) order(concurrent) - allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[omp::sequence (omp::directive (target parallel for simd @@ -206,22 +206,23 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) order(concurrent) - allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]] + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda)))]] for (int i = 0; i < 64; i++) ll++; [[omp::sequence (directive (target teams device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) nowait depend(inout: dd[0]) - allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]] + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda)))]] ; [[omp::sequence (directive (target device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) - nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]] + nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda)))]] ; [[omp::sequence (omp::directive (target teams distribute device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) order(concurrent) - collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]] + collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) + has_device_addr (hda)))]] for (int i = 0; i < 64; i++) ; [[omp::directive (target teams distribute parallel for @@ -230,7 +231,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, collapse(1) dist_schedule(static, 16) if (parallel: i2) num_threads (nth) proc_bind(spread) lastprivate (l) schedule(static, 4) nowait depend(inout: dd[0]) order(concurrent) - allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[omp::directive (target teams distribute parallel for simd @@ -240,7 +241,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, if (parallel: i2) num_threads (nth) proc_bind(spread) lastprivate (l) schedule(static, 4) order(concurrent) safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) - allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[omp::directive (target teams distribute simd @@ -248,14 +249,14 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) collapse(1) dist_schedule(static, 16) order(concurrent) safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) - allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[omp::directive (target simd device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) safelen(8) simdlen(4) lastprivate (l) linear(ll: 1) aligned(q: 32) reduction(+:r) nowait depend(inout: dd[0]) nontemporal(ntm) if(simd:i3) order(concurrent) - allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[omp::sequence (directive (taskgroup task_reduction(+:r2) allocate (r2)), @@ -515,28 +516,28 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) nowait depend(inout: dd[0]) lastprivate (l) bind(parallel) order(concurrent) collapse(1) - allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]] + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]] for (l = 0; l < 64; ++l) ; [[omp::directive (target parallel loop device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) nowait depend(inout: dd[0]) lastprivate (l) order(concurrent) collapse(1) - allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]] + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]] for (l = 0; l < 64; ++l) ; [[omp::directive (target teams loop device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0]) lastprivate (l) bind(teams) collapse(1) - allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]] + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]] for (l = 0; l < 64; ++l) ; [[omp::directive (target teams loop device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) shared(s) default(shared) reduction(+:r) num_teams(nte - 1 : nte) thread_limit(tl) nowait depend(inout: dd[0]) lastprivate (l) order(concurrent) collapse(1) - allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]] + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]] for (l = 0; l < 64; ++l) ; [[omp::directive (critical)]] { diff --git a/gcc/testsuite/g++.dg/gomp/attrs-2.C b/gcc/testsuite/g++.dg/gomp/attrs-2.C index 955b2dd04c7..cc91fa28307 100644 --- a/gcc/testsuite/g++.dg/gomp/attrs-2.C +++ b/gcc/testsuite/g++.dg/gomp/attrs-2.C @@ -121,7 +121,7 @@ baz (int d, int m, int i1, int i2, int p, int *idp, int s, } void -bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, +bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int hda, int s, int nte, int tl, int nth, int g, int nta, int fi, int pp, int *q, int *dd, int ntm, const char *msg) { @@ -185,20 +185,20 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, [[omp::directive (target parallel, device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp), if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread) - nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]] + nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]] ; [[omp::directive (target parallel for, device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp), if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread), lastprivate (l),linear (ll:1),ordered schedule(static, 4),collapse(1),nowait depend(inout: dd[0]), - allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[using omp:directive (target parallel for, device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp), if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread), lastprivate (l),linear (ll:1),schedule(static, 4),collapse(1),nowait depend(inout: dd[0]),order(concurrent), - allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[omp::sequence (omp::directive (target parallel for simd, @@ -206,22 +206,23 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread), lastprivate (l),linear (ll:1),schedule(static, 4),collapse(1), safelen(8),simdlen(4),aligned(q: 32),nowait depend(inout: dd[0]),nontemporal(ntm),if (simd: i3),order(concurrent), - allocate (omp_default_mem_alloc:f),in_reduction(+:r2)))]] + allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda)))]] for (int i = 0; i < 64; i++) ll++; [[using omp:sequence (directive (target teams, device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp), - shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),nowait, depend(inout: dd[0]), - allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]] + shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),nowait,depend(inout: dd[0]), + allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda)))]] ; [[using omp:sequence (directive (target, device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp), - nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2)))]] + nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr(hda)))]] ; [[omp::sequence (omp::directive (target teams distribute, device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp), shared(s),default(shared),reduction(+:r),num_teams(nte-1:nte),thread_limit(tl),order(concurrent), - collapse(1),dist_schedule(static, 16),nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2)))]] + collapse(1),dist_schedule(static, 16),nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2), + has_device_addr (hda)))]] for (int i = 0; i < 64; i++) ; [[omp::directive (target teams distribute parallel for, @@ -230,7 +231,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, collapse(1),dist_schedule(static, 16), if (parallel: i2),num_threads (nth),proc_bind(spread), lastprivate (l),schedule(static, 4),nowait depend(inout: dd[0]),order(concurrent), - allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[omp::directive (target teams distribute parallel for simd, @@ -240,7 +241,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, if (parallel: i2),num_threads (nth),proc_bind(spread), lastprivate (l),schedule(static, 4),order(concurrent), safelen(8),simdlen(4),aligned(q: 32),nowait depend(inout: dd[0]),nontemporal(ntm),if (simd: i3), - allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[omp::directive (target teams distribute simd, @@ -248,14 +249,14 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl), collapse(1),dist_schedule(static, 16),order(concurrent), safelen(8),simdlen(4),aligned(q: 32),nowait depend(inout: dd[0]),nontemporal(ntm), - allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[omp::directive (target simd, device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp), safelen(8),simdlen(4),lastprivate (l),linear(ll: 1),aligned(q: 32),reduction(+:r), nowait depend(inout: dd[0]),nontemporal(ntm),if(simd:i3),order(concurrent), - allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]] + allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]] for (int i = 0; i < 64; i++) ll++; [[omp::sequence (directive (taskgroup, task_reduction(+:r2), allocate (r2)), @@ -515,28 +516,28 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s, device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp), if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread), nowait depend(inout: dd[0]),lastprivate (l),bind(parallel),order(concurrent),collapse(1), - allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]] + allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]] for (l = 0; l < 64; ++l) ; [[omp::directive (target parallel loop, device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp), if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread), nowait depend(inout: dd[0]),lastprivate (l),order(concurrent),collapse(1), - allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]] + allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]] for (l = 0; l < 64; ++l) ; [[omp::directive (target teams loop, device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp), shared(s),default(shared),reduction(+:r),num_teams(nte-1:nte),thread_limit(tl),nowait,depend(inout: dd[0]), lastprivate (l),bind(teams),collapse(1), - allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]] + allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]] for (l = 0; l < 64; ++l) ; [[omp::directive (target teams loop, device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp), shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),nowait,depend(inout: dd[0]), lastprivate (l),order(concurrent),collapse(1) - allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]] + allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]] for (l = 0; l < 64; ++l) ; [[omp::directive (critical)]] { diff --git a/gcc/testsuite/gfortran.dg/gomp/is_device_ptr-3.f90 b/gcc/testsuite/gfortran.dg/gomp/is_device_ptr-3.f90 new file mode 100644 index 00000000000..c3de7726e88 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/is_device_ptr-3.f90 @@ -0,0 +1,27 @@ +! Test to ensure that IS_DEVICE_PTR is removed for non-used variables. + +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +program main + use iso_c_binding + implicit none + + integer :: x, y + call foo (x, y) + +contains + subroutine foo (a, b) + integer, target :: a, b + + !$omp target data map(a, b) use_device_ptr(a, b) + !$omp target is_device_ptr(a, b) + a = 42 + !$omp end target + !$omp end target data + end subroutine foo + +end program main + +! { dg-final { scan-tree-dump "is_device_ptr\\(a\\)" "gimple" } } +! { dg-final { scan-tree-dump-not "is_device_ptr\\(b\\)" "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-1.f90 new file mode 100644 index 00000000000..db3fa46b4a1 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-1.f90 @@ -0,0 +1,36 @@ +! { dg-do compile } + +implicit none + +integer, target :: x +integer, pointer :: ptr +integer :: a(5) + +!$omp target has_device_addr(x) +!$omp end target +!$omp target has_device_addr(ptr) +!$omp end target +!$omp target has_device_addr(a) +!$omp end target +!$omp target has_device_addr(a(2:3)) +!$omp end target +!$omp target has_device_addr(a(:3)) +!$omp end target +!$omp target has_device_addr(a(2:)) +!$omp end target +!$omp target has_device_addr(a(2)) +!$omp end target + +!$omp target has_device_addr(x) has_device_addr(x) ! { dg-error "'x' present on multiple clauses" } +!$omp end target + +!$omp target private(x) has_device_addr(x) ! { dg-error "'x' present on multiple clauses" } +!$omp end target +!$omp target has_device_addr(x) private(x) ! { dg-error "'x' present on multiple clauses" } +!$omp end target +!$omp target firstprivate(x) has_device_addr(x) ! { dg-error "'x' present on multiple clauses" } +!$omp end target +!$omp target has_device_addr(x) firstprivate(x) ! { dg-error "'x' present on multiple clauses" } +!$omp end target + +end \ No newline at end of file diff --git a/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-2.f90 new file mode 100644 index 00000000000..7fc92b3cb6e --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-2.f90 @@ -0,0 +1,27 @@ +! Test to ensure that HAS_DEVICE_ADDR is removed for non-used variables. + +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +program main + use iso_c_binding + implicit none + + integer :: x, y + call foo (x, y) + +contains + subroutine foo (a, b) + integer :: a, b + + !$omp target data map(a) use_device_addr(a) + !$omp target has_device_addr(a) + a = 42 + !$omp end target + !$omp end target data + end subroutine foo + +end program main + +! { dg-final { scan-tree-dump "has_device_addr\\(a\\)" "gimple" } } +! { dg-final { scan-tree-dump-not "has_device_addr\\(b\\)" "gimple" } } diff --git a/gcc/tree-core.h b/gcc/tree-core.h index bf2efa61330..01a1ce499da 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -342,6 +342,9 @@ enum omp_clause_code { OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ OMP_CLAUSE_MAP, + /* OpenMP clause: has_device_addr (variable-list). */ + OMP_CLAUSE_HAS_DEVICE_ADDR, + /* Internal structure to hold OpenACC cache directive's variable-list. #pragma acc cache (variable-list). */ OMP_CLAUSE__CACHE_, diff --git a/gcc/tree-nested.cc b/gcc/tree-nested.cc index b7e9a3b472f..078ceab3ca3 100644 --- a/gcc/tree-nested.cc +++ b/gcc/tree-nested.cc @@ -1339,6 +1339,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_LINK: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_ADDR: + case OMP_CLAUSE_HAS_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_DETACH: do_decl_clause: @@ -2123,6 +2124,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_LINK: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_ADDR: + case OMP_CLAUSE_HAS_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_DETACH: do_decl_clause: diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 666b7a70ea2..99af977979d 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -493,6 +493,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case OMP_CLAUSE_USE_DEVICE_ADDR: name = "use_device_addr"; goto print_remap; + case OMP_CLAUSE_HAS_DEVICE_ADDR: + name = "has_device_addr"; + goto print_remap; case OMP_CLAUSE_IS_DEVICE_PTR: name = "is_device_ptr"; goto print_remap; diff --git a/gcc/tree.cc b/gcc/tree.cc index dfcdf6822f1..aeb3b4cebcf 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -289,6 +289,7 @@ unsigned const char omp_clause_num_ops[] = 2, /* OMP_CLAUSE_FROM */ 2, /* OMP_CLAUSE_TO */ 2, /* OMP_CLAUSE_MAP */ + 1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */ 2, /* OMP_CLAUSE__CACHE_ */ 2, /* OMP_CLAUSE_GANG */ 1, /* OMP_CLAUSE_ASYNC */ @@ -378,6 +379,7 @@ const char * const omp_clause_code_name[] = "from", "to", "map", + "has_device_addr", "_cache_", "gang", "async", diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 424459f4442..161a423ac7c 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -294,7 +294,7 @@ The OpenMP 4.5 specification is fully supported. @item @code{align} clause/modifier in @code{allocate} directive/clause and @code{allocator} directive @tab P @tab C/C++ on clause only @item @code{thread_limit} clause to @code{target} construct @tab Y @tab -@item @code{has_device_addr} clause to @code{target} construct @tab N @tab +@item @code{has_device_addr} clause to @code{target} construct @tab Y @tab @item iterators in @code{target update} motion clauses and @code{map} clauses @tab N @tab @item indirect calls to the device version of a procedure or function in diff --git a/libgomp/target.c b/libgomp/target.c index 698ff14a05f..9017458885e 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -2510,7 +2510,7 @@ copy_firstprivate_data (char *tgt, size_t mapnum, void **hostaddrs, tgt_size = 0; size_t i; for (i = 0; i < mapnum; i++) - if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE) + if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE && hostaddrs[i] != NULL) { size_t align = (size_t) 1 << (kinds[i] >> 8); tgt_size = (tgt_size + align - 1) & ~(align - 1); diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-2.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-2.C new file mode 100644 index 00000000000..d9a309d7af4 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-2.C @@ -0,0 +1,23 @@ +/* Testing 'has_device_addr' clause on the target construct with reference. */ + +#include + +int +main () +{ + int *dp = (int*)omp_target_alloc (sizeof(int), 0); + + #pragma omp target is_device_ptr(dp) + *dp = 42; + + int &x = *dp; + + #pragma omp target has_device_addr(x) + x = 24; + + #pragma omp target has_device_addr(x) + if (x != 24) + __builtin_abort (); + + omp_target_free(dp, 0); +} diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-4.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-4.C new file mode 100644 index 00000000000..6468c6c8433 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-4.C @@ -0,0 +1,33 @@ +#include + +int +main () +{ + int *dp = (int*)omp_target_alloc (30*sizeof(int), 0); + + #pragma omp target is_device_ptr(dp) + for (int i = 0; i < 30; i++) + dp[i] = i; + + int (&x)[30] = *static_cast(static_cast(dp)); + + #pragma omp target has_device_addr(x) + for (int i = 0; i < 30; i++) + x[i] = 2 * i; + + #pragma omp target has_device_addr(x) + for (int i = 0; i < 30; i++) + if (x[i] != 2 * i) + __builtin_abort (); + + #pragma omp target has_device_addr(x[1:5]) + for (int i = 1; i < 6; i++) + x[i] = 3 * i; + + #pragma omp target has_device_addr(x[1:5]) + for (int i = 1; i < 6; i++) + if (x[i] != 3 * i) + __builtin_abort (); + + omp_target_free (dp, 0); +} diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-5.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-5.C new file mode 100644 index 00000000000..e847cdceb44 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-5.C @@ -0,0 +1,33 @@ +/* Testing 'has_device_addr' clause on the target construct with reference. */ + +#include + +int +main () +{ + int *dpx = (int*)omp_target_alloc (sizeof(int), 0); + int **dpy = (int**)omp_target_alloc (sizeof(int*), 0); + + #pragma omp target is_device_ptr(dpx, dpy) + { + *dpx = 42; + int z = 77; + *dpy = &z; + } + + int& x = *dpx; + int*& y = *dpy; + + #pragma omp target has_device_addr(x, y) + { + x = 24; + y = &x; + } + + #pragma omp target has_device_addr(x, y) + if (x != 24 || y != &x) + __builtin_abort (); + + omp_target_free(dpx, 0); + omp_target_free(dpy, 0); +} diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-6.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-6.C new file mode 100644 index 00000000000..141edb14dec --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-6.C @@ -0,0 +1,32 @@ +/* Testing 'has_device_addr' clause on the target construct with reference. */ + +#include + +int +main () +{ + int *dpx = (int*)omp_target_alloc (sizeof(int), 0); + double *dpy = (double*)omp_target_alloc (sizeof(double), 0); + + #pragma omp target is_device_ptr(dpx, dpy) + { + *dpx = 42; + *dpy = 43.5; + } + + int &x = *dpx; + double &y = *dpy; + + #pragma omp target has_device_addr(x, y) + { + x = 24; + y = 25.7; + } + + #pragma omp target has_device_addr(y, x) + if (x != 24 || y != 25.7) + __builtin_abort (); + + omp_target_free(dpx, 0); + omp_target_free(dpy, 0); +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-1.c new file mode 100644 index 00000000000..fcc5c9e8553 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-1.c @@ -0,0 +1,73 @@ +/* Testing the 'has_device_addr' clause on the target construct with + enclosing 'target data' construct. */ + +#define N 40 + +int +main () +{ + int x = 24; + + #pragma omp target data map(x) use_device_addr(x) + #pragma omp target has_device_addr(x) + x = 42; + if (x != 42) + __builtin_abort (); + + int y[N]; + + for (int i = 0; i < N; i++) + y[i] = 42; + #pragma omp target data map(y) use_device_addr(y) + #pragma omp target has_device_addr(y) + for (int i = 0; i < N; i++) + y[i] = i; + for (int i = 0; i < N; i++) + if (y[i] != i) + __builtin_abort (); + + #pragma omp target data map(y[:N]) use_device_addr(y) + #pragma omp target has_device_addr(y[:N]) + for (int i = 0; i < N; i++) + y[i] = i + 2; + for (int i = 0; i < N; i++) + if (y[i] != i + 2) + __builtin_abort (); + + #pragma omp target data map(y[:N]) use_device_addr(y) + #pragma omp target has_device_addr(y[24]) + y[24] = 42; + if (y[24] != 42) + __builtin_abort (); + + #pragma omp target data map(y[:N]) use_device_addr(y) + #pragma omp target has_device_addr(y[24:]) + for (int i = 24; i < N; i++) + y[i] = i + 3; + for (int i = 24; i < N; i++) + if (y[i] != i + 3) + __builtin_abort (); + + #pragma omp target data map(y[:N]) use_device_addr(y) + #pragma omp target has_device_addr(y[12:24]) + for (int i = 12; i < 24; i++) + y[i] = i + 4; + for (int i = 12; i < 24; i++) + if (y[i] != i + 4) + __builtin_abort (); + + int u[0]; + #pragma omp target data map(u) use_device_addr(u) + #pragma omp target has_device_addr(u) + ; + + struct S { int m; } s; + s.m = 42; + #pragma omp target data map (s) use_device_addr (s) + #pragma omp target has_device_addr (s) + ++s.m; + if (s.m != 43) + __builtin_abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-has-device-addr-3.c b/libgomp/testsuite/libgomp.c/target-has-device-addr-3.c new file mode 100644 index 00000000000..fd99a82f66a --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-has-device-addr-3.c @@ -0,0 +1,33 @@ +/* Testing 'has_device_addr' clause with variable sized array. */ + +int +foo (int size) +{ + int x[size]; + + #pragma omp target data map(x[:size]) use_device_addr(x) + #pragma omp target has_device_addr(x) + for (int i = 0; i < size; i++) + x[i] = i; + for (int i = 0; i < size; i++) + if (x[i] != i) + __builtin_abort (); + + #pragma omp target data map(x) use_device_addr(x) + #pragma omp target has_device_addr(x[2:3]) + for (int i = 0; i < size; i++) + x[i] = i; + for (int i = 0; i < size; i++) + if (x[i] != i) + __builtin_abort (); + + return 0; +} + +int +main () +{ + foo (40); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-1.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-1.f90 new file mode 100644 index 00000000000..2945864fa53 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-1.f90 @@ -0,0 +1,50 @@ +program main + use omp_lib + use iso_c_binding + implicit none + + integer, parameter :: N = 40 + integer :: x, i + integer :: y (N) + integer :: u (0) + + x = 24 + !$omp target data map(x) use_device_addr(x) + !$omp target has_device_addr(x) + x = 42; + !$omp end target + !$omp end target data + if (x /= 42) stop 1 + + y = 42 + !$omp target data map(y) use_device_addr(y) + !$omp target has_device_addr(y) + y = [(i, i=1, N)] + !$omp end target + !$omp end target data + if (any (y /= [(i, i = 1, N)])) stop 2 + + !$omp target data map(y(:N)) use_device_addr(y) + !$omp target has_device_addr(y(:N)) + y = [(i+2, i=1, N)] + !$omp end target + !$omp end target data + if (any (y /= [(i+2, i = 1, N)])) stop 3 + + !$omp target data map(y) use_device_addr(y) + !$omp target has_device_addr(y(24:)) + do i = 24, N + y(i) = i + 3 + end do + !$omp end target + !$omp end target data + do i = 24, N + if (y(i) /= i + 3) stop 5 + end do + + !$omp target data map(u) use_device_addr(u) + !$omp target has_device_addr(u) + !$omp end target + !$omp end target data + +end program main diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-2.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-2.f90 new file mode 100644 index 00000000000..a8d78a75af3 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-2.f90 @@ -0,0 +1,40 @@ +program main + use omp_lib + use iso_c_binding + implicit none + + integer, parameter :: N = 5 + integer :: i, x(N), y(N), z(N:2*N-1) + target :: z + + x = 42 + y = 43 + z = 44 + + call foo (x, y, z) + if (any (x /= [(i, i = 1, N)])) stop 1 + if (any (y /= [(2*i, i = 1, N)])) stop 2 + if (any (z /= [(3*i, i = 1, N)])) stop 3 + + contains + subroutine foo(a, b, c) + integer :: a(:) + integer :: b(*) + integer, pointer, intent(in) :: c(:) + + !$omp target data map(a,b(:N),c) use_device_addr(a,b(:N),c) + !$omp target has_device_addr(A,B(:N),C) + if (lbound(a,dim=1) /= 1 .or. ubound(a,dim=1) /= N) stop 10 + if (lbound(b,dim=1) /= 1) stop 11 + if (lbound(c,dim=1) /= N .or. ubound(c,dim=1) /= 2*N-1) stop 12 + if (any (a /= 42)) stop 13 + if (any (b(:N) /= 43)) stop 14 + if (any (c /= 44)) stop 15 + a = [(i, i=1, N)] + b(:N) = [(2*i, i = 1, N)] + c = [(3*i, i = 1, N)] + !$omp end target + !$omp end target data + end subroutine foo + +end program main diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-3.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-3.f90 new file mode 100644 index 00000000000..c6293b4de2e --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-3.f90 @@ -0,0 +1,90 @@ +! Test optional dummy arguments in HAS_DEVICE_ADDR. + +program main + use omp_lib + use iso_c_binding + implicit none + + integer, target :: x + integer, pointer :: ptr + integer, parameter :: N=7 + real :: y1(N), y2(N) + integer, target :: y3(N:2*N-1) + integer :: i + + x = 24 + ptr => x + y1 = 42.24 + y2 = 42.24 + y3 = 42 + + call optional_scalar (is_present=.false.) + if (x /= 24) stop 1 + + call optional_scalar (x, is_present=.true.) + if (x /= 42) stop 2 + + call optional_ptr (is_present=.false.) + if (x /= 42) stop 3 + if (ptr /= 42) stop 4 + + call optional_ptr (ptr, is_present=.true.) + if (x /= 84) stop 5 + if (ptr /= 84) stop 6 + + call optional_array (is_present=.false.) + if (any (y1 /= [(42.24, i=1, N)])) stop 7 + if (any (y2 /= [(42.24, i=1, N)])) stop 8 + if (any (y3 /= [(42, i=1, N)])) stop 9 + + call optional_array (y1, y2, y3, is_present=.true.) + if (any (y1 /= [(42.24+i, i=1, N)])) stop 10 + if (any (y2 /= [(42.24+2*i, i=1, N)])) stop 11 + if (any (y3 /= [(42+3*i, i=1, N)])) stop 12 + +contains + subroutine optional_scalar (a, is_present) + integer, optional :: a + logical, value :: is_present + + !$omp target data map(a) use_device_addr(a) + !$omp target has_device_addr(a) + if (is_present) a = 42 + !$omp end target + !$omp end target data + end subroutine optional_scalar + + subroutine optional_ptr (a, is_present) + integer, pointer, optional :: a + logical, value :: is_present + !$omp target data map(a) use_device_addr(a) + !$omp target has_device_addr(a) + if (is_present) a = 84 + !$omp end target + !$omp end target data + end subroutine optional_ptr + + subroutine optional_array (a, b, c, is_present) + real, optional :: a(:), b(*) + integer, optional, pointer, intent(in) :: c(:) + logical, value :: is_present + integer :: i + + !$omp target data map(a, b(:N), c) use_device_addr(a, b, c) + !$omp target has_device_addr(a, b, c) + if (is_present) then + if (lbound(a,dim=1) /= 1 .or. ubound(a,dim=1) /= N) stop 21 + if (lbound(b,dim=1) /= 1) stop 22 + if (lbound(c,dim=1) /= N .or. ubound(c,dim=1) /= 2*N-1) stop 23 + if (any (a /= [(42.24, i = 1, N)])) stop 24 + if (any (b(:N) /= [(42.24, i = 1, N)])) stop 25 + if (any (c /= [(42, i = 1, N)])) stop 26 + a = [(42.24+i, i=1, N)] + b(:N) = [(42.24+2*i, i=1, N)] + c = [(42+3*i, i=1, N)] + end if + !$omp end target + !$omp end target data + end subroutine optional_array + +end program main diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-4.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-4.f90 new file mode 100644 index 00000000000..59d3e3d31dd --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-4.f90 @@ -0,0 +1,71 @@ +! Test allocatables in HAS_DEVICE_ADDR. + +program main + use omp_lib + use iso_c_binding + implicit none + + integer, parameter :: N = 5 + integer, allocatable :: x + integer, allocatable :: y(:) + call scalar_dummy (x) + call array_dummy (y) + call array_dummy_optional (y) + call array_dummy_optional () + +contains + subroutine scalar_dummy (a) + integer, allocatable :: a + + allocate (a) + a = 24 + + !$omp target data map(a) use_device_addr(a) + !$omp target has_device_addr(a) + a = 42 + !$omp end target + !$omp end target data + if (a /= 42) stop 1 + + deallocate (a) + end subroutine scalar_dummy + + subroutine array_dummy (a) + integer, allocatable :: a(:) + integer :: i + + allocate (a(N)) + a = 42 + + !$omp target data map(a) use_device_addr(a) + !$omp target has_device_addr(a) + a = [(i, i=1, N)] + !$omp end target + !$omp end target data + if (any (a /= [(i, i=1, N)])) stop 2 + + deallocate (a) + end subroutine array_dummy + + subroutine array_dummy_optional (a) + integer, optional, allocatable :: a(:) + integer :: i + + if (present (a)) then + allocate (a(N)) + a = 42 + end if + + !$omp target data map(a) use_device_addr(a) + !$omp target has_device_addr(a) + if (present (a)) a = [(i, i=1, N)] + !$omp end target + !$omp end target data + + if (present (a)) then + if (any (a /= [(i, i=1, N)])) stop 2 + deallocate (a) + end if + end subroutine array_dummy_optional + +end program main