aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--gcc/ChangeLog17
-rw-r--r--gcc/c-family/ChangeLog7
-rw-r--r--gcc/c-family/c-pragma.h5
-rw-r--r--gcc/c/ChangeLog13
-rw-r--r--gcc/c/c-parser.c25
-rw-r--r--gcc/c/c-typeck.c28
-rw-r--r--gcc/cp/ChangeLog14
-rw-r--r--gcc/cp/parser.c16
-rw-r--r--gcc/cp/pt.c2
-rw-r--r--gcc/cp/semantics.c39
-rw-r--r--gcc/gimplify.c4
-rw-r--r--gcc/omp-low.c31
-rw-r--r--gcc/testsuite/ChangeLog6
-rw-r--r--gcc/testsuite/c-c++-common/gomp/target-data-1.c30
-rw-r--r--gcc/tree-core.h5
-rw-r--r--gcc/tree-nested.c2
-rw-r--r--gcc/tree-pretty-print.c3
-rw-r--r--gcc/tree.c3
-rw-r--r--libgomp/ChangeLog12
-rw-r--r--libgomp/testsuite/libgomp.c++/target-9.C43
-rw-r--r--libgomp/testsuite/libgomp.c/target-18.c30
21 files changed, 290 insertions, 45 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index b85a1db..23837ab 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,20 @@
+2019-08-07 Jakub Jelinek <jakub@redhat.com>
+
+ * 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 <linkw@gcc.gnu.org>
* config/rs6000/vector.md (vrotr<mode>3): New define_expand.
diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog
index 693a546..873b9ad 100644
--- a/gcc/c-family/ChangeLog
+++ b/gcc/c-family/ChangeLog
@@ -1,3 +1,10 @@
+2019-08-07 Jakub Jelinek <jakub@redhat.com>
+
+ * 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 <polacek@redhat.com>
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 e8a509f..803dc1b 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 18c4e41..18c1ce0 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,16 @@
+2019-08-07 Jakub Jelinek <jakub@redhat.com>
+
+ * 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 <jakub@redhat.com>
PR c/91192
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index c7c13a6..9b5cfcf 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 %<map%> or %<use_device_ptr%> clause");
+ "one %<map%>, %<use_device_ptr%> or %<use_device_addr%> "
+ "clause");
return NULL_TREE;
}
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 9a1a910..bda90fd 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 9b0a34c..c1e977b 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,17 @@
+2019-08-07 Jakub Jelinek <jakub@redhat.com>
+
+ * 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 <jason@redhat.com>
PR c++/91378 - ICE with noexcept and auto return type.
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 79da7b5..4d07a6a 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 %<map%> or %<use_device_ptr%> clause");
+ "one %<map%>, %<use_device_ptr%> or %<use_device_addr%> "
+ "clause");
return NULL_TREE;
}
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index b71fbaad..b1ad99d 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 fa69624..77e7a6d 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 10b9b68..978df86 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 4a6ea0a..ca8edde 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 a372d7f..d524f72 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2019-08-07 Jakub Jelinek <jakub@redhat.com>
+
+ * 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 <linkw@gcc.gnu.org>
* 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 0d4975b..7aa111a 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 4719562..8216a60 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 5d18957..4df07ed 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 9bea132..8091a0a 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 efa49e9..1f2aac4 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 0331c51..01440f7 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,15 @@
+2019-08-07 Jakub Jelinek <jakub@redhat.com>
+
+ * 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 <jakub@redhat.com>
* 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 a5d171b..83a61cf 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 cbacaf6..dd511fb 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
generated by cgit v1.1 at 2024-07-19 20:30:27 +0000