From patchwork Thu Nov 6 15:18:22 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 407580 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 A196514007D for ; Fri, 7 Nov 2014 02:18:49 +1100 (AEDT) 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:subject:in-reply-to:references:date:message-id:mime-version :content-type; q=dns; s=default; b=I7usGmwnuf6vvmn0wcbnndbtZUOgF 8jLGmupWAgYe9mB95o2/t3gXCxn8/9fi+lrplVBi4HIF4EoWH8J/9NEpqv6mupEk LmSVvR4BxNMGviENXFc+zXd9i4Y9lj36zx6mW9RDSIpJwQOFerXSlM4wuOZje5YC nCNOJ0IVtY9cZc= 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:subject:in-reply-to:references:date:message-id:mime-version :content-type; s=default; bh=0HrIG92AjC+OmvQ7QLCRFzPPxjE=; b=da5 FXcjIPg/+khO1t3Q2eVp+UCgy06PdY2zDhxcCixKpg97rUJ2vCGRAJQk0oRrgrhe cSv9Z7y91AGyiZPwnAlUFqbSvNV6x2xI5gAQyqsJ4n4M47LDNC8XRydPvTDt36Wx Dv1KfEgnTbXbac/m/ulxYX4W29+fxEJDw4NKeDRs= Received: (qmail 9752 invoked by alias); 6 Nov 2014 15:18:36 -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 9661 invoked by uid 89); 6 Nov 2014 15:18:35 -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_NONE 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, 06 Nov 2014 15:18:33 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1XmOpQ-0004e5-Q9 from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Thu, 06 Nov 2014 07:18:29 -0800 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.3.181.6; Thu, 6 Nov 2014 15:18:27 +0000 From: Thomas Schwinge To: James Norris , Subject: [gomp4] OpenACC Executable Directives In-Reply-To: <545934A3.4040405@mentor.com> References: <545934A3.4040405@mentor.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.3.1 (x86_64-pc-linux-gnu) Date: Thu, 6 Nov 2014 16:18:22 +0100 Message-ID: <87egtgs6gx.fsf@schwinge.name> MIME-Version: 1.0 Hi Jim! On Tue, 4 Nov 2014 14:18:43 -0600, James Norris wrote: > void > f (int flag,, int *a) > { > if (flag) > #pragma acc update host (a[0:10]) > > return; > } > > I read bullet one under "Restrictions" on page 43, to mean the > placement of the update directive statement in the above example > is wrong. Given the poor grammar, this could be read a different way. That's also my interpretation, and I think it should be that way for all executable directives (as listed 2.12 Executable Directives, plus 2.8 Cache Directive -- am I missing additional ones?), for the reason we just talked about: if -fopenacc is not active, this code will turn into: if (flag) return; ..., which probably was not intended. I think this is spelled out clearly in the OpenMP specification (but didn't verify right now), and see also the discussion in . As to fixing this, see how I'm handling PRAGMA_OACC_UPDATE in gcc/c/c-parser:c_parser_pragma; modelled after PRAGMA_OMP_* executable directives. Plus, I started testing this in gcc/testsuite/c-c++-common/goacc/pragma_context.c. We should also be adding appropritate documentation to the C/C++ parser source code (or generally GCC user documentation?) -- it's easy to understand once you got it, but it also took me a too long time to... As you're saying, this is a different kind of C/C++ pragma usage that what they're "normally" used for. This is not a priority right now, but if you or somebody wants to pick it up, here is my WIP patch for the C PRAGMA_OACC_CACHE: commit a294b3299d4214f2cbe611cc71ef190b247716cf Author: Thomas Schwinge Date: Wed Nov 5 18:57:41 2014 +0100 OpenACC cache: pragma context TODO: gcc/cp/ --- gcc/c/c-parser.c | 33 +++++++++++++---------- gcc/testsuite/c-c++-common/goacc/pragma_context.c | 13 +++++++-- 2 files changed, 30 insertions(+), 16 deletions(-) Grüße, Thomas diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 9f4b013..57d29e5 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -1247,6 +1247,7 @@ static vec *c_parser_expr_list (c_parser *, bool, bool, static tree c_parser_oacc_loop (location_t, c_parser *, char *); static void c_parser_omp_construct (c_parser *); static void c_parser_omp_threadprivate (c_parser *); +static void c_parser_oacc_cache (c_parser *); static void c_parser_oacc_enter_exit_data (c_parser *, bool); static void c_parser_oacc_update (c_parser *); static void c_parser_omp_barrier (c_parser *); @@ -9575,6 +9576,17 @@ c_parser_pragma (c_parser *parser, enum pragma_context context) switch (id) { + case PRAGMA_OACC_CACHE: + if (context != pragma_compound) + { + if (context == pragma_stmt) + c_parser_error (parser, "%<#pragma acc cache%> may only be " + "used in compound statements"); + goto bad_stmt; + } + c_parser_oacc_cache (parser); + return false; + case PRAGMA_OACC_ENTER_DATA: c_parser_oacc_enter_exit_data (parser, true); return false; @@ -11926,27 +11938,24 @@ c_parser_omp_structured_block (c_parser *parser) /* OpenACC 2.0: # pragma acc cache (variable-list) new-line - - LOC is the location of the #pragma token. */ -static tree -c_parser_oacc_cache (location_t loc, c_parser *parser) +static void +c_parser_oacc_cache (c_parser *parser) { - tree stmt, clauses; + location_t loc = c_parser_peek_token (parser)->location; - clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL); + c_parser_consume_pragma (parser); + + tree clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL); clauses = c_finish_omp_clauses (clauses); - c_parser_skip_to_pragma_eol (parser); - stmt = make_node (OACC_CACHE); + tree stmt = make_node (OACC_CACHE); TREE_TYPE (stmt) = void_type_node; OACC_CACHE_CLAUSES (stmt) = clauses; SET_EXPR_LOCATION (stmt, loc); add_stmt (stmt); - - return stmt; } /* OpenACC 2.0: @@ -14749,10 +14758,6 @@ c_parser_omp_construct (c_parser *parser) switch (p_kind) { - case PRAGMA_OACC_CACHE: - strcpy (p_name, "#pragma acc"); - stmt = c_parser_oacc_cache (loc, parser); - break; case PRAGMA_OACC_DATA: stmt = c_parser_oacc_data (loc, parser); break; diff --git gcc/testsuite/c-c++-common/goacc/pragma_context.c gcc/testsuite/c-c++-common/goacc/pragma_context.c index ad33d92..e76d2d6 100644 --- gcc/testsuite/c-c++-common/goacc/pragma_context.c +++ gcc/testsuite/c-c++-common/goacc/pragma_context.c @@ -1,15 +1,18 @@ // pragma_external +#pragma acc cache /* { dg-error "expected declaration specifiers before '#pragma'" } */ #pragma acc update /* { dg-error "expected declaration specifiers before '#pragma'" } */ // pragma_struct struct s_pragma_struct { +#pragma acc cache /* { dg-error "expected declaration specifiers before '#pragma'" } */ #pragma acc update /* { dg-error "expected declaration specifiers before '#pragma'" } */ }; // pragma_param void f_pragma_param ( +#pragma acc cache /* { dg-error "expected declaration specifiers before '#pragma'" } */ #pragma acc update /* { dg-error "expected declaration specifiers before '#pragma'" } */ void) { @@ -20,13 +23,19 @@ void f2 (void) { if (0) +#pragma acc cache /* { dg-error "'#pragma acc cache' may only be used in compound statements before '#pragma'" } */ + ; + if (0) #pragma acc update /* { dg-error "'#pragma acc update' may only be used in compound statements before '#pragma'" } */ + ; } // pragma_compound void f3 (void) { - int i = 0; -#pragma acc update device(i) + int i[10] = { 0 }; + +#pragma acc cache (i[1:3]) +#pragma acc update device(i[3:2]) }