Hi Thomas,
On 11/04/2015 10:49 AM, Thomas Schwinge wrote:
Hi Jim!
On Tue, 3 Nov 2015 10:31:32 -0600, James Norris <jnor...@codesourcery.com>
wrote:
On 10/27/2015 03:18 PM, James Norris wrote:
This patch adds the processing of OpenACC declare directive in C
and C++. (Note: Support in Fortran is already in trunk.)
..., and a patch adjusting some Fortran front end things is awaiting
review,
<http://news.gmane.org/find-root.php?message_id=%3C5637692F.7050306%40codesourcery.com%3E>.
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.
Oh, wow, you could remove a lot of code!
Yes, I missed that patch when it entered into the code base. My bad.
Just a superficial review on your patch; patch re-ordered a bit for
review.
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
+DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_UINT, BT_VOID, BT_PTR, BT_INT, BT_UINT)
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
+DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_UINT, BT_VOID, BT_PTR, BT_INT, BT_UINT)
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_STATIC, "GOACC_register_static",
+ BT_FN_VOID_PTR_INT_UINT, ATTR_NOTHROW_LIST)
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
+ GOACC_register_static;
I think these changes can be dropped -- assuming you have not
unintentionally dropped the GOACC_register_static function/usage in your
v2 patch.
Will fix.
--- 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[] =
+ { "oacc declare", 0, -1, true, false, false, NULL, false },
As far as I can tell, nothing is setting this attribute anymore in your
v2 patch, so I guess this and all handling code ("lookup_attribute",
"remove_attribute") can also be dropped?
Will fix.
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
/* 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) \
For uniformity, please use/add a new alias "PRAGMA_OACC_CLAUSE_* =
PRAGMA_OMP_CLAUSE_LINK" instead of using PRAGMA_OMP_CLAUSE_* here, and
also in c_parser_oacc_data_clause, c_parser_oacc_all_clauses, and in the
C++ front end.
Will fix.
+ | (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)
+{
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
+static tree
+cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
+{
+ [...]
+ tree prev_attr = lookup_attribute ("oacc declare",
+ DECL_ATTRIBUTES (decl));
Per my comment above, this would always be NULL_TREE. The C front end is
different?
Will fix.
+
+ 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;
+ }
+ }
--- 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;
Note to Jakub et al.: code for handling OACC_* is generally missing here,
also for other constructs and clauses; we'll be adding that.
--- 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,
Need to update gcc/gimple-pretty-print.c:dump_gimple_omp_target.
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
+/* 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));
Will always be NULL_TREE, as far as I can tell, so...
Will fix.
+
+ if (!attr)
+ return false;
+
... will always return "false" here, and this is dead code:
Will fix.
+ 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;
+}
@@ -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;
Unreachable condition if device_resident_p always returns "false".
Will fix.
/* 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)
+/* 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);
As above, obsolete "oacc declare" attribute.
Will fix.
--- 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. */
Confused -- I don't see these two getting handled in libgomp?
These won't be 'seen' by libgomp. So should these
be defined by some other means?
--- 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:
(As far as I can tell, these three hunks are not related to OpenACC
declare, but a bug fix for OpenACC enter/exit data. Will submit that
later on, with test cases.)
I'll eliminate the three hunks from the patch.
Thank you for taking the time to review the patch.
Jim
