From patchwork Fri Feb 26 15:59:42 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Martin Jambor X-Patchwork-Id: 589166 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 BEE061402D6 for ; Sat, 27 Feb 2016 02:59:56 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=v37IwI9V; 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:date :from:to:subject:message-id:mime-version:content-type; q=dns; s= default; b=WRIT2i2d5gFeMih8KhdzrGJTdlPDP9c1lUasONZpILe/sdEx8PvOj nH3HNi4N6kkdD51Qzn1Eje2l1Lpu7zGWyGaMa7aiWFZc8Hp1HBtCRmPVrOC847VU SuKnJPDD5F4dgo4bWPemlnbYMLg7bF34x+j681QPnDJshDiG8mLel0= 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:date :from:to:subject:message-id:mime-version:content-type; s= default; bh=afJGv+UZwhBCD+wdmuuC4vH/98A=; b=v37IwI9V/2+k9F+yrk06 D5rXJVJzeA6XI14S2tTMuRMzihXTJrXRmAtKftNXC9x3vcjmeSN0QvrLmejla45d yd1QqmtgX1SFz+hPX5cLXGPeisGYZDC1BRqhA56/asSSkJGN8GasNZklhz8dtmrz p+/Fqt4clEEV3lKYufQzy20= Received: (qmail 3209 invoked by alias); 26 Feb 2016 15:59:47 -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 2547 invoked by uid 89); 26 Feb 2016 15:59:46 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=vary 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 (CAMELLIA256-SHA encrypted) ESMTPS; Fri, 26 Feb 2016 15:59:45 +0000 Received: from relay1.suse.de (charybdis-ext.suse.de [195.135.220.254]) by mx2.suse.de (Postfix) with ESMTP id 0DE20ACDF for ; Fri, 26 Feb 2016 15:59:42 +0000 (UTC) Date: Fri, 26 Feb 2016 16:59:42 +0100 From: Martin Jambor To: GCC Patches Subject: [hsa, testsuite] Introduce offload_device_shared_as effective target Message-ID: <20160226155941.GI3094@virgil.suse.cz> Mail-Followup-To: GCC Patches MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.5.24 (2015-08-30) X-IsSubscribed: yes Hello, this patch has been written by Keith McDaniel when he was working at AMD (and so it should be covered by their blanket copyright assignment) and adds a libgomp testsuite effective target predicate offload_device_shared_as that allows us to run tests only when target constructs are run on a device with shared memory (this includes the host). Keith included a C++ test to illustrate the use. I have tested this thoroughly, both with and without HSA (enabled or present). OK for trunk? Thanks, Martin 2016-02-10 Keith McDaniel Martin Jambor * testsuite/lib/libgomp.exp (check_effective_target_offload_device_shared_as): New proc. * testsuite/libgomp.c++/declare_target-1.C: New test. --- libgomp/testsuite/lib/libgomp.exp | 13 ++++++++ libgomp/testsuite/libgomp.c++/declare_target-1.C | 38 ++++++++++++++++++++++++ 2 files changed, 51 insertions(+) create mode 100644 libgomp/testsuite/libgomp.c++/declare_target-1.C diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index a4c9d83..154a447 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -343,6 +343,19 @@ proc check_effective_target_offload_device_nonshared_as { } { } } ] } + +# Return 1 if offload device is available and it has shared address space. +proc check_effective_target_offload_device_shared_as { } { + return [check_runtime_nocache offload_device_shared_as { + int main () + { + int x = 10; + #pragma omp target map(to: x) + x++; + return x == 10; + } + } ] +} # Return 1 if at least one nvidia board is present. diff --git a/libgomp/testsuite/libgomp.c++/declare_target-1.C b/libgomp/testsuite/libgomp.c++/declare_target-1.C new file mode 100644 index 0000000..4394bb1 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare_target-1.C @@ -0,0 +1,38 @@ +// { dg-do run } +// { dg-require-effective-target offload_device_shared_as } + +#include + +struct typeX +{ + int a; +}; + +class typeY +{ +public: + int foo () { return a^0x01; } + int a; +}; + +#pragma omp declare target +struct typeX varX; +class typeY varY; +#pragma omp end declare target + +int main () +{ + varX.a = 0; + varY.a = 0; + + #pragma omp target + { + varX.a = 100; + varY.a = 100; + } + + if (varX.a != 100 || varY.a != 100) + abort (); + + return 0; +}