[OpenACC] C, C++: bind and nohost clauses

Message ID	87fv097zaj.fsf@hertz.schwinge.homeip.net
State	New
Headers	show Return-Path: <gcc-patches-return-414091-incoming=patchwork.ozlabs.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 :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= Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk Sender: gcc-patches-owner@gcc.gnu.org From: Thomas Schwinge <thomas@codesourcery.com> To: Jakub Jelinek <jakub@redhat.com>, <gcc-patches@gcc.gnu.org> CC: Cesar Philippidis <cesar@codesourcery.com>, Joseph Myers <joseph@codesourcery.com> 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 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable

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)

[OpenACC] C, C++: bind and nohost clauses

Commit Message

Patch