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! 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. > --- 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? > --- 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. > + | (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? > + > + 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... > + > + if (!attr) > + return false; > + ... will always return "false" here, and this is dead code: > + 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". > /* 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. > --- 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? > --- 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.) Grüße Thomas
signature.asc
Description: PGP signature
