This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.
| Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
|---|---|---|
| Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |
| Other format: | [Raw text] | |
On 10/27/2015 03:18 PM, James Norris wrote:
Hi!
This patch adds the processing of OpenACC declare directive in C
and C++. (Note: Support in Fortran is already in trunk.)
Commentary on the changes is included as an attachment (NOTES).
All of the code is in the gomp-4_0-branch.
Regtested on x86_64-linux.
Thanks!
Jim
Ping!
I've revised the patch since I originally submitted it for review
(https://gcc.gnu.org/ml/gcc-patches/2015-10/msg02967.html). The
revision is due to Jakub and et al OpenMP 4.5 work in the area of
'omp declare target'. I now exploit that functionality and have
revised the patch accordingly.
Updated ChangeLog, patch, and commentary (NOTES) are attached.
Regtested on x86_64-linux
Thanks!
Jim
Attachment:
NOTES
Description: Text document
Attachment:
ChangeLog
Description: Text document
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index b561436..a109806 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -450,6 +450,7 @@ DEF_FUNCTION_TYPE_3 (BT_FN_BOOL_ULONG_ULONG_ULONGPTR, BT_BOOL, BT_ULONG,
BT_ULONG, BT_PTR_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_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/c-family/c-common.c b/gcc/c-family/c-common.c
index c87704b..53f92f7 100644
--- a/gcc/c-family/c-common.c
+++ b/gcc/c-family/c-common.c
@@ -830,6 +830,7 @@ const struct attribute_spec c_common_attribute_table[] =
handle_bnd_legacy, false },
{ "bnd_instrument", 0, 0, true, false, false,
handle_bnd_instrument, false },
+ { "oacc declare", 0, -1, true, false, false, NULL, false },
{ NULL, 0, 0, false, false, false, NULL, false }
};
diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c
index d99c2af..ad8cdbf 100644
--- a/gcc/c-family/c-pragma.c
+++ b/gcc/c-family/c-pragma.c
@@ -1206,6 +1206,7 @@ struct omp_pragma_def { const char *name; unsigned int id; };
static const struct omp_pragma_def oacc_pragmas[] = {
{ "cache", PRAGMA_OACC_CACHE },
{ "data", PRAGMA_OACC_DATA },
+ { "declare", PRAGMA_OACC_DECLARE },
{ "enter", PRAGMA_OACC_ENTER_DATA },
{ "exit", PRAGMA_OACC_EXIT_DATA },
{ "kernels", PRAGMA_OACC_KERNELS },
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index cec920f..dcba221 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -29,6 +29,7 @@ enum pragma_kind {
PRAGMA_OACC_CACHE,
PRAGMA_OACC_DATA,
+ PRAGMA_OACC_DECLARE,
PRAGMA_OACC_ENTER_DATA,
PRAGMA_OACC_EXIT_DATA,
PRAGMA_OACC_KERNELS,
@@ -150,6 +151,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_CREATE,
PRAGMA_OACC_CLAUSE_DELETE,
PRAGMA_OACC_CLAUSE_DEVICEPTR,
+ PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_HOST,
PRAGMA_OACC_CLAUSE_NUM_GANGS,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 90038d5..c21a274 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -1229,6 +1229,7 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *, bool, bool,
vec<tree, va_gc> **, location_t *,
tree *, vec<location_t> *,
unsigned int * = NULL);
+static void c_parser_oacc_declare (c_parser *);
static void c_parser_oacc_enter_exit_data (c_parser *, bool);
static void c_parser_oacc_update (c_parser *);
static void c_parser_omp_construct (c_parser *);
@@ -9695,6 +9696,10 @@ c_parser_pragma (c_parser *parser, enum pragma_context context)
switch (id)
{
+ case PRAGMA_OACC_DECLARE:
+ c_parser_oacc_declare (parser);
+ return false;
+
case PRAGMA_OACC_ENTER_DATA:
c_parser_oacc_enter_exit_data (parser, true);
return false;
@@ -9980,6 +9985,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_DEVICE;
else if (!strcmp ("deviceptr", p))
result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
+ else if (!strcmp ("device_resident", p))
+ result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT;
else if (!strcmp ("dist_schedule", p))
result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
break;
@@ -10412,10 +10419,16 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_DEVICE:
kind = GOMP_MAP_FORCE_TO;
break;
+ case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
+ kind = GOMP_MAP_DEVICE_RESIDENT;
+ break;
case PRAGMA_OACC_CLAUSE_HOST:
case PRAGMA_OACC_CLAUSE_SELF:
kind = GOMP_MAP_FORCE_FROM;
break;
+ case PRAGMA_OMP_CLAUSE_LINK:
+ kind = GOMP_MAP_LINK;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@@ -12584,6 +12597,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses);
c_name = "deviceptr";
break;
+ case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "device_resident";
+ break;
case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
clauses = c_parser_omp_clause_firstprivate (parser, clauses);
c_name = "firstprivate";
@@ -12601,6 +12618,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_omp_clause_if (parser, clauses, false);
c_name = "if";
break;
+ case PRAGMA_OMP_CLAUSE_LINK:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "link";
+ break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
clauses = c_parser_omp_clause_num_gangs (parser, clauses);
c_name = "num_gangs";
@@ -13054,6 +13075,220 @@ c_parser_oacc_data (location_t loc, c_parser *parser)
}
/* OpenACC 2.0:
+ # pragma acc declare oacc-data-clause[optseq] new-line
+*/
+
+#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_OMP_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 void
+c_parser_oacc_declare (c_parser *parser)
+{
+ location_t pragma_loc = c_parser_peek_token (parser)->location;
+ tree c, clauses, ret_clauses, stmt, t;
+
+ bool error = false;
+
+ c_parser_consume_pragma (parser);
+
+ clauses = c_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
+ "#pragma acc declare");
+ if (!clauses)
+ {
+ error_at (pragma_loc,
+ "no valid clauses specified in %<#pragma acc declare%>");
+ return;
+ }
+
+ for (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;
+ }
+
+ 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;
+ }
+ else if (TREE_PUBLIC (decl))
+ {
+ error_at (loc,
+ "invalid use of %<global%> variable %qD "
+ "in %<#pragma acc declare%>", decl);
+ error = true;
+ continue;
+ }
+ break;
+ }
+
+ tree id;
+ tree prev_attr = lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (decl));
+
+ if (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_LINK)
+ id = get_identifier ("omp declare target link");
+ else
+ id = get_identifier ("omp declare target");
+
+ 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 (cl), "previous directive was here");
+ error = true;
+ continue;
+ }
+ }
+
+ if (!error)
+ DECL_ATTRIBUTES (decl) =
+ tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl));
+ }
+
+ if (error || global_bindings_p ())
+ return;
+
+ 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;
+ }
+ }
+
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DECLARE_CLAUSES (stmt) = clauses;
+ OACC_DECLARE_RETURN_CLAUSES (stmt) = ret_clauses;
+ SET_EXPR_LOCATION (stmt, pragma_loc);
+
+ add_stmt (stmt);
+
+ return;
+}
+
+/* OpenACC 2.0:
# pragma acc enter data oacc-enter-data-clause[optseq] new-line
or
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 24cb47f..6d47352 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -29094,6 +29094,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_DEVICE;
else if (!strcmp ("deviceptr", p))
result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
+ else if (!strcmp ("device_resident", p))
+ result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT;
else if (!strcmp ("dist_schedule", p))
result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
break;
@@ -29503,10 +29505,16 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_DEVICE:
kind = GOMP_MAP_FORCE_TO;
break;
+ case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
+ kind = GOMP_MAP_DEVICE_RESIDENT;
+ break;
case PRAGMA_OACC_CLAUSE_HOST:
case PRAGMA_OACC_CLAUSE_SELF:
kind = GOMP_MAP_FORCE_FROM;
break;
+ case PRAGMA_OMP_CLAUSE_LINK:
+ kind = GOMP_MAP_LINK;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@@ -31475,6 +31483,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses);
c_name = "deviceptr";
break;
+ case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "device_resident";
+ break;
case PRAGMA_OACC_CLAUSE_GANG:
c_name = "gang";
clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG,
@@ -31488,6 +31500,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_omp_clause_if (parser, clauses, here, false);
c_name = "if";
break;
+ case PRAGMA_OMP_CLAUSE_LINK:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "link";
+ break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
clauses = cp_parser_omp_clause_num_gangs (parser, clauses);
c_name = "num_gangs";
@@ -34380,6 +34396,221 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok)
}
/* OpenACC 2.0:
+ # pragma acc declare oacc-data-clause[optseq] new-line
+*/
+
+#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_OMP_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 c, clauses, ret_clauses, stmt, t;
+ bool error = false;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
+ "#pragma acc declare", pragma_tok, true);
+
+
+ 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;
+ }
+ else if (TREE_PUBLIC (decl))
+ {
+ error_at (loc,
+ "invalid use of %<global%> variable %qD "
+ "in %<#pragma acc declare%>", decl);
+ error = true;
+ continue;
+ }
+ break;
+ }
+
+ tree id;
+ tree prev_attr = lookup_attribute ("oacc declare",
+ DECL_ATTRIBUTES (decl));
+
+ if (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_LINK)
+ id = get_identifier ("omp declare target link");
+ else
+ id = get_identifier ("omp declare target");
+
+ 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;
+ }
+ }
+
+ if (!error)
+ DECL_ATTRIBUTES (decl) =
+ tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl));
+ }
+
+ if (error || global_bindings_p ())
+ return NULL_TREE;
+
+ 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;
+ }
+ }
+
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DECLARE_CLAUSES (stmt) = clauses;
+ OACC_DECLARE_RETURN_CLAUSES (stmt) = ret_clauses;
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+
+ add_stmt (stmt);
+
+ return NULL_TREE;
+}
+
+/* OpenACC 2.0:
# pragma acc enter data oacc-enter-data-clause[optseq] new-line
or
@@ -36040,6 +36271,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_CACHE:
case PRAGMA_OACC_DATA:
case PRAGMA_OACC_ENTER_DATA:
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index e836ec7..0a6b190 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -15314,6 +15314,17 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
add_stmt (t);
break;
+ case OACC_DECLARE:
+ t = copy_node (t);
+ tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false, false,
+ args, complain, in_decl);
+ OACC_DECLARE_CLAUSES (t) = tmp;
+ tmp = tsubst_omp_clauses (OACC_DECLARE_RETURN_CLAUSES (t), false, false,
+ args, complain, in_decl);
+ OACC_DECLARE_RETURN_CLAUSES (t) = tmp;
+ add_stmt (t);
+ break;
+
case OMP_TARGET_UPDATE:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index ca75654..6d993db 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -145,6 +145,7 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I2_INT, BT_VOID, BT_VOLATILE_PTR, BT_I2, BT
DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I4_INT, BT_VOID, BT_VOLATILE_PTR, BT_I4, BT_INT)
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_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)
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 781801b..e45162d 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -170,6 +170,7 @@ enum gf_mask {
GF_OMP_TARGET_KIND_OACC_DATA = 7,
GF_OMP_TARGET_KIND_OACC_UPDATE = 8,
GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 9,
+ GF_OMP_TARGET_KIND_OACC_DECLARE = 10,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
@@ -6004,6 +6005,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 03203c0..0b685b9 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -154,6 +154,7 @@ struct gimplify_omp_ctx
bool target_map_scalars_firstprivate;
bool target_map_pointers_as_0len_arrays;
bool target_firstprivatize_array_bases;
+ gomp_target *declare_returns;
};
static struct gimplify_ctx *gimplify_ctxp;
@@ -373,6 +374,7 @@ new_omp_context (enum omp_region_type region_type)
c->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
else
c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
+ c->declare_returns = NULL;
return c;
}
@@ -5789,6 +5791,26 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
return false;
}
+/* Return true if global var DECL is device resident. */
+
+static bool
+device_resident_p (tree decl)
+{
+ tree attr = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
+
+ if (!attr)
+ return false;
+
+ for (tree t = TREE_VALUE (attr); t; t = TREE_PURPOSE (t))
+ {
+ tree c = TREE_VALUE (t);
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DEVICE_RESIDENT)
+ return true;
+ }
+
+ return false;
+}
+
/* Determine outer default flags for DECL mentioned in an OMP region
but not declared in an enclosing clause.
@@ -5838,6 +5860,8 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
flags |= GOVD_FIRSTPRIVATE;
break;
case OMP_CLAUSE_DEFAULT_UNSPECIFIED:
+ if (is_global_var (decl) && device_resident_p (decl))
+ flags |= GOVD_MAP_TO_ONLY | GOVD_MAP;
/* decl will be either GOVD_FIRSTPRIVATE or GOVD_SHARED. */
gcc_assert ((ctx->region_type & ORT_TASK) != 0);
if (struct gimplify_omp_ctx *octx = ctx->outer_context)
@@ -7520,6 +7544,62 @@ gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p)
*expr_p = NULL_TREE;
}
+/* Gimplify OACC_DECLARE. */
+
+static void
+gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
+{
+ tree expr = *expr_p;
+ gomp_target *stmt;
+ tree clauses, t;
+
+ clauses = OACC_DECLARE_CLAUSES (expr);
+
+ gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE);
+
+ for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+ {
+ tree attrs, decl = OMP_CLAUSE_DECL (t);
+
+ if (TREE_CODE (decl) == MEM_REF)
+ continue;
+
+ omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN);
+
+ attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
+ if (attrs)
+ DECL_ATTRIBUTES (decl) = remove_attribute ("oacc declare", attrs);
+ }
+
+ stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
+ clauses);
+
+ gimplify_seq_add_stmt (pre_p, stmt);
+
+ clauses = OACC_DECLARE_RETURN_CLAUSES (expr);
+ if (clauses)
+ {
+ struct gimplify_omp_ctx *c;
+
+ /* Any clauses that affect the state of a variable prior
+ to return are saved and dealt with after the body has
+ been gimplified. */
+
+ gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA,
+ OACC_DECLARE);
+
+ c = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = c->outer_context;
+ delete_omp_context (c);
+
+ stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
+ clauses);
+ gimplify_omp_ctxp->declare_returns = stmt;
+ }
+
+ *expr_p = NULL_TREE;
+}
+
/* Gimplify the contents of an OMP_PARALLEL statement. This involves
gimplification of the body, as well as scanning the body for used
variables. We need to do this scan now, because variable-sized
@@ -9586,11 +9666,15 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
break;
case OACC_HOST_DATA:
- case OACC_DECLARE:
sorry ("directive not yet implemented");
ret = GS_ALL_DONE;
break;
+ case OACC_DECLARE:
+ gimplify_oacc_declare (expr_p, pre_p);
+ ret = GS_ALL_DONE;
+ break;
+
case OACC_DATA:
case OACC_KERNELS:
case OACC_PARALLEL:
@@ -10251,6 +10335,28 @@ gimplify_body (tree fndecl, bool do_parms)
gimplify_seq_add_stmt (&seq, outer_stmt);
}
+ if (flag_openacc && gimplify_omp_ctxp)
+ {
+ while (gimplify_omp_ctxp)
+ {
+ struct gimplify_omp_ctx *c;
+
+ if (gimplify_omp_ctxp->declare_returns)
+ {
+ /* Clauses are present that affect the state of a
+ variable, insert the statment to handle this
+ as the very last statement. */
+
+ gimplify_seq_add_stmt (&seq, gimplify_omp_ctxp->declare_returns);
+ gimplify_omp_ctxp->declare_returns = NULL;
+ }
+
+ c = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = c->outer_context;
+ delete_omp_context (c);
+ }
+ }
+
/* The body must contain exactly one statement, a GIMPLE_BIND. If this is
not the case, wrap everything in a GIMPLE_BIND to make it so. */
if (gimple_code (outer_stmt) == GIMPLE_BIND
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index ea9cf0d..4af3640 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -317,3 +317,7 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA,
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_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 d0264e9..215adfa 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -12120,6 +12120,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:
@@ -12351,6 +12352,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 ();
}
@@ -12473,6 +12477,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_DATA:
break;
case BUILT_IN_GOMP_TARGET:
@@ -12755,6 +12760,7 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
case GF_OMP_TARGET_KIND_EXIT_DATA:
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;
@@ -14968,6 +14974,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:
@@ -15042,6 +15049,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_FORCE_PRESENT:
case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_LINK:
gcc_assert (is_gimple_omp_oacc (stmt));
break;
default:
@@ -16710,6 +16719,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
case GF_OMP_TARGET_KIND_EXIT_DATA:
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/c-c++-common/goacc/declare-1.c b/gcc/testsuite/c-c++-common/goacc/declare-1.c
new file mode 100644
index 0000000..b036c63
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/declare-1.c
@@ -0,0 +1,83 @@
+/* Test valid uses of declare directive. */
+/* { dg-do compile } */
+
+int v0;
+#pragma acc declare create(v0)
+
+int v1;
+#pragma acc declare copyin(v1)
+
+int *v2;
+#pragma acc declare deviceptr(v2)
+
+int v3;
+#pragma acc declare device_resident(v3)
+
+int v4;
+#pragma acc declare link(v4)
+
+int v5, v6, v7, v8;
+#pragma acc declare create(v5, v6) copyin(v7, v8)
+
+void
+f (void)
+{
+ int va0;
+#pragma acc declare create(va0)
+
+ int va1;
+#pragma acc declare copyin(va1)
+
+ int *va2;
+#pragma acc declare deviceptr(va2)
+
+ int va3;
+#pragma acc declare device_resident(va3)
+
+ extern int ve0;
+#pragma acc declare create(ve0)
+
+ extern int ve1;
+#pragma acc declare copyin(ve1)
+
+ extern int *ve2;
+#pragma acc declare deviceptr(ve2)
+
+ extern int ve3;
+#pragma acc declare device_resident(ve3)
+
+ extern int ve4;
+#pragma acc declare link(ve4)
+
+ int va5;
+#pragma acc declare copy(va5)
+
+ int va6;
+#pragma acc declare copyout(va6)
+
+ int va7;
+#pragma acc declare present(va7)
+
+ int va8;
+#pragma acc declare present_or_copy(va8)
+
+ int va9;
+#pragma acc declare present_or_copyin(va9)
+
+ int va10;
+#pragma acc declare present_or_copyout(va10)
+
+ int va11;
+#pragma acc declare present_or_create(va11)
+
+ a:
+ {
+ int va0;
+#pragma acc declare create(va0)
+ if (v1)
+ goto a;
+ else
+ goto b;
+ }
+ b:;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c
new file mode 100644
index 0000000..7979f0c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c
@@ -0,0 +1,68 @@
+/* Test invalid uses of declare directive. */
+/* { dg-do compile } */
+
+#pragma acc declare /* { dg-error "no valid clauses" } */
+
+#pragma acc declare create(undeclared) /* { dg-error "undeclared" } */
+/* { dg-error "no valid clauses" "second error" { target *-*-* } 6 } */
+
+int v0[10];
+#pragma acc declare create(v0[1:3]) /* { dg-error "subarray" } */
+
+int v1;
+#pragma acc declare create(v1, v1) /* { dg-error "more than once" } */
+
+int v2;
+#pragma acc declare create(v2) /* { dg-message "previous directive" } */
+#pragma acc declare copyin(v2) /* { dg-error "more than once" } */
+
+int v3;
+#pragma acc declare copy(v3) /* { dg-error "at file scope" } */
+
+int v4;
+#pragma acc declare copyout(v4) /* { dg-error "at file scope" } */
+
+int v5;
+#pragma acc declare present(v5) /* { dg-error "at file scope" } */
+
+int v6;
+#pragma acc declare present_or_copy(v6) /* { dg-error "at file scope" } */
+
+int v7;
+#pragma acc declare present_or_copyin(v7) /* { dg-error "at file scope" } */
+
+int v8;
+#pragma acc declare present_or_copyout(v8) /* { dg-error "at file scope" } */
+
+int v9;
+#pragma acc declare present_or_create(v9) /* { dg-error "at file scope" } */
+
+void
+f (void)
+{
+ int va0;
+#pragma acc declare link(va0) /* { dg-error "global variable" } */
+
+ extern int ve0;
+#pragma acc declare copy(ve0) /* { dg-error "invalid use of" } */
+
+ extern int ve1;
+#pragma acc declare copyout(ve1) /* { dg-error "invalid use of" } */
+
+ extern int ve2;
+#pragma acc declare present(ve2) /* { dg-error "invalid use of" } */
+
+ extern int ve3;
+#pragma acc declare present_or_copy(ve3) /* { dg-error "invalid use of" } */
+
+ extern int ve4;
+#pragma acc declare present_or_copyin(ve4) /* { dg-error "invalid use of" } */
+
+ extern int ve5;
+#pragma acc declare present_or_copyout(ve5) /* { dg-error "invalid use of" } */
+
+ extern int ve6;
+#pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */
+
+#pragma acc declare present (v9) /* { dg-error "invalid use of" } */
+}
diff --git a/gcc/tree.def b/gcc/tree.def
index fc7490a..d6269c0 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1186,8 +1186,9 @@ DEFTREECODE (OMP_TASKGROUP, "omp_taskgroup", tcc_statement, 1)
DEFTREECODE (OACC_CACHE, "oacc_cache", tcc_statement, 1)
/* OpenACC - #pragma acc declare [clause1 ... clauseN]
- Operand 0: OACC_DECLARE_CLAUSES: List of clauses. */
-DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 1)
+ Operand 0: OACC_DECLARE_CLAUSES: List of clauses.
+ Operand 1: OACC_DECLARE_RETURN_CLAUSES: List of clauses for returns. */
+DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 2)
/* OpenACC - #pragma acc enter data [clause1 ... clauseN]
Operand 0: OACC_ENTER_DATA_CLAUSES: List of clauses. */
diff --git a/gcc/tree.h b/gcc/tree.h
index 65c3117..66b97bc 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1232,6 +1232,8 @@ extern void protected_set_expr_location (tree, location_t);
#define OACC_DECLARE_CLAUSES(NODE) \
TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 0)
+#define OACC_DECLARE_RETURN_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 1)
#define OACC_ENTER_DATA_CLAUSES(NODE) \
TREE_OPERAND (OACC_ENTER_DATA_CHECK (NODE), 0)
diff --git a/gcc/varpool.c b/gcc/varpool.c
index 478f365..a8cdb1c 100644
--- a/gcc/varpool.c
+++ b/gcc/varpool.c
@@ -137,18 +137,9 @@ varpool_node::create_empty (void)
return node;
}
-/* Return varpool node assigned to DECL. Create new one when needed. */
-varpool_node *
-varpool_node::get_create (tree decl)
+static void
+make_offloadable (varpool_node *node, tree decl)
{
- varpool_node *node = varpool_node::get (decl);
- gcc_checking_assert (TREE_CODE (decl) == VAR_DECL);
- if (node)
- 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)))
{
@@ -160,6 +151,25 @@ varpool_node::get_create (tree decl)
node->force_output = 1;
#endif
}
+}
+
+/* Return varpool node assigned to DECL. Create new one when needed. */
+varpool_node *
+varpool_node::get_create (tree decl)
+{
+ varpool_node *node = varpool_node::get (decl);
+ gcc_checking_assert (TREE_CODE (decl) == VAR_DECL);
+ if (node)
+ {
+ if (!node->offloadable)
+ make_offloadable (node, decl);
+ return node;
+ }
+
+ node = varpool_node::create_empty ();
+ node->decl = decl;
+
+ make_offloadable (node, decl);
node->register_symbol ();
return node;
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index f834dec..4128912 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -73,6 +73,11 @@ enum gomp_map_kind
POINTER_SIZE_UNITS. */
GOMP_MAP_FORCE_DEVICEPTR = (GOMP_MAP_FLAG_SPECIAL_1 | 0),
/* Do not map, copy bits for firstprivate instead. */
+ /* OpenACC device_resident. */
+ GOMP_MAP_DEVICE_RESIDENT = (GOMP_MAP_FLAG_SPECIAL_1 | 1),
+ /* OpenACC link. */
+ GOMP_MAP_LINK = (GOMP_MAP_FLAG_SPECIAL_1 | 2),
+ /* Allocate. */
GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
/* Similarly, but store the value in the pointer rather than
pointed by the pointer. */
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 2153661..e96f929 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -378,10 +378,12 @@ GOACC_2.0 {
GOACC_wait;
GOACC_get_thread_num;
GOACC_get_num_threads;
+ GOACC_register_static;
};
GOACC_2.0.1 {
global:
+ GOACC_declare;
GOACC_parallel_keyed;
} GOACC_2.0;
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index b150106..fd9348c 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -297,7 +297,9 @@ GOACC_enter_exit_data (int device, size_t mapnum,
if (kind == GOMP_MAP_FORCE_ALLOC
|| kind == GOMP_MAP_FORCE_PRESENT
- || kind == GOMP_MAP_FORCE_TO)
+ || kind == GOMP_MAP_FORCE_TO
+ || kind == GOMP_MAP_TO
+ || kind == GOMP_MAP_ALLOC)
{
data_enter = true;
break;
@@ -324,6 +326,9 @@ GOACC_enter_exit_data (int device, size_t mapnum,
{
switch (kind)
{
+ case GOMP_MAP_ALLOC:
+ acc_present_or_create (hostaddrs[i], sizes[i]);
+ break;
case GOMP_MAP_POINTER:
gomp_acc_insert_pointer (1, &hostaddrs[i], &sizes[i],
&kinds[i]);
@@ -332,6 +337,7 @@ GOACC_enter_exit_data (int device, size_t mapnum,
acc_create (hostaddrs[i], sizes[i]);
break;
case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_TO:
acc_present_or_copyin (hostaddrs[i], sizes[i]);
break;
case GOMP_MAP_FORCE_TO:
@@ -501,3 +507,62 @@ GOACC_get_thread_num (void)
{
return 0;
}
+
+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/declare-1.c b/libgomp/testsuite/declare-1.c
new file mode 100644
index 0000000..8fbec4d
--- /dev/null
+++ b/libgomp/testsuite/declare-1.c
@@ -0,0 +1,122 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <openacc.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+#define N 8
+
+void
+subr2 (int *a)
+{
+ int i;
+ int f[N];
+#pragma acc declare copyout (f)
+
+#pragma acc parallel copy (a[0:N])
+ {
+ for (i = 0; i < N; i++)
+ {
+ f[i] = a[i];
+ a[i] = f[i] + f[i] + f[i];
+ }
+ }
+}
+
+void
+subr1 (int *a)
+{
+ int f[N];
+#pragma acc declare copy (f)
+
+#pragma acc parallel copy (a[0:N])
+ {
+ int i;
+
+ for (i = 0; i < N; i++)
+ {
+ f[i] = a[i];
+ a[i] = f[i] + f[i];
+ }
+ }
+}
+
+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)
+{
+ 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 ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 1234;
+ }
+
+ subr1 (&a[0]);
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 1234 * 2)
+ abort ();
+ }
+
+ subr2 (&a[0]);
+
+ for (i = 0; i < 1; i++)
+ {
+ if (a[i] != 1234 * 6)
+ abort ();
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/declare-5.c b/libgomp/testsuite/declare-5.c
new file mode 100644
index 0000000..1e2f6ce
--- /dev/null
+++ b/libgomp/testsuite/declare-5.c
@@ -0,0 +1,13 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+int
+main (int argc, char **argv)
+{
+ int a[8] __attribute__((unused));
+
+ __builtin_printf ("CheCKpOInT\n");
+#pragma acc declare present (a)
+}
+
+/* { dg-output "CheCKpOInT" } */
+/* { dg-shouldfail "" } */
| Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
|---|---|---|
| Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |