From patchwork Wed Jun 24 15:23:50 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: James Norris X-Patchwork-Id: 488097 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 9E920140316 for ; Thu, 25 Jun 2015 01:24:06 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=ZlaYnxtc; 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:content-type; q=dns; s=default; b=Du6ZppZl45v7WE3qsjo34iV6smSGsyATsCCbkS6vAGd RR8h2nQ2CxtfyZKYKjAL8gLsSztdDw3jWUWKd0XW4CEx52aJq8yA5jn7sDb+fFQo cSfWqQU7ueggImtqwIOBP6PTFAERl6QTcjma66YderXzEOWpbCW2YARaCVYiu3bE = 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:content-type; s=default; bh=IzUqR5r0kaOjnFipDF47F3emdhY=; b=ZlaYnxtcU3o6z98jC El7hy+WtX9FBiL6qDo7Pb+ibbR6zaxpy72AxmslYCzJejfhyvEKfUWAGKjYJIgPt d4tEUSoGyvrsgFtiGoqFA7QyjIEVYQPtaSdPtBu9gDtZo6HVl2Ic4NQCuDdjtIHz eP8vS+C2CK10QfYO79XNIkgfoM= Received: (qmail 88827 invoked by alias); 24 Jun 2015 15:23:58 -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 88812 invoked by uid 89); 24 Jun 2015 15:23:58 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.5 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; Wed, 24 Jun 2015 15:23:56 +0000 Received: from svr-orw-fem-04.mgc.mentorg.com ([147.34.97.41]) by relay1.mentorg.com with esmtp id 1Z7mWn-0000CH-61 from James_Norris@mentor.com for gcc-patches@gcc.gnu.org; Wed, 24 Jun 2015 08:23:53 -0700 Received: from [172.30.80.204] (147.34.91.1) by svr-orw-fem-04.mgc.mentorg.com (147.34.97.41) with Microsoft SMTP Server id 14.3.224.2; Wed, 24 Jun 2015 08:23:52 -0700 Message-ID: <558ACB86.3020608@codesourcery.com> Date: Wed, 24 Jun 2015 10:23:50 -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: CC: Thomas Schwinge Subject: [gomp4] Additional tests for declare directive and fixes. Hi! The following patch adds additional testing of the declare directive and fixes for issues that arose from the testing. Committed to gomp-4_0-branch. Jim diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index e7df751..bcbd163 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -1767,12 +1767,15 @@ finish_oacc_declare (tree fnbody, tree decls) break; } - stmt = make_node (OACC_DECLARE); - TREE_TYPE (stmt) = void_type_node; - OACC_DECLARE_CLAUSES (stmt) = ret_clauses; - SET_EXPR_LOCATION (stmt, loc); + if (ret_clauses) + { + stmt = make_node (OACC_DECLARE); + TREE_TYPE (stmt) = void_type_node; + OACC_DECLARE_CLAUSES (stmt) = ret_clauses; + SET_EXPR_LOCATION (stmt, loc); - tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); + tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); + } DECL_ATTRIBUTES (fndecl) = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl)); @@ -12812,6 +12815,14 @@ c_parser_oacc_declare (c_parser *parser) error = true; continue; } + else if (TREE_PUBLIC (decl)) + { + error_at (loc, + "invalid use of % variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } break; } diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c index 15da51e..a35f599 100644 --- a/gcc/cp/decl.c +++ b/gcc/cp/decl.c @@ -14343,7 +14343,17 @@ finish_oacc_declare (tree fndecl, tree decls) { t = tsi_stmt (i); if (TREE_CODE (t) == BIND_EXPR) - list = BIND_EXPR_BODY (t); + { + list = BIND_EXPR_BODY (t); + if (TREE_CODE (list) != STATEMENT_LIST) + { + stmt = list; + list = alloc_stmt_list (); + BIND_EXPR_BODY (t) = list; + i = tsi_start (list); + tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING); + } + } } if (clauses) @@ -14371,11 +14381,11 @@ finish_oacc_declare (tree fndecl, tree decls) } } - if (!found) - { - i = tsi_start (list); - tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); - } + if (!found) + { + i = tsi_start (list); + tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); + } } while (oacc_returns) @@ -14405,18 +14415,21 @@ finish_oacc_declare (tree fndecl, tree decls) free (r); } - for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i)) + if (ret_clauses) { - if (tsi_end_p (i)) - break; - } + for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i)) + { + if (tsi_end_p (i)) + break; + } - stmt = make_node (OACC_DECLARE); - TREE_TYPE (stmt) = void_type_node; - OMP_STANDALONE_CLAUSES (stmt) = ret_clauses; - SET_EXPR_LOCATION (stmt, loc); + stmt = make_node (OACC_DECLARE); + TREE_TYPE (stmt) = void_type_node; + OMP_STANDALONE_CLAUSES (stmt) = ret_clauses; + SET_EXPR_LOCATION (stmt, loc); - tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); + tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); + } DECL_ATTRIBUTES (fndecl) = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl)); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 78bcb0a1..41fb35e 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -32123,6 +32123,14 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) error = true; continue; } + else if (TREE_PUBLIC (decl)) + { + error_at (loc, + "invalid use of % variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } break; } diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c index ce12463..7979f0c 100644 --- a/gcc/testsuite/c-c++-common/goacc/declare-2.c +++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c @@ -63,4 +63,6 @@ f (void) extern int ve6; #pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */ + +#pragma acc declare present (v9) /* { dg-error "invalid use of" } */ } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c index 59cfe51..584b921 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c @@ -4,6 +4,26 @@ #include #include +#define N 8 + +void +subr1 (int *a) +{ + int f[N]; +#pragma acc declare copy (f) + +#pragma acc parallel copy (a[0:N]) + { + int i; + + for (i = 0; i < N; i++) + { + f[i] = a[i]; + a[i] = f[i] + f[i]; + } + } +} + int b[8]; #pragma acc declare create (b) @@ -13,7 +33,6 @@ int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 }; int main (int argc, char **argv) { - const int N = 8; int a[N]; int e[N]; #pragma acc declare create (e) @@ -61,5 +80,18 @@ main (int argc, char **argv) abort (); } + for (i = 0; i < N; i++) + { + a[i] = 1234; + } + + subr1 (&a[0]); + + for (i = 0; i < N; i++) + { + if (a[i] != 1234 * 2) + abort (); + } + return 0; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c new file mode 100644 index 0000000..8bb7dec --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c @@ -0,0 +1,16 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include +#include +#include + +#define N 8 + +int +main (int argc, char **argv) +{ + int a[N] __attribute__((unused)); +#pragma acc declare present (a) +} + +/* { dg-shouldfail "" } */