From patchwork Thu Jun 12 10:58:42 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 359082 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 822121400B0 for ; Thu, 12 Jun 2014 20:59:08 +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:from :to:cc:subject:date:message-id:mime-version:content-type; q=dns; s=default; b=Aex8eKACijsIBixkrqm0XkwKnIUFT6Xup5xHwYVOvRPLRxlge/ FW0qC6g1lxLoyhZrXNmCYWHZhogW7AVRsyr4+XvGyI+dysoaNgFaygy6vusXUxlI 62T/eqTQ2YYGg1/TCO9OJtvBKq06tepjVm02EYHthLnIzDIzs0huVSyEU= 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; s= default; bh=/SgouFc8KtV/aE/OjpalZte+muI=; b=XlfloGfwKt8rzyGaZPHb x3aYolWtxLmGsNp2+Ne7B1PFd50OAP/dc2I3Cij8fYYWel6GdxnML9/rKzk9T6pO TBLEHmU4X7fW1DZads0JuFyQI7+WoPs+wlMLMmi2YUVeVur11cdmTTq5TeJVFKIW uREutaCjzswguZtL0rVFmRA= Received: (qmail 26290 invoked by alias); 12 Jun 2014 10:59:01 -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 26141 invoked by uid 89); 12 Jun 2014 10:58:59 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL, BAYES_00 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; Thu, 12 Jun 2014 10:58:58 +0000 Received: from svr-orw-fem-01.mgc.mentorg.com ([147.34.98.93]) by relay1.mentorg.com with esmtp id 1Wv2ib-0004Dp-8o from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Thu, 12 Jun 2014 03:58:53 -0700 Received: from SVR-IES-FEM-01.mgc.mentorg.com ([137.202.0.104]) by svr-orw-fem-01.mgc.mentorg.com over TLS secured channel with Microsoft SMTPSVC(6.0.3790.4675); Thu, 12 Jun 2014 03:58:53 -0700 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.2.247.3; Thu, 12 Jun 2014 11:58:51 +0100 From: Thomas Schwinge To: CC: Subject: [GOMP4, COMMITTED] OpenACC if clause. Date: Thu, 12 Jun 2014 12:58:42 +0200 Message-ID: <1402570722-17540-1-git-send-email-thomas@codesourcery.com> MIME-Version: 1.0 From: tschwinge gcc/c/ * c-parser.c (c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_IF. (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK) (OACC_PARALLEL_CLAUSE_MASK, OACC_UPDATE_CLAUSE_MASK): Add it. gcc/ * omp-low.c (scan_sharing_clauses): Allow OMP_CLAUSE_IF. (expand_oacc_offload, expand_omp_target): Handle it. gcc/testsuite/ * c-c++-common/goacc/if-clause-1.c: New file. * c-c++-common/goacc/if-clause-2.c: Likewise. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@211510 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 6 ++ gcc/c/ChangeLog.gomp | 8 +++ gcc/c/c-parser.c | 10 +++- gcc/omp-low.c | 81 +++++++++++++++++++++----- gcc/testsuite/ChangeLog.gomp | 5 ++ gcc/testsuite/c-c++-common/goacc/if-clause-1.c | 8 +++ gcc/testsuite/c-c++-common/goacc/if-clause-2.c | 11 ++++ 7 files changed, 112 insertions(+), 17 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/if-clause-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/if-clause-2.c diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index be1aa16..2abe179 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,3 +1,9 @@ +2014-06-12 Thomas Schwinge + James Norris + + * omp-low.c (scan_sharing_clauses): Allow OMP_CLAUSE_IF. + (expand_oacc_offload, expand_omp_target): Handle it. + 2014-06-06 Thomas Schwinge James Norris diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp index f1e45f3..108ce3e 100644 --- gcc/c/ChangeLog.gomp +++ gcc/c/ChangeLog.gomp @@ -1,3 +1,11 @@ +2014-06-12 Thomas Schwinge + James Norris + + * c-parser.c (c_parser_oacc_all_clauses): Handle + PRAGMA_OMP_CLAUSE_IF. + (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK) + (OACC_PARALLEL_CLAUSE_MASK, OACC_UPDATE_CLAUSE_MASK): Add it. + 2014-06-06 Thomas Schwinge James Norris diff --git gcc/c/c-parser.c gcc/c/c-parser.c index bf4bad62..6269923 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -10203,7 +10203,7 @@ c_parser_omp_clause_final (c_parser *parser, tree list) return list; } -/* OpenMP 2.5: +/* OpenACC, OpenMP 2.5: if ( expression ) */ static tree @@ -11295,6 +11295,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "host"; break; + case PRAGMA_OMP_CLAUSE_IF: + clauses = c_parser_omp_clause_if (parser, clauses); + c_name = "if"; + break; case PRAGMA_OMP_CLAUSE_NUM_GANGS: clauses = c_parser_omp_clause_num_gangs (parser, clauses); c_name = "num_gangs"; @@ -11614,6 +11618,7 @@ c_parser_omp_structured_block (c_parser *parser) | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \ @@ -11649,6 +11654,7 @@ c_parser_oacc_data (location_t loc, c_parser *parser) | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \ @@ -11727,6 +11733,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name) | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_GANGS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_WORKERS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \ @@ -11775,6 +11782,7 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name) #define OACC_UPDATE_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HOST) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SELF) ) static void diff --git gcc/omp-low.c gcc/omp-low.c index 6fd0f30..958f116 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1634,12 +1634,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c); break; - case OMP_CLAUSE_IF: - if (is_gimple_omp_oacc_specifically (ctx->stmt)) - { - sorry ("clause not supported yet"); - break; - } case OMP_CLAUSE_FINAL: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_NUM_TEAMS: @@ -1649,6 +1643,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_DIST_SCHEDULE: case OMP_CLAUSE_DEPEND: gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt)); + case OMP_CLAUSE_IF: case OMP_CLAUSE_NUM_GANGS: case OMP_CLAUSE_NUM_WORKERS: case OMP_CLAUSE_VECTOR_LENGTH: @@ -1916,12 +1911,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) } break; - case OMP_CLAUSE_IF: - if (is_gimple_omp_oacc_specifically (ctx->stmt)) - { - sorry ("clause not supported yet"); - break; - } case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_COPYIN: case OMP_CLAUSE_DEFAULT: @@ -1945,6 +1934,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt)); + case OMP_CLAUSE_IF: case OMP_CLAUSE_NUM_GANGS: case OMP_CLAUSE_NUM_WORKERS: case OMP_CLAUSE_VECTOR_LENGTH: @@ -5115,7 +5105,7 @@ expand_oacc_offload (struct omp_region *region) /* Emit a library call to launch CHILD_FN. */ tree t1, t2, t3, t4, t_num_gangs, t_num_workers, t_vector_length, - device, c, clauses; + device, cond, c, clauses; enum built_in_function start_ix; location_t clause_loc; tree (*gimple_omp_clauses) (const_gimple); @@ -5165,9 +5155,15 @@ expand_oacc_offload (struct omp_region *region) break; } - /* By default, the value of DEVICE is -1 (let runtime library choose). */ + /* By default, the value of DEVICE is -1 (let runtime library choose) + and there is no conditional. */ + cond = NULL_TREE; device = build_int_cst (integer_type_node, -1); + c = find_omp_clause (clauses, OMP_CLAUSE_IF); + if (c) + cond = OMP_CLAUSE_IF_EXPR (c); + c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE); gcc_assert (c == NULL); if (c) @@ -5181,6 +5177,61 @@ expand_oacc_offload (struct omp_region *region) /* Ensure 'device' is of the correct type. */ device = fold_convert_loc (clause_loc, integer_type_node, device); + /* If we found the clause 'if (cond)', build + (cond ? device : -2). */ + if (cond) + { + cond = gimple_boolify (cond); + + basic_block cond_bb, then_bb, else_bb; + edge e; + tree tmp_var; + + tmp_var = create_tmp_var (TREE_TYPE (device), NULL); + /* Preserve indentation of expand_omp_target. */ + if (0) + { + gsi = gsi_last_bb (new_bb); + gsi_prev (&gsi); + e = split_block (new_bb, gsi_stmt (gsi)); + } + else + e = split_block (new_bb, NULL); + cond_bb = e->src; + new_bb = e->dest; + remove_edge (e); + + then_bb = create_empty_bb (cond_bb); + else_bb = create_empty_bb (then_bb); + set_immediate_dominator (CDI_DOMINATORS, then_bb, cond_bb); + set_immediate_dominator (CDI_DOMINATORS, else_bb, cond_bb); + + stmt = gimple_build_cond_empty (cond); + gsi = gsi_last_bb (cond_bb); + gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING); + + gsi = gsi_start_bb (then_bb); + stmt = gimple_build_assign (tmp_var, device); + gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING); + + gsi = gsi_start_bb (else_bb); + stmt = gimple_build_assign (tmp_var, + build_int_cst (integer_type_node, -2)); + gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING); + + make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE); + make_edge (cond_bb, else_bb, EDGE_FALSE_VALUE); + if (current_loops) + { + add_bb_to_loop (then_bb, cond_bb->loop_father); + add_bb_to_loop (else_bb, cond_bb->loop_father); + } + make_edge (then_bb, new_bb, EDGE_FALLTHRU); + make_edge (else_bb, new_bb, EDGE_FALLTHRU); + + device = tmp_var; + } + gsi = gsi_last_bb (new_bb); t = gimple_omp_data_arg (entry_stmt); if (t == NULL) @@ -8661,7 +8712,6 @@ expand_omp_target (struct omp_region *region) device = build_int_cst (integer_type_node, -1); c = find_omp_clause (clauses, OMP_CLAUSE_IF); - gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA); if (c) cond = OMP_CLAUSE_IF_EXPR (c); @@ -8684,7 +8734,6 @@ expand_omp_target (struct omp_region *region) (cond ? device : -2). */ if (cond) { - gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA); cond = gimple_boolify (cond); basic_block cond_bb, then_bb, else_bb; diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 17493ce..9b65cc9 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-06-12 Thomas Schwinge + + * c-c++-common/goacc/if-clause-1.c: New file. + * c-c++-common/goacc/if-clause-2.c: Likewise. + 2014-06-06 Thomas Schwinge * c-c++-common/goacc/pragma_context.c: New file. diff --git gcc/testsuite/c-c++-common/goacc/if-clause-1.c gcc/testsuite/c-c++-common/goacc/if-clause-1.c new file mode 100644 index 0000000..c078596 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/if-clause-1.c @@ -0,0 +1,8 @@ +void +f (void) +{ + struct { int i; } *p; +#pragma acc data copyout(p) if(1) if(1) /* { dg-error "too many 'if' clauses" } */ + ; +#pragma acc update device(p) if(*p) /* { dg-error "used struct type value where scalar is required" } */ +} diff --git gcc/testsuite/c-c++-common/goacc/if-clause-2.c gcc/testsuite/c-c++-common/goacc/if-clause-2.c new file mode 100644 index 0000000..5ab8459 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/if-clause-2.c @@ -0,0 +1,11 @@ +void +f (short c) +{ +#pragma acc parallel if(c) + ; +#pragma acc kernels if(c) + ; +#pragma acc data if(c) + ; +#pragma acc update device(c) if(c) +}