From patchwork Tue Nov 22 20:00:47 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 83485 Delivered-To: patch@linaro.org Received: by 10.140.97.165 with SMTP id m34csp2288083qge; Tue, 22 Nov 2016 12:01:21 -0800 (PST) X-Received: by 10.84.210.46 with SMTP id z43mr3043497plh.115.1479844881448; Tue, 22 Nov 2016 12:01:21 -0800 (PST) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id s136si29892648pgc.91.2016.11.22.12.01.21 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Tue, 22 Nov 2016 12:01:21 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-return-442283-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-442283-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-442283-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 :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=DRsr7P64Z8xK5DORg WwG6L4jML0SlVFSusbyI/YFFSNCtE1W7ip3lsiQnbkMG0ETblgKbJeEp8WW2C03C ZvWG97ndM0mpjkLS7dN+rS+zgm5zmH4vAkgDUKUGjx7sb81/t/YzUuoHwCHFp+gZ dYAOnr6P7tf/K4I785Czn0PBcI= 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 :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=21S4IMGrJOOHVa3xalOVSAt G1a8=; b=juBa01+DEKTkn0fFZEBxwWhPBf/noGyncmru/gi4Xh601yiXLkSBHTj T/h3uj1u/8GvKooZrlFqwlnEpix1FHfoHFY2FU/3LwFEnqp0A/XrJHZTpyi7LCFy eZJKAQmeNxuaj/mgrYCUsADGXCFSEUehWMCQ5ZBpXxVGGtl5IeEQ= Received: (qmail 118533 invoked by alias); 22 Nov 2016 20:01:04 -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 118276 invoked by uid 89); 22 Nov 2016 20:01:02 -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=20a, offload, Discarding, discarding 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; Tue, 22 Nov 2016 20:00:52 +0000 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtp id 1c9HFK-0003YT-QY from Cesar_Philippidis@mentor.com ; Tue, 22 Nov 2016 12:00:50 -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; Tue, 22 Nov 2016 12:00:48 -0800 Subject: Re: [PATCH] OpenACC routines -- middle end To: Jakub Jelinek References: <5aa1ff0a-f5a6-ef64-aaf1-0666eed4fd42@codesourcery.com> <20161118121418.GX3541@tucnak.redhat.com> <5a660e56-dfd1-cb7a-644d-9b8e688bdae8@codesourcery.com> <20161122195816.GD3541@tucnak.redhat.com> CC: "gcc-patches@gcc.gnu.org" , Thomas Schwinge From: Cesar Philippidis Message-ID: Date: Tue, 22 Nov 2016 12:00:47 -0800 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.3.0 MIME-Version: 1.0 In-Reply-To: <20161122195816.GD3541@tucnak.redhat.com> X-ClientProxiedBy: svr-orw-mbx-03.mgc.mentorg.com (147.34.90.203) To svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) On 11/22/2016 11:58 AM, Jakub Jelinek wrote: > On Tue, Nov 22, 2016 at 11:53:50AM -0800, Cesar Philippidis wrote: >> I've incorporated those changes in this patch. Is it ok for trunk? > > The ChangeLog mentions omp-low.[ch] changes, but the patch doesn't include > them. > Have they been dropped, or moved to another patch? No, sorry I forgot to include them in the diff. This patch should contain all of the middle end changes. Cesar >> 2016-11-22 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. > > Jakub > 2016-11-22 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 964efe9..49999b8 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 8611060..04b591e 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 7c58c03..b8a414b 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), + "previous %qs clause", + 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), + "previous %qs clause", + 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, "previous %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 a3d220d..bd0e254 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 096eefd..5494441 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -1022,6 +1022,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 a4c5b1b..521fb22 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", }; @@ -11871,6 +11875,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 */ @@ -11892,6 +11897,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 b4ec3fd..a17606f 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)