From patchwork Thu Oct 22 19:15:43 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: James Norris X-Patchwork-Id: 534547 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 B30F2141303 for ; Fri, 23 Oct 2015 06:15:59 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=L903gTxX; 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 :message-id:date:from:mime-version:to:cc:subject:references :in-reply-to:content-type; q=dns; s=default; b=X0MvBErl5s8vLMBOB yAqY5TnMDIkJs4fBSJITR9geOqqE+y/LoYe3N75olkfmTbD2EWgRNuS8MvNTengJ x6gaK+VPb8f8fFWZ3UkKpnr1NQmDCDlaDXmGwTtnJRGkP+WLCGy9vDGeiAX8SgZJ 229UPAeKEnjBS0bV9OmknvvZek= 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 :message-id:date:from:mime-version:to:cc:subject:references :in-reply-to:content-type; s=default; bh=rwrKMkfl+p80Vdho7gOauL3 JhuI=; b=L903gTxXOyxxf+rxx+rQdnsQnJgiAx1OaQXm9bshLs30ZIaw1GOqbqp Pcr7yh8EFgemU4iW68Ep/ATXdyc/jiRpRIfeFhzH81cAe0z2EbbNLKN7ktyCP3jB h1Flui17uM2PDP84eZ2Qmng6+LoMYdxfSTTM1CiPqscOgeQ+6xGY= Received: (qmail 43957 invoked by alias); 22 Oct 2015 19:15:50 -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 43946 invoked by uid 89); 22 Oct 2015 19:15:49 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.6 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_LOW, SPF_PASS, T_FROM_12LTRDOM, UNWANTED_LANGUAGE_BODY 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:15:48 +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 1ZpLKz-0003IL-TX from James_Norris@mentor.com ; Thu, 22 Oct 2015 12:15:45 -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:15:45 -0700 Message-ID: <562935DF.9040503@codesourcery.com> Date: Thu, 22 Oct 2015 14:15:43 -0500 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 3/7] host_data construct (C front-end) References: <56293476.5020801@codesourcery.com> In-Reply-To: <56293476.5020801@codesourcery.com> X-TagToolbar-Keys: D20151022141543460 gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h - Add definition for c_finish_oacc_host_data(). gcc/cp/parser.c b/gcc/cp/parser.c - Add handling of use_device clause in cp_parser_omp_clause_name(). - Add handling of use_device clause in cp_parser_oacc_all_clauses(). - Add new macro OACC_HOST_DATA_CLAUSE_MASK. - Add new function cp_parser_oacc_host_data() to handle host_data. - Add handling of host_data pragma to cp_parser_omp_construct(). - Add handling of host_data pragma to cp_parser_pragma(). gcc/cp/semantics.c b/gcc/cp/semantics.c - Add handling of use_device clause to finish_omp_clauses(). - Add new function finish_oacc_host_data(). diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index 16db41f..76ece42 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -6318,6 +6318,7 @@ extern void finish_omp_threadprivate (tree); extern tree begin_omp_structured_block (void); extern tree finish_omp_structured_block (tree); extern tree finish_oacc_data (tree, tree); +extern tree finish_oacc_host_data (tree, tree); extern tree finish_oacc_kernels (tree, tree); extern tree finish_oacc_parallel (tree, tree); extern tree begin_omp_parallel (void); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index f07a5e4..714e69c 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -29235,6 +29235,8 @@ cp_parser_omp_clause_name (cp_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_length", p)) @@ -31381,6 +31383,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "self"; break; + case PRAGMA_OACC_CLAUSE_USE_DEVICE: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE, + clauses); + c_name = "use_device"; + break; case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH: clauses = cp_parser_oacc_clause_vector_length (parser, clauses); c_name = "vector_length"; @@ -34221,6 +34228,30 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok) return stmt; } +#define OACC_HOST_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) ) + +/* OpenACC 2.0: + # pragma acc host_data new-line + structured-block */ + +static tree +cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok) +{ + tree stmt, clauses, block; + unsigned int save; + + clauses = cp_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, + "#pragma acc host_data", pragma_tok); + + block = begin_omp_parallel (); + save = cp_parser_begin_omp_structured_block (parser); + cp_parser_statement (parser, NULL_TREE, false, NULL); + cp_parser_end_omp_structured_block (parser, save); + stmt = finish_oacc_host_data (clauses, block); + return stmt; +} + /* OpenACC 2.0: # pragma acc enter data oacc-enter-data-clause[optseq] new-line @@ -35288,6 +35319,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok) case PRAGMA_OACC_EXIT_DATA: stmt = cp_parser_oacc_enter_exit_data (parser, pragma_tok, false); break; + case PRAGMA_OACC_HOST_DATA: + stmt = cp_parser_oacc_host_data (parser, pragma_tok); + break; case PRAGMA_OACC_KERNELS: stmt = cp_parser_oacc_kernels (parser, pragma_tok); break; @@ -35856,6 +35890,7 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context) case PRAGMA_OACC_DATA: case PRAGMA_OACC_ENTER_DATA: case PRAGMA_OACC_EXIT_DATA: + case PRAGMA_OACC_HOST_DATA: case PRAGMA_OACC_KERNELS: case PRAGMA_OACC_PARALLEL: case PRAGMA_OACC_LOOP: diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index c0a8b32..25482e7 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -6689,6 +6689,7 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) case OMP_CLAUSE_SIMD: case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE__CILK_FOR_COUNT_: + case OMP_CLAUSE_USE_DEVICE: break; case OMP_CLAUSE_INBRANCH: @@ -7119,6 +7120,24 @@ finish_oacc_data (tree clauses, tree block) return add_stmt (stmt); } +/* Generate OACC_HOST_DATA, with CLAUSES and BLOCK as its compound + statement. */ + +tree +finish_oacc_host_data (tree clauses, tree block) +{ + tree stmt; + + block = finish_omp_structured_block (block); + + stmt = make_node (OACC_HOST_DATA); + TREE_TYPE (stmt) = void_type_node; + OACC_HOST_DATA_CLAUSES (stmt) = clauses; + OACC_HOST_DATA_BODY (stmt) = block; + + return add_stmt (stmt); +} + /* Generate OACC_KERNELS, with CLAUSES and BLOCK as its compound statement. LOC is the location of the OACC_KERNELS. */