From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 62225 invoked by alias); 7 Jan 2016 18:57:41 -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 62206 invoked by uid 89); 7 Jan 2016 18:57:40 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.3 required=5.0 tests=AWL,BAYES_00,KAM_ASCII_DIVIDERS,RCVD_IN_DNSWL_NONE,SPF_PASS autolearn=no version=3.3.2 spammy=func2, UD:openacc.h, openacc.h, openacch 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, 07 Jan 2016 18:57:38 +0000 Received: from svr-orw-fem-03.mgc.mentorg.com ([147.34.97.39]) by relay1.mentorg.com with esmtp id 1aHFkc-0003RK-Fp from James_Norris@mentor.com for gcc-patches@gcc.gnu.org; Thu, 07 Jan 2016 10:57:34 -0800 Received: from [172.30.80.57] (147.34.91.1) by svr-orw-fem-03.mgc.mentorg.com (147.34.97.39) with Microsoft SMTP Server id 14.3.224.2; Thu, 7 Jan 2016 10:57:34 -0800 To: GCC Patches From: James Norris Subject: [gomp4] Fix use of declare'd vars by routine procedures. Message-ID: <568EB51C.9010405@codesourcery.com> Date: Thu, 07 Jan 2016 18:57:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.4.0 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------010400040609000909000909" X-SW-Source: 2016-01/txt/msg00352.txt.bz2 --------------010400040609000909000909 Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit Content-length: 397 Hi! The checking of variables declared by OpenACC declare directives used within an OpenACC routine procedure was not being done correctly. This fix adds the checking required and also adds to the testing. This fix resolves the issue pointed out by Cesar with declare-4.c (https://gcc.gnu.org/ml/gcc-patches/2016-01/msg00339.html). This patch has been applied to gomp-4_0-branch. Thanks, Jim --------------010400040609000909000909 Content-Type: text/x-patch; name="declare.patch" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="declare.patch" Content-length: 6278 Index: gcc/c/ChangeLog.gomp =================================================================== --- gcc/c/ChangeLog.gomp (revision 232138) +++ gcc/c/ChangeLog.gomp (working copy) @@ -1,3 +1,8 @@ +2016-01-07 James Norris + + * c-parser.c (c_finish_oacc_routine): Add new attribute. + * c-typeck.c (build_external_ref): Add usage check. + 2015-12-08 Thomas Schwinge * c-parser.c (c_parser_oacc_clause_bind, c_parser_oacc_routine) Index: gcc/c/c-parser.c =================================================================== --- gcc/c/c-parser.c (revision 232138) +++ gcc/c/c-parser.c (working copy) @@ -14115,6 +14115,10 @@ /* Also add an "omp declare target" attribute, with clauses. */ DECL_ATTRIBUTES (fndecl) = tree_cons (get_identifier ("omp declare target"), clauses, DECL_ATTRIBUTES (fndecl)); + + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("oacc routine"), + clauses, DECL_ATTRIBUTES (fndecl)); } /* OpenACC 2.0: Index: gcc/c/c-typeck.c =================================================================== --- gcc/c/c-typeck.c (revision 232138) +++ gcc/c/c-typeck.c (working copy) @@ -2664,6 +2664,26 @@ tree ref; tree decl = lookup_name (id); + if (decl + && decl != error_mark_node + && current_function_decl + && TREE_CODE (decl) == VAR_DECL + && is_global_var (decl) + && lookup_attribute ("oacc routine", + DECL_ATTRIBUTES (current_function_decl))) + { + if (lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (decl)) + || ((!lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (decl)) + && ((TREE_STATIC (decl) && !DECL_EXTERNAL (decl)) + || (!TREE_STATIC (decl) && DECL_EXTERNAL (decl)))))) + { + error_at (loc, "invalid use in % function"); + return error_mark_node; + } + } + /* In Objective-C, an instance variable (ivar) may be preferred to whatever lookup_name() found. */ decl = objc_lookup_ivar (decl, id); Index: gcc/cp/ChangeLog.gomp =================================================================== --- gcc/cp/ChangeLog.gomp (revision 232138) +++ gcc/cp/ChangeLog.gomp (working copy) @@ -1,3 +1,8 @@ +2016-01-07 James Norris + + * parser.c (cp_finalize_oacc_routine): Add new attribute. + * semantics.c (finish_id_expression): Add usage check. + 2016-01-07 Cesar Philippidis * cp-tree.h (bind_decls_match): Declare. Index: gcc/cp/parser.c =================================================================== --- gcc/cp/parser.c (revision 232138) +++ gcc/cp/parser.c (working copy) @@ -36732,6 +36732,10 @@ DECL_ATTRIBUTES (fndecl) = tree_cons (get_identifier ("omp declare target"), clauses, DECL_ATTRIBUTES (fndecl)); + + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("oacc routine"), + NULL_TREE, DECL_ATTRIBUTES (fndecl)); } } Index: gcc/cp/semantics.c =================================================================== --- gcc/cp/semantics.c (revision 232138) +++ gcc/cp/semantics.c (working copy) @@ -3700,6 +3700,25 @@ decl = convert_from_reference (decl); } + + if (decl != error_mark_node + && current_function_decl + && TREE_CODE (decl) == VAR_DECL + && is_global_var (decl) + && lookup_attribute ("oacc routine", + DECL_ATTRIBUTES (current_function_decl))) + { + if (lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (decl)) + || ((!lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (decl)) + && ((TREE_STATIC (decl) && !DECL_EXTERNAL (decl)) + || (!TREE_STATIC (decl) && DECL_EXTERNAL (decl)))))) + { + *error_msg = "invalid use in % function"; + return error_mark_node; + } + } } return cp_expr (decl, location); Index: gcc/testsuite/ChangeLog.gomp =================================================================== --- gcc/testsuite/ChangeLog.gomp (revision 232138) +++ gcc/testsuite/ChangeLog.gomp (working copy) @@ -1,3 +1,7 @@ +2016-01-07 James Norris + + * c-c++-common/goacc/routine-5.c: Additional tests. + 2016-01-07 Cesar Philippidis * g++.dg/goacc/routine-2.C: Add more coverage. Index: gcc/testsuite/c-c++-common/goacc/routine-5.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/routine-5.c (revision 232138) +++ gcc/testsuite/c-c++-common/goacc/routine-5.c (working copy) @@ -59,3 +59,49 @@ #pragma acc routine (Foo) gang // { dg-error "must be applied before definition" } #pragma acc routine (Baz) // { dg-error "not been declared" } + +float vb1; + +#pragma acc routine +int +func1 (int a) +{ + vb1 = a + 1; /* { dg-error "invalid use in" } */ + + return vb1; /* { dg-error "invalid use in" } */ +} + +#pragma acc routine +int +func2 (int a) +{ + static int vb2; + + vb2 = a + 1; /* { dg-error "invalid use in" } */ + + return vb2; /* { dg-error "invalid use in" } */ +} + +float vb3; +#pragma acc declare link (vb3) + +#pragma acc routine +int +func3 (int a) +{ + vb3 = a + 1; /* { dg-error "invalid use in" } */ + + return vb3; /* { dg-error "invalid use in" } */ +} + +float vb4; +#pragma acc declare create (vb4) + +#pragma acc routine +int +func4 (int a) +{ + vb4 = a + 1; + + return vb4; +} Index: libgomp/ChangeLog.gomp =================================================================== --- libgomp/ChangeLog.gomp (revision 232138) +++ libgomp/ChangeLog.gomp (working copy) @@ -1,3 +1,7 @@ +2016-01-07 James Norris + + * testsuite/libgomp.oacc-c-c++-common/declare-4.c: Fix test. + 2016-01-06 Cesar Philippidis * testsuite/libgomp.oacc-fortran/pr68813.f90: New test. Index: libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c (revision 232138) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c (working copy) @@ -4,7 +4,7 @@ #include float b; -#pragma acc declare link (b) +#pragma acc declare create (b) #pragma acc routine int --------------010400040609000909000909--