* [gomp4] declare directive [1/5]
2015-06-08 14:59 [gomp4] declare directive [0/5] James Norris
@ 2015-06-08 15:01 ` James Norris
2015-06-08 15:04 ` [gomp4] declare directive [2/5] James Norris
` (3 subsequent siblings)
4 siblings, 0 replies; 10+ messages in thread
From: James Norris @ 2015-06-08 15:01 UTC (permalink / raw)
To: gcc-patches; +Cc: Thomas Schwinge, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 1 bytes --]
[-- Attachment #2: c.patch --]
[-- Type: text/x-patch, Size: 13058 bytes --]
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index f508b91..83c1432 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -82,6 +82,7 @@ along with GCC; see the file COPYING3. If not see
#include "omp-low.h"
#include "builtins.h"
#include "gomp-constants.h"
+#include "tree-iterator.h"
\f
/* Initialization routine for this file. */
@@ -1472,6 +1473,316 @@ c_parser_external_declaration (c_parser *parser)
}
}
+static tree
+check_oacc_vars_1 (tree *tp, int *, void *l)
+{
+ if (TREE_CODE (*tp) == VAR_DECL && TREE_PUBLIC (*tp))
+ {
+ location_t loc = DECL_SOURCE_LOCATION (*tp);
+ tree attrs;
+ attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (*tp));
+ if (attrs)
+ {
+ tree t;
+
+ for (t = TREE_VALUE (attrs); t; t = TREE_CHAIN (t))
+ {
+ loc = EXPR_LOCATION ((tree) l);
+
+ if (OMP_CLAUSE_MAP_KIND (TREE_VALUE (t)) == GOMP_MAP_LINK)
+ {
+ error_at (loc, "%<link%> clause cannot be used with %qE",
+ *tp);
+ break;
+ }
+ }
+ }
+ else
+ error_at (loc, "no %<#pragma acc declare%> for %qE", *tp);
+ }
+ return NULL_TREE;
+}
+
+static tree
+check_oacc_vars (tree *tp, int *, void *)
+{
+ if (TREE_CODE (*tp) == STATEMENT_LIST)
+ {
+ tree_stmt_iterator i;
+
+ for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
+ {
+ tree t = tsi_stmt (i);
+ walk_tree_without_duplicates (&t, check_oacc_vars_1, t);
+ }
+ }
+
+ return NULL_TREE;
+}
+
+static struct oacc_return
+{
+ tree_stmt_iterator iter;
+ tree stmt;
+ int op;
+ struct oacc_return *next;
+} *oacc_returns;
+
+static tree
+find_oacc_return (tree *tp, int *, void *)
+{
+ if (TREE_CODE (*tp) == STATEMENT_LIST)
+ {
+ tree_stmt_iterator i;
+
+ for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
+ {
+ tree t;
+ struct oacc_return *r;
+
+ t = tsi_stmt (i);
+
+ if (TREE_CODE (t) == RETURN_EXPR)
+ {
+ r = XNEW (struct oacc_return);
+ r->iter = i;
+ r->stmt = NULL_TREE;
+ r->op = 1;
+ r->next = NULL;
+
+ if (oacc_returns)
+ r->next = oacc_returns;
+
+ oacc_returns = r;
+ }
+ else if (TREE_CODE (t) == COND_EXPR)
+ {
+ bool op1, op2;
+ tree op;
+
+ op1 = op2 = false;
+
+ op = TREE_OPERAND (t, 1);
+ op1 = (op && TREE_CODE (op) == RETURN_EXPR);
+
+ op = TREE_OPERAND (t, 2);
+ op2 = (op && TREE_CODE (op) == RETURN_EXPR);
+
+ if (op1 || op2)
+ {
+ r = XNEW (struct oacc_return);
+ r->stmt = t;
+ r->op = op1 ? 1 : 2;
+ r->next = NULL;
+
+ if (oacc_returns)
+ r->next = oacc_returns;
+
+ oacc_returns = r;
+ }
+ }
+ }
+ }
+
+ return NULL_TREE;
+}
+
+static void
+finish_oacc_declare (tree fnbody, tree decls)
+{
+ tree t, stmt, body, c, ret_clauses, clauses;
+ location_t loc;
+ tree_stmt_iterator i;
+ tree fndecl = current_function_decl;
+
+ if (lookup_attribute ("oacc function", DECL_ATTRIBUTES (fndecl)))
+ {
+ if (lookup_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl)))
+ {
+ location_t loc = DECL_SOURCE_LOCATION (fndecl);
+ error_at (loc, "%<#pragma acc declare%> not allowed in %qE", fndecl);
+ }
+
+ walk_tree_without_duplicates (&fnbody, check_oacc_vars, NULL);
+ return;
+ }
+
+ if (!decls)
+ return;
+
+ body = BIND_EXPR_BODY (fnbody);
+
+ if (TREE_CODE (body) != STATEMENT_LIST)
+ {
+ tree list;
+
+ list = alloc_stmt_list ();
+ append_to_statement_list (body, &list);
+ BIND_EXPR_BODY (fnbody) = list;
+ body = list;
+ }
+
+ walk_tree_without_duplicates (&body, find_oacc_return, NULL);
+
+ clauses = NULL_TREE;
+
+ for (t = decls; t; t = TREE_CHAIN (t))
+ {
+ c = TREE_VALUE (TREE_VALUE (t));
+
+ if (clauses)
+ OMP_CLAUSE_CHAIN (c) = clauses;
+ else
+ loc = OMP_CLAUSE_LOCATION (c);
+
+ clauses = c;
+ }
+
+ ret_clauses = NULL_TREE;
+
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ bool ret = false;
+ HOST_WIDE_INT kind, new_op;
+
+ kind = OMP_CLAUSE_MAP_KIND (c);
+
+ switch (kind)
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_TO:
+ new_op = GOMP_MAP_FORCE_DEALLOC;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FORCE_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+ new_op = GOMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FORCE_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO);
+ new_op = GOMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+ new_op = GOMP_MAP_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+ new_op = GOMP_MAP_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_LINK:
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_TO:
+ break;
+
+ default:
+ gcc_unreachable ();
+ break;
+ }
+
+ if (ret)
+ {
+ t = copy_node (c);
+
+ OMP_CLAUSE_SET_MAP_KIND (t, new_op);
+
+ if (ret_clauses)
+ OMP_CLAUSE_CHAIN (t) = ret_clauses;
+
+ ret_clauses = t;
+ }
+ }
+
+ if (clauses)
+ {
+ bool found = false;
+
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DECLARE_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ c = OMP_CLAUSE_DECL (TREE_VALUE (TREE_VALUE (decls)));
+
+ for (i = tsi_start (body); !tsi_end_p (i); tsi_next (&i))
+ {
+ tree it;
+
+ it = tsi_stmt (i);
+
+ if ((TREE_CODE (it) == DECL_EXPR) && (DECL_EXPR_DECL (it) == c))
+ {
+ tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING);
+ found = true;
+ break;
+ }
+ }
+
+ if (!found)
+ {
+ i = tsi_start (body);
+ tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
+ }
+ }
+
+ while (oacc_returns)
+ {
+ struct oacc_return *r;
+
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DECLARE_CLAUSES (stmt) = ret_clauses;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ r = oacc_returns;
+ if (r->stmt)
+ {
+ tree l;
+
+ l = alloc_stmt_list ();
+ append_to_statement_list (stmt, &l);
+ stmt = TREE_OPERAND (r->stmt, r->op);
+ append_to_statement_list (stmt, &l);
+ TREE_OPERAND (r->stmt, r->op) = l;
+ }
+ else
+ tsi_link_before (&r->iter, stmt, TSI_CONTINUE_LINKING);
+
+ oacc_returns = r->next;
+ free (r);
+ }
+
+ for (i = tsi_start (body); !tsi_end_p (i); tsi_next (&i))
+ {
+ if (tsi_end_p (i))
+ break;
+ }
+
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DECLARE_CLAUSES (stmt) = ret_clauses;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
+
+ DECL_ATTRIBUTES (fndecl)
+ = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
+}
+
+
static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec<c_token>);
static void c_finish_oacc_routine (c_parser *, tree, tree, bool);
@@ -2019,6 +2330,9 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
fnbody = c_parser_compound_statement (parser);
if (flag_cilkplus && contains_array_notation_expr (fnbody))
fnbody = expand_array_notation_exprs (fnbody);
+ tree decls = lookup_attribute ("oacc declare",
+ DECL_ATTRIBUTES (current_function_decl));
+ finish_oacc_declare (fnbody, decls);
if (nested)
{
tree decl = current_function_decl;
@@ -12426,6 +12740,8 @@ c_parser_oacc_data (location_t loc, c_parser *parser)
# pragma acc declare oacc-data-clause[optseq] new-line
*/
+static int oacc_dcl_idx = 0;
+
#define OACC_DECLARE_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
@@ -12445,6 +12761,7 @@ c_parser_oacc_declare (c_parser *parser)
{
location_t pragma_loc = c_parser_peek_token (parser)->location;
tree clauses;
+ bool error = false;
c_parser_consume_pragma (parser);
@@ -12460,18 +12777,23 @@ c_parser_oacc_declare (c_parser *parser)
{
location_t loc = OMP_CLAUSE_LOCATION (t);
tree decl = OMP_CLAUSE_DECL (t);
+ tree devres = NULL_TREE;
if (!DECL_P (decl))
{
error_at (loc, "subarray in %<#pragma acc declare%>");
+ error = true;
continue;
}
- gcc_assert (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_MAP);
+
switch (OMP_CLAUSE_MAP_KIND (t))
{
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_DEVICEPTR:
+ break;
+
case GOMP_MAP_DEVICE_RESIDENT:
+ devres = t;
break;
case GOMP_MAP_POINTER:
@@ -12483,8 +12805,10 @@ c_parser_oacc_declare (c_parser *parser)
if (!global_bindings_p () && !DECL_EXTERNAL (decl))
{
error_at (loc,
- "invalid variable %qD in %<#pragma acc declare link%>",
+ "%qD must be a global variable in"
+ "%<#pragma acc declare link%>",
decl);
+ error = true;
continue;
}
break;
@@ -12493,6 +12817,7 @@ c_parser_oacc_declare (c_parser *parser)
if (global_bindings_p ())
{
error_at (loc, "invalid OpenACC clause at file scope");
+ error = true;
continue;
}
if (DECL_EXTERNAL (decl))
@@ -12500,6 +12825,7 @@ c_parser_oacc_declare (c_parser *parser)
error_at (loc,
"invalid use of %<extern%> variable %qD "
"in %<#pragma acc declare%>", decl);
+ error = true;
continue;
}
break;
@@ -12516,17 +12842,23 @@ c_parser_oacc_declare (c_parser *parser)
if (prev_attr)
{
tree p = TREE_VALUE (prev_attr);
- error_at (loc,
- "variable %qD used more than once with "
- "%<#pragma acc declare%>", decl);
- inform (OMP_CLAUSE_LOCATION (TREE_VALUE (p)),
- "previous directive was here");
- continue;
+ tree cl = TREE_VALUE (p);
+
+ if (!devres
+ && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
+ {
+ error_at (loc,
+ "variable %qD used more than once with "
+ "%<#pragma acc declare%>", decl);
+ inform (OMP_CLAUSE_LOCATION (cl),
+ "previous directive was here");
+ error = true;
+ continue;
+ }
}
}
else
{
- bool ok = true;
decl_for_attr = current_function_decl;
tree prev_attr = lookup_attribute ("oacc declare",
DECL_ATTRIBUTES (decl_for_attr));
@@ -12544,17 +12876,82 @@ c_parser_oacc_declare (c_parser *parser)
"%<#pragma acc declare%>", decl);
inform (OMP_CLAUSE_LOCATION (cl),
"previous directive was here");
- ok = false;
+ error = true;
break;
}
}
- if (!ok)
- continue;
}
- tree attr = tree_cons (NULL_TREE, t, NULL_TREE);
- tree attrs = tree_cons (get_identifier ("oacc declare"),
- attr, NULL_TREE);
- decl_attributes (&decl_for_attr, attrs, 0);
+
+ if (!error)
+ {
+ tree attr = tree_cons (NULL_TREE, t, NULL_TREE);
+ tree attrs = tree_cons (get_identifier ("oacc declare"),
+ attr, NULL_TREE);
+ decl_attributes (&decl_for_attr, attrs, 0);
+ }
+ }
+
+ if (error)
+ return;
+
+ if (global_bindings_p ())
+ {
+ char buf[128];
+ struct c_declarator *target;
+ tree stmt, attrs;
+ c_arg_info *arg_info = build_arg_info ();
+ struct c_declarator *declarator;
+ struct c_declspecs *specs;
+ struct c_typespec spec;
+ location_t loc = UNKNOWN_LOCATION;
+ tree f, t, fnbody, call_fn;
+
+ sprintf (buf, "__openacc_c_constructor__%d", oacc_dcl_idx++);
+ target = build_id_declarator (get_identifier (buf));
+ arg_info->types = void_list_node;
+ declarator = build_function_declarator (arg_info, target);
+
+ specs = build_null_declspecs ();
+ spec.kind = ctsk_resword;
+ spec.spec = get_identifier ("void");
+ spec.expr = NULL_TREE;
+ spec.expr_const_operands = true;
+
+ declspecs_add_type (pragma_loc, specs, spec);
+ finish_declspecs (specs);
+
+ attrs = tree_cons (get_identifier ("constructor") , NULL_TREE, NULL_TREE);
+ start_function (specs, declarator, attrs);
+ store_parm_decls ();
+ f = c_begin_compound_stmt (true);
+ TREE_USED (current_function_decl) = 1;
+ call_fn = builtin_decl_explicit (BUILT_IN_GOACC_STATIC);
+ TREE_SIDE_EFFECTS (call_fn) = 1;
+
+ for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+ {
+ tree d, a1, a2, a3;
+ vec<tree, va_gc> *args;
+ vec_alloc (args, 3);
+
+ d = OMP_CLAUSE_DECL (t);
+
+ a1 = build_unary_op (loc, ADDR_EXPR, d, 0);
+ a2 = DECL_SIZE_UNIT (d);
+ a3 = build_int_cst (unsigned_type_node, OMP_CLAUSE_MAP_KIND (t));
+
+ args->quick_push (a1);
+ args->quick_push (a2);
+ args->quick_push (a3);
+
+ stmt = build_function_call_vec (loc, vNULL, call_fn, args, NULL);
+ add_stmt (stmt);
+ }
+
+ fnbody = c_end_compound_stmt (loc, f, true);
+ add_stmt (fnbody);
+
+ finish_function ();
}
}
^ permalink raw reply [flat|nested] 10+ messages in thread* Re: [gomp4] declare directive [2/5]
2015-06-08 14:59 [gomp4] declare directive [0/5] James Norris
2015-06-08 15:01 ` [gomp4] declare directive [1/5] James Norris
@ 2015-06-08 15:04 ` James Norris
2015-06-08 15:05 ` [gomp4] declare directive [3/5] James Norris
` (2 subsequent siblings)
4 siblings, 0 replies; 10+ messages in thread
From: James Norris @ 2015-06-08 15:04 UTC (permalink / raw)
To: gcc-patches; +Cc: Thomas Schwinge, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 1 bytes --]
[-- Attachment #2: cp.patch --]
[-- Type: text/x-patch, Size: 14733 bytes --]
diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c
index 261a12d..15da51e 100644
--- a/gcc/cp/decl.c
+++ b/gcc/cp/decl.c
@@ -78,6 +78,7 @@ along with GCC; see the file COPYING3. If not see
#include "cilk.h"
#include "wide-int.h"
#include "builtins.h"
+#include "gomp-constants.h"
/* Possible cases of bad specifiers type used by bad_specifiers. */
enum bad_spec_place {
@@ -14113,6 +14114,314 @@ maybe_save_function_definition (tree fun)
register_constexpr_fundef (fun, DECL_SAVED_TREE (fun));
}
+static tree
+check_oacc_vars_1 (tree *tp, int *, void *l)
+{
+ if (TREE_CODE (*tp) == VAR_DECL && TREE_PUBLIC (*tp))
+ {
+ location_t loc = DECL_SOURCE_LOCATION (*tp);
+ tree attrs;
+ attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (*tp));
+ if (attrs)
+ {
+ tree t;
+
+ for (t = TREE_VALUE (attrs); t; t = TREE_CHAIN (t))
+ {
+ loc = EXPR_LOCATION ((tree) l);
+
+ if (OMP_CLAUSE_MAP_KIND (TREE_VALUE (t)) == GOMP_MAP_LINK)
+ {
+ error_at (loc, "%<link%> clause cannot be used with %qE",
+ *tp);
+ break;
+ }
+ }
+ }
+ else
+ error_at (loc, "no %<#pragma acc declare%> for %qE", *tp);
+ }
+ return NULL_TREE;
+}
+
+static tree
+check_oacc_vars (tree *tp, int *, void *)
+{
+ if (TREE_CODE (*tp) == STATEMENT_LIST)
+ {
+ tree_stmt_iterator i;
+
+ for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
+ {
+ tree t = tsi_stmt (i);
+ walk_tree_without_duplicates (&t, check_oacc_vars_1, t);
+ }
+ }
+
+ return NULL_TREE;
+}
+
+static struct oacc_return
+{
+ tree_stmt_iterator iter;
+ tree stmt;
+ int op;
+ struct oacc_return *next;
+} *oacc_returns;
+
+static tree
+find_oacc_return (tree *tp, int *, void *)
+{
+ if (TREE_CODE (*tp) == STATEMENT_LIST)
+ {
+ tree_stmt_iterator i;
+
+ for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
+ {
+ tree t;
+ struct oacc_return *r;
+
+ t = tsi_stmt (i);
+
+ if (TREE_CODE (t) == RETURN_EXPR)
+ {
+ r = XNEW (struct oacc_return);
+ r->iter = i;
+ r->stmt = NULL_TREE;
+ r->op = 1;
+ r->next = NULL;
+
+ if (oacc_returns)
+ r->next = oacc_returns;
+
+ oacc_returns = r;
+ }
+ else if (TREE_CODE (t) == IF_STMT)
+ {
+ bool op1, op2;
+ tree op;
+
+ op1 = op2 = false;
+
+ op = TREE_OPERAND (t, 1);
+ op1 = (op && TREE_CODE (op) == RETURN_EXPR);
+
+ op = TREE_OPERAND (t, 2);
+ op2 = (op && TREE_CODE (op) == RETURN_EXPR);
+
+ if (op1 || op2)
+ {
+ r = XNEW (struct oacc_return);
+ r->stmt = t;
+ r->op = op1 ? 1 : 2;
+ r->next = NULL;
+
+ if (oacc_returns)
+ r->next = oacc_returns;
+
+ oacc_returns = r;
+ }
+ }
+ }
+ }
+
+ return NULL_TREE;
+}
+
+static void
+finish_oacc_declare (tree fndecl, tree decls)
+{
+ tree t, stmt, list, c, ret_clauses, clauses;
+ location_t loc;
+ tree_stmt_iterator i;
+
+ list = cur_stmt_list;
+
+ if (lookup_attribute ("oacc function", DECL_ATTRIBUTES (fndecl)))
+ {
+ if (lookup_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl)))
+ {
+ location_t loc = DECL_SOURCE_LOCATION (fndecl);
+ error_at (loc, "%<#pragma acc declare%> not allowed in %qE", fndecl);
+ }
+
+ walk_tree_without_duplicates (&list, check_oacc_vars, NULL);
+ return;
+ }
+
+ if (!decls)
+ return;
+
+ walk_tree_without_duplicates (&list, find_oacc_return, NULL);
+
+ clauses = NULL_TREE;
+
+ for (t = decls; t; t = TREE_CHAIN (t))
+ {
+ c = TREE_VALUE (TREE_VALUE (t));
+
+ if (clauses)
+ OMP_CLAUSE_CHAIN (c) = clauses;
+ else
+ loc = OMP_CLAUSE_LOCATION (c);
+
+ clauses = c;
+ }
+
+ ret_clauses = NULL_TREE;
+
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ bool ret = false;
+ HOST_WIDE_INT kind, new_op;
+
+ kind = OMP_CLAUSE_MAP_KIND (c);
+
+ switch (kind)
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_TO:
+ new_op = GOMP_MAP_FORCE_DEALLOC;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FORCE_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+ new_op = GOMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FORCE_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO);
+ new_op = GOMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+ new_op = GOMP_MAP_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+ new_op = GOMP_MAP_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_TO:
+ break;
+
+ case GOMP_MAP_LINK:
+ continue;
+
+ default:
+ gcc_unreachable ();
+ break;
+ }
+
+ if (ret)
+ {
+ t = copy_node (c);
+
+ OMP_CLAUSE_SET_MAP_KIND (t, new_op);
+
+ if (ret_clauses)
+ OMP_CLAUSE_CHAIN (t) = ret_clauses;
+
+ ret_clauses = t;
+ }
+ }
+
+ i = tsi_start (list);
+ if (!tsi_end_p (i))
+ {
+ t = tsi_stmt (i);
+ if (TREE_CODE (t) == BIND_EXPR)
+ list = BIND_EXPR_BODY (t);
+ }
+
+ if (clauses)
+ {
+ bool found = false;
+
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OMP_STANDALONE_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ c = OMP_CLAUSE_DECL (TREE_VALUE (TREE_VALUE (decls)));
+
+ for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i))
+ {
+ tree it;
+
+ it = tsi_stmt (i);
+
+ if ((TREE_CODE (it) == DECL_EXPR) && (DECL_EXPR_DECL (it) == c))
+ {
+ tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING);
+ found = true;
+ break;
+ }
+ }
+
+ if (!found)
+ {
+ i = tsi_start (list);
+ tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
+ }
+ }
+
+ while (oacc_returns)
+ {
+ struct oacc_return *r;
+
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OMP_STANDALONE_CLAUSES (stmt) = ret_clauses;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ r = oacc_returns;
+ if (r->stmt)
+ {
+ tree l;
+
+ l = alloc_stmt_list ();
+ append_to_statement_list (stmt, &l);
+ stmt = TREE_OPERAND (r->stmt, r->op);
+ append_to_statement_list (stmt, &l);
+ TREE_OPERAND (r->stmt, r->op) = l;
+ }
+ else
+ tsi_link_before (&r->iter, stmt, TSI_CONTINUE_LINKING);
+
+ oacc_returns = r->next;
+ free (r);
+ }
+
+ for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i))
+ {
+ if (tsi_end_p (i))
+ break;
+ }
+
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OMP_STANDALONE_CLAUSES (stmt) = ret_clauses;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
+
+ DECL_ATTRIBUTES (fndecl)
+ = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
+}
+
/* Finish up a function declaration and compile that function
all the way to assembler language output. The free the storage
for the function definition.
@@ -14141,6 +14450,9 @@ finish_function (int flags)
gcc_assert (!defer_mark_used_calls);
defer_mark_used_calls = true;
+ tree decls = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
+ finish_oacc_declare (fndecl, decls);
+
record_key_method_defined (fndecl);
fntype = TREE_TYPE (fndecl);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 2947bf4..fb6b7ed 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -61,6 +61,7 @@ along with GCC; see the file COPYING3. If not see
#include "type-utils.h"
#include "omp-low.h"
#include "gomp-constants.h"
+#include "tree-iterator.h"
\f
/* The lexer. */
@@ -32035,6 +32036,221 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok)
return stmt;
}
+/* OpenACC 2.0:
+ # pragma acc declare oacc-data-clause[optseq] new-line
+*/
+
+static int oacc_dcl_idx = 0;
+
+#define OACC_DECLARE_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE))
+
+static tree
+cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree clauses;
+ bool error = false;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
+ "#pragma acc declare", pragma_tok);
+
+ if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+ {
+ error_at (pragma_tok->location,
+ "no valid clauses specified in %<#pragma acc declare%>");
+ return NULL_TREE;
+ }
+
+ for (tree t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+ {
+ location_t loc = OMP_CLAUSE_LOCATION (t);
+ tree decl = OMP_CLAUSE_DECL (t);
+ tree devres = NULL_TREE;
+ if (!DECL_P (decl))
+ {
+ error_at (loc, "subarray in %<#pragma acc declare%>");
+ error = true;
+ continue;
+ }
+ gcc_assert (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_MAP);
+ switch (OMP_CLAUSE_MAP_KIND (t))
+ {
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ break;
+
+ case GOMP_MAP_DEVICE_RESIDENT:
+ devres = t;
+ break;
+
+ case GOMP_MAP_POINTER:
+ /* Generated by c_finish_omp_clauses from array sections;
+ avoid spurious diagnostics. */
+ break;
+
+ case GOMP_MAP_LINK:
+ if (!global_bindings_p () && !DECL_EXTERNAL (decl))
+ {
+ error_at (loc,
+ "%qD must be a global variable in"
+ "%<#pragma acc declare link%>",
+ decl);
+ error = true;
+ continue;
+ }
+ break;
+
+ default:
+ if (global_bindings_p ())
+ {
+ error_at (loc, "invalid OpenACC clause at file scope");
+ error = true;
+ continue;
+ }
+ if (DECL_EXTERNAL (decl))
+ {
+ error_at (loc,
+ "invalid use of %<extern%> variable %qD "
+ "in %<#pragma acc declare%>", decl);
+ error = true;
+ continue;
+ }
+ break;
+ }
+
+ /* Store the clause in an attribute on the variable, at file
+ scope, or the function, at block scope. */
+ tree decl_for_attr;
+ if (global_bindings_p ())
+ {
+ decl_for_attr = decl;
+ tree prev_attr = lookup_attribute ("oacc declare",
+ DECL_ATTRIBUTES (decl));
+ if (prev_attr)
+ {
+ tree p = TREE_VALUE (prev_attr);
+ tree cl = TREE_VALUE (p);
+
+ if (!devres
+ && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
+ {
+ error_at (loc,
+ "variable %qD used more than once with "
+ "%<#pragma acc declare%>", decl);
+ inform (OMP_CLAUSE_LOCATION (TREE_VALUE (p)),
+ "previous directive was here");
+ error = true;
+ continue;
+ }
+ }
+ }
+ else
+ {
+ decl_for_attr = current_function_decl;
+ tree prev_attr = lookup_attribute ("oacc declare",
+ DECL_ATTRIBUTES (decl_for_attr));
+ for (;
+ prev_attr;
+ prev_attr = lookup_attribute ("oacc declare",
+ TREE_CHAIN (prev_attr)))
+ {
+ tree p = TREE_VALUE (prev_attr);
+ tree cl = TREE_VALUE (p);
+ if (OMP_CLAUSE_DECL (cl) == decl)
+ {
+ error_at (loc,
+ "variable %qD used more than once with "
+ "%<#pragma acc declare%>", decl);
+ inform (OMP_CLAUSE_LOCATION (cl),
+ "previous directive was here");
+ error = true;
+ break;
+ }
+ }
+ }
+
+ if (!error)
+ {
+ tree attr = tree_cons (NULL_TREE, t, NULL_TREE);
+ tree attrs = tree_cons (get_identifier ("oacc declare"),
+ attr, NULL_TREE);
+ decl_attributes (&decl_for_attr, attrs, 0);
+ }
+ }
+
+ if (error)
+ return NULL_TREE;
+
+ if (global_bindings_p ())
+ {
+ char buf[128];
+ cp_decl_specifier_seq decl_specifiers;
+ cp_declarator *declarator;
+ tree attrs, parms;
+ tree f, t, call_fn, stmt;
+ location_t loc = UNKNOWN_LOCATION;
+ void *p;
+
+ p = obstack_alloc (&declarator_obstack, 0);
+ clear_decl_specs (&decl_specifiers);
+ decl_specifiers.type = void_type_node;
+ sprintf (buf, "__openacc_cp_constructor__%d", oacc_dcl_idx++);
+
+ declarator = make_id_declarator (NULL_TREE, get_identifier (buf),
+ sfk_none);
+ parms = void_list_node;
+ declarator = make_call_declarator (declarator, parms,
+ TYPE_UNQUALIFIED,
+ VIRT_SPEC_UNSPECIFIED,
+ REF_QUAL_NONE,
+ NULL_TREE,
+ NULL_TREE);
+ attrs = tree_cons (get_identifier ("constructor") , NULL_TREE, NULL_TREE);
+ start_function (&decl_specifiers, declarator, attrs);
+ f = begin_compound_stmt (0);
+ call_fn = builtin_decl_explicit (BUILT_IN_GOACC_STATIC);
+ TREE_SIDE_EFFECTS (call_fn) = 1;
+
+ for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+ {
+ tree d, a1, a2, a3;
+ vec<tree, va_gc> *args;
+ vec_alloc (args, 3);
+
+ d = OMP_CLAUSE_DECL (t);
+
+ a1 = build_unary_op (loc, ADDR_EXPR, d, 0);
+ a2 = DECL_SIZE_UNIT (d);
+ a3 = build_int_cst (unsigned_type_node, OMP_CLAUSE_MAP_KIND (t));
+
+ args->quick_push (a1);
+ args->quick_push (a2);
+ args->quick_push (a3);
+
+ stmt = build_function_call_vec (loc, vNULL, call_fn, args, NULL);
+ finish_expr_stmt (stmt);
+ }
+
+ finish_compound_stmt (f);
+ expand_or_defer_fn (finish_function (0));
+ obstack_free (&declarator_obstack, p);
+ }
+
+ return NULL_TREE;
+}
+
#define OACC_HOST_DATA_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
@@ -33903,6 +34119,10 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
cp_parser_omp_declare (parser, pragma_tok, context);
return false;
+ case PRAGMA_OACC_DECLARE:
+ cp_parser_oacc_declare (parser, pragma_tok);
+ return false;
+
case PRAGMA_OACC_ENTER_DATA:
if (context == pragma_stmt)
{
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index caafb43..f6e5c3b 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -14230,6 +14230,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
break;
case OMP_TARGET_UPDATE:
+ case OACC_DECLARE:
case OACC_ENTER_DATA:
case OACC_EXIT_DATA:
case OACC_UPDATE:
^ permalink raw reply [flat|nested] 10+ messages in thread* [gomp4] declare directive [3/5]
2015-06-08 14:59 [gomp4] declare directive [0/5] James Norris
2015-06-08 15:01 ` [gomp4] declare directive [1/5] James Norris
2015-06-08 15:04 ` [gomp4] declare directive [2/5] James Norris
@ 2015-06-08 15:05 ` James Norris
2015-06-17 10:04 ` Thomas Schwinge
2015-10-30 13:30 ` Thomas Schwinge
2015-06-08 15:06 ` [gomp4] declare directive [5/5] James Norris
2015-06-08 15:06 ` [gomp4] declare directive [4/5] James Norris
4 siblings, 2 replies; 10+ messages in thread
From: James Norris @ 2015-06-08 15:05 UTC (permalink / raw)
To: gcc-patches; +Cc: Thomas Schwinge, Tobias Burnus
[-- Attachment #1: Type: text/plain, Size: 1 bytes --]
[-- Attachment #2: fortran.patch --]
[-- Type: text/x-patch, Size: 25559 bytes --]
diff --git a/gcc/fortran/f95-lang.c b/gcc/fortran/f95-lang.c
index 5003581..a889342 100644
--- a/gcc/fortran/f95-lang.c
+++ b/gcc/fortran/f95-lang.c
@@ -119,6 +119,8 @@ static const struct attribute_spec gfc_attribute_table[] =
affects_type_identity } */
{ "omp declare target", 0, 0, true, false, false,
gfc_handle_omp_declare_target_attribute, false },
+ { "oacc declare", 0, 0, true, false, false,
+ gfc_handle_omp_declare_target_attribute, false },
{ "oacc function", 0, 0, true, false, false,
gfc_handle_omp_declare_target_attribute, false },
{ NULL, 0, 0, false, false, false, NULL, false }
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index e73c269..a90b0f8 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -861,6 +861,13 @@ typedef struct
/* Mentioned in OMP DECLARE TARGET. */
unsigned omp_declare_target:1;
+ /* Mentioned in OACC DECLARE. */
+ unsigned oacc_declare_create:1;
+ unsigned oacc_declare_copyin:1;
+ unsigned oacc_declare_deviceptr:1;
+ unsigned oacc_declare_device_resident:1;
+ unsigned oacc_declare_link:1;
+
/* This is an OpenACC acclerator function. */
unsigned oacc_function:1;
@@ -1132,6 +1139,8 @@ typedef enum
OMP_MAP_FORCE_TOFROM,
OMP_MAP_FORCE_PRESENT,
OMP_MAP_FORCE_DEVICEPTR,
+ OMP_MAP_DEVICE_RESIDENT,
+ OMP_MAP_LINK,
OMP_MAP_FORCE_TO_GANGLOCAL
}
gfc_omp_map_op;
@@ -1174,6 +1183,7 @@ enum
OMP_LIST_FROM,
OMP_LIST_REDUCTION,
OMP_LIST_DEVICE_RESIDENT,
+ OMP_LIST_LINK,
OMP_LIST_USE_DEVICE,
OMP_LIST_CACHE,
OMP_LIST_NUM
@@ -1269,6 +1279,7 @@ typedef struct gfc_oacc_declare
{
struct gfc_oacc_declare *next;
locus where;
+ bool module_var;
gfc_omp_clauses *clauses;
}
gfc_oacc_declare;
@@ -3276,6 +3287,6 @@ void gfc_convert_mpz_to_signed (mpz_t, int);
/* trans-decl.c */
-void insert_oacc_declare (gfc_namespace *);
+void finish_oacc_declare (gfc_namespace *, enum sym_flavor);
#endif /* GCC_GFORTRAN_H */
diff --git a/gcc/fortran/module.c b/gcc/fortran/module.c
index 1abfc46..c174902 100644
--- a/gcc/fortran/module.c
+++ b/gcc/fortran/module.c
@@ -1894,7 +1894,9 @@ typedef enum
AB_IS_CLASS, AB_PROCEDURE, AB_PROC_POINTER, AB_ASYNCHRONOUS, AB_CODIMENSION,
AB_COARRAY_COMP, AB_VTYPE, AB_VTAB, AB_CONTIGUOUS, AB_CLASS_POINTER,
AB_IMPLICIT_PURE, AB_ARTIFICIAL, AB_UNLIMITED_POLY, AB_OMP_DECLARE_TARGET,
- AB_ARRAY_OUTER_DEPENDENCY
+ AB_ARRAY_OUTER_DEPENDENCY, AB_OACC_DECLARE_CREATE, AB_OACC_DECLARE_COPYIN,
+ AB_OACC_DECLARE_DEVICEPTR, AB_OACC_DECLARE_DEVICE_RESIDENT,
+ AB_OACC_DECLARE_LINK
}
ab_attribute;
@@ -1951,6 +1953,11 @@ static const mstring attr_bits[] =
minit ("UNLIMITED_POLY", AB_UNLIMITED_POLY),
minit ("OMP_DECLARE_TARGET", AB_OMP_DECLARE_TARGET),
minit ("ARRAY_OUTER_DEPENDENCY", AB_ARRAY_OUTER_DEPENDENCY),
+ minit ("OACC_DECLARE_CREATE", AB_OACC_DECLARE_CREATE),
+ minit ("OACC_DECLARE_COPYIN", AB_OACC_DECLARE_COPYIN),
+ minit ("OACC_DECLARE_DEVICEPTR", AB_OACC_DECLARE_DEVICEPTR),
+ minit ("OACC_DECLARE_DEVICE_RESIDENT", AB_OACC_DECLARE_DEVICE_RESIDENT),
+ minit ("OACC_DECLARE_LINK", AB_OACC_DECLARE_LINK),
minit (NULL, -1)
};
@@ -2133,6 +2140,16 @@ mio_symbol_attribute (symbol_attribute *attr)
MIO_NAME (ab_attribute) (AB_OMP_DECLARE_TARGET, attr_bits);
if (attr->array_outer_dependency)
MIO_NAME (ab_attribute) (AB_ARRAY_OUTER_DEPENDENCY, attr_bits);
+ if (attr->oacc_declare_create)
+ MIO_NAME (ab_attribute) (AB_OACC_DECLARE_CREATE, attr_bits);
+ if (attr->oacc_declare_copyin)
+ MIO_NAME (ab_attribute) (AB_OACC_DECLARE_COPYIN, attr_bits);
+ if (attr->oacc_declare_deviceptr)
+ MIO_NAME (ab_attribute) (AB_OACC_DECLARE_DEVICEPTR, attr_bits);
+ if (attr->oacc_declare_device_resident)
+ MIO_NAME (ab_attribute) (AB_OACC_DECLARE_DEVICE_RESIDENT, attr_bits);
+ if (attr->oacc_declare_link)
+ MIO_NAME (ab_attribute) (AB_OACC_DECLARE_LINK, attr_bits);
mio_rparen ();
@@ -2302,6 +2319,21 @@ mio_symbol_attribute (symbol_attribute *attr)
case AB_ARRAY_OUTER_DEPENDENCY:
attr->array_outer_dependency =1;
break;
+ case AB_OACC_DECLARE_CREATE:
+ attr->oacc_declare_create = 1;
+ break;
+ case AB_OACC_DECLARE_COPYIN:
+ attr->oacc_declare_copyin = 1;
+ break;
+ case AB_OACC_DECLARE_DEVICEPTR:
+ attr->oacc_declare_deviceptr = 1;
+ break;
+ case AB_OACC_DECLARE_DEVICE_RESIDENT:
+ attr->oacc_declare_device_resident = 1;
+ break;
+ case AB_OACC_DECLARE_LINK:
+ attr->oacc_declare_link = 1;
+ break;
}
}
}
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index fc16d8c..46bf865 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -475,6 +475,7 @@ match_oacc_clause_gang (gfc_omp_clauses *cp)
#define OMP_CLAUSE_BIND ((uint64_t) 1 << 58)
#define OMP_CLAUSE_NOHOST ((uint64_t) 1 << 59)
#define OMP_CLAUSE_DEVICE_TYPE ((uint64_t) 1 << 60)
+#define OMP_CLAUSE_LINK ((uint64_t) 1 << 61)
/* Helper function for OpenACC and OpenMP clauses involving memory
mapping. */
@@ -749,6 +750,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
true)
== MATCH_YES)
continue;
+ if ((mask & OMP_CLAUSE_LINK)
+ && gfc_match_omp_variable_list ("link (",
+ &c->lists[OMP_LIST_LINK],
+ true)
+ == MATCH_YES)
+ continue;
if ((mask & OMP_CLAUSE_OACC_DEVICE)
&& gfc_match ("device ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
@@ -1352,7 +1359,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
| OMP_CLAUSE_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_DEVICE_RESIDENT \
| OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \
| OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
- | OMP_CLAUSE_PRESENT_OR_CREATE)
+ | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_LINK)
#define OACC_UPDATE_CLAUSES \
(OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST \
| OMP_CLAUSE_OACC_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
@@ -1501,12 +1508,18 @@ gfc_match_oacc_declare (void)
gfc_omp_namelist *n;
gfc_namespace *ns = gfc_current_ns;
gfc_oacc_declare *new_oc, *oc;
- locus where = gfc_current_locus;
+ bool module_var = false;
if (gfc_match_omp_clauses (&c, OACC_DECLARE_CLAUSES, 0, false, false, true)
!= MATCH_YES)
return MATCH_ERROR;
+ for (n = c->lists[OMP_LIST_DEVICE_RESIDENT]; n != NULL; n = n->next)
+ n->sym->attr.oacc_declare_device_resident = 1;
+
+ for (n = c->lists[OMP_LIST_LINK]; n != NULL; n = n->next)
+ n->sym->attr.oacc_declare_link = 1;
+
for (n = c->lists[OMP_LIST_MAP]; n != NULL; n = n->next)
{
gfc_symbol *s = n->sym;
@@ -1520,6 +1533,14 @@ gfc_match_oacc_declare (void)
"$!ACC DECLARE at %C");
return MATCH_ERROR;
}
+
+ module_var = true;
+ }
+
+ if (ns->proc_name->attr.oacc_function)
+ {
+ gfc_error ("Invalid declare in routine with " "$!ACC DECLARE at %C");
+ return MATCH_ERROR;
}
if (s->attr.in_common)
@@ -1543,12 +1564,31 @@ gfc_match_oacc_declare (void)
"$!ACC DECLARE at %C");
return MATCH_ERROR;
}
+
+ switch (n->u.map_op)
+ {
+ case OMP_MAP_FORCE_ALLOC:
+ s->attr.oacc_declare_create = 1;
+ break;
+
+ case OMP_MAP_FORCE_TO:
+ s->attr.oacc_declare_copyin = 1;
+ break;
+
+ case OMP_MAP_FORCE_DEVICEPTR:
+ s->attr.oacc_declare_deviceptr = 1;
+ break;
+
+ default:
+ break;
+ }
}
new_oc = gfc_get_oacc_declare ();
new_oc->next = ns->oacc_declare;
- new_oc->where = where;
+ new_oc->module_var = module_var;
new_oc->clauses = c;
+ new_oc->where = gfc_current_locus;
for (oc = new_oc; oc; oc = oc->next)
{
@@ -4961,6 +5001,33 @@ gfc_resolve_oacc_declare (gfc_namespace *ns)
n->sym->name, &loc);
}
}
+
+ for (oc = ns->oacc_declare; oc; oc = oc->next)
+ {
+ for (list = OMP_LIST_LINK; list <= OMP_LIST_LINK; list++)
+ for (n = oc->clauses->lists[list]; n; n = n->next)
+ n->sym->mark = 0;
+ }
+
+ for (oc = ns->oacc_declare; oc; oc = oc->next)
+ {
+ for (list = OMP_LIST_LINK; list <= OMP_LIST_LINK; list++)
+ for (n = oc->clauses->lists[list]; n; n = n->next)
+ {
+ if (n->sym->mark)
+ gfc_error ("Symbol %qs present on multiple clauses at %L",
+ n->sym->name, &loc);
+ else
+ n->sym->mark = 1;
+ }
+ }
+
+ for (oc = ns->oacc_declare; oc; oc = oc->next)
+ {
+ for (list = OMP_LIST_LINK; list <= OMP_LIST_LINK; list++)
+ for (n = oc->clauses->lists[list]; n; n = n->next)
+ n->sym->mark = 0;
+ }
}
diff --git a/gcc/fortran/symbol.c b/gcc/fortran/symbol.c
index b18608b..1ecc16d 100644
--- a/gcc/fortran/symbol.c
+++ b/gcc/fortran/symbol.c
@@ -375,6 +375,11 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
*contiguous = "CONTIGUOUS", *generic = "GENERIC";
static const char *threadprivate = "THREADPRIVATE";
static const char *omp_declare_target = "OMP DECLARE TARGET";
+ static const char *oacc_declare_copyin = "OACC DECLARE COPYIN";
+ static const char *oacc_declare_create = "OACC DECLARE CREATE";
+ static const char *oacc_declare_deviceptr = "OACC DECLARE DEVICEPTR";
+ static const char *oacc_declare_device_resident =
+ "OACC DECLARE DEVICE_RESIDENT";
const char *a1, *a2;
int standard;
@@ -506,6 +511,10 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
conf (in_equivalence, allocatable);
conf (in_equivalence, threadprivate);
conf (in_equivalence, omp_declare_target);
+ conf (in_equivalence, oacc_declare_create);
+ conf (in_equivalence, oacc_declare_copyin);
+ conf (in_equivalence, oacc_declare_deviceptr);
+ conf (in_equivalence, oacc_declare_device_resident);
conf (dummy, result);
conf (entry, result);
@@ -555,6 +564,10 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
conf (cray_pointee, in_equivalence);
conf (cray_pointee, threadprivate);
conf (cray_pointee, omp_declare_target);
+ conf (cray_pointee, oacc_declare_create);
+ conf (cray_pointee, oacc_declare_copyin);
+ conf (cray_pointee, oacc_declare_deviceptr);
+ conf (cray_pointee, oacc_declare_device_resident);
conf (data, dummy);
conf (data, function);
@@ -609,6 +622,10 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
conf (proc_pointer, abstract)
conf (entry, omp_declare_target)
+ conf (entry, oacc_declare_create)
+ conf (entry, oacc_declare_copyin)
+ conf (entry, oacc_declare_deviceptr)
+ conf (entry, oacc_declare_device_resident)
a1 = gfc_code2string (flavors, attr->flavor);
@@ -646,6 +663,10 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
conf2 (subroutine);
conf2 (threadprivate);
conf2 (omp_declare_target);
+ conf2 (oacc_declare_create);
+ conf2 (oacc_declare_copyin);
+ conf2 (oacc_declare_deviceptr);
+ conf2 (oacc_declare_device_resident);
if (attr->access == ACCESS_PUBLIC || attr->access == ACCESS_PRIVATE)
{
@@ -728,6 +749,10 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
conf2 (threadprivate);
conf2 (result);
conf2 (omp_declare_target);
+ conf2 (oacc_declare_create);
+ conf2 (oacc_declare_copyin);
+ conf2 (oacc_declare_deviceptr);
+ conf2 (oacc_declare_device_resident);
if (attr->intent != INTENT_UNKNOWN)
{
@@ -1239,6 +1264,62 @@ gfc_add_omp_declare_target (symbol_attribute *attr, const char *name,
bool
+gfc_add_oacc_declare_create (symbol_attribute *attr, const char *name, locus *where)
+{
+ if (check_used (attr, name, where))
+ return false;
+
+ if (attr->oacc_declare_create)
+ return true;
+
+ attr->oacc_declare_create = 1;
+ return check_conflict (attr, name, where);
+}
+
+
+bool
+gfc_add_oacc_declare_copyin (symbol_attribute *attr, const char *name, locus *where)
+{
+ if (check_used (attr, name, where))
+ return false;
+
+ if (attr->oacc_declare_copyin)
+ return true;
+
+ attr->oacc_declare_copyin = 1;
+ return check_conflict (attr, name, where);
+}
+
+
+bool
+gfc_add_oacc_declare_deviceptr (symbol_attribute *attr, const char *name, locus *where)
+{
+ if (check_used (attr, name, where))
+ return false;
+
+ if (attr->oacc_declare_deviceptr)
+ return true;
+
+ attr->oacc_declare_deviceptr = 1;
+ return check_conflict (attr, name, where);
+}
+
+
+bool
+gfc_add_oacc_declare_device_resident (symbol_attribute *attr, const char *name, locus *where)
+{
+ if (check_used (attr, name, where))
+ return false;
+
+ if (attr->oacc_declare_device_resident)
+ return true;
+
+ attr->oacc_declare_device_resident = 1;
+ return check_conflict (attr, name, where);
+}
+
+
+bool
gfc_add_target (symbol_attribute *attr, locus *where)
{
@@ -1796,6 +1877,18 @@ gfc_copy_attr (symbol_attribute *dest, symbol_attribute *src, locus *where)
if (src->omp_declare_target
&& !gfc_add_omp_declare_target (dest, NULL, where))
goto fail;
+ if (src->oacc_declare_create
+ && !gfc_add_oacc_declare_create (dest, NULL, where))
+ goto fail;
+ if (src->oacc_declare_copyin
+ && !gfc_add_oacc_declare_copyin (dest, NULL, where))
+ goto fail;
+ if (src->oacc_declare_deviceptr
+ && !gfc_add_oacc_declare_deviceptr (dest, NULL, where))
+ goto fail;
+ if (src->oacc_declare_device_resident
+ && !gfc_add_oacc_declare_device_resident (dest, NULL, where))
+ goto fail;
if (src->target && !gfc_add_target (dest, where))
goto fail;
if (src->dummy && !gfc_add_dummy (dest, NULL, where))
diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
index 6cdc472..77fdc8b 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -1309,6 +1309,16 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list)
list = tree_cons (get_identifier ("omp declare target"),
NULL_TREE, list);
+ if (sym_attr.oacc_declare_create
+ || sym_attr.oacc_declare_copyin
+ || sym_attr.oacc_declare_deviceptr
+ || sym_attr.oacc_declare_device_resident
+ || sym_attr.oacc_declare_link)
+ {
+ list = tree_cons (get_identifier ("oacc declare"),
+ NULL_TREE, list);
+ }
+
if (sym_attr.oacc_function)
list = tree_cons (get_identifier ("oacc function"),
NULL_TREE, list);
@@ -5754,14 +5764,49 @@ is_ieee_module_used (gfc_namespace *ns)
}
+static struct oacc_return
+{
+ gfc_code *code;
+ struct oacc_return *next;
+} *oacc_returns;
+
+
+static void
+find_oacc_return (gfc_code *code)
+{
+ if (code->next)
+ {
+ if (code->next->op == EXEC_RETURN)
+ {
+ struct oacc_return *r;
+
+ r = XCNEW (struct oacc_return);
+ r->code = code;
+ r->next = NULL;
+
+ if (oacc_returns)
+ r->next = oacc_returns;
+
+ oacc_returns = r;
+ }
+ else
+ {
+ find_oacc_return (code->next);
+ }
+ }
+
+ if (code->block)
+ find_oacc_return (code->block);
+
+ return;
+}
+
+
static gfc_code *
find_end (gfc_code *code)
{
gcc_assert (code);
- if (code->op == EXEC_END_PROCEDURE)
- return code;
-
if (code->next)
{
if (code->next->op == EXEC_END_PROCEDURE)
@@ -5774,38 +5819,284 @@ find_end (gfc_code *code)
}
+static gfc_omp_clauses *module_oacc_clauses;
+
+
+static void
+add_clause (gfc_symbol *sym, gfc_omp_map_op map_op)
+{
+ gfc_omp_namelist *n;
+
+ n = gfc_get_omp_namelist ();
+ n->sym = sym;
+ n->u.map_op = map_op;
+
+ if (!module_oacc_clauses)
+ module_oacc_clauses = gfc_get_omp_clauses ();
+
+ if (module_oacc_clauses->lists[OMP_LIST_MAP])
+ n->next = module_oacc_clauses->lists[OMP_LIST_MAP];
+
+ module_oacc_clauses->lists[OMP_LIST_MAP] = n;
+}
+
+
+static void
+find_module_oacc_declare_clauses (gfc_symbol *sym)
+{
+ if (sym->attr.use_assoc)
+ {
+ gfc_omp_map_op map_op;
+
+ sym->attr.referenced = sym->attr.oacc_declare_create
+ | sym->attr.oacc_declare_copyin
+ | sym->attr.oacc_declare_deviceptr
+ | sym->attr.oacc_declare_device_resident;
+
+ if (sym->attr.oacc_declare_create)
+ map_op = OMP_MAP_FORCE_ALLOC;
+
+ if (sym->attr.oacc_declare_copyin)
+ map_op = OMP_MAP_FORCE_TO;
+
+ if (sym->attr.oacc_declare_deviceptr)
+ map_op = OMP_MAP_FORCE_DEVICEPTR;
+
+ if (sym->attr.oacc_declare_device_resident)
+ map_op = OMP_MAP_DEVICE_RESIDENT;
+
+ if (sym->attr.referenced)
+ add_clause (sym, map_op);
+ }
+}
+
+
void
-insert_oacc_declare (gfc_namespace *ns)
+finish_oacc_declare (gfc_namespace *ns, enum sym_flavor flavor)
{
- gfc_code *code;
+ gfc_code *code, *end_c, *code2;
+ gfc_oacc_declare *oc;
+ gfc_omp_clauses *omp_clauses = NULL, *ret_clauses = NULL;
+ gfc_omp_namelist *n;
+ locus where = gfc_current_locus;
+
+ gfc_traverse_ns (ns, find_module_oacc_declare_clauses);
+
+ if (module_oacc_clauses && flavor == FL_PROGRAM)
+ {
+ gfc_oacc_declare *new_oc;
+
+ new_oc = gfc_get_oacc_declare ();
+ new_oc->next = ns->oacc_declare;
+ new_oc->clauses = module_oacc_clauses;
+
+ ns->oacc_declare = new_oc;
+ module_oacc_clauses = NULL;
+ }
+
+ if (!ns->oacc_declare)
+ return;
+
+ for (oc = ns->oacc_declare; oc; oc = oc->next)
+ {
+ if (oc->module_var)
+ continue;
+
+ if (oc->clauses)
+ {
+ if (omp_clauses)
+ {
+ gfc_omp_namelist *p;
+
+ for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
+ p = n;
+
+ p->next = oc->clauses->lists[OMP_LIST_MAP];
+ }
+ else
+ {
+ omp_clauses = oc->clauses;
+ }
+ }
+ }
+
+ while (ns->oacc_declare)
+ {
+ oc = ns->oacc_declare;
+ ns->oacc_declare = oc->next;
+ free (oc);
+ }
code = XCNEW (gfc_code);
code->op = EXEC_OACC_DECLARE;
- code->loc = ns->oacc_declare->where;
+ code->loc = where;
+ code->ext.omp_clauses = omp_clauses;
+
+ for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
+ {
+ bool ret = false;
+ gfc_omp_map_op new_op;
+
+ switch (n->u.map_op)
+ {
+ case OMP_MAP_ALLOC:
+ case OMP_MAP_FORCE_ALLOC:
+ new_op = OMP_MAP_FORCE_DEALLOC;
+ ret = true;
+ break;
+
+ case OMP_MAP_DEVICE_RESIDENT:
+ n->u.map_op = OMP_MAP_FORCE_ALLOC;
+ new_op = OMP_MAP_FORCE_DEALLOC;
+ ret = true;
+ break;
+
+ case OMP_MAP_FORCE_FROM:
+ n->u.map_op = OMP_MAP_FORCE_ALLOC;
+ new_op = OMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case OMP_MAP_FORCE_TO:
+ new_op = OMP_MAP_FORCE_DEALLOC;
+ ret = true;
+ break;
+
+ case OMP_MAP_FORCE_TOFROM:
+ n->u.map_op = OMP_MAP_FORCE_TO;
+ new_op = OMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case OMP_MAP_FROM:
+ n->u.map_op = OMP_MAP_FORCE_ALLOC;
+ new_op = OMP_MAP_FROM;
+ ret = true;
+ break;
+
+ case OMP_MAP_FORCE_DEVICEPTR:
+ case OMP_MAP_FORCE_PRESENT:
+ case OMP_MAP_LINK:
+ case OMP_MAP_TO:
+ break;
+
+ case OMP_MAP_TOFROM:
+ n->u.map_op = OMP_MAP_TO;
+ new_op = OMP_MAP_FROM;
+ ret = true;
+ break;
+
+ default:
+ gcc_unreachable ();
+ break;
+ }
+
+ if (ret)
+ {
+ gfc_omp_namelist *new_n;
+
+ new_n = gfc_get_omp_namelist ();
+ new_n->sym = n->sym;
+ new_n->u.map_op = new_op;
+
+ if (!ret_clauses)
+ ret_clauses = gfc_get_omp_clauses ();
+
+ if (ret_clauses->lists[OMP_LIST_MAP])
+ new_n->next = ret_clauses->lists[OMP_LIST_MAP];
+
+ ret_clauses->lists[OMP_LIST_MAP] = new_n;
+ ret = false;
+ }
+ }
- code->ext.oacc_declare = ns->oacc_declare;
+ if (!ret_clauses)
+ {
+ code->next = ns->code;
+ ns->code = code;
+ return;
+ }
- code->block = XCNEW (gfc_code);
- code->block->op = EXEC_OACC_DECLARE;
- code->block->loc = ns->oacc_declare->where;
+ code2 = XCNEW (gfc_code);
+ code2->op = EXEC_OACC_DECLARE;
+ code2->loc = where;
+ code2->ext.omp_clauses = ret_clauses;
if (ns->code)
{
- gfc_code *c;
+ find_oacc_return (ns->code);
+
+ if (ns->code->op == EXEC_END_PROCEDURE)
+ {
+ code2->next = ns->code;
+ code->next = code2;
+ }
+ else
+ {
+ end_c = find_end (ns->code);
+ if (end_c)
+ {
+ code2->next = end_c->next;
+ end_c->next = code2;
+ code->next = ns->code;
+ }
+ else
+ {
+ gfc_code *last;
+
+ last = ns->code;
+
+ while (last->next)
+ last = last->next;
+
+ last->next = code2;
+ code->next = ns->code;
+ }
+ }
+ }
+ else
+ {
+ code->next = code2;
+ }
+
+ while (oacc_returns)
+ {
+ struct oacc_return *r;
+
+ r = oacc_returns;
- c = find_end (ns->code);
- if (c)
+ ret_clauses = gfc_get_omp_clauses ();
+
+ for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
{
- code->next = c->next;
- c->next = NULL;
+ if (n->u.map_op == OMP_MAP_FORCE_ALLOC
+ || n->u.map_op == OMP_MAP_FORCE_TO)
+ {
+ gfc_omp_namelist *new_n;
+
+ new_n = gfc_get_omp_namelist ();
+ new_n->sym = n->sym;
+ new_n->u.map_op = OMP_MAP_FORCE_DEALLOC;
+
+ if (ret_clauses->lists[OMP_LIST_MAP])
+ new_n->next = ret_clauses->lists[OMP_LIST_MAP];
+
+ ret_clauses->lists[OMP_LIST_MAP] = new_n;
+ }
}
- code->block->next = ns->code;
- code->block->ext.oacc_declare = NULL;
+ code2 = XCNEW (gfc_code);
+ code2->op = EXEC_OACC_DECLARE;
+ code2->loc = where;
+ code2->ext.omp_clauses = ret_clauses;
+ code2->next = r->code->next;
+ r->code->next = code2;
+
+ oacc_returns = r->next;
+ free (r);
}
- ns->code = code;
- ns->oacc_declare = NULL;
+ ns->code = code;
}
@@ -5946,8 +6237,7 @@ gfc_generate_function_code (gfc_namespace * ns)
add_argument_checking (&body, sym);
/* Generate !$ACC DECLARE directive. */
- if (ns->oacc_declare)
- insert_oacc_declare (ns);
+ finish_oacc_declare (ns, sym->attr.flavor);
tmp = gfc_trans_code (ns->code);
gfc_add_expr_to_block (&body, tmp);
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 1aa33c0..f73e366 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1784,12 +1784,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_LIST_USE_DEVICE:
clause_code = OMP_CLAUSE_USE_DEVICE;
goto add_clause;
- case OMP_LIST_DEVICE_RESIDENT:
- clause_code = OMP_CLAUSE_DEVICE_RESIDENT;
- goto add_clause;
case OMP_LIST_CACHE:
clause_code = OMP_CLAUSE__CACHE_;
goto add_clause;
+ case OMP_LIST_DEVICE_RESIDENT:
+ case OMP_LIST_LINK:
+ continue;
add_clause:
omp_clauses
@@ -1937,6 +1937,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
if (!n->sym->attr.referenced)
continue;
+ if (n->sym->attr.use_assoc && n->sym->attr.oacc_declare_link)
+ continue;
+
tree node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
tree node2 = NULL_TREE;
tree node3 = NULL_TREE;
@@ -2160,6 +2163,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_MAP_FORCE_TO_GANGLOCAL:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO_GANGLOCAL);
break;
+ case OMP_MAP_DEVICE_RESIDENT:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DEVICE_RESIDENT);
+ break;
default:
gcc_unreachable ();
}
@@ -4391,25 +4397,18 @@ tree
gfc_trans_oacc_declare (gfc_code *code)
{
stmtblock_t block;
- struct gfc_oacc_declare *d;
- tree stmt, clauses = NULL_TREE;
+ tree stmt, oacc_clauses;
+ enum tree_code construct_code;
gfc_start_block (&block);
- for (d = code->ext.oacc_declare; d; d = d->next)
- {
- tree t;
-
- t = gfc_trans_omp_clauses (&block, d->clauses, d->clauses->loc);
+ construct_code = OACC_DECLARE;
- if (clauses)
- OMP_CLAUSE_CHAIN (clauses) = t;
- else
- clauses = t;
- }
-
- stmt = gfc_trans_omp_code (code->block->next, true);
- stmt = build2_loc (input_location, OACC_DATA, void_type_node, stmt, clauses);
+ gfc_start_block (&block);
+ oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
+ code->loc);
+ stmt = build1_loc (input_location, construct_code, void_type_node,
+ oacc_clauses);
gfc_add_expr_to_block (&block, stmt);
return gfc_finish_block (&block);
}
diff --git a/gcc/fortran/trans-stmt.c b/gcc/fortran/trans-stmt.c
index c6be9ad..352b383 100644
--- a/gcc/fortran/trans-stmt.c
+++ b/gcc/fortran/trans-stmt.c
@@ -1588,8 +1588,7 @@ gfc_trans_block_construct (gfc_code* code)
code->exit_label = exit_label;
/* Generate !$ACC DECLARE directive. */
- if (ns->oacc_declare)
- insert_oacc_declare (ns);
+ finish_oacc_declare (ns, FL_UNKNOWN);
gfc_add_expr_to_block (&body, gfc_trans_code (ns->code));
gfc_add_expr_to_block (&body, build1_v (LABEL_EXPR, exit_label));
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index 067882f..cc11d11 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -147,6 +147,7 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I4_INT, BT_VOID, BT_VOLATILE_PTR, BT_I4, BT
DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I8_INT, BT_VOID, BT_VOLATILE_PTR, BT_I8, BT_INT)
DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I16_INT, BT_VOID, BT_VOLATILE_PTR, BT_I16, BT_INT)
DEF_FUNCTION_TYPE_3 (BT_FN_INT_INT_INT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_UINT, BT_VOID, BT_PTR, BT_INT, BT_UINT)
DEF_FUNCTION_TYPE_4 (BT_FN_VOID_OMPFN_PTR_UINT_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT)
^ permalink raw reply [flat|nested] 10+ messages in thread* Re: [gomp4] declare directive [3/5]
2015-06-08 15:05 ` [gomp4] declare directive [3/5] James Norris
@ 2015-06-17 10:04 ` Thomas Schwinge
2015-06-18 16:21 ` James Norris
2015-10-30 13:30 ` Thomas Schwinge
1 sibling, 1 reply; 10+ messages in thread
From: Thomas Schwinge @ 2015-06-17 10:04 UTC (permalink / raw)
To: James Norris; +Cc: Tobias Burnus, gcc-patches, fortran
[-- Attachment #1: Type: text/plain, Size: 3749 bytes --]
Hi Jim!
I had mentioned that the Fortran front end changes cause regressions in a
few libgomp execution tests, if configured for Intel MIC (emulation)
offloading. I have now located *where* this is coming from, but would
you please work on figuring out *why*?
Fortunately, you'll be able to work on the problem even without Intel MIC
(emulation) offloading configured: to reproduce, just look at the
difference in -fdump-tree-original without and with your patch applied.
You'll notice that clauses are getting "lost" from OpenMP target update
directives; for example, for
libgomp/testsuite/libgomp.fortran/declare-target-1.f90 I see:
--- GOOD/declare-target-2.f90.003t.original 2015-06-16 18:16:07.472763339 +0200
+++ ./declare-target-2.f90.003t.original 2015-06-16 19:28:22.706845250 +0200
@@ -3,14 +3,14 @@
extern integer(kind=4) var_x;
var_x = 10;
- #pragma omp target update to(var_x)
+ #pragma omp target update
#pragma omp target
{
{
var_x = var_x * 2;
}
}
- #pragma omp target update from(var_x)
+ #pragma omp target update
if (var_x != 20)
{
_gfortran_abort ();
(This is the only test case that I looked at, so far.)
I tracked this down to:
On Mon, 8 Jun 2015 10:04:11 -0500, James Norris <jnorris@codesourcery.com> wrote:
> --- a/gcc/fortran/trans-decl.c
> +++ b/gcc/fortran/trans-decl.c
> +static void
> +find_module_oacc_declare_clauses (gfc_symbol *sym)
> +{
> + if (sym->attr.use_assoc)
> + {
> + gfc_omp_map_op map_op;
> +
> + sym->attr.referenced = sym->attr.oacc_declare_create
> + | sym->attr.oacc_declare_copyin
> + | sym->attr.oacc_declare_deviceptr
> + | sym->attr.oacc_declare_device_resident;
> +
> + if (sym->attr.oacc_declare_create)
> + map_op = OMP_MAP_FORCE_ALLOC;
> +
> + if (sym->attr.oacc_declare_copyin)
> + map_op = OMP_MAP_FORCE_TO;
> +
> + if (sym->attr.oacc_declare_deviceptr)
> + map_op = OMP_MAP_FORCE_DEVICEPTR;
> +
> + if (sym->attr.oacc_declare_device_resident)
> + map_op = OMP_MAP_DEVICE_RESIDENT;
> +
> + if (sym->attr.referenced)
> + add_clause (sym, map_op);
> + }
> +}
... this function apparently doing "something inappropriate". It gets
(unconditionally) called from:
> +finish_oacc_declare (gfc_namespace *ns, enum sym_flavor flavor)
> {
> [...]
> + gfc_traverse_ns (ns, find_module_oacc_declare_clauses);
... here, which in turn gets (unconditionally) called from:
> @@ -5946,8 +6237,7 @@ gfc_generate_function_code (gfc_namespace * ns)
> add_argument_checking (&body, sym);
>
> /* Generate !$ACC DECLARE directive. */
> - if (ns->oacc_declare)
> - insert_oacc_declare (ns);
> + finish_oacc_declare (ns, sym->attr.flavor);
>
> tmp = gfc_trans_code (ns->code);
> gfc_add_expr_to_block (&body, tmp);
... here, and:
> --- a/gcc/fortran/trans-stmt.c
> +++ b/gcc/fortran/trans-stmt.c
> @@ -1588,8 +1588,7 @@ gfc_trans_block_construct (gfc_code* code)
> code->exit_label = exit_label;
>
> /* Generate !$ACC DECLARE directive. */
> - if (ns->oacc_declare)
> - insert_oacc_declare (ns);
> + finish_oacc_declare (ns, FL_UNKNOWN);
>
> gfc_add_expr_to_block (&body, gfc_trans_code (ns->code));
> gfc_add_expr_to_block (&body, build1_v (LABEL_EXPR, exit_label));
... here.
Is that sufficient information for you to reproduce the problem?
As soon as you have a patch to bring back the lost clauses in the
-fdump-tree-original, I'll be happy to test it in my Intel MIC (emulated)
offloading build.
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 10+ messages in thread* Re: [gomp4] declare directive [3/5]
2015-06-17 10:04 ` Thomas Schwinge
@ 2015-06-18 16:21 ` James Norris
0 siblings, 0 replies; 10+ messages in thread
From: James Norris @ 2015-06-18 16:21 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: Tobias Burnus, gcc-patches, fortran
[-- Attachment #1: Type: text/plain, Size: 1170 bytes --]
Hi Thomas!
On 06/17/2015 04:59 AM, Thomas Schwinge wrote:
> Hi Jim!
>
> I had mentioned that the Fortran front end changes cause regressions in a
> few libgomp execution tests, if configured for Intel MIC (emulation)
> offloading. I have now located *where* this is coming from, but would
> you please work on figuring out *why*?
>
There are actually two bugs in find_module_oacc_declare_clauses which
are causing the issues you are seeing.
With the first bug, if none of the 'oacc_declare_*' bits were asserted,
then the referenced field within the attribute structure was set to
zero. If the referenced field was already set to one prior to
find_module_oacc_declare_clauses being called, then the field gets
incorrectly set to zero, if none of the 'oacc_declare_*' bits were
asserted.
With the second bug, if the referenced field within the attribute
structure is already set to one prior to
find_module_oacc_declare_clauses being called, then the subroutine
add_clause was called. The subroutine add_clause should only be
called if one of the 'oacc_declare_*' bits are asserted.
The attached patch resolves the above issues.
Committed to gomp-4_0-branch
Jim
[-- Attachment #2: declare.patch --]
[-- Type: text/x-patch, Size: 1013 bytes --]
diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
index 77fdc8b..7387a80 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -5848,11 +5848,6 @@ find_module_oacc_declare_clauses (gfc_symbol *sym)
{
gfc_omp_map_op map_op;
- sym->attr.referenced = sym->attr.oacc_declare_create
- | sym->attr.oacc_declare_copyin
- | sym->attr.oacc_declare_deviceptr
- | sym->attr.oacc_declare_device_resident;
-
if (sym->attr.oacc_declare_create)
map_op = OMP_MAP_FORCE_ALLOC;
@@ -5865,8 +5860,14 @@ find_module_oacc_declare_clauses (gfc_symbol *sym)
if (sym->attr.oacc_declare_device_resident)
map_op = OMP_MAP_DEVICE_RESIDENT;
- if (sym->attr.referenced)
- add_clause (sym, map_op);
+ if (sym->attr.oacc_declare_create
+ || sym->attr.oacc_declare_copyin
+ || sym->attr.oacc_declare_deviceptr
+ || sym->attr.oacc_declare_device_resident)
+ {
+ sym->attr.referenced = 1;
+ add_clause (sym, map_op);
+ }
}
}
^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [gomp4] declare directive [3/5]
2015-06-08 15:05 ` [gomp4] declare directive [3/5] James Norris
2015-06-17 10:04 ` Thomas Schwinge
@ 2015-10-30 13:30 ` Thomas Schwinge
1 sibling, 0 replies; 10+ messages in thread
From: Thomas Schwinge @ 2015-10-30 13:30 UTC (permalink / raw)
To: James Norris, gcc-patches, fortran; +Cc: Tobias Burnus
[-- Attachment #1: Type: text/plain, Size: 2245 bytes --]
Hi!
On Mon, 8 Jun 2015 10:04:11 -0500, James Norris <jnorris@codesourcery.com> wrote:
> --- a/gcc/fortran/gfortran.h
> +++ b/gcc/fortran/gfortran.h
> @@ -1174,6 +1183,7 @@ enum
> OMP_LIST_FROM,
> OMP_LIST_REDUCTION,
> OMP_LIST_DEVICE_RESIDENT,
> + OMP_LIST_LINK,
> OMP_LIST_USE_DEVICE,
> OMP_LIST_CACHE,
> OMP_LIST_NUM
I noticed (my means of hitting a segmentation fault) that this was
missing an update to the clause_names in
gcc/fortran/openmp.c:resolve_omp_clauses. (Yes, I agree that is a
strange, non-obvious dependency that this function needs to be updated
for OMP_LIST_* changes...) Fixed on gomp-4_0-branch in r229576:
commit a5246d7b6c91e0800eeb6355bf5e4c63d27aafb2
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri Oct 30 13:24:35 2015 +0000
Fix OMP_LIST_LINK handling
gcc/fortran/
* openmp.c (resolve_omp_clauses): Add "LINK" to clause_names.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@229576 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/fortran/ChangeLog.gomp | 4 ++++
gcc/fortran/openmp.c | 2 +-
2 files changed, 5 insertions(+), 1 deletion(-)
diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp
index 7fe3eac..592dd8d 100644
--- gcc/fortran/ChangeLog.gomp
+++ gcc/fortran/ChangeLog.gomp
@@ -1,3 +1,7 @@
+2015-10-30 Thomas Schwinge <thomas@codesourcery.com>
+
+ * openmp.c (resolve_omp_clauses): Add "LINK" to clause_names.
+
2015-10-29 Thomas Schwinge <thomas@codesourcery.com>
* openmp.c (gfc_match_omp_map_clause): Remove allow_sections
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index a2c5105..32779f7 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -3197,7 +3197,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
static const char *clause_names[]
= { "PRIVATE", "FIRSTPRIVATE", "LASTPRIVATE", "COPYPRIVATE", "SHARED",
"COPYIN", "UNIFORM", "ALIGNED", "LINEAR", "DEPEND", "MAP",
- "TO", "FROM", "REDUCTION", "DEVICE_RESIDENT", "USE_DEVICE",
+ "TO", "FROM", "REDUCTION", "DEVICE_RESIDENT", "LINK", "USE_DEVICE",
"CACHE" };
if (omp_clauses == NULL)
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 10+ messages in thread
* [gomp4] declare directive [5/5]
2015-06-08 14:59 [gomp4] declare directive [0/5] James Norris
` (2 preceding siblings ...)
2015-06-08 15:05 ` [gomp4] declare directive [3/5] James Norris
@ 2015-06-08 15:06 ` James Norris
2015-07-13 11:56 ` Thomas Schwinge
2015-06-08 15:06 ` [gomp4] declare directive [4/5] James Norris
4 siblings, 1 reply; 10+ messages in thread
From: James Norris @ 2015-06-08 15:06 UTC (permalink / raw)
To: gcc-patches; +Cc: Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 1 bytes --]
[-- Attachment #2: libgomp.patch --]
[-- Type: text/x-patch, Size: 13780 bytes --]
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index fe38dc6..663c27c 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -318,6 +318,7 @@ GOACC_2.0 {
global:
GOACC_data_end;
GOACC_data_start;
+ GOACC_declare;
GOACC_enter_exit_data;
GOACC_parallel;
GOACC_update;
@@ -331,6 +332,7 @@ GOACC_2.0.GOMP_4_BRANCH {
GOACC_deviceptr;
GOACC_get_ganglocal_ptr;
GOACC_kernels;
+ GOACC_register_static;
} GOACC_2.0;
GOMP_PLUGIN_1.0 {
diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c
index 9f24dc3..e772f48 100644
--- a/libgomp/oacc-init.c
+++ b/libgomp/oacc-init.c
@@ -205,6 +205,8 @@ acc_shutdown_1 (acc_device_t d)
if (!base_dev)
gomp_fatal ("device %s not supported", name_of_acc_device_t (d));
+ goacc_deallocate_static (d);
+
gomp_mutex_lock (&goacc_thread_lock);
/* Free target-specific TLS data and close all devices. */
@@ -373,7 +375,9 @@ goacc_attach_host_thread_to_device (int ord)
void
acc_init (acc_device_t d)
{
- if (!cached_base_dev)
+ bool init = !cached_base_dev;
+
+ if (init)
gomp_init_targets_once ();
gomp_mutex_lock (&acc_device_lock);
@@ -381,6 +385,9 @@ acc_init (acc_device_t d)
cached_base_dev = acc_init_1 (d);
gomp_mutex_unlock (&acc_device_lock);
+
+ if (init)
+ goacc_allocate_static (d);
goacc_attach_host_thread_to_device (-1);
}
diff --git a/libgomp/oacc-int.h b/libgomp/oacc-int.h
index 0ace737..8f4938e 100644
--- a/libgomp/oacc-int.h
+++ b/libgomp/oacc-int.h
@@ -98,6 +98,9 @@ void goacc_save_and_set_bind (acc_device_t);
void goacc_restore_bind (void);
void goacc_lazy_initialize (void);
+void goacc_allocate_static (acc_device_t);
+void goacc_deallocate_static (acc_device_t);
+
#ifdef HAVE_ATTRIBUTE_VISIBILITY
# pragma GCC visibility pop
#endif
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 513d0bc..70758bc 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -109,6 +109,68 @@ alloc_ganglocal_addrs (size_t mapnum, void **hostaddrs, size_t *sizes,
}
}
+static struct oacc_static
+{
+ void *addr;
+ size_t size;
+ unsigned short mask;
+ bool free;
+ struct oacc_static *next;
+} *oacc_statics;
+
+static bool alloc_done = false;
+
+void
+goacc_allocate_static (acc_device_t d)
+{
+ struct oacc_static *s;
+
+ if (alloc_done)
+ assert (0);
+
+ for (s = oacc_statics; s; s = s->next)
+ {
+ void *d;
+
+ switch (s->mask)
+ {
+ case GOMP_MAP_FORCE_ALLOC:
+ break;
+
+ case GOMP_MAP_FORCE_TO:
+ d = acc_deviceptr (s->addr);
+ acc_memcpy_to_device (d, s->addr, s->size);
+ break;
+
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_LINK:
+ break;
+
+ default:
+ assert (0);
+ break;
+ }
+ }
+
+ alloc_done = true;
+}
+
+void
+goacc_deallocate_static (acc_device_t d)
+{
+ struct oacc_static *s;
+ unsigned short mask = GOMP_MAP_FORCE_DEALLOC;
+
+ if (!alloc_done)
+ return;
+
+ for (s = oacc_statics; s; s = s->next)
+ GOACC_enter_exit_data (d, 1, &s->addr, &s->size, &mask, 0, 0);
+
+ alloc_done = false;
+}
+
static void goacc_wait (int async, int num_waits, va_list ap);
void
@@ -592,3 +654,82 @@ GOACC_get_thread_num (int gang, int worker, int vector)
{
return 0;
}
+
+void
+GOACC_register_static (void *addr, int size, unsigned int mask)
+{
+ struct oacc_static *s;
+
+ s = (struct oacc_static *) malloc (sizeof (struct oacc_static));
+ s->addr = addr;
+ s->size = (size_t) size;
+ s->mask = mask;
+ s->free = false;
+ s->next = NULL;
+
+ if (oacc_statics)
+ s->next = oacc_statics;
+
+ oacc_statics = s;
+}
+
+#include <stdio.h>
+
+void
+GOACC_declare (int device, size_t mapnum,
+ void **hostaddrs, size_t *sizes, unsigned short *kinds)
+{
+ int i;
+
+ for (i = 0; i < mapnum; i++)
+ {
+ unsigned char kind = kinds[i] & 0xff;
+
+ if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET)
+ continue;
+
+ switch (kind)
+ {
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_DEALLOC:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_POINTER:
+ GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
+ &kinds[i], 0, 0);
+ break;
+
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ break;
+
+ case GOMP_MAP_ALLOC:
+ if (!acc_is_present (hostaddrs[i], sizes[i]))
+ {
+ GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
+ &kinds[i], 0, 0);
+ }
+ break;
+
+ case GOMP_MAP_TO:
+ GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
+ &kinds[i], 0, 0);
+
+ break;
+
+ case GOMP_MAP_FROM:
+ kinds[i] = GOMP_MAP_FORCE_FROM;
+ GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
+ &kinds[i], 0, 0);
+ break;
+
+ case GOMP_MAP_FORCE_PRESENT:
+ if (!acc_is_present (hostaddrs[i], sizes[i]))
+ gomp_fatal ("[%p,%zd] is not mapped", hostaddrs[i], sizes[i]);
+ break;
+
+ default:
+ assert (0);
+ break;
+ }
+ }
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-1.C b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C
new file mode 100644
index 0000000..268809b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C
@@ -0,0 +1,24 @@
+
+template<class T>
+T foo()
+{
+ T a;
+ #pragma acc declare create (a)
+
+ #pragma acc parallel
+ {
+ a = 5;
+ }
+
+ return a;
+}
+
+int
+main (void)
+{
+ int rc;
+
+ rc = foo<int>();
+
+ return rc;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
new file mode 100644
index 0000000..59cfe51
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
@@ -0,0 +1,65 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <openacc.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+int b[8];
+#pragma acc declare create (b)
+
+int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+#pragma acc declare copyin (d)
+
+int
+main (int argc, char **argv)
+{
+ const int N = 8;
+ int a[N];
+ int e[N];
+#pragma acc declare create (e)
+ int i;
+
+ for (i = 0; i < N; i++)
+ a[i] = i + 1;
+
+ if (!acc_is_present (&b, sizeof (b)))
+ abort ();
+
+ if (!acc_is_present (&d, sizeof (d)))
+ abort ();
+
+ if (!acc_is_present (&e, sizeof (e)))
+ abort ();
+
+#pragma acc parallel copyin (a[0:N])
+ {
+ for (i = 0; i < N; i++)
+ {
+ b[i] = a[i];
+ a[i] = b[i];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != i + 1)
+ abort ();
+ }
+
+#pragma acc parallel copy (a[0:N])
+ {
+ for (i = 0; i < N; i++)
+ {
+ e[i] = a[i] + d[i];
+ a[i] = e[i];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != (i + 1) * 2)
+ abort ();
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c
new file mode 100644
index 0000000..2078a33
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c
@@ -0,0 +1,64 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <stdlib.h>
+
+#define N 16
+
+float c[N];
+#pragma acc declare device_resident (c)
+
+#pragma acc routine
+float
+subr2 (float a)
+{
+ int i;
+
+ for (i = 0; i < N; i++)
+ c[i] = 2.0;
+
+ for (i = 0; i < N; i++)
+ a += c[i];
+
+ return a;
+}
+
+float b[N];
+#pragma acc declare copyin (b)
+
+#pragma acc routine
+float
+subr1 (float a)
+{
+ int i;
+
+ for (i = 0; i < N; i++)
+ a += b[i];
+
+ return a;
+}
+
+int
+main (int argc, char **argv)
+{
+ float a;
+ int i;
+
+ for (i = 0; i < 16; i++)
+ b[i] = 1.0;
+
+ a = 0.0;
+
+ a = subr1 (a);
+
+ if (a != 16.0)
+ abort ();
+
+ a = 0.0;
+
+ a = subr2 (a);
+
+ if (a != 32.0)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-3.c
new file mode 100644
index 0000000..c3a2187
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-3.c
@@ -0,0 +1,61 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <stdlib.h>
+#include <openacc.h>
+
+float *b;
+#pragma acc declare deviceptr (b)
+
+#pragma acc routine
+float *
+subr2 (void)
+{
+ return b;
+}
+
+float
+subr1 (float a)
+{
+ float b;
+#pragma acc declare present_or_copy (b)
+ float c;
+#pragma acc declare present_or_copyin (c)
+ float d;
+#pragma acc declare present_or_create (d)
+ float e;
+#pragma acc declare present_or_copyout (e)
+
+#pragma acc parallel copy (a)
+ {
+ b = a;
+ c = b;
+ d = c;
+ e = d;
+ a = e;
+ }
+
+ return a;
+}
+
+int
+main (int argc, char **argv)
+{
+ float a;
+ float *c;
+
+ a = 2.0;
+
+ a = subr1 (a);
+
+ if (a != 2.0)
+ abort ();
+
+ b = (float *) acc_malloc (sizeof (float));
+
+ c = subr2 ();
+
+ if (b != c)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c
new file mode 100644
index 0000000..84ec64f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c
@@ -0,0 +1,27 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <stdlib.h>
+#include <openacc.h>
+
+float b;
+#pragma acc declare link (b)
+
+int
+main (int argc, char **argv)
+{
+ float a;
+
+ a = 2.0;
+
+#pragma acc parallel copy (a)
+ {
+ b = a;
+ a = 1.0;
+ a = a + b;
+ }
+
+ if (a != 3.0)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
index 0bab5bd..4d58e70 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
@@ -1,5 +1,10 @@
! { dg-do run { target openacc_nvidia_accel_selected } }
+module vars
+ integer z
+ !$acc declare create (z)
+end module vars
+
subroutine subr6 (a, d)
integer, parameter :: N = 8
integer :: i
@@ -200,6 +205,7 @@ subroutine subr0 (a, b, c, d)
end subroutine
program main
+ use vars
use openacc
integer, parameter :: N = 8
integer :: a(N)
@@ -212,6 +218,8 @@ program main
c(:) = 4
d(:) = 5
+ if (acc_is_present (z) .neqv. .true.) call abort
+
call subr0 (a, b, c, d)
call test (a, .false.)
@@ -226,4 +234,5 @@ program main
if (d(i) .ne. 16) call abort
end do
+
end program
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-2.f90
new file mode 100644
index 0000000..9b75aa1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-2.f90
@@ -0,0 +1,14 @@
+! { dg-do run { target openacc_nvidia_accel_selected } }
+
+module globalvars
+ integer a
+ !$acc declare create (a)
+end module globalvars
+
+program test
+ use globalvars
+ use openacc
+
+ if (acc_is_present (a) .neqv. .true.) call abort
+
+end program test
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-3.f90
new file mode 100644
index 0000000..79fc011
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-3.f90
@@ -0,0 +1,65 @@
+! { dg-do run { target openacc_nvidia_accel_selected } }
+
+module globalvars
+ real b
+ !$acc declare link (b)
+end module globalvars
+
+program test
+ use openacc
+
+ real a
+ real c
+ !$acc declare link (c)
+
+ if (acc_is_present (b) .neqv. .false.) call abort
+ if (acc_is_present (c) .neqv. .false.) call abort
+
+ a = 0.0
+ b = 1.0
+
+ !$acc parallel copy (a) copyin (b)
+ b = b + 4.0
+ a = b
+ !$acc end parallel
+
+ if (a .ne. 5.0) call abort
+
+ if (acc_is_present (b) .neqv. .false.) call abort
+
+ a = 0.0
+
+ !$acc parallel copy (a) create (b)
+ b = 4.0
+ a = b
+ !$acc end parallel
+
+ if (a .ne. 4.0) call abort
+
+ if (acc_is_present (b) .neqv. .false.) call abort
+
+ a = 0.0
+
+ !$acc parallel copy (a) copy (b)
+ b = 4.0
+ a = b
+ !$acc end parallel
+
+ if (a .ne. 4.0) call abort
+ if (b .ne. 4.0) call abort
+
+ if (acc_is_present (b) .neqv. .false.) call abort
+
+ a = 0.0
+
+ !$acc parallel copy (a) copy (b) copy (c)
+ b = 4.0
+ c = b
+ a = c
+ !$acc end parallel
+
+ if (a .ne. 4.0) call abort
+
+ if (acc_is_present (b) .neqv. .false.) call abort
+
+end program test
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-4.f90
new file mode 100644
index 0000000..997c8ac
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-4.f90
@@ -0,0 +1,27 @@
+! { dg-do run { target openacc_nvidia_accel_selected } }
+
+module vars
+ real b
+ !$acc declare create (b)
+end module vars
+
+program test
+ use vars
+ use openacc
+ real a
+
+ if (acc_is_present (b) .neqv. .true.) call abort
+
+ a = 2.0
+
+ !$acc parallel copy (a)
+ b = a
+ a = 1.0
+ a = a + b
+ !$acc end parallel
+
+ if (acc_is_present (b) .neqv. .true.) call abort
+
+ if (a .ne. 3.0) call abort
+
+end program test
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-5.f90
new file mode 100644
index 0000000..d7c9bac
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-5.f90
@@ -0,0 +1,28 @@
+! { dg-do run { target openacc_nvidia_accel_selected } }
+
+module vars
+ implicit none
+ real b
+ !$acc declare device_resident (b)
+end module vars
+
+program test
+ use vars
+ use openacc
+ real a
+
+ if (acc_is_present (b) .neqv. .true.) call abort
+
+ a = 2.0
+
+ !$acc parallel copy (a)
+ b = a
+ a = 1.0
+ a = a + b
+ !$acc end parallel
+
+ if (acc_is_present (b) .neqv. .true.) call abort
+
+ if (a .ne. 3.0) call abort
+
+end program test
^ permalink raw reply [flat|nested] 10+ messages in thread* Re: [gomp4] declare directive [5/5]
2015-06-08 15:06 ` [gomp4] declare directive [5/5] James Norris
@ 2015-07-13 11:56 ` Thomas Schwinge
0 siblings, 0 replies; 10+ messages in thread
From: Thomas Schwinge @ 2015-07-13 11:56 UTC (permalink / raw)
To: James Norris, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 2330 bytes --]
Hi Jim!
On Mon, 8 Jun 2015 10:06:21 -0500, James Norris <jnorris@codesourcery.com> wrote:
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C
> @@ -0,0 +1,24 @@
> +
> +template<class T>
> +T foo()
> +{
> + T a;
> + #pragma acc declare create (a)
> +
> + #pragma acc parallel
> + {
> + a = 5;
> + }
> +
> + return a;
> +}
> +
> +int
> +main (void)
> +{
> + int rc;
> +
> + rc = foo<int>();
> +
> + return rc;
> +}
I wonder, in a shared-memory setting (say, host-fallback because of the
OpenACC if clause, or acc_device_host, or acc_device_host_nonshm),
shouldn't the original and "declare"d objects of variable a be the same
(just like with the other data clauses), and thus the function foo return
the value 5 instead of 0?
Anyway, as-is, this test case FAILed in 32-bit x86 GNU/Linux testing
(acc_device_host, acc_device_host_nonshm), which I fixed in r225734 as
follows:
commit 016e15e94b8511f2041646c43d4344e1ea424e62
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon Jul 13 11:48:33 2015 +0000
libgomp testsuite: Don't read from uninitialized variables
libgomp/
* testsuite/libgomp.oacc-c++/declare-1.C (foo): Initialize a.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@225734 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog.gomp | 2 ++
libgomp/testsuite/libgomp.oacc-c++/declare-1.C | 2 +-
2 files changed, 3 insertions(+), 1 deletion(-)
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index fd7887a..7d1e9ad 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,5 +1,7 @@
2015-07-13 Thomas Schwinge <thomas@codesourcery.com>
+ * testsuite/libgomp.oacc-c++/declare-1.C (foo): Initialize a.
+
* testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-5.c:
Add XFAIL.
diff --git libgomp/testsuite/libgomp.oacc-c++/declare-1.C libgomp/testsuite/libgomp.oacc-c++/declare-1.C
index 268809b..6618b10 100644
--- libgomp/testsuite/libgomp.oacc-c++/declare-1.C
+++ libgomp/testsuite/libgomp.oacc-c++/declare-1.C
@@ -2,7 +2,7 @@
template<class T>
T foo()
{
- T a;
+ T a = 0;
#pragma acc declare create (a)
#pragma acc parallel
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 10+ messages in thread
* [gomp4] declare directive [4/5]
2015-06-08 14:59 [gomp4] declare directive [0/5] James Norris
` (3 preceding siblings ...)
2015-06-08 15:06 ` [gomp4] declare directive [5/5] James Norris
@ 2015-06-08 15:06 ` James Norris
4 siblings, 0 replies; 10+ messages in thread
From: James Norris @ 2015-06-08 15:06 UTC (permalink / raw)
To: gcc-patches; +Cc: Thomas Schwinge, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 1 bytes --]
[-- Attachment #2: gcc.patch --]
[-- Type: text/x-patch, Size: 12780 bytes --]
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 7c3273f..0774da5 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -451,6 +451,7 @@ DEF_FUNCTION_TYPE_3 (BT_FN_BOOL_ULONG_ULONG_ULONGPTR, BT_BOOL, BT_ULONG,
DEF_FUNCTION_TYPE_3 (BT_FN_BOOL_ULONGLONG_ULONGLONG_ULONGLONGPTR, BT_BOOL,
BT_ULONGLONG, BT_ULONGLONG, BT_PTR_ULONGLONG)
DEF_FUNCTION_TYPE_3 (BT_FN_INT_INT_INT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_UINT, BT_VOID, BT_PTR, BT_INT, BT_UINT)
DEF_FUNCTION_TYPE_4 (BT_FN_SIZE_CONST_PTR_SIZE_SIZE_FILEPTR,
BT_SIZE, BT_CONST_PTR, BT_SIZE, BT_SIZE, BT_FILEPTR)
diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c
index a640a96..f447af6 100644
--- a/gcc/gimple-pretty-print.c
+++ b/gcc/gimple-pretty-print.c
@@ -1365,6 +1365,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs,
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
kind = " oacc_enter_exit_data";
break;
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
+ kind = " oacc_declare";
+ break;
default:
gcc_unreachable ();
}
diff --git a/gcc/gimple.h b/gcc/gimple.h
index bf048e6..bd92c96 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -100,7 +100,7 @@ enum gf_mask {
GF_OMP_FOR_KIND_CILKSIMD = GF_OMP_FOR_SIMD | 1,
GF_OMP_FOR_COMBINED = 1 << 3,
GF_OMP_FOR_COMBINED_INTO = 1 << 4,
- GF_OMP_TARGET_KIND_MASK = (1 << 3) - 1,
+ GF_OMP_TARGET_KIND_MASK = (1 << 4) - 1,
GF_OMP_TARGET_KIND_REGION = 0,
GF_OMP_TARGET_KIND_DATA = 1,
GF_OMP_TARGET_KIND_UPDATE = 2,
@@ -109,6 +109,7 @@ enum gf_mask {
GF_OMP_TARGET_KIND_OACC_DATA = 5,
GF_OMP_TARGET_KIND_OACC_UPDATE = 6,
GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 7,
+ GF_OMP_TARGET_KIND_OACC_DECLARE = 8,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
@@ -5663,6 +5664,7 @@ is_gimple_omp_oacc (const_gimple stmt)
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
return true;
default:
return false;
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index c85b424..b1f768f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -5819,10 +5819,26 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
splay_tree_node n;
unsigned flags = in_code ? GOVD_SEEN : 0;
bool ret = false, shared;
+ bool device_resident = false;
if (error_operand_p (decl))
return false;
+ if (flag_openacc && is_global_var (decl))
+ {
+ tree attr = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
+ if (attr)
+ {
+ tree t, c;
+ for (t = TREE_VALUE (attr); t; t = TREE_PURPOSE (t))
+ {
+ c = TREE_VALUE (t);
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DEVICE_RESIDENT)
+ device_resident = true;
+ }
+ }
+ }
+
/* Threadprivate variables are predetermined. */
if (is_global_var (decl))
{
@@ -5899,7 +5915,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
by default are firstprivate (gang-local) in parallel. */
if (!n2 && !AGGREGATE_TYPE_P (type))
{
- if (ctx->acc_region_kind == ARK_PARALLEL)
+ if (device_resident)
+ flags |= GOVD_MAP_TO_ONLY;
+ else if (ctx->acc_region_kind == ARK_PARALLEL)
flags |= (GOVD_GANGLOCAL | GOVD_MAP_TO_ONLY);
/* Scalars under kernels are default 'copy'. */
else if (ctx->acc_region_kind == ARK_KERNELS)
@@ -7729,6 +7747,10 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
+ case OACC_DECLARE:
+ kind = GF_OMP_TARGET_KIND_OACC_DECLARE;
+ ork = ORK_OACC;
+ break;
case OACC_ENTER_DATA:
kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
ork = ORK_OACC;
@@ -8707,11 +8729,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = gimplify_oacc_host_data (expr_p, pre_p);
break;
- case OACC_DECLARE:
- sorry ("directive not yet implemented");
- ret = GS_ALL_DONE;
- break;
-
case OACC_KERNELS:
case OACC_PARALLEL:
case OACC_DATA:
@@ -8724,6 +8741,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
+ case OACC_DECLARE:
case OACC_ENTER_DATA:
case OACC_EXIT_DATA:
case OACC_UPDATE:
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index 6e70d0b..b31cb2d 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -299,3 +299,7 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update",
BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams",
BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_STATIC, "GOACC_register_static",
+ BT_FN_VOID_PTR_INT_UINT, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DECLARE, "GOACC_declare",
+ BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 0b31992..e1c9db4 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -9519,6 +9519,7 @@ expand_omp_target (struct omp_region *region)
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
data_region = false;
break;
case GF_OMP_TARGET_KIND_DATA:
@@ -9825,6 +9826,9 @@ expand_omp_target (struct omp_region *region)
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
break;
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
+ start_ix = BUILT_IN_GOACC_DECLARE;
+ break;
default:
gcc_unreachable ();
}
@@ -9944,6 +9948,7 @@ expand_omp_target (struct omp_region *region)
args.quick_push (build_zero_cst (ptr_type_node));
break;
case BUILT_IN_GOACC_DATA_START:
+ case BUILT_IN_GOACC_DECLARE:
case BUILT_IN_GOACC_ENTER_EXIT_DATA:
case BUILT_IN_GOACC_KERNELS:
case BUILT_IN_GOACC_KERNELS_INTERNAL:
@@ -9960,6 +9965,7 @@ expand_omp_target (struct omp_region *region)
switch (start_ix)
{
case BUILT_IN_GOACC_DATA_START:
+ case BUILT_IN_GOACC_DECLARE:
case BUILT_IN_GOMP_TARGET:
case BUILT_IN_GOMP_TARGET_DATA:
case BUILT_IN_GOMP_TARGET_UPDATE:
@@ -10268,6 +10274,7 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
case GF_OMP_TARGET_KIND_UPDATE:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
/* ..., other than for those stand-alone directives... */
region = NULL;
break;
@@ -12771,6 +12778,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
data_region = false;
break;
case GF_OMP_TARGET_KIND_DATA:
@@ -12835,6 +12843,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_FORCE_PRESENT:
case GOMP_MAP_FORCE_DEALLOC:
case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_LINK:
gcc_assert (is_gimple_omp_oacc (stmt));
break;
default:
@@ -13888,6 +13898,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
case GF_OMP_TARGET_KIND_UPDATE:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
cur_region = cur_region->outer;
break;
default:
diff --git a/gcc/testsuite/ChangeLog.gomp b/gcc/testsuite/ChangeLog.gomp
index fb480cf..649740c 100644
--- a/gcc/testsuite/ChangeLog.gomp
+++ b/gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,10 @@
+
+2015-06-04 James Norris <jnorris@codesourcery.com>
+
+ * c-c++-common/goacc/declare-1.c: Update tests.
+ * c-c++-common/goacc/declare-2.c: Likewise.
+ * gfortran.dg/goacc/declare-1.f95: Update tests.
+
2015-06-01 Tom de Vries <tom@codesourcery.com>
Revert:
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-1.c b/gcc/testsuite/c-c++-common/goacc/declare-1.c
index cf50f02..b036c63 100644
--- a/gcc/testsuite/c-c++-common/goacc/declare-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/declare-1.c
@@ -1,6 +1,5 @@
/* Test valid uses of declare directive. */
/* { dg-do compile } */
-/* { dg-skip-if "not yet" { c++ } } */
int v0;
#pragma acc declare create(v0)
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c
index a2b5d6f..ce12463 100644
--- a/gcc/testsuite/c-c++-common/goacc/declare-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c
@@ -1,11 +1,10 @@
/* Test invalid uses of declare directive. */
/* { dg-do compile } */
-/* { dg-skip-if "not yet" { c++ } } */
#pragma acc declare /* { dg-error "no valid clauses" } */
#pragma acc declare create(undeclared) /* { dg-error "undeclared" } */
-/* { dg-error "no valid clauses" "second error" { target *-*-* } 7 } */
+/* { dg-error "no valid clauses" "second error" { target *-*-* } 6 } */
int v0[10];
#pragma acc declare create(v0[1:3]) /* { dg-error "subarray" } */
@@ -42,7 +41,7 @@ void
f (void)
{
int va0;
-#pragma acc declare link(va0) /* { dg-error "invalid variable" } */
+#pragma acc declare link(va0) /* { dg-error "global variable" } */
extern int ve0;
#pragma acc declare copy(ve0) /* { dg-error "invalid use of" } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-1.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-1.f95
index 14190a7..50f75dc 100644
--- a/gcc/testsuite/gfortran.dg/goacc/declare-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/declare-1.f95
@@ -15,5 +15,6 @@ contains
END BLOCK
end function foo
end program test
-! { dg-final { scan-tree-dump-times "pragma acc data map\\(force_tofrom:i\\)" 2 "original" } }
+! { dg-final { scan-tree-dump-times "pragma acc declare map\\(force_to:i\\)" 2 "original" } }
+! { dg-final { scan-tree-dump-times "pragma acc declare map\\(force_from:i\\)" 2 "original" } }
! { dg-final { cleanup-tree-dump "original" } }
diff --git a/gcc/varpool.c b/gcc/varpool.c
index 76148a5..070d1c3 100644
--- a/gcc/varpool.c
+++ b/gcc/varpool.c
@@ -57,6 +57,7 @@ along with GCC; see the file COPYING3. If not see
#include "lto-streamer.h"
#include "context.h"
#include "omp-low.h"
+#include "gomp-constants.h"
const char * const tls_model_names[]={"none", "emulated",
"global-dynamic", "local-dynamic",
@@ -161,6 +162,58 @@ varpool_node::create_empty (void)
return node;
}
+static void
+make_offloadable_1 (varpool_node *node, tree decl)
+{
+ node->offloadable = 1;
+#ifdef ENABLE_OFFLOADING
+ g->have_offload = true;
+ if (!in_lto_p)
+ vec_safe_push (offload_vars, decl);
+ node->force_output = 1;
+#endif
+}
+
+void
+make_offloadable (varpool_node *node, tree decl)
+{
+ tree attrs;
+
+ if (node->offloadable)
+ return;
+
+ if (flag_openmp)
+ {
+ make_offloadable_1 (node, decl);
+ return;
+ }
+
+ attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
+ if (attrs)
+ {
+ tree *t;
+ int total = 0, skip = 0;
+
+ gcc_assert (&TREE_VALUE (attrs));
+
+ for (t = &TREE_VALUE (attrs); *t; t = &TREE_CHAIN (*t))
+ {
+ HOST_WIDE_INT kind = OMP_CLAUSE_MAP_KIND (TREE_VALUE (*t));
+
+ total++;
+
+ if (kind == GOMP_MAP_LINK)
+ skip++;
+ }
+
+ if (total - skip > 0)
+ make_offloadable_1 (node, decl);
+
+ DECL_ATTRIBUTES (decl)
+ = remove_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
+ }
+}
+
/* Return varpool node assigned to DECL. Create new one when needed. */
varpool_node *
varpool_node::get_create (tree decl)
@@ -168,22 +221,18 @@ varpool_node::get_create (tree decl)
varpool_node *node = varpool_node::get (decl);
gcc_checking_assert (TREE_CODE (decl) == VAR_DECL);
if (node)
- return node;
+ {
+ if (flag_openacc && !DECL_EXTERNAL (decl))
+ make_offloadable (node, decl);
+ return node;
+ }
node = varpool_node::create_empty ();
node->decl = decl;
if ((flag_openacc || flag_openmp) && !DECL_EXTERNAL (decl)
&& lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)))
- {
- node->offloadable = 1;
-#ifdef ENABLE_OFFLOADING
- g->have_offload = true;
- if (!in_lto_p)
- vec_safe_push (offload_vars, decl);
- node->force_output = 1;
-#endif
- }
+ make_offloadable (node, decl);
node->register_symbol ();
return node;
^ permalink raw reply [flat|nested] 10+ messages in thread