From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 41956 invoked by alias); 22 Oct 2015 19:15:01 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 41936 invoked by uid 89); 22 Oct 2015 19:15:01 -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,RCVD_IN_DNSWL_LOW,SPF_PASS,T_FROM_12LTRDOM 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, 22 Oct 2015 19:14:59 +0000 Received: from svr-orw-fem-02x.mgc.mentorg.com ([147.34.96.206] helo=SVR-ORW-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1ZpLKC-0002yX-Nj from James_Norris@mentor.com ; Thu, 22 Oct 2015 12:14:56 -0700 Received: from [172.30.80.51] (147.34.91.1) by svr-orw-fem-02.mgc.mentorg.com (147.34.96.168) with Microsoft SMTP Server id 14.3.224.2; Thu, 22 Oct 2015 12:14:55 -0700 Message-ID: <562935AE.2010109@codesourcery.com> Date: Thu, 22 Oct 2015 19:15:00 -0000 From: James Norris User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:31.0) Gecko/20100101 Thunderbird/31.7.0 MIME-Version: 1.0 To: GCC Patches CC: "Joseph S. Myers" , Nathan Sidwell , Jakub Jelinek Subject: [OpenACC 2/7] host_data construct (C FE) References: <56293476.5020801@codesourcery.com> In-Reply-To: <56293476.5020801@codesourcery.com> X-TagToolbar-Keys: D20151022141454406 Content-Type: multipart/mixed; boundary="------------060609090906020704010506" X-SW-Source: 2015-10/txt/msg02332.txt.bz2 --------------060609090906020704010506 Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit Content-length: 710 gcc/c/c-parser.c b/gcc/c/c-parser.c - Add handling of use_device clause in c_parser_omp_clause_name(). - Add new function c_parser_oacc_clause_use_device() to handle use_device clause with host_data. - Add handling of use_device clause in c_parser_oacc_all_clauses(). - Add new macro OACC_HOST_DATA_CLAUSE_MASK. - Add new function c_parser_oacc_host_data() to handle host_data. - Add handling of host_data pragma to c_parser_omp_construct(). gcc/c/c-tree.h b/gcc/c/c-tree.h - Add definition for c_finish_oacc_host_data(). gcc/c/c-typeck.c b/gcc/c/c-typeck.c - Add new function c_finish_oacc_host_data(). --------------060609090906020704010506 Content-Type: text/x-patch; name="p2.patch" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="p2.patch" Content-length: 4133 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 704ebc6..ead98b9 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -10116,6 +10116,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_UNTIED; else if (!strcmp ("use_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; + else if (!strcmp ("use_device", p)) + result = PRAGMA_OACC_CLAUSE_USE_DEVICE; break; case 'v': if (!strcmp ("vector", p)) @@ -11219,6 +11221,15 @@ c_parser_oacc_clause_async (c_parser *parser, tree list) return list; } +/* OpenACC 2.0: + use_device ( variable-list ) */ + +static tree +c_parser_oacc_clause_use_device (c_parser *parser, tree list) +{ + return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE, list); +} + /* OpenACC: wait ( int-expr-list ) */ @@ -12474,6 +12485,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 = "self"; break; + case PRAGMA_OACC_CLAUSE_USE_DEVICE: + clauses = c_parser_oacc_clause_use_device (parser, clauses); + c_name = "use_device"; + break; case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH: clauses = c_parser_omp_clause_vector_length (parser, clauses); c_name = "vector_length"; @@ -13003,6 +13018,29 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter) /* OpenACC 2.0: + # pragma acc host_data oacc-data-clause[optseq] new-line + structured-block +*/ + +#define OACC_HOST_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) ) + +static tree +c_parser_oacc_host_data (location_t loc, c_parser *parser) +{ + tree stmt, clauses, block; + + clauses = c_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, + "#pragma acc host_data"); + + block = c_begin_omp_parallel (); + add_stmt (c_parser_omp_structured_block (parser)); + stmt = c_finish_oacc_host_data (loc, clauses, block); + return stmt; +} + + +/* OpenACC 2.0: # pragma acc loop oacc-loop-clause[optseq] new-line structured-block @@ -16075,6 +16113,9 @@ c_parser_omp_construct (c_parser *parser) case PRAGMA_OACC_DATA: stmt = c_parser_oacc_data (loc, parser); break; + case PRAGMA_OACC_HOST_DATA: + stmt = c_parser_oacc_host_data (loc, parser); + break; case PRAGMA_OACC_KERNELS: strcpy (p_name, "#pragma acc"); stmt = c_parser_oacc_kernels (loc, parser, p_name); diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h index bee03d3..a9c5975 100644 --- a/gcc/c/c-tree.h +++ b/gcc/c/c-tree.h @@ -643,6 +643,7 @@ extern tree c_expr_to_decl (tree, bool *, bool *); extern tree c_finish_oacc_parallel (location_t, tree, tree); extern tree c_finish_oacc_kernels (location_t, tree, tree); extern tree c_finish_oacc_data (location_t, tree, tree); +extern tree c_finish_oacc_host_data (location_t, tree, tree); extern tree c_begin_omp_parallel (void); extern tree c_finish_omp_parallel (location_t, tree, tree); extern tree c_begin_omp_task (void); diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index bc43602..a5e2a4a 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -11510,6 +11510,25 @@ c_finish_oacc_data (location_t loc, tree clauses, tree block) return add_stmt (stmt); } +/* Generate OACC_HOST_DATA, with CLAUSES and BLOCK as its compound + statement. LOC is the location of the OACC_HOST_DATA. */ + +tree +c_finish_oacc_host_data (location_t loc, tree clauses, tree block) +{ + tree stmt; + + block = c_end_compound_stmt (loc, block, true); + + stmt = make_node (OACC_HOST_DATA); + TREE_TYPE (stmt) = void_type_node; + OACC_HOST_DATA_CLAUSES (stmt) = clauses; + OACC_HOST_DATA_BODY (stmt) = block; + SET_EXPR_LOCATION (stmt, loc); + + return add_stmt (stmt); +} + /* Like c_begin_compound_stmt, except force the retention of the BLOCK. */ tree @@ -12942,6 +12961,7 @@ 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_USE_DEVICE: pc = &OMP_CLAUSE_CHAIN (c); continue; --------------060609090906020704010506--