From patchwork Thu Feb 8 12:57:22 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Martin Jambor X-Patchwork-Id: 870869 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-472847-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="BLfCD/3Z"; dkim-atps=neutral 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 3zcdWz3TB1z9s7F for ; Thu, 8 Feb 2018 23:57:35 +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:cc:subject:date:message-id:mime-version:content-type; q=dns; s=default; b=Klb1pIbnPMn6f6ueeMDrI8+4+ZXrrtyO1JfexYyQ2RDBOW0RYt 1nQoB/qXvDqdIP+f5c4Hlp0qHB8CenC0fY+wsnhzz2HdRPFT+/+2MuQdFR+tCv3X l3g7jPjPj9F9K0BaxMG1ULO/64Qt0Wi4sXZO8am2/Yc2XgmX1xdOzgWrE= 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:cc:subject:date:message-id:mime-version:content-type; s= default; bh=W+gOR8c77nTGKmkY5IAKdgzVSS0=; b=BLfCD/3ZmCEY04tXWlhZ rJw9KKoMtWbaBadP8Bi7o/yVJPsl6rASjJ5mQSWvB4o7Fn3r0ZxDBSA7nYSMEXVp L2xZ61MNDLeTDkYSm1XMdBBAdGCTGs/58RFGyh2o786ErZhMH2bmPZQ2vwPNIMgY Elfa5JeRmMlowRTQzdy/Q7A= Received: (qmail 93552 invoked by alias); 8 Feb 2018 12:57:28 -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 93543 invoked by uid 89); 8 Feb 2018 12:57:27 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.2 spammy=H*UA:https, Hx-languages-length:1925 X-HELO: mx2.suse.de Received: from mx2.suse.de (HELO mx2.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 08 Feb 2018 12:57:26 +0000 Received: from relay1.suse.de (charybdis-ext.suse.de [195.135.220.254]) by mx2.suse.de (Postfix) with ESMTP id 1073CAC57 for ; Thu, 8 Feb 2018 12:57:23 +0000 (UTC) From: Martin Jambor To: GCC Patches Cc: Subject: [hsa] Set program allocation for static local variables User-Agent: Notmuch/0.25.1 (https://notmuchmail.org) Emacs/25.3.1 (x86_64-suse-linux-gnu) Date: Thu, 08 Feb 2018 13:57:22 +0100 Message-ID: MIME-Version: 1.0 X-IsSubscribed: yes Hi, it has been brought to my attention that libgomp.c/target-28.c testcase fails to finalize because the static variable s has illegal hsa allocation. Fixed by the patch below, which I am about to commit to trunk (and will commit to the gcc-7-branch after testing there). Thanks, Martin 2018-02-08 Martin Jambor * hsa-gen.c (get_symbol_for_decl): Set program allocation for static local variables. libgomp/ * testsuite/libgomp.hsa.c/staticvar.c: New test. Added testcase --- gcc/hsa-gen.c | 10 +++++++--- libgomp/testsuite/libgomp.hsa.c/staticvar.c | 23 +++++++++++++++++++++++ 2 files changed, 30 insertions(+), 3 deletions(-) create mode 100644 libgomp/testsuite/libgomp.hsa.c/staticvar.c diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c index af0b33d658f..55a46b5a16a 100644 --- a/gcc/hsa-gen.c +++ b/gcc/hsa-gen.c @@ -932,9 +932,13 @@ get_symbol_for_decl (tree decl) else if (lookup_attribute ("hsa_group_segment", DECL_ATTRIBUTES (decl))) segment = BRIG_SEGMENT_GROUP; - else if (TREE_STATIC (decl) - || lookup_attribute ("hsa_global_segment", - DECL_ATTRIBUTES (decl))) + else if (TREE_STATIC (decl)) + { + segment = BRIG_SEGMENT_GLOBAL; + allocation = BRIG_ALLOCATION_PROGRAM; + } + else if (lookup_attribute ("hsa_global_segment", + DECL_ATTRIBUTES (decl))) segment = BRIG_SEGMENT_GLOBAL; else segment = BRIG_SEGMENT_PRIVATE; diff --git a/libgomp/testsuite/libgomp.hsa.c/staticvar.c b/libgomp/testsuite/libgomp.hsa.c/staticvar.c new file mode 100644 index 00000000000..6d20c9aa328 --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/staticvar.c @@ -0,0 +1,23 @@ +extern void abort (void); + +#pragma omp declare target +int +foo (void) +{ + static int s; + return ++s; +} +#pragma omp end declare target + +int +main () +{ + int r; + #pragma omp target map(from:r) + { + r = foo (); + } + if (r != 1) + abort (); + return 0; +}