From patchwork Fri Nov 11 23:43:02 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 81905 Delivered-To: patch@linaro.org Received: by 10.140.97.165 with SMTP id m34csp9422qge; Fri, 11 Nov 2016 15:43:35 -0800 (PST) X-Received: by 10.98.214.157 with SMTP id a29mr11703221pfl.162.1478907814130; Fri, 11 Nov 2016 15:43:34 -0800 (PST) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id n65si12639677pga.6.2016.11.11.15.43.33 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Fri, 11 Nov 2016 15:43:34 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-return-441208-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) client-ip=209.132.180.131; Authentication-Results: mx.google.com; dkim=pass header.i=@gcc.gnu.org; spf=pass (google.com: domain of gcc-patches-return-441208-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-441208-patch=linaro.org@gcc.gnu.org DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :subject:to:cc:message-id:date:mime-version:content-type; q=dns; s=default; b=QUF+tVT/LQvPP2fQS/lCWMv97VLZAgSdtxQcoohbrj5aHYHuYc 4A7H7apDojw4nznOxzYBwzy4u8LfAFBVK7kFZwGiNaywDO24Ej1AyNWh10jDRwJb qaVl6E55RvMjttUK0khRc1Pfz7fovuFAlUL7dQ3KvuSfwH+f70oeGekjA= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :subject:to:cc:message-id:date:mime-version:content-type; s= default; bh=Hn/uWOgnt8ApF9Jg1QKuJer3n/E=; b=PUGJVPItjN2txJqEr5qX ufmzeVs2GWgP1y3a5DHjWQ8kpXViik6xjwsAQrpZz3hx1f6Hk2Y2HTSeb1I5WMJy 3BS+STWM12IzwIG2n8CZBYuGBeY3bCRd/+nv/TrIO0orF+W8e/XrUMuSWHl5/+SQ Wo+qwu9I+43QvXwZKtHuvy8= Received: (qmail 50067 invoked by alias); 11 Nov 2016 23:43:20 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 49121 invoked by uid 89); 11 Nov 2016 23:43:19 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=sk:c155d56, FNDECL, sk:TREE_ST, sk:tree_st X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 11 Nov 2016 23:43:08 +0000 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtp id 1c5LTO-0002e0-1R from Cesar_Philippidis@mentor.com ; Fri, 11 Nov 2016 15:43:06 -0800 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Fri, 11 Nov 2016 15:43:03 -0800 From: Cesar Philippidis Subject: [PATCH] OpenACC routines -- middle end To: "gcc-patches@gcc.gnu.org" , Jakub Jelinek CC: Thomas Schwinge Message-ID: <5aa1ff0a-f5a6-ef64-aaf1-0666eed4fd42@codesourcery.com> Date: Fri, 11 Nov 2016 15:43:02 -0800 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.3.0 MIME-Version: 1.0 X-ClientProxiedBy: svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) To svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) Currently GCC lacks support for the bind and nohost clauses in OpenACC routine. Furthermore, none of the FEs preform much error handling to detect incompatible acc loops inside those functions. This patch adds the common middle end components, namely tree codes for the clauses, and OMP lowering and gimplification code for diagnostics and code generation. I've also included the changes to c-family, because I wanted to break out the FE changes into separate patches. Is this patch OK for trunk? Cesar 2016-11-11 Cesar Philippidis Thomas Schwinge gcc/c-family/ * c-attribs.c (c_common_attribute_table): Adjust "omp declare target". * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_BIND and PRAGMA_OACC_CLAUSE_NOHOST. gcc/ * gimplify.c (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_BIND and OMP_CLAUSE_NOHOST. (gimplify_adjust_omp_clauses): Likewise. * omp-low.c (scan_sharing_clauses): Likewise. (verify_oacc_routine_clauses): New function. (maybe_discard_oacc_function): New function. (execute_oacc_device_lower): Don't generate code for NOHOST. * omp-low.h (verify_oacc_routine_clauses): Declare. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_BIND and OMP_CLAUSE_NOHOST. * tree-pretty-print.c (dump_omp_clause): Likewise. * tree.c (omp_clause_num_ops): Likewise. (omp_clause_code_name): Likewise. (walk_tree_1): Handle OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST. * tree.h (OMP_CLAUSE_BIND_NAME): Define. diff --git a/gcc/c-family/c-attribs.c b/gcc/c-family/c-attribs.c index 925f1b2..55c53ea 100644 --- a/gcc/c-family/c-attribs.c +++ b/gcc/c-family/c-attribs.c @@ -322,7 +322,7 @@ const struct attribute_spec c_common_attribute_table[] = handle_omp_declare_simd_attribute, false }, { "simd", 0, 1, true, false, false, handle_simd_attribute, false }, - { "omp declare target", 0, 0, true, false, false, + { "omp declare target", 0, -1, true, false, false, handle_omp_declare_target_attribute, false }, { "omp declare target link", 0, 0, true, false, false, handle_omp_declare_target_attribute, false }, diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 6d9cb08..dd2722a 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -149,6 +149,7 @@ enum pragma_omp_clause { /* Clauses for OpenACC. */ PRAGMA_OACC_CLAUSE_ASYNC = PRAGMA_CILK_CLAUSE_VECTORLENGTH + 1, PRAGMA_OACC_CLAUSE_AUTO, + PRAGMA_OACC_CLAUSE_BIND, PRAGMA_OACC_CLAUSE_COPY, PRAGMA_OACC_CLAUSE_COPYOUT, PRAGMA_OACC_CLAUSE_CREATE, @@ -158,6 +159,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST, PRAGMA_OACC_CLAUSE_INDEPENDENT, + PRAGMA_OACC_CLAUSE_NOHOST, PRAGMA_OACC_CLAUSE_NUM_GANGS, PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT, diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 16573dd..c1d24fc 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8373,6 +8373,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c); break; + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: default: gcc_unreachable (); } @@ -9112,6 +9114,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, remove = true; break; + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: default: gcc_unreachable (); } diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 331da6a..13f186e 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -2201,6 +2201,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, install_var_local (decl, ctx); break; + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_TILE: case OMP_CLAUSE__CACHE_: default: @@ -2365,6 +2367,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, case OMP_CLAUSE__GRIDDIM_: break; + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_TILE: case OMP_CLAUSE__CACHE_: default: @@ -12684,9 +12688,192 @@ set_oacc_fn_attrib (tree fn, tree clauses, bool is_kernel, vec *args) } } -/* Process the routine's dimension clauess to generate an attribute - value. Issue diagnostics as appropriate. We default to SEQ - (OpenACC 2.5 clarifies this). All dimensions have a size of zero +/* Verify OpenACC routine clauses. + + Returns 0 if FNDECL should be marked as an accelerator routine, 1 if it has + already been marked in compatible way, and -1 if incompatible. Upon + returning, the chain of clauses will contain exactly one clause specifying + the level of parallelism. */ + +int +verify_oacc_routine_clauses (tree fndecl, tree *clauses, location_t loc, + const char *routine_str) +{ + tree c_level = NULL_TREE; + tree c_bind = NULL_TREE; + tree c_nohost = NULL_TREE; + tree c_p = NULL_TREE; + for (tree c = *clauses; c; c_p = c, c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_SEQ: + if (c_level == NULL_TREE) + c_level = c; + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_CODE (c_level)) + { + /* This has already been diagnosed in the front ends. */ + /* Drop the duplicate clause. */ + gcc_checking_assert (c_p != NULL_TREE); + OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c); + c = c_p; + } + else + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qs specifies a conflicting level of parallelism", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + inform (OMP_CLAUSE_LOCATION (c_level), + "... to the previous %qs clause here", + omp_clause_code_name[OMP_CLAUSE_CODE (c_level)]); + /* Drop the conflicting clause. */ + gcc_checking_assert (c_p != NULL_TREE); + OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c); + c = c_p; + } + break; + case OMP_CLAUSE_BIND: + /* Don't bother with duplicate clauses at this point. */ + c_bind = c; + break; + case OMP_CLAUSE_NOHOST: + /* Don't bother with duplicate clauses at this point. */ + c_nohost = c; + break; + default: + gcc_unreachable (); + } + if (c_level == NULL_TREE) + { + /* OpenACC 2.5 makes this an error; for the current OpenACC 2.0a + implementation add an implicit "seq" clause. */ + c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ); + OMP_CLAUSE_CHAIN (c_level) = *clauses; + *clauses = c_level; + } + /* In *clauses, we now have exactly one clause specifying the level of + parallelism. */ + + /* Still got some work to do for Fortran... */ + if (fndecl == NULL_TREE) + return 0; + + tree attr + = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl)); + if (attr != NULL_TREE) + { + /* If a "#pragma acc routine" has already been applied, just verify + this one for compatibility. */ + /* Collect previous directive's clauses. */ + tree c_level_p = NULL_TREE; + tree c_bind_p = NULL_TREE; + tree c_nohost_p = NULL_TREE; + for (tree c = TREE_VALUE (attr); c; c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_SEQ: + gcc_checking_assert (c_level_p == NULL_TREE); + c_level_p = c; + break; + case OMP_CLAUSE_BIND: + /* Don't bother with duplicate clauses at this point. */ + c_bind_p = c; + break; + case OMP_CLAUSE_NOHOST: + /* Don't bother with duplicate clauses at this point. */ + c_nohost_p = c; + break; + default: + gcc_unreachable (); + } + gcc_checking_assert (c_level_p != NULL_TREE); + /* ..., and compare to current directive's, which we've already collected + above. */ + tree c_diag; + tree c_diag_p; + /* Matching level of parallelism? */ + if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p)) + { + c_diag = c_level; + c_diag_p = c_level_p; + goto incompatible; + } + /* Matching bind clauses? */ + if ((c_bind == NULL_TREE) != (c_bind_p == NULL_TREE)) + { + c_diag = c_bind; + c_diag_p = c_bind_p; + goto incompatible; + } + /* Matching bind clauses' names? */ + if ((c_bind != NULL_TREE) && (c_bind_p != NULL_TREE)) + { + tree c_bind_name = OMP_CLAUSE_BIND_NAME (c_bind); + tree c_bind_name_p = OMP_CLAUSE_BIND_NAME (c_bind_p); + /* TODO: will/should actually be the trees/strings/string pointers be + identical? */ + if (strcmp (TREE_STRING_POINTER (c_bind_name), + TREE_STRING_POINTER (c_bind_name_p)) != 0) + { + c_diag = c_bind; + c_diag_p = c_bind_p; + goto incompatible; + } + } + /* Matching nohost clauses? */ + if ((c_nohost == NULL_TREE) != (c_nohost_p == NULL_TREE)) + { + c_diag = c_nohost; + c_diag_p = c_nohost_p; + goto incompatible; + } + /* Compatible. */ + return 1; + + incompatible: + if (c_diag != NULL_TREE) + error_at (OMP_CLAUSE_LOCATION (c_diag), + "incompatible %qs clause when applying" + " %<%s%> to %qD, which has already been" + " marked as an accelerator routine", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)], + routine_str, fndecl); + else if (c_diag_p != NULL_TREE) + error_at (loc, + "missing %qs clause when applying" + " %<%s%> to %qD, which has already been" + " marked as an accelerator routine", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)], + routine_str, fndecl); + else + gcc_unreachable (); + if (c_diag_p != NULL_TREE) + inform (OMP_CLAUSE_LOCATION (c_diag_p), + "... with %qs clause here", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)]); + else + { + /* In the front ends, we don't preserve location information for the + OpenACC routine directive itself. However, that of c_level_p + should be close. */ + location_t loc_routine = OMP_CLAUSE_LOCATION (c_level_p); + inform (loc_routine, "... without %qs clause near to here", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]); + } + /* Incompatible. */ + return -1; + } + + return 0; +} + +/* Process the OpenACC routine's clauses to generate an attribute + for the level of parallelism. All dimensions have a size of zero (dynamic). TREE_PURPOSE is set to indicate whether that dimension can have a loop partitioned on it. non-zero indicates yes, zero indicates no. By construction once a non-zero has been @@ -19694,6 +19881,28 @@ default_goacc_reduction (gcall *call) gsi_replace_with_seq (&gsi, seq, true); } +/* Determine whether DECL should be discarded in this offload + compilation. */ + +static bool +maybe_discard_oacc_function (tree decl) +{ + tree attr = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)); + + if (!attr) + return false; + + enum omp_clause_code kind = OMP_CLAUSE_NOHOST; + +#ifdef ACCEL_COMPILER + kind = OMP_CLAUSE_BIND; +#endif + if (find_omp_clause (TREE_VALUE (attr), kind)) + return true; + + return false; +} + /* Main entry point for oacc transformations which run on the device compiler after LTO, so we know what the target device is at this point (including the host fallback). */ @@ -19707,6 +19916,14 @@ execute_oacc_device_lower () /* Not an offloaded function. */ return 0; + if (maybe_discard_oacc_function (current_function_decl)) + { + if (dump_file) + fprintf (dump_file, "Discarding function\n"); + TREE_ASM_WRITTEN (current_function_decl) = 1; + return TODO_discard_function; + } + /* Parse the default dim argument exactly once. */ if ((const void *)flag_openacc_dims != &flag_openacc_dims) { diff --git a/gcc/omp-low.h b/gcc/omp-low.h index b1f7885..2602a12 100644 --- a/gcc/omp-low.h +++ b/gcc/omp-low.h @@ -31,6 +31,7 @@ extern bool make_gimple_omp_edges (basic_block, struct omp_region **, int *); extern void omp_finish_file (void); extern tree omp_member_access_dummy_var (tree); extern void replace_oacc_fn_attrib (tree, tree); +extern int verify_oacc_routine_clauses (tree, tree *, location_t, const char *); extern tree build_oacc_routine_dims (tree); extern tree get_oacc_fn_attrib (tree); extern void set_oacc_fn_attrib (tree, tree, bool, vec *); diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 3e3f31e..5871849 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -465,7 +465,13 @@ enum omp_clause_code { /* OpenMP internal-only clause to specify grid dimensions of a gridified kernel. */ - OMP_CLAUSE__GRIDDIM_ + OMP_CLAUSE__GRIDDIM_, + + /* OpenACC clause: bind (string). */ + OMP_CLAUSE_BIND, + + /* OpenACC clause: nohost. */ + OMP_CLAUSE_NOHOST }; #undef DEFTREESTRUCT diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index ebbf606..00f4ba7 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -1021,6 +1021,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags) spc, flags, false); pp_right_paren (pp); break; + case OMP_CLAUSE_NOHOST: + pp_string (pp, "nohost"); + break; + case OMP_CLAUSE_BIND: + pp_string (pp, "bind("); + dump_generic_node (pp, OMP_CLAUSE_BIND_NAME (clause), + spc, flags, false); + pp_string (pp, ")"); + break; case OMP_CLAUSE__GRIDDIM_: pp_string (pp, "_griddim_("); diff --git a/gcc/tree.c b/gcc/tree.c index c155d56..f51df11 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -329,6 +329,8 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_VECTOR_LENGTH */ 1, /* OMP_CLAUSE_TILE */ 2, /* OMP_CLAUSE__GRIDDIM_ */ + 1, /* OMP_CLAUSE_BIND */ + 0, /* OMP_CLAUSE_NOHOST */ }; const char * const omp_clause_code_name[] = @@ -399,7 +401,9 @@ const char * const omp_clause_code_name[] = "num_workers", "vector_length", "tile", - "_griddim_" + "_griddim_", + "bind", + "nohost", }; @@ -11869,6 +11873,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE__LOOPTEMP_: case OMP_CLAUSE__SIMDUID_: case OMP_CLAUSE__CILK_FOR_COUNT_: + case OMP_CLAUSE_BIND: WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 0)); /* FALLTHRU */ @@ -11890,6 +11895,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_TILE: WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp)); diff --git a/gcc/tree.h b/gcc/tree.h index 6a98b6e..7757a9a 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1526,6 +1526,9 @@ extern void protected_set_expr_location (tree, location_t); #define OMP_CLAUSE_VECTOR_LENGTH_EXPR(NODE) \ OMP_CLAUSE_OPERAND ( \ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_VECTOR_LENGTH), 0) +#define OMP_CLAUSE_BIND_NAME(NODE) \ + OMP_CLAUSE_OPERAND ( \ + OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_BIND), 0) #define OMP_CLAUSE_DEPEND_KIND(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEPEND)->omp_clause.subcode.depend_kind)