From patchwork Fri May 24 13:07:56 2013 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Jakub Jelinek X-Patchwork-Id: 246139 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 with cipher DHE-RSA-AES256-SHA (256/256 bits)) (Client CN "localhost", Issuer "www.qmailtoaster.com" (not verified)) by ozlabs.org (Postfix) with ESMTPS id 325BE2C00AC for ; Fri, 24 May 2013 23:08:29 +1000 (EST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:reply-to:mime-version :content-type; q=dns; s=default; b=Y27jHlGCo5jUP0kMi1fvGnORuJKQA TEUWcgO2+H6Rrztv8N+udTrD4tqQpxZwtJhnEPTVp+HCjpjb54kCQPMeJiN8RWm3 TuS/a/yCQVQLEes+ZTs1BEVyH3wI/wxVO73CCBEPAKnhytNUyAZ3nrL1dmXSipBy RLv8JkzNkeEcZc= 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:date :from:to:cc:subject:message-id:reply-to:mime-version :content-type; s=default; bh=Z2UvUFX5UiCJmoA1H91lHtEIq8k=; b=a8/ 1zTWy5jpPjyBi47RPVCpvylW+x4ojZR23ojUuw5Zk0bcwragoDBVGn4Rem74pWU6 G5Rg3K6Ew476ZBA35Bvu//nr6d6B5oObFJCPWXJW44Azl2+kd4OcOYVzILIqDERW V7bSyZV5xzpyaBSu9GQwVZ3d4qiSRolRGe16QR3c= Received: (qmail 18212 invoked by alias); 24 May 2013 13:08: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 18202 invoked by uid 89); 24 May 2013 13:08:20 -0000 X-Spam-SWARE-Status: No, score=-6.8 required=5.0 tests=AWL, BAYES_00, RCVD_IN_HOSTKARMA_W, RCVD_IN_HOSTKARMA_WL, RP_MATCHES_RCVD, SPF_HELO_PASS, SPF_PASS, TW_TM autolearn=ham version=3.3.1 Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.84/v0.84-167-ge50287c) with ESMTP; Fri, 24 May 2013 13:08:04 +0000 Received: from int-mx12.intmail.prod.int.phx2.redhat.com (int-mx12.intmail.prod.int.phx2.redhat.com [10.5.11.25]) by mx1.redhat.com (8.14.4/8.14.4) with ESMTP id r4OD83Xm007747 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-SHA bits=256 verify=OK) for ; Fri, 24 May 2013 09:08:03 -0400 Received: from zalov.cz (vpn-49-71.rdu2.redhat.com [10.10.49.71]) by int-mx12.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id r4OD80Ej028366 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-SHA bits=256 verify=NO); Fri, 24 May 2013 09:08:02 -0400 Received: from zalov.cz (localhost [127.0.0.1]) by zalov.cz (8.14.5/8.14.5) with ESMTP id r4OD7x8O030076; Fri, 24 May 2013 15:08:00 +0200 Received: (from jakub@localhost) by zalov.cz (8.14.5/8.14.5/Submit) id r4OD7wpr030075; Fri, 24 May 2013 15:07:58 +0200 Date: Fri, 24 May 2013 15:07:56 +0200 From: Jakub Jelinek To: Richard Henderson , Jason Merrill , Torvald Riegel Cc: gcc-patches@gcc.gnu.org Subject: [gomp4] Initial support for accelerator support parsing Message-ID: <20130524130756.GR1377@tucnak.redhat.com> Reply-To: Jakub Jelinek MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.5.21 (2010-09-15) Hi! This patch contains C++ parser changes etc. to handle #pragma omp {teams,target {,data,update},distribute} parsing all the way through till omp lowering (it bombs badly in omp expansion, but already omp lowering will need to be tought out). Things not handled yet are #pragma omp declare target/#pragma omp end declare target (I assume we can handle it by automatically adding "omp target" attribute to vars/functions in there) and there is no support for array sections yet (also relevant for OMP_CLAUSE_DEPEND). Say: void baz (float *, float *, int); float foo (int x) { float b[1024], c[1024], s = 0; int i; baz (b, c, x); #pragma omp target map(to: b, c) #pragma omp parallel for reduction(+:s) for (i = 0; i < x; i++) s += b[i] * c[i]; return s; } float bar (int x, int y, int z) { float b[1024], c[1024], s = 0; int i, j; baz (b, c, x); #pragma omp target data map(to: b) { #pragma omp target map(to: c) #pragma omp teams num_teams(y) num_threads(z) reduction(+:s) #pragma omp distribute dist_schedule(static, 4) collapse(1) for (j=0; j < x; j += y) #pragma omp parallel for reduction(+:s) for (i = j; i < j + y; i++) s += b[i] * c[i]; #pragma omp target update from(b, c) } return s; } now parses with g++ -S -fopenmp -fdump-tree-all and shows dump up to *.cfg (crash during ompexp). Comments on this before I commit it to gomp-4_0-branch? 2013-05-24 Jakub Jelinek * tree.def (OMP_TEAMS, OMP_TARGET_DATA, OMP_TARGET, OMP_TARGET_UPDATE): New tree codes. * tree-cfg.c (make_edges): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS. * omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_DIST_SCHEDULE. * gimple-low.c (lower_stmt): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS. * tree.h (OMP_TEAMS_BODY, OMP_TEAMS_CLAUSES, OMP_TARGET_DATA_BODY, OMP_TARGET_DATA_CLAUSES, OMP_TARGET_BODY, OMP_TARGET_CLAUSES, OMP_TARGET_UPDATE_CLAUSES): Define. * tree-nested.c (convert_nonlocal_reference_stmt, convert_local_reference_stmt, convert_gimple_call): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS. * tree-inline.c (estimate_num_insns): Likewise. (remap_gimple_stmt): Likewise. Adjust gimple_build_omp_for caller. * gimple.def: Adjust comments describing OMP_CLAUSEs. (GIMPLE_OMP_TARGET, GIMPLE_OMP_TEAMS): New GIMPLE stmts. * tree-parloops.c (create_parallel_loop): Adjust gimple_build_omp_for caller. * tree-pretty-print.c (dump_generic_node): Handle OMP_TEAMS, OMP_TARGET, OMP_TARGET_DATA and OMP_TARGET_UPDATE. * gimple.h (GF_OMP_TARGET_KIND_MASK, GF_OMP_TARGET_KIND_REGION, GF_OMP_TARGET_KIND_DATA, GF_OMP_TARGET_KIND_UPDATE): New. (gimple_build_omp_for): Add kind argument to prototype. (gimple_build_omp_target, gimple_build_omp_teams): New prototypes. (gimple_has_substatements): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS. (gimple_omp_subcode): Change GIMPLE_OMP_SINGLE to GIMPLE_OMP_TEAMS. (gimple_omp_target_clauses, gimple_omp_target_clauses_ptr, gimple_omp_target_set_clauses, gimple_omp_target_kind, gimple_omp_target_set_kind, gimple_omp_teams_clauses, gimple_omp_teams_clauses_ptr, gimple_omp_teams_set_clauses): New inline functions. (gimple_return_set_retval): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS. * gimple.c (gimple_build_omp_for): Add kind argument, call gimple_omp_for_set_kind. (gimple_build_omp_target, gimple_build_omp_teams): New functions. (walk_gimple_op, walk_gimple_stmt, gimple_copy): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS. * gimple-pretty-print.c (dump_gimple_omp_target, dump_gimple_omp_teams): New functions. (pp_gimple_stmt_1): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS. * gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP. (enum omp_region_type): Add ORT_TEAMS, ORT_TARGET and ORT_TARGET_DATA. (omp_add_variable): Add temporary assertions. (omp_notice_threadprivate_variable): Complain if threadprivate vars appear in target region. (omp_notice_variable): ORT_TARGET, ORT_TARGET_DATA and ORT_TEAMS handling. (omp_check_private): Ignore ORT_TARGET and ORT_TARGET_DATA regions. (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_MAP, OMP_CLAUSE_TO, OMP_CLAUSE_FROM, OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_DIST_SCHEDULE and OMP_CLAUSE_DEVICE. (gimplify_adjust_omp_clauses): Likewise. (gimplify_adjust_omp_clauses_1): Handle GOVD_MAP. Fix up check for privatization by also testing for GOVD_LINEAR. (gimplify_omp_for): Adjust gimple_build_omp_for caller. Clear *expr_p. (gimplify_omp_workshare): Handle also OMP_TARGET, OMP_TARGET_DATA and OMP_TEAMS. Clear *expr_p. (gimplify_omp_target_update): New function. (gimplify_expr): Handle OMP_TARGET, OMP_TARGET_DATA, OMP_TARGET_UPDATE and OMP_TEAMS. cp/ * parser.c (cp_parser_omp_clause_cancelkind): Remove diagnostics. (cp_parser_omp_all_clauses): Require that OMP_CLAUSE_{TO,FROM} and OMP_CLAUSE_{PARALLEL,FOR,SECTIONS,TASKGROUP} must be first in the list of clauses. (OMP_TEAMS_CLAUSE_MASK, OMP_TARGET_CLAUSE_MASK, OMP_TARGET_DATA_CLAUSE_MASK, OMP_TARGET_UPDATE_CLAUSE_MASK, OMP_DISTRIBUTE_CLAUSE_MASK): Define. (cp_parser_omp_teams, cp_parser_omp_target, cp_parser_omp_target_data, cp_parser_omp_target_update, cp_parser_omp_distribute): New functions. (cp_parser_omp_construct): Handle PRAGMA_OMP_DISTRIBUTE and PRAGMA_OMP_TEAMS. (cp_parser_pragma): Handle PRAGMA_OMP_DISTRIBUTE, PRAGMA_OMP_TEAMS and PRAGMA_OMP_TARGET. * pt.c (tsubst_expr): Handle OMP_TEAMS, OMP_TARGET, OMP_TARGET_DATA and OMP_TARGET_UPDATE. Jakub --- gcc/tree.def.jj 2013-04-30 18:03:33.000000000 +0200 +++ gcc/tree.def 2013-05-22 15:22:45.143759788 +0200 @@ -1042,6 +1042,21 @@ DEFTREECODE (OMP_FOR_SIMD, "omp_for_simd Operands like for OMP_FOR. */ DEFTREECODE (OMP_DISTRIBUTE, "omp_distribute", tcc_statement, 6) +/* OpenMP - #pragma omp teams [clause1 ... clauseN] + Operand 0: OMP_TEAMS_BODY: Teams body. + Operand 1: OMP_TEAMS_CLAUSES: List of clauses. */ +DEFTREECODE (OMP_TEAMS, "omp_teams", tcc_statement, 2) + +/* OpenMP - #pragma omp target data [clause1 ... clauseN] + Operand 0: OMP_TARGET_DATA_BODY: Target data construct body. + Operand 1: OMP_TARGET_DATA_CLAUSES: List of clauses. */ +DEFTREECODE (OMP_TARGET_DATA, "omp_target_data", tcc_statement, 2) + +/* OpenMP - #pragma omp target [clause1 ... clauseN] + Operand 0: OMP_TARGET_BODY: Target construct body. + Operand 1: OMP_TARGET_CLAUSES: List of clauses. */ +DEFTREECODE (OMP_TARGET, "omp_target", tcc_statement, 2) + /* OpenMP - #pragma omp sections [clause1 ... clauseN] Operand 0: OMP_SECTIONS_BODY: Sections body. Operand 1: OMP_SECTIONS_CLAUSES: List of clauses. */ @@ -1069,6 +1084,10 @@ DEFTREECODE (OMP_ORDERED, "omp_ordered", Operand 1: OMP_CRITICAL_NAME: Identifier for critical section. */ DEFTREECODE (OMP_CRITICAL, "omp_critical", tcc_statement, 2) +/* OpenMP - #pragma omp target update [clause1 ... clauseN] + Operand 0: OMP_TARGET_UPDATE_CLAUSES: List of clauses. */ +DEFTREECODE (OMP_TARGET_UPDATE, "omp_target_update", tcc_statement, 1) + /* OMP_ATOMIC through OMP_ATOMIC_CAPTURE_NEW must be consecutive, or OMP_ATOMIC_SEQ_CST needs adjusting. */ --- gcc/tree-cfg.c.jj 2013-05-20 13:21:43.000000000 +0200 +++ gcc/tree-cfg.c 2013-05-24 13:36:50.784061934 +0200 @@ -592,6 +592,8 @@ make_edges (void) case GIMPLE_OMP_TASK: case GIMPLE_OMP_FOR: case GIMPLE_OMP_SINGLE: + case GIMPLE_OMP_TARGET: + case GIMPLE_OMP_TEAMS: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_ORDERED: case GIMPLE_OMP_CRITICAL: --- gcc/omp-low.c.jj 2013-05-20 15:07:59.000000000 +0200 +++ gcc/omp-low.c 2013-05-24 13:47:32.465024619 +0200 @@ -1483,6 +1483,7 @@ scan_sharing_clauses (tree clauses, omp_ case OMP_CLAUSE_IF: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_SCHEDULE: + case OMP_CLAUSE_DIST_SCHEDULE: if (ctx->outer) scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer); break; @@ -1548,6 +1549,7 @@ scan_sharing_clauses (tree clauses, omp_ case OMP_CLAUSE_IF: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_SCHEDULE: + case OMP_CLAUSE_DIST_SCHEDULE: case OMP_CLAUSE_NOWAIT: case OMP_CLAUSE_ORDERED: case OMP_CLAUSE_COLLAPSE: --- gcc/gimple-low.c.jj 2013-03-20 10:07:24.000000000 +0100 +++ gcc/gimple-low.c 2013-05-24 13:40:19.382080371 +0200 @@ -444,6 +444,8 @@ lower_stmt (gimple_stmt_iterator *gsi, s case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: + case GIMPLE_OMP_TARGET: + case GIMPLE_OMP_TEAMS: data->cannot_fallthru = false; lower_omp_directive (gsi, data); data->cannot_fallthru = false; --- gcc/tree.h.jj 2013-05-20 13:18:33.000000000 +0200 +++ gcc/tree.h 2013-05-22 15:07:13.696048592 +0200 @@ -1863,6 +1863,20 @@ extern void protected_set_expr_location #define OMP_CRITICAL_BODY(NODE) TREE_OPERAND (OMP_CRITICAL_CHECK (NODE), 0) #define OMP_CRITICAL_NAME(NODE) TREE_OPERAND (OMP_CRITICAL_CHECK (NODE), 1) +#define OMP_TEAMS_BODY(NODE) TREE_OPERAND (OMP_TEAMS_CHECK (NODE), 0) +#define OMP_TEAMS_CLAUSES(NODE) TREE_OPERAND (OMP_TEAMS_CHECK (NODE), 1) + +#define OMP_TARGET_DATA_BODY(NODE) \ + TREE_OPERAND (OMP_TARGET_DATA_CHECK (NODE), 0) +#define OMP_TARGET_DATA_CLAUSES(NODE)\ + TREE_OPERAND (OMP_TARGET_DATA_CHECK (NODE), 1) + +#define OMP_TARGET_BODY(NODE) TREE_OPERAND (OMP_TARGET_CHECK (NODE), 0) +#define OMP_TARGET_CLAUSES(NODE) TREE_OPERAND (OMP_TARGET_CHECK (NODE), 1) + +#define OMP_TARGET_UPDATE_CLAUSES(NODE)\ + TREE_OPERAND (OMP_TARGET_UPDATE_CHECK (NODE), 0) + #define OMP_CLAUSE_CHAIN(NODE) TREE_CHAIN (OMP_CLAUSE_CHECK (NODE)) #define OMP_CLAUSE_DECL(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \ --- gcc/tree-nested.c.jj 2013-03-20 10:08:27.000000000 +0100 +++ gcc/tree-nested.c 2013-05-23 13:23:48.785572843 +0200 @@ -1291,6 +1291,22 @@ convert_nonlocal_reference_stmt (gimple_ info->suppress_expansion = save_suppress; break; + case GIMPLE_OMP_TARGET: + save_suppress = info->suppress_expansion; + convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi); + walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op, + info, gimple_omp_body_ptr (stmt)); + info->suppress_expansion = save_suppress; + break; + + case GIMPLE_OMP_TEAMS: + save_suppress = info->suppress_expansion; + convert_nonlocal_omp_clauses (gimple_omp_teams_clauses_ptr (stmt), wi); + walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op, + info, gimple_omp_body_ptr (stmt)); + info->suppress_expansion = save_suppress; + break; + case GIMPLE_OMP_SECTION: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_ORDERED: @@ -1714,6 +1730,22 @@ convert_local_reference_stmt (gimple_stm info->suppress_expansion = save_suppress; break; + case GIMPLE_OMP_TARGET: + save_suppress = info->suppress_expansion; + convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi); + walk_body (convert_local_reference_stmt, convert_local_reference_op, + info, gimple_omp_body_ptr (stmt)); + info->suppress_expansion = save_suppress; + break; + + case GIMPLE_OMP_TEAMS: + save_suppress = info->suppress_expansion; + convert_local_omp_clauses (gimple_omp_teams_clauses_ptr (stmt), wi); + walk_body (convert_local_reference_stmt, convert_local_reference_op, + info, gimple_omp_body_ptr (stmt)); + info->suppress_expansion = save_suppress; + break; + case GIMPLE_OMP_SECTION: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_ORDERED: @@ -2071,6 +2103,8 @@ convert_gimple_call (gimple_stmt_iterato case GIMPLE_OMP_SECTIONS: case GIMPLE_OMP_SECTION: case GIMPLE_OMP_SINGLE: + case GIMPLE_OMP_TARGET: + case GIMPLE_OMP_TEAMS: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_ORDERED: case GIMPLE_OMP_CRITICAL: --- gcc/tree-inline.c.jj 2013-05-13 16:49:40.000000000 +0200 +++ gcc/tree-inline.c 2013-05-23 13:53:52.258166517 +0200 @@ -1298,7 +1298,8 @@ remap_gimple_stmt (gimple stmt, copy_bod case GIMPLE_OMP_FOR: s1 = remap_gimple_seq (gimple_omp_body (stmt), id); s2 = remap_gimple_seq (gimple_omp_for_pre_body (stmt), id); - copy = gimple_build_omp_for (s1, gimple_omp_for_clauses (stmt), + copy = gimple_build_omp_for (s1, gimple_omp_for_kind (stmt), + gimple_omp_for_clauses (stmt), gimple_omp_for_collapse (stmt), s2); { size_t i; @@ -1345,6 +1346,19 @@ remap_gimple_stmt (gimple stmt, copy_bod (s1, gimple_omp_single_clauses (stmt)); break; + case GIMPLE_OMP_TARGET: + s1 = remap_gimple_seq (gimple_omp_body (stmt), id); + copy = gimple_build_omp_target + (s1, gimple_omp_target_kind (stmt), + gimple_omp_target_clauses (stmt)); + break; + + case GIMPLE_OMP_TEAMS: + s1 = remap_gimple_seq (gimple_omp_body (stmt), id); + copy = gimple_build_omp_teams + (s1, gimple_omp_teams_clauses (stmt)); + break; + case GIMPLE_OMP_CRITICAL: s1 = remap_gimple_seq (gimple_omp_body (stmt), id); copy @@ -3716,6 +3730,8 @@ estimate_num_insns (gimple stmt, eni_wei case GIMPLE_OMP_SECTION: case GIMPLE_OMP_SECTIONS: case GIMPLE_OMP_SINGLE: + case GIMPLE_OMP_TARGET: + case GIMPLE_OMP_TEAMS: return (weights->omp_cost + estimate_num_insns_seq (gimple_omp_body (stmt), weights)); --- gcc/gimple.def.jj 2013-03-20 10:05:01.000000000 +0100 +++ gcc/gimple.def 2013-05-23 12:45:29.902586194 +0200 @@ -287,7 +287,7 @@ DEFGSCODE(GIMPLE_OMP_ORDERED, "gimple_om BODY is a the sequence of statements to be executed by all threads. - CLAUSES is a TREE_LIST node with all the clauses. + CLAUSES is an OMP_CLAUSE chain with all the clauses. CHILD_FN is set when outlining the body of the parallel region. All the statements in BODY are moved into this newly created @@ -306,7 +306,7 @@ DEFGSCODE(GIMPLE_OMP_PARALLEL, "gimple_o BODY is a the sequence of statements to be executed by all threads. - CLAUSES is a TREE_LIST node with all the clauses. + CLAUSES is an OMP_CLAUSE chain with all the clauses. CHILD_FN is set when outlining the body of the explicit task region. All the statements in BODY are moved into this newly created @@ -334,7 +334,7 @@ DEFGSCODE(GIMPLE_OMP_SECTION, "gimple_om /* OMP_SECTIONS represents #pragma omp sections. BODY is the sequence of statements in the sections body. - CLAUSES is a TREE_LIST node holding the list of associated clauses. + CLAUSES is an OMP_CLAUSE chain holding the list of associated clauses. CONTROL is a VAR_DECL used for deciding which of the sections to execute. */ DEFGSCODE(GIMPLE_OMP_SECTIONS, "gimple_omp_sections", GSS_OMP_SECTIONS) @@ -346,9 +346,21 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "g /* GIMPLE_OMP_SINGLE represents #pragma omp single BODY is the sequence of statements inside the single section. - CLAUSES is a TREE_LIST node holding the associated clauses. */ + CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */ DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE) +/* GIMPLE_OMP_TARGET represents + #pragma omp target {,data,update} + BODY is the sequence of statements inside the target construct + (NULL for target update). + CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */ +DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_SINGLE) + +/* GIMPLE_OMP_TEAMS represents #pragma omp teams + BODY is the sequence of statements inside the single section. + CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */ +DEFGSCODE(GIMPLE_OMP_TEAMS, "gimple_omp_teams", GSS_OMP_SINGLE) + /* GIMPLE_PREDICT specifies a hint for branch prediction. PREDICT is one of the predictors from predict.def. --- gcc/cp/parser.c.jj 2013-05-20 13:21:27.000000000 +0200 +++ gcc/cp/parser.c 2013-05-22 19:12:25.523182096 +0200 @@ -26823,22 +26823,7 @@ cp_parser_omp_clause_cancelkind (cp_pars enum omp_clause_code code, tree list, location_t location) { - tree c; - - for (c = list; c; c = OMP_CLAUSE_CHAIN (c)) - switch (OMP_CLAUSE_CODE (c)) - { - case OMP_CLAUSE_PARALLEL: - case OMP_CLAUSE_FOR: - case OMP_CLAUSE_SECTIONS: - case OMP_CLAUSE_TASKGROUP: - error_at (location, "only one of %, %, % " - "and % clauses can be specified"); - break; - default: - break; - } - c = build_omp_clause (location, code); + tree c = build_omp_clause (location, code); OMP_CLAUSE_CHAIN (c) = list; return c; } @@ -27260,7 +27245,6 @@ cp_parser_omp_all_clauses (cp_parser *pa token = cp_lexer_peek_token (parser->lexer); c_kind = cp_parser_omp_clause_name (parser); - first = false; switch (c_kind) { @@ -27359,31 +27343,48 @@ cp_parser_omp_all_clauses (cp_parser *pa clauses = cp_parser_omp_clause_cancelkind (parser, OMP_CLAUSE_PARALLEL, clauses, token->location); c_name = "parallel"; + if (!first) + { + clause_not_first: + error_at (token->location, "%qs must be the first clause of %qs", + c_name, where); + clauses = prev; + } break; case PRAGMA_OMP_CLAUSE_FOR: clauses = cp_parser_omp_clause_cancelkind (parser, OMP_CLAUSE_FOR, clauses, token->location); c_name = "for"; + if (!first) + goto clause_not_first; break; case PRAGMA_OMP_CLAUSE_SECTIONS: clauses = cp_parser_omp_clause_cancelkind (parser, OMP_CLAUSE_SECTIONS, clauses, token->location); c_name = "sections"; + if (!first) + goto clause_not_first; break; case PRAGMA_OMP_CLAUSE_TASKGROUP: clauses = cp_parser_omp_clause_cancelkind (parser, OMP_CLAUSE_TASKGROUP, clauses, token->location); c_name = "taskgroup"; + if (!first) + goto clause_not_first; break; case PRAGMA_OMP_CLAUSE_TO: clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses); c_name = "to"; + if (!first) + goto clause_not_first; break; case PRAGMA_OMP_CLAUSE_FROM: clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses); c_name = "from"; + if (!first) + goto clause_not_first; break; case PRAGMA_OMP_CLAUSE_UNIFORM: clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_UNIFORM, @@ -27441,6 +27442,8 @@ cp_parser_omp_all_clauses (cp_parser *pa goto saw_error; } + first = false; + if (((mask >> c_kind) & 1) == 0) { /* Remove the invalid clause(s) from the list to avoid @@ -29013,6 +29016,180 @@ cp_parser_omp_cancellation_point (cp_par } /* OpenMP 4.0: + # pragma omp teams teams-clause[optseq] new-line + structured-block */ + +#define OMP_TEAMS_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SHARED) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_TEAMS) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_THREADS) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULT)) + +static tree +cp_parser_omp_teams (cp_parser *parser, cp_token *pragma_tok) +{ + tree stmt = make_node (OMP_TEAMS); + TREE_TYPE (stmt) = void_type_node; + + OMP_TEAMS_CLAUSES (stmt) + = cp_parser_omp_all_clauses (parser, OMP_TEAMS_CLAUSE_MASK, + "#pragma omp teams", pragma_tok); + OMP_TEAMS_BODY (stmt) = cp_parser_omp_structured_block (parser); + + return add_stmt (stmt); +} + +/* OpenMP 4.0: + # pragma omp target data target-data-clause[optseq] new-line + structured-block */ + +#define OMP_TARGET_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)) + +static tree +cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok) +{ + tree stmt = make_node (OMP_TARGET_DATA); + TREE_TYPE (stmt) = void_type_node; + + OMP_TARGET_DATA_CLAUSES (stmt) + = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, + "#pragma omp target data", pragma_tok); + OMP_TARGET_DATA_BODY (stmt) = cp_parser_omp_structured_block (parser); + + SET_EXPR_LOCATION (stmt, pragma_tok->location); + return add_stmt (stmt); +} + +/* OpenMP 4.0: + # pragma omp target update target-update-clause[optseq] new-line */ + +#define OMP_TARGET_UPDATE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FROM) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_TO) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)) + +static bool +cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok, + enum pragma_context context) +{ + if (context == pragma_stmt) + { + error_at (pragma_tok->location, + "%<#pragma omp target update%> may only be " + "used in compound statements"); + cp_parser_skip_to_pragma_eol (parser, pragma_tok); + return false; + } + + tree clauses + = cp_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK, + "#pragma omp target update", pragma_tok); + if (find_omp_clause (clauses, OMP_CLAUSE_TO) == NULL_TREE + && find_omp_clause (clauses, OMP_CLAUSE_FROM) == NULL_TREE) + { + error_at (pragma_tok->location, + "%<#pragma omp target update must contain either " + "% or % clauses"); + return false; + } + + tree stmt = make_node (OMP_TARGET_UPDATE); + TREE_TYPE (stmt) = void_type_node; + OMP_TARGET_UPDATE_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, pragma_tok->location); + add_stmt (stmt); + return false; +} + +/* OpenMP 4.0: + # pragma omp target target-clause[optseq] new-line + structured-block */ + +#define OMP_TARGET_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)) + +static bool +cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, + enum pragma_context context) +{ + if (context != pragma_stmt && context != pragma_compound) + { + cp_parser_error (parser, "expected declaration specifiers"); + cp_parser_skip_to_pragma_eol (parser, pragma_tok); + return false; + } + + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + tree id = cp_lexer_peek_token (parser->lexer)->u.value; + const char *p = IDENTIFIER_POINTER (id); + + if (strcmp (p, "data") == 0) + { + cp_lexer_consume_token (parser->lexer); + cp_parser_omp_target_data (parser, pragma_tok); + return true; + } + else if (strcmp (p, "update") == 0) + { + cp_lexer_consume_token (parser->lexer); + return cp_parser_omp_target_update (parser, pragma_tok, context); + } + } + + tree stmt = make_node (OMP_TARGET); + TREE_TYPE (stmt) = void_type_node; + + OMP_TARGET_CLAUSES (stmt) + = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, + "#pragma omp target", pragma_tok); + OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser); + + SET_EXPR_LOCATION (stmt, pragma_tok->location); + add_stmt (stmt); + return true; +} + +/* OpenMP 4.0: + #pragma omp distribute distribute-clause[optseq] new-line + for-loop */ + +#define OMP_DISTRIBUTE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)\ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE)) + +static tree +cp_parser_omp_distribute (cp_parser *parser, cp_token *pragma_tok) +{ + tree clauses, sb, ret; + unsigned int save; + + clauses = cp_parser_omp_all_clauses (parser, OMP_DISTRIBUTE_CLAUSE_MASK, + "#pragma omp distribute", pragma_tok); + + sb = begin_omp_structured_block (); + save = cp_parser_begin_omp_structured_block (parser); + + ret = cp_parser_omp_for_loop (parser, OMP_DISTRIBUTE, clauses, NULL); + + cp_parser_end_omp_structured_block (parser, save); + add_stmt (finish_omp_structured_block (sb)); + + return ret; +} + +/* OpenMP 4.0: # pragma omp declare simd declare-simd-clauses[optseq] new-line */ #define OMP_DECLARE_SIMD_CLAUSE_MASK \ @@ -29112,6 +29289,9 @@ cp_parser_omp_construct (cp_parser *pars case PRAGMA_OMP_CRITICAL: stmt = cp_parser_omp_critical (parser, pragma_tok); break; + case PRAGMA_OMP_DISTRIBUTE: + stmt = cp_parser_omp_distribute (parser, pragma_tok); + break; case PRAGMA_OMP_FOR: stmt = cp_parser_omp_for (parser, pragma_tok); break; @@ -29139,6 +29319,9 @@ cp_parser_omp_construct (cp_parser *pars case PRAGMA_OMP_TASKGROUP: cp_parser_omp_taskgroup (parser, pragma_tok); return; + case PRAGMA_OMP_TEAMS: + stmt = cp_parser_omp_teams (parser, pragma_tok); + break; default: gcc_unreachable (); } @@ -29609,6 +29792,7 @@ cp_parser_pragma (cp_parser *parser, enu case PRAGMA_OMP_ATOMIC: case PRAGMA_OMP_CRITICAL: + case PRAGMA_OMP_DISTRIBUTE: case PRAGMA_OMP_FOR: case PRAGMA_OMP_MASTER: case PRAGMA_OMP_ORDERED: @@ -29618,11 +29802,15 @@ cp_parser_pragma (cp_parser *parser, enu case PRAGMA_OMP_SINGLE: case PRAGMA_OMP_TASK: case PRAGMA_OMP_TASKGROUP: + case PRAGMA_OMP_TEAMS: if (context != pragma_stmt && context != pragma_compound) goto bad_stmt; cp_parser_omp_construct (parser, pragma_tok); return true; + case PRAGMA_OMP_TARGET: + return cp_parser_omp_target (parser, pragma_tok, context); + case PRAGMA_OMP_SECTION: error_at (pragma_tok->location, "%<#pragma omp section%> may only be used in " --- gcc/cp/pt.c.jj 2013-05-20 13:21:25.000000000 +0200 +++ gcc/cp/pt.c 2013-05-22 18:44:31.959390265 +0200 @@ -13330,6 +13330,9 @@ tsubst_expr (tree t, tree args, tsubst_f case OMP_SECTIONS: case OMP_SINGLE: + case OMP_TEAMS: + case OMP_TARGET_DATA: + case OMP_TARGET: tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, args, complain, in_decl); stmt = push_stmt_list (); @@ -13341,6 +13344,14 @@ tsubst_expr (tree t, tree args, tsubst_f OMP_CLAUSES (t) = tmp; add_stmt (t); break; + + case OMP_TARGET_UPDATE: + tmp = tsubst_omp_clauses (OMP_TARGET_UPDATE_CLAUSES (t), false, + args, complain, in_decl); + t = copy_node (t); + OMP_CLAUSES (t) = tmp; + add_stmt (t); + break; case OMP_SECTION: case OMP_CRITICAL: --- gcc/tree-parloops.c.jj 2013-05-13 16:46:37.000000000 +0200 +++ gcc/tree-parloops.c 2013-05-23 13:54:21.854439781 +0200 @@ -1686,7 +1686,7 @@ create_parallel_loop (struct loop *loop, t = build_omp_clause (loc, OMP_CLAUSE_SCHEDULE); OMP_CLAUSE_SCHEDULE_KIND (t) = OMP_CLAUSE_SCHEDULE_STATIC; - for_stmt = gimple_build_omp_for (NULL, t, 1, NULL); + for_stmt = gimple_build_omp_for (NULL, GF_OMP_FOR_KIND_FOR, t, 1, NULL); gimple_set_location (for_stmt, loc); gimple_omp_for_set_index (for_stmt, 0, initvar); gimple_omp_for_set_initial (for_stmt, 0, cvar_init); --- gcc/tree-pretty-print.c.jj 2013-05-20 13:18:24.000000000 +0200 +++ gcc/tree-pretty-print.c 2013-05-22 19:00:26.349842964 +0200 @@ -2347,6 +2347,27 @@ dump_generic_node (pretty_printer *buffe pp_string (buffer, "#pragma omp distribute"); goto dump_omp_loop; + case OMP_TEAMS: + pp_string (buffer, "#pragma omp teams"); + dump_omp_clauses (buffer, OMP_TEAMS_CLAUSES (node), spc, flags); + goto dump_omp_body; + + case OMP_TARGET_DATA: + pp_string (buffer, "#pragma omp target data"); + dump_omp_clauses (buffer, OMP_TARGET_DATA_CLAUSES (node), spc, flags); + goto dump_omp_body; + + case OMP_TARGET: + pp_string (buffer, "#pragma omp target"); + dump_omp_clauses (buffer, OMP_TARGET_CLAUSES (node), spc, flags); + goto dump_omp_body; + + case OMP_TARGET_UPDATE: + pp_string (buffer, "#pragma omp target update"); + dump_omp_clauses (buffer, OMP_TARGET_UPDATE_CLAUSES (node), spc, flags); + is_expr = false; + break; + dump_omp_loop: dump_omp_clauses (buffer, OMP_FOR_CLAUSES (node), spc, flags); --- gcc/gimple.h.jj 2013-05-13 16:49:47.000000000 +0200 +++ gcc/gimple.h 2013-05-23 14:06:41.729631141 +0200 @@ -115,6 +115,10 @@ enum gf_mask { GF_OMP_FOR_KIND_SIMD = 1 << 0, GF_OMP_FOR_KIND_FOR_SIMD = 2 << 0, GF_OMP_FOR_KIND_DISTRIBUTE = 3 << 0, + GF_OMP_TARGET_KIND_MASK = 3 << 0, + GF_OMP_TARGET_KIND_REGION = 0 << 0, + GF_OMP_TARGET_KIND_DATA = 1 << 0, + GF_OMP_TARGET_KIND_UPDATE = 2 << 0, /* True on an GIMPLE_OMP_RETURN statement if the return does not require a thread synchronization via some sort of barrier. The exact barrier @@ -618,7 +622,7 @@ struct GTY(()) gimple_statement_omp_cont tree control_use; }; -/* GIMPLE_OMP_SINGLE */ +/* GIMPLE_OMP_SINGLE, GIMPLE_OMP_TARGET, GIMPLE_OMP_TEAMS */ struct GTY(()) gimple_statement_omp_single { /* [ WORD 1-7 ] */ @@ -805,7 +809,7 @@ gimple gimple_build_switch_nlabels (unsi gimple gimple_build_switch (tree, tree, vec ); gimple gimple_build_omp_parallel (gimple_seq, tree, tree, tree); gimple gimple_build_omp_task (gimple_seq, tree, tree, tree, tree, tree, tree); -gimple gimple_build_omp_for (gimple_seq, tree, size_t, gimple_seq); +gimple gimple_build_omp_for (gimple_seq, int, tree, size_t, gimple_seq); gimple gimple_build_omp_critical (gimple_seq, tree); gimple gimple_build_omp_section (gimple_seq); gimple gimple_build_omp_continue (tree, tree); @@ -815,6 +819,8 @@ gimple gimple_build_omp_ordered (gimple_ gimple gimple_build_omp_sections (gimple_seq, tree); gimple gimple_build_omp_sections_switch (void); gimple gimple_build_omp_single (gimple_seq, tree); +gimple gimple_build_omp_target (gimple_seq, int, tree); +gimple gimple_build_omp_teams (gimple_seq, tree); gimple gimple_build_cdt (tree, tree); gimple gimple_build_omp_atomic_load (tree, tree); gimple gimple_build_omp_atomic_store (tree); @@ -1264,6 +1270,8 @@ gimple_has_substatements (gimple g) case GIMPLE_OMP_TASK: case GIMPLE_OMP_SECTIONS: case GIMPLE_OMP_SINGLE: + case GIMPLE_OMP_TARGET: + case GIMPLE_OMP_TEAMS: case GIMPLE_OMP_CRITICAL: case GIMPLE_WITH_CLEANUP_EXPR: case GIMPLE_TRANSACTION: @@ -1691,7 +1699,7 @@ static inline unsigned gimple_omp_subcode (const_gimple s) { gcc_gimple_checking_assert (gimple_code (s) >= GIMPLE_OMP_ATOMIC_LOAD - && gimple_code (s) <= GIMPLE_OMP_SINGLE); + && gimple_code (s) <= GIMPLE_OMP_TEAMS); return s->gsbase.subcode; } @@ -4604,6 +4612,87 @@ gimple_omp_single_set_clauses (gimple gs } +/* Return the clauses associated with OMP_TARGET GS. */ + +static inline tree +gimple_omp_target_clauses (const_gimple gs) +{ + GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET); + return gs->gimple_omp_single.clauses; +} + + +/* Return a pointer to the clauses associated with OMP_TARGET GS. */ + +static inline tree * +gimple_omp_target_clauses_ptr (gimple gs) +{ + GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET); + return &gs->gimple_omp_single.clauses; +} + + +/* Set CLAUSES to be the clauses associated with OMP_TARGET GS. */ + +static inline void +gimple_omp_target_set_clauses (gimple gs, tree clauses) +{ + GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET); + gs->gimple_omp_single.clauses = clauses; +} + + +/* Return the kind of OMP target statemement. */ + +static inline int +gimple_omp_target_kind (const_gimple g) +{ + GIMPLE_CHECK (g, GIMPLE_OMP_TARGET); + return (gimple_omp_subcode (g) & GF_OMP_TARGET_KIND_MASK); +} + + +/* Set the OMP target kind. */ + +static inline void +gimple_omp_target_set_kind (gimple g, int kind) +{ + GIMPLE_CHECK (g, GIMPLE_OMP_TARGET); + g->gsbase.subcode = (g->gsbase.subcode & ~GF_OMP_TARGET_KIND_MASK) + | (kind & GF_OMP_TARGET_KIND_MASK); +} + + +/* Return the clauses associated with OMP_TEAMS GS. */ + +static inline tree +gimple_omp_teams_clauses (const_gimple gs) +{ + GIMPLE_CHECK (gs, GIMPLE_OMP_TEAMS); + return gs->gimple_omp_single.clauses; +} + + +/* Return a pointer to the clauses associated with OMP_TEAMS GS. */ + +static inline tree * +gimple_omp_teams_clauses_ptr (gimple gs) +{ + GIMPLE_CHECK (gs, GIMPLE_OMP_TEAMS); + return &gs->gimple_omp_single.clauses; +} + + +/* Set CLAUSES to be the clauses associated with OMP_TEAMS GS. */ + +static inline void +gimple_omp_teams_set_clauses (gimple gs, tree clauses) +{ + GIMPLE_CHECK (gs, GIMPLE_OMP_TEAMS); + gs->gimple_omp_single.clauses = clauses; +} + + /* Return the clauses associated with OMP_SECTIONS GS. */ static inline tree @@ -4946,6 +5035,8 @@ gimple_return_set_retval (gimple gs, tre case GIMPLE_OMP_SECTIONS: \ case GIMPLE_OMP_SECTIONS_SWITCH: \ case GIMPLE_OMP_SINGLE: \ + case GIMPLE_OMP_TARGET: \ + case GIMPLE_OMP_TEAMS: \ case GIMPLE_OMP_SECTION: \ case GIMPLE_OMP_MASTER: \ case GIMPLE_OMP_ORDERED: \ --- gcc/gimple.c.jj 2013-05-13 16:49:46.000000000 +0200 +++ gcc/gimple.c 2013-05-23 13:40:12.487789257 +0200 @@ -908,13 +908,14 @@ gimple_build_omp_critical (gimple_seq bo PRE_BODY is the sequence of statements that are loop invariant. */ gimple -gimple_build_omp_for (gimple_seq body, tree clauses, size_t collapse, +gimple_build_omp_for (gimple_seq body, int kind, tree clauses, size_t collapse, gimple_seq pre_body) { gimple p = gimple_alloc (GIMPLE_OMP_FOR, 0); if (body) gimple_omp_set_body (p, body); gimple_omp_for_set_clauses (p, clauses); + gimple_omp_for_set_kind (p, kind); p->gimple_omp_for.collapse = collapse; p->gimple_omp_for.iter = ggc_alloc_cleared_vec_gimple_omp_for_iter (collapse); @@ -1094,6 +1095,41 @@ gimple_build_omp_single (gimple_seq body } +/* Build a GIMPLE_OMP_TARGET statement. + + BODY is the sequence of statements that will be executed. + CLAUSES are any of the OMP target construct's clauses. */ + +gimple +gimple_build_omp_target (gimple_seq body, int kind, tree clauses) +{ + gimple p = gimple_alloc (GIMPLE_OMP_TARGET, 0); + if (body) + gimple_omp_set_body (p, body); + gimple_omp_target_set_clauses (p, clauses); + gimple_omp_target_set_kind (p, kind); + + return p; +} + + +/* Build a GIMPLE_OMP_TEAMS statement. + + BODY is the sequence of statements that will be executed. + CLAUSES are any of the OMP teams construct's clauses. */ + +gimple +gimple_build_omp_teams (gimple_seq body, tree clauses) +{ + gimple p = gimple_alloc (GIMPLE_OMP_TEAMS, 0); + if (body) + gimple_omp_set_body (p, body); + gimple_omp_teams_set_clauses (p, clauses); + + return p; +} + + /* Build a GIMPLE_OMP_ATOMIC_LOAD statement. */ gimple @@ -1610,6 +1646,20 @@ walk_gimple_op (gimple stmt, walk_tree_f return ret; break; + case GIMPLE_OMP_TARGET: + ret = walk_tree (gimple_omp_target_clauses_ptr (stmt), callback_op, wi, + pset); + if (ret) + return ret; + break; + + case GIMPLE_OMP_TEAMS: + ret = walk_tree (gimple_omp_teams_clauses_ptr (stmt), callback_op, wi, + pset); + if (ret) + return ret; + break; + case GIMPLE_OMP_ATOMIC_LOAD: ret = walk_tree (gimple_omp_atomic_load_lhs_ptr (stmt), callback_op, wi, pset); @@ -1786,6 +1836,8 @@ walk_gimple_stmt (gimple_stmt_iterator * case GIMPLE_OMP_TASK: case GIMPLE_OMP_SECTIONS: case GIMPLE_OMP_SINGLE: + case GIMPLE_OMP_TARGET: + case GIMPLE_OMP_TEAMS: ret = walk_gimple_seq_mod (gimple_omp_body_ptr (stmt), callback_stmt, callback_op, wi); if (ret) @@ -2308,6 +2360,8 @@ gimple_copy (gimple stmt) /* FALLTHRU */ case GIMPLE_OMP_SINGLE: + case GIMPLE_OMP_TARGET: + case GIMPLE_OMP_TEAMS: case GIMPLE_OMP_SECTION: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_ORDERED: --- gcc/gimple-pretty-print.c.jj 2013-05-13 16:49:01.000000000 +0200 +++ gcc/gimple-pretty-print.c 2013-05-23 13:17:37.589649980 +0200 @@ -1264,6 +1264,78 @@ dump_gimple_omp_single (pretty_printer * } } +/* Dump a GIMPLE_OMP_TARGET tuple on the pretty_printer BUFFER. */ + +static void +dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags) +{ + const char *kind; + switch (gimple_omp_target_kind (gs)) + { + case GF_OMP_TARGET_KIND_REGION: + kind = ""; + break; + case GF_OMP_TARGET_KIND_DATA: + kind = " data"; + break; + case GF_OMP_TARGET_KIND_UPDATE: + kind = " update"; + break; + default: + gcc_unreachable (); + } + if (flags & TDF_RAW) + { + dump_gimple_fmt (buffer, spc, flags, "%G%s <%+BODY <%S>%nCLAUSES <", gs, + kind, gimple_omp_body (gs)); + dump_omp_clauses (buffer, gimple_omp_target_clauses (gs), spc, flags); + dump_gimple_fmt (buffer, spc, flags, " >"); + } + else + { + pp_string (buffer, "#pragma omp target"); + pp_string (buffer, kind); + dump_omp_clauses (buffer, gimple_omp_target_clauses (gs), spc, flags); + if (!gimple_seq_empty_p (gimple_omp_body (gs))) + { + newline_and_indent (buffer, spc + 2); + pp_character (buffer, '{'); + pp_newline (buffer); + dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags); + newline_and_indent (buffer, spc + 2); + pp_character (buffer, '}'); + } + } +} + +/* Dump a GIMPLE_OMP_TEAMS tuple on the pretty_printer BUFFER. */ + +static void +dump_gimple_omp_teams (pretty_printer *buffer, gimple gs, int spc, int flags) +{ + if (flags & TDF_RAW) + { + dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S>%nCLAUSES <", gs, + gimple_omp_body (gs)); + dump_omp_clauses (buffer, gimple_omp_teams_clauses (gs), spc, flags); + dump_gimple_fmt (buffer, spc, flags, " >"); + } + else + { + pp_string (buffer, "#pragma omp teams"); + dump_omp_clauses (buffer, gimple_omp_teams_clauses (gs), spc, flags); + if (!gimple_seq_empty_p (gimple_omp_body (gs))) + { + newline_and_indent (buffer, spc + 2); + pp_character (buffer, '{'); + pp_newline (buffer); + dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags); + newline_and_indent (buffer, spc + 2); + pp_character (buffer, '}'); + } + } +} + /* Dump a GIMPLE_OMP_SECTIONS tuple on the pretty_printer BUFFER. */ static void @@ -2038,6 +2110,14 @@ pp_gimple_stmt_1 (pretty_printer *buffer dump_gimple_omp_single (buffer, gs, spc, flags); break; + case GIMPLE_OMP_TARGET: + dump_gimple_omp_target (buffer, gs, spc, flags); + break; + + case GIMPLE_OMP_TEAMS: + dump_gimple_omp_teams (buffer, gs, spc, flags); + break; + case GIMPLE_OMP_RETURN: dump_gimple_omp_return (buffer, gs, spc, flags); break; --- gcc/gimplify.c.jj 2013-05-13 16:49:10.000000000 +0200 +++ gcc/gimplify.c 2013-05-24 14:04:01.534225316 +0200 @@ -57,10 +57,11 @@ enum gimplify_omp_var_data GOVD_LASTPRIVATE = 32, GOVD_REDUCTION = 64, GOVD_LOCAL = 128, - GOVD_DEBUG_PRIVATE = 256, - GOVD_PRIVATE_OUTER_REF = 512, - GOVD_LINEAR = 1024, - GOVD_ALIGNED = 2048, + GOVD_MAP = 256, + GOVD_DEBUG_PRIVATE = 512, + GOVD_PRIVATE_OUTER_REF = 1024, + GOVD_LINEAR = 2048, + GOVD_ALIGNED = 4096, GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -74,7 +75,10 @@ enum omp_region_type ORT_PARALLEL = 2, ORT_COMBINED_PARALLEL = 3, ORT_TASK = 4, - ORT_UNTIED_TASK = 5 + ORT_UNTIED_TASK = 5, + ORT_TEAMS = 8, + ORT_TARGET_DATA = 16, + ORT_TARGET = 32 }; struct gimplify_omp_ctx @@ -5829,6 +5833,9 @@ omp_add_variable (struct gimplify_omp_ct the parameters of the type. */ if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { + /* To be handled later. */ + gcc_assert ((flags & GOVD_MAP) == 0); + /* Add the pointer replacement variable as PRIVATE if the variable replacement is private, else FIRSTPRIVATE since we'll need the address of the original variable either for SHARED, or for the @@ -5870,6 +5877,9 @@ omp_add_variable (struct gimplify_omp_ct } else if (lang_hooks.decls.omp_privatize_by_reference (decl)) { + /* To be handled later. */ + gcc_assert ((flags & GOVD_MAP) == 0); + gcc_assert ((flags & GOVD_LOCAL) == 0); omp_firstprivatize_type_sizes (ctx, TREE_TYPE (decl)); @@ -5896,6 +5906,22 @@ omp_notice_threadprivate_variable (struc tree decl2) { splay_tree_node n; + struct gimplify_omp_ctx *octx; + + for (octx = ctx; octx; octx = octx->outer_context) + if (octx->region_type == ORT_TARGET) + { + n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); + if (n == NULL) + { + error ("threadprivate variable %qE used in target region", + DECL_NAME (decl)); + error_at (octx->location, "enclosing target region"); + splay_tree_insert (octx->variables, (splay_tree_key)decl, 0); + } + if (decl2) + splay_tree_insert (octx->variables, (splay_tree_key)decl2, 0); + } if (ctx->region_type != ORT_UNTIED_TASK) return false; @@ -5944,13 +5970,24 @@ omp_notice_variable (struct gimplify_omp } n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); + if (ctx->region_type == ORT_TARGET) + { + if (n == NULL) + omp_add_variable (ctx, decl, GOVD_MAP | flags); + else + n->value |= flags; + ret = lang_hooks.decls.omp_disregard_value_expr (decl, true); + goto do_outer; + } + if (n == NULL) { enum omp_clause_default_kind default_kind, kind; struct gimplify_omp_ctx *octx; if (ctx->region_type == ORT_WORKSHARE - || ctx->region_type == ORT_SIMD) + || ctx->region_type == ORT_SIMD + || ctx->region_type == ORT_TARGET_DATA) goto do_outer; /* ??? Some compiler-generated variables (like SAVE_EXPRs) could be @@ -5964,12 +6001,24 @@ omp_notice_variable (struct gimplify_omp switch (default_kind) { case OMP_CLAUSE_DEFAULT_NONE: - error ("%qE not specified in enclosing parallel", - DECL_NAME (lang_hooks.decls.omp_report_decl (decl))); if ((ctx->region_type & ORT_TASK) != 0) - error_at (ctx->location, "enclosing task"); + { + error ("%qE not specified in enclosing task", + DECL_NAME (lang_hooks.decls.omp_report_decl (decl))); + error_at (ctx->location, "enclosing task"); + } + else if (ctx->region_type == ORT_TEAMS) + { + error ("%qE not specified in enclosing teams construct", + DECL_NAME (lang_hooks.decls.omp_report_decl (decl))); + error_at (ctx->location, "enclosing teams construct"); + } else - error_at (ctx->location, "enclosing parallel"); + { + error ("%qE not specified in enclosing parallel", + DECL_NAME (lang_hooks.decls.omp_report_decl (decl))); + error_at (ctx->location, "enclosing parallel"); + } /* FALLTHRU */ case OMP_CLAUSE_DEFAULT_SHARED: flags |= GOVD_SHARED; @@ -5989,13 +6038,15 @@ omp_notice_variable (struct gimplify_omp { splay_tree_node n2; + if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0) + continue; n2 = splay_tree_lookup (octx->variables, (splay_tree_key) decl); if (n2 && (n2->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED) { flags |= GOVD_FIRSTPRIVATE; break; } - if ((octx->region_type & ORT_PARALLEL) != 0) + if ((octx->region_type & (ORT_PARALLEL | ORT_TEAMS)) != 0) break; } if (flags & GOVD_FIRSTPRIVATE) @@ -6137,6 +6188,9 @@ omp_check_private (struct gimplify_omp_c /* References might be private, but might be shared too. */ || lang_hooks.decls.omp_privatize_by_reference (decl)); + if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0) + continue; + n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); if (n != NULL) return (n->value & GOVD_SHARED) == 0; @@ -6204,6 +6258,20 @@ gimplify_scan_omp_clauses (tree *list_p, } flags = GOVD_LINEAR | GOVD_EXPLICIT; goto do_add; + case OMP_CLAUSE_MAP: + flags = GOVD_MAP | GOVD_EXPLICIT; + notice_outer = false; + goto do_add; + + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: + decl = OMP_CLAUSE_DECL (c); + if (error_operand_p (decl)) + { + remove = true; + break; + } + goto do_notice; do_add: decl = OMP_CLAUSE_DECL (c); @@ -6292,6 +6360,9 @@ gimplify_scan_omp_clauses (tree *list_p, case OMP_CLAUSE_SCHEDULE: case OMP_CLAUSE_NUM_THREADS: + case OMP_CLAUSE_NUM_TEAMS: + case OMP_CLAUSE_DIST_SCHEDULE: + case OMP_CLAUSE_DEVICE: if (gimplify_expr (&OMP_CLAUSE_OPERAND (c, 0), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) remove = true; @@ -6357,12 +6428,40 @@ gimplify_adjust_omp_clauses_1 (splay_tre gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_PRIVATE); private_debug = true; } + else if (flags & GOVD_MAP) + private_debug = false; else private_debug = lang_hooks.decls.omp_private_debug_clause (decl, !!(flags & GOVD_SHARED)); if (private_debug) code = OMP_CLAUSE_PRIVATE; + else if (flags & GOVD_MAP) + { + /* If decl is already in the enclosing device data environment, + the spec says that it should just be used and no init/assignment + should be done. If there was any privatization in between though, + it means that original decl might be in the enclosing device data + environment, but the privatized might not. */ + struct gimplify_omp_ctx *ctx; + for (ctx = gimplify_omp_ctxp->outer_context; + ctx; ctx = ctx->outer_context) + { + n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); + if (n == NULL) + continue; + if (ctx->region_type == ORT_TARGET_DATA) + { + if ((n->value & GOVD_MAP) != 0) + return 0; + } + else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE + | GOVD_PRIVATE | GOVD_REDUCTION + | GOVD_LINEAR)) != 0) + break; + } + code = OMP_CLAUSE_MAP; + } else if (flags & GOVD_SHARED) { if (is_global_var (decl)) @@ -6373,7 +6472,8 @@ gimplify_adjust_omp_clauses_1 (splay_tre splay_tree_node on = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); if (on && (on->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE - | GOVD_PRIVATE | GOVD_REDUCTION)) != 0) + | GOVD_PRIVATE | GOVD_REDUCTION + | GOVD_LINEAR)) != 0) break; ctx = ctx->outer_context; } @@ -6400,6 +6500,8 @@ gimplify_adjust_omp_clauses_1 (splay_tre OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1; else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF)) OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1; + else if (code == OMP_CLAUSE_MAP) + OMP_CLAUSE_MAP_KIND (clause) = OMP_CLAUSE_MAP_TOFROM; *list_p = clause; lang_hooks.decls.omp_finish_clause (clause); @@ -6517,11 +6619,47 @@ gimplify_adjust_omp_clauses (tree *list_ } break; + case OMP_CLAUSE_MAP: + decl = OMP_CLAUSE_DECL (c); + n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); + remove = false; + if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN)) + remove = true; + else + { + /* If decl is already in the enclosing device data environment, + the spec says that it should just be used and no init/assignment + should be done. If there was any privatization in between though, + it means that original decl might be in the enclosing device data + environment, but the privatized might not. */ + struct gimplify_omp_ctx *octx; + for (octx = ctx->outer_context; octx; octx = octx->outer_context) + { + n = splay_tree_lookup (octx->variables, + (splay_tree_key) decl); + if (n == NULL) + continue; + if (octx->region_type == ORT_TARGET_DATA) + { + if ((n->value & GOVD_MAP) != 0) + remove = true; + } + else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE + | GOVD_PRIVATE | GOVD_REDUCTION + | GOVD_LINEAR)) != 0) + break; + } + } + break; + case OMP_CLAUSE_REDUCTION: case OMP_CLAUSE_COPYIN: case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_IF: case OMP_CLAUSE_NUM_THREADS: + case OMP_CLAUSE_NUM_TEAMS: + case OMP_CLAUSE_DIST_SCHEDULE: + case OMP_CLAUSE_DEVICE: case OMP_CLAUSE_SCHEDULE: case OMP_CLAUSE_NOWAIT: case OMP_CLAUSE_ORDERED: @@ -6532,6 +6670,8 @@ gimplify_adjust_omp_clauses (tree *list_ case OMP_CLAUSE_MERGEABLE: case OMP_CLAUSE_PROC_BIND: case OMP_CLAUSE_SAFELEN: + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: break; default: @@ -6847,25 +6987,19 @@ gimplify_omp_for (tree *expr_p, gimple_s gimplify_adjust_omp_clauses (&OMP_FOR_CLAUSES (for_stmt)); - gfor = gimple_build_omp_for (for_body, OMP_FOR_CLAUSES (for_stmt), - TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)), - for_pre_body); + int kind; switch (TREE_CODE (for_stmt)) { - case OMP_FOR: - break; - case OMP_SIMD: - gimple_omp_for_set_kind (gfor, GF_OMP_FOR_KIND_SIMD); - break; - case OMP_FOR_SIMD: - gimple_omp_for_set_kind (gfor, GF_OMP_FOR_KIND_FOR_SIMD); - break; - case OMP_DISTRIBUTE: - gimple_omp_for_set_kind (gfor, GF_OMP_FOR_KIND_DISTRIBUTE); - break; + case OMP_FOR: kind = GF_OMP_FOR_KIND_FOR; break; + case OMP_SIMD: kind = GF_OMP_FOR_KIND_SIMD; break; + case OMP_FOR_SIMD: kind = GF_OMP_FOR_KIND_FOR_SIMD; break; + case OMP_DISTRIBUTE: kind = GF_OMP_FOR_KIND_DISTRIBUTE; break; default: gcc_unreachable (); } + gfor = gimple_build_omp_for (for_body, kind, OMP_FOR_CLAUSES (for_stmt), + TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)), + for_pre_body); for (i = 0; i < TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)); i++) { @@ -6880,11 +7014,15 @@ gimplify_omp_for (tree *expr_p, gimple_s } gimplify_seq_add_stmt (pre_p, gfor); - return ret == GS_ALL_DONE ? GS_ALL_DONE : GS_ERROR; + if (ret != GS_ALL_DONE) + return GS_ERROR; + *expr_p = NULL_TREE; + return GS_ALL_DONE; } -/* Gimplify the gross structure of other OpenMP worksharing constructs. - In particular, OMP_SECTIONS and OMP_SINGLE. */ +/* Gimplify the gross structure of other OpenMP constructs. + In particular, OMP_SECTIONS, OMP_SINGLE, OMP_TARGET, OMP_TARGET_DATA + and OMP_TEAMS. */ static void gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) @@ -6892,19 +7030,72 @@ gimplify_omp_workshare (tree *expr_p, gi tree expr = *expr_p; gimple stmt; gimple_seq body = NULL; + enum omp_region_type ort = ORT_WORKSHARE; - gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ORT_WORKSHARE); + switch (TREE_CODE (expr)) + { + case OMP_SECTIONS: + case OMP_SINGLE: + break; + case OMP_TARGET: + ort = ORT_TARGET; + break; + case OMP_TARGET_DATA: + ort = ORT_TARGET_DATA; + break; + case OMP_TEAMS: + ort = ORT_TEAMS; + break; + default: + gcc_unreachable (); + } + gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort); gimplify_and_add (OMP_BODY (expr), &body); gimplify_adjust_omp_clauses (&OMP_CLAUSES (expr)); - if (TREE_CODE (expr) == OMP_SECTIONS) - stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr)); - else if (TREE_CODE (expr) == OMP_SINGLE) - stmt = gimple_build_omp_single (body, OMP_CLAUSES (expr)); - else - gcc_unreachable (); + switch (TREE_CODE (expr)) + { + case OMP_SECTIONS: + stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr)); + break; + case OMP_SINGLE: + stmt = gimple_build_omp_single (body, OMP_CLAUSES (expr)); + break; + case OMP_TARGET: + stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_REGION, + OMP_CLAUSES (expr)); + break; + case OMP_TARGET_DATA: + stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_DATA, + OMP_CLAUSES (expr)); + break; + case OMP_TEAMS: + stmt = gimple_build_omp_teams (body, OMP_CLAUSES (expr)); + break; + default: + gcc_unreachable (); + } + + gimplify_seq_add_stmt (pre_p, stmt); + *expr_p = NULL_TREE; +} + +/* Gimplify the gross structure of OpenMP target update construct. */ + +static void +gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) +{ + tree expr = *expr_p; + gimple stmt; + + gimplify_scan_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr), pre_p, + ORT_WORKSHARE); + gimplify_adjust_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr)); + stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_UPDATE, + OMP_TARGET_UPDATE_CLAUSES (expr)); gimplify_seq_add_stmt (pre_p, stmt); + *expr_p = NULL_TREE; } /* A subroutine of gimplify_omp_atomic. The front end is supposed to have @@ -7811,10 +8002,18 @@ gimplify_expr (tree *expr_p, gimple_seq case OMP_SECTIONS: case OMP_SINGLE: + case OMP_TARGET: + case OMP_TARGET_DATA: + case OMP_TEAMS: gimplify_omp_workshare (expr_p, pre_p); ret = GS_ALL_DONE; break; + case OMP_TARGET_UPDATE: + gimplify_omp_target_update (expr_p, pre_p); + ret = GS_ALL_DONE; + break; + case OMP_SECTION: case OMP_MASTER: case OMP_ORDERED: