From patchwork Sat Nov 14 08:36:36 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 544742 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 0AD88141173 for ; Sat, 14 Nov 2015 19:37:12 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=dHbkda0D; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:mime-version:content-type :content-transfer-encoding; q=dns; s=default; b=Z60dlNKCtiokwxpx 10BKPeuAN+nxWvIF1K+vqAIbMqTA2MKSfdsSCRzU5Eyw8U7Gj9vBlLJrCziKLnyb RZIXEo1BmX/98puuWM+w58WqTuzVwr3UXKShCu/KTuCK24Pxs3mTFA3BlD3gKMl1 iaA+iEm+e4OM3os5r+shXz5RUTc= 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 :to:cc:subject:date:message-id:mime-version:content-type :content-transfer-encoding; s=default; bh=rWz5dDw+qyo2It9dRy4K3A fdQkQ=; b=dHbkda0DY5KPuxSabVQXWCmrMx0yY7Fxuqkv+Xa/WYdB5FN9taHBez YRZBE5Ahibl0XHV4GYIOJOGYqHDkN7OeKwAsJ/td2P7XucKuLFLeEBdc4cqtqfPx iBxM1OdLBt0w9K7qiSH+hW1Zy1GO/C8UogcmAgs4yIFDjvkqM9Lwo= Received: (qmail 40475 invoked by alias); 14 Nov 2015 08:37:05 -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 40465 invoked by uid 89); 14 Nov 2015 08:37:04 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.7 required=5.0 tests=AWL, BAYES_50, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=ham version=3.3.2 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; Sat, 14 Nov 2015 08:37:00 +0000 Received: from svr-orw-fem-05.mgc.mentorg.com ([147.34.97.43]) by relay1.mentorg.com with esmtp id 1ZxWKP-0005Ul-47 from Thomas_Schwinge@mentor.com ; Sat, 14 Nov 2015 00:36:57 -0800 Received: from tftp-cs (147.34.91.1) by svr-orw-fem-05.mgc.mentorg.com (147.34.97.43) with Microsoft SMTP Server id 14.3.224.2; Sat, 14 Nov 2015 00:36:56 -0800 Received: by tftp-cs (Postfix, from userid 49978) id 34263C23D2; Sat, 14 Nov 2015 00:36:56 -0800 (PST) From: Thomas Schwinge To: Jakub Jelinek , CC: Cesar Philippidis , Joseph Myers Subject: [OpenACC] C, C++: bind and nohost clauses User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (x86_64-pc-linux-gnu) Date: Sat, 14 Nov 2015 09:36:36 +0100 Message-ID: <87fv097zaj.fsf@hertz.schwinge.homeip.net> MIME-Version: 1.0 Hi! Initial support for the OpenACC bind and nohost clauses (routine directive) for C, C++. Fortran to follow. Middle end handling and more complete testsuite coverage also to follow once we got a few details clarified. OK for trunk? commit 96ca3e7fd643a67282a05750799bb78d4c634a71 Author: Thomas Schwinge Date: Fri Nov 13 14:36:57 2015 +0100 [OpenACC] C, C++: bind and nohost clauses YYYY-MM-DD Thomas Schwinge Joseph Myers gcc/ * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_BIND and OMP_CLAUSE_NOHOST. * tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries for these. (walk_tree_1): Handle these. * gimplify.c (gimplify_scan_omp_clauses) (gimplify_adjust_omp_clauses): Likewise. * omp-low.c (scan_sharing_clauses): Likewise. * tree-pretty-print.c (dump_omp_clause): Likewise. * tree.h (OMP_CLAUSE_BIND_NAME): New macro. YYYY-MM-DD Thomas Schwinge Joseph Myers gcc/c-family/ * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_BIND and PRAGMA_OACC_CLAUSE_NOHOST. YYYY-MM-DD Thomas Schwinge Joseph Myers gcc/c/ * c-parser.c (c_parser_omp_clause_name): Handle "bind" and "nohost". (c_parser_oacc_clause_bind): New function. (c_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_BIND and PRAGMA_OACC_CLAUSE_NOHOST. (OACC_ROUTINE_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_BIND and PRAGMA_OACC_CLAUSE_NOHOST. * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_BIND and OMP_CLAUSE_NOHOST. YYYY-MM-DD Thomas Schwinge Cesar Philippidis gcc/cp/ * parser.c (cp_parser_omp_clause_name): Handle "bind" and "nohost". (cp_parser_oacc_clause_bind): New function. (cp_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_BIND and PRAGMA_OACC_CLAUSE_NOHOST. (OACC_ROUTINE_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_BIND and PRAGMA_OACC_CLAUSE_NOHOST. * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_BIND and OMP_CLAUSE_NOHOST. * semantics.c (finish_omp_clauses): Likewise. YYYY-MM-DD Thomas Schwinge Joseph Myers gcc/testsuite/ * c-c++-common/goacc/routine-1.c: Extend. --- gcc/c-family/c-pragma.h | 2 + gcc/c/c-parser.c | 55 ++++++++++++++++++++-- gcc/c/c-typeck.c | 2 + gcc/cp/parser.c | 70 ++++++++++++++++++++++++++-- gcc/cp/pt.c | 2 + gcc/cp/semantics.c | 2 + gcc/gimplify.c | 4 ++ gcc/omp-low.c | 4 ++ gcc/testsuite/c-c++-common/goacc/routine-1.c | 40 +++++++++++++++- gcc/tree-core.h | 6 +++ gcc/tree-pretty-print.c | 9 ++++ gcc/tree.c | 6 +++ gcc/tree.h | 3 ++ 13 files changed, 198 insertions(+), 7 deletions(-) Grüße Thomas diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h index afeceff..ccb8ddf 100644 --- gcc/c-family/c-pragma.h +++ gcc/c-family/c-pragma.h @@ -147,6 +147,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, @@ -155,6 +156,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 gcc/c/c-parser.c gcc/c/c-parser.c index 2484b92..384b395 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -9978,6 +9978,10 @@ c_parser_omp_clause_name (c_parser *parser) else if (!strcmp ("async", p)) result = PRAGMA_OACC_CLAUSE_ASYNC; break; + case 'b': + if (!strcmp ("bind", p)) + result = PRAGMA_OACC_CLAUSE_BIND; + break; case 'c': if (!strcmp ("collapse", p)) result = PRAGMA_OMP_CLAUSE_COLLAPSE; @@ -10053,6 +10057,10 @@ c_parser_omp_clause_name (c_parser *parser) case 'n': if (!strcmp ("nogroup", p)) result = PRAGMA_OMP_CLAUSE_NOGROUP; + else if (!strcmp ("nohost", p)) + result = PRAGMA_OACC_CLAUSE_NOHOST; + else if (flag_cilkplus && !strcmp ("nomask", p)) + result = PRAGMA_CILK_CLAUSE_NOMASK; else if (!strcmp ("notinbranch", p)) result = PRAGMA_OMP_CLAUSE_NOTINBRANCH; else if (!strcmp ("nowait", p)) @@ -10067,8 +10075,6 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_NUM_THREADS; else if (!strcmp ("num_workers", p)) result = PRAGMA_OACC_CLAUSE_NUM_WORKERS; - else if (flag_cilkplus && !strcmp ("nomask", p)) - result = PRAGMA_CILK_CLAUSE_NOMASK; break; case 'o': if (!strcmp ("ordered", p)) @@ -11413,6 +11419,38 @@ c_parser_oacc_clause_async (c_parser *parser, tree list) } /* OpenACC 2.0: + bind ( identifier ) + bind ( string-literal ) */ + +static tree +c_parser_oacc_clause_bind (c_parser *parser, tree list) +{ + location_t loc = c_parser_peek_token (parser)->location; + + parser->lex_untranslated_string = true; + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + { + parser->lex_untranslated_string = false; + return list; + } + if (c_parser_next_token_is (parser, CPP_NAME) + || c_parser_next_token_is (parser, CPP_STRING)) + { + tree t = c_parser_peek_token (parser)->value; + c_parser_consume_token (parser); + tree c = build_omp_clause (loc, OMP_CLAUSE_BIND); + OMP_CLAUSE_BIND_NAME (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + } + else + c_parser_error (parser, "expected identifier or character string literal"); + parser->lex_untranslated_string = false; + c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"); + return list; +} + +/* OpenACC 2.0: tile ( size-expr-list ) */ static tree @@ -12688,6 +12726,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses); c_name = "auto"; break; + case PRAGMA_OACC_CLAUSE_BIND: + clauses = c_parser_oacc_clause_bind (parser, clauses); + c_name = "bind"; + break; case PRAGMA_OACC_CLAUSE_COLLAPSE: clauses = c_parser_omp_clause_collapse (parser, clauses); c_name = "collapse"; @@ -12746,6 +12788,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses); c_name = "independent"; break; + case PRAGMA_OACC_CLAUSE_NOHOST: + clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_NOHOST, + clauses); + c_name = "nohost"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: clauses = c_parser_omp_clause_num_gangs (parser, clauses); c_name = "num_gangs"; @@ -13430,7 +13477,9 @@ c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser, ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) ) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_BIND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST) ) /* Parse an OpenACC routine directive. For named directives, we apply immediately to the named function. For unnamed ones we then parse diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index 4335a87..033aa9b 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -13043,6 +13043,8 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) case OMP_CLAUSE_GANG: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_TILE: pc = &OMP_CLAUSE_CHAIN (c); continue; diff --git gcc/cp/parser.c gcc/cp/parser.c index a87675e..db8fd60 100644 --- gcc/cp/parser.c +++ gcc/cp/parser.c @@ -29105,6 +29105,10 @@ cp_parser_omp_clause_name (cp_parser *parser) else if (!strcmp ("async", p)) result = PRAGMA_OACC_CLAUSE_ASYNC; break; + case 'b': + if (!strcmp ("bind", p)) + result = PRAGMA_OACC_CLAUSE_BIND; + break; case 'c': if (!strcmp ("collapse", p)) result = PRAGMA_OMP_CLAUSE_COLLAPSE; @@ -29178,12 +29182,14 @@ cp_parser_omp_clause_name (cp_parser *parser) case 'n': if (!strcmp ("nogroup", p)) result = PRAGMA_OMP_CLAUSE_NOGROUP; + else if (!strcmp ("nohost", p)) + result = PRAGMA_OACC_CLAUSE_NOHOST; + else if (flag_cilkplus && !strcmp ("nomask", p)) + result = PRAGMA_CILK_CLAUSE_NOMASK; else if (!strcmp ("notinbranch", p)) result = PRAGMA_OMP_CLAUSE_NOTINBRANCH; else if (!strcmp ("nowait", p)) result = PRAGMA_OMP_CLAUSE_NOWAIT; - else if (flag_cilkplus && !strcmp ("nomask", p)) - result = PRAGMA_CILK_CLAUSE_NOMASK; else if (!strcmp ("num_gangs", p)) result = PRAGMA_OACC_CLAUSE_NUM_GANGS; else if (!strcmp ("num_tasks", p)) @@ -31473,6 +31479,53 @@ cp_parser_oacc_clause_async (cp_parser *parser, tree list) return list; } +/* OpenACC 2.0: + bind ( identifier ) + bind ( string-literal ) */ + +static tree +cp_parser_oacc_clause_bind (cp_parser *parser, tree list) +{ + location_t loc = cp_lexer_peek_token (parser->lexer)->location; + bool save_translate_strings_p = parser->translate_strings_p; + + parser->translate_strings_p = false; + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + { + parser->translate_strings_p = save_translate_strings_p; + return list; + } + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME) + || cp_lexer_next_token_is (parser->lexer, CPP_STRING)) + { + tree t; + + if (cp_lexer_peek_token (parser->lexer)->type == CPP_STRING) + { + t = cp_lexer_peek_token (parser->lexer)->u.value; + cp_lexer_consume_token (parser->lexer); + } + else + t = cp_parser_id_expression (parser, /*template_p=*/false, + /*check_dependency_p=*/true, + /*template_p=*/NULL, + /*declarator_p=*/false, + /*optional_p=*/false); + if (t == error_mark_node) + return t; + + tree c = build_omp_clause (loc, OMP_CLAUSE_BIND); + OMP_CLAUSE_BIND_NAME (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + } + else + cp_parser_error (parser, "expected identifier or character string literal"); + parser->translate_strings_p = save_translate_strings_p; + cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN); + return list; +} + /* Parse all OpenACC clauses. The set clauses allowed by the directive is a bitmask in MASK. Return the list of clauses found. */ @@ -31509,6 +31562,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses, here); c_name = "auto"; break; + case PRAGMA_OACC_CLAUSE_BIND: + clauses = cp_parser_oacc_clause_bind (parser, clauses); + c_name = "bind"; + break; case PRAGMA_OACC_CLAUSE_COLLAPSE: clauses = cp_parser_omp_clause_collapse (parser, clauses, here); c_name = "collapse"; @@ -31569,6 +31626,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses, here); c_name = "independent"; break; + case PRAGMA_OACC_CLAUSE_NOHOST: + clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_NOHOST, + clauses, here); + c_name = "nohost"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: code = OMP_CLAUSE_NUM_GANGS; c_name = "num_gangs"; @@ -35634,7 +35696,9 @@ cp_parser_omp_taskloop (cp_parser *parser, cp_token *pragma_tok, ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_BIND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST)) /* Finalize #pragma acc routine clauses after direct declarator has been parsed, and put that into "omp declare target" attribute. */ diff --git gcc/cp/pt.c gcc/cp/pt.c index 62659ec..4264d33 100644 --- gcc/cp/pt.c +++ gcc/cp/pt.c @@ -14388,6 +14388,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, case OMP_CLAUSE_VECTOR: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: + case OMP_CLAUSE_BIND: OMP_CLAUSE_OPERAND (nc, 0) = tsubst_expr (OMP_CLAUSE_OPERAND (oc, 0), args, complain, in_decl, /*integral_constant_expression_p=*/false); @@ -14458,6 +14459,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_NOHOST: break; case OMP_CLAUSE_TILE: { diff --git gcc/cp/semantics.c gcc/cp/semantics.c index db37e85..fc94b00 100644 --- gcc/cp/semantics.c +++ gcc/cp/semantics.c @@ -6850,6 +6850,8 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) case OMP_CLAUSE_AUTO: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: break; case OMP_CLAUSE_TILE: diff --git gcc/gimplify.c gcc/gimplify.c index 66e5168..e8bb750 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -7178,6 +7178,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_VECTOR_LENGTH: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_BIND: if (gimplify_expr (&OMP_CLAUSE_OPERAND (c, 0), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) remove = true; @@ -7221,6 +7222,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_NOGROUP: case OMP_CLAUSE_THREADS: case OMP_CLAUSE_SIMD: + case OMP_CLAUSE_NOHOST: break; case OMP_CLAUSE_DEFAULTMAP: @@ -7724,6 +7726,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p, case OMP_CLAUSE_VECTOR: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_TILE: break; diff --git gcc/omp-low.c gcc/omp-low.c index 51b471c..8f59fe3 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -2123,6 +2123,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_GANG: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_TILE: break; @@ -2298,6 +2300,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_GANG: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_TILE: break; diff --git gcc/testsuite/c-c++-common/goacc/routine-1.c gcc/testsuite/c-c++-common/goacc/routine-1.c index a5e0d69..4cb3cff 100644 --- gcc/testsuite/c-c++-common/goacc/routine-1.c +++ gcc/testsuite/c-c++-common/goacc/routine-1.c @@ -1,3 +1,4 @@ +/* Test valid use of clauses with routine. */ #pragma acc routine gang void gang (void) @@ -19,15 +20,52 @@ void seq (void) { } +#pragma acc routine bind (bind_f_1) +void bind_f_1 (void) +{ +} + +#pragma acc routine bind (bind_f_1) +void bind_f_1_1 (void) +{ +} + +#pragma acc routine bind ("bind_f_2") +void bind_f_2 (void) +{ +} + +#pragma acc routine bind ("bind_f_2") +void bind_f_2_1 (void) +{ +} + +typedef int T; + +#pragma acc routine bind (T) +void T_1 (void) +{ +} + +#pragma acc routine nohost +void nohost (void) +{ +} + int main () { - #pragma acc parallel num_gangs (32) num_workers (32) vector_length (32) { gang (); worker (); vector (); seq (); + bind_f_1 (); + bind_f_1_1 (); + bind_f_2 (); + bind_f_2_1 (); + T_1 (); + nohost (); } return 0; diff --git gcc/tree-core.h gcc/tree-core.h index ff061ef..8ddb9a6 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -430,6 +430,12 @@ enum omp_clause_code { /* OpenACC clause: vector_length (integer-expression). */ OMP_CLAUSE_VECTOR_LENGTH, + /* OpenACC clause: bind ( identifer | string ). */ + OMP_CLAUSE_BIND, + + /* OpenACC clause: nohost. */ + OMP_CLAUSE_NOHOST, + /* OpenACC clause: tile ( size-expr-list ). */ OMP_CLAUSE_TILE }; diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c index 3f0a4e6..3cd389e 100644 --- gcc/tree-pretty-print.c +++ gcc/tree-pretty-print.c @@ -932,6 +932,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags) case OMP_CLAUSE_INDEPENDENT: pp_string (pp, "independent"); 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_NOHOST: + pp_string (pp, "nohost"); + break; case OMP_CLAUSE_TILE: pp_string (pp, "tile("); dump_generic_node (pp, OMP_CLAUSE_TILE_LIST (clause), diff --git gcc/tree.c gcc/tree.c index 50e1db0..38da31e 100644 --- gcc/tree.c +++ gcc/tree.c @@ -328,6 +328,8 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_NUM_GANGS */ 1, /* OMP_CLAUSE_NUM_WORKERS */ 1, /* OMP_CLAUSE_VECTOR_LENGTH */ + 1, /* OMP_CLAUSE_BIND */ + 0, /* OMP_CLAUSE_NOHOST */ 1, /* OMP_CLAUSE_TILE */ }; @@ -400,6 +402,8 @@ const char * const omp_clause_code_name[] = "num_gangs", "num_workers", "vector_length", + "bind", + "nohost", "tile" }; @@ -11577,6 +11581,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 */ @@ -11598,6 +11603,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 gcc/tree.h gcc/tree.h index 1bb59f2..d6f29b4 100644 --- gcc/tree.h +++ gcc/tree.h @@ -1454,6 +1454,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)