From patchwork Fri Nov 13 20:06:56 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Nathan Sidwell X-Patchwork-Id: 544440 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 3FFC8141453 for ; Sat, 14 Nov 2015 07:07:09 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=MS6/a1BX; 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 :subject:to:references:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=QuH2ZxYeDG64lLtig mX0Ox2SyAOBmOdqgwemMVoVedBhPM4vHA6LSQjVeeziW7Q97ugJG1CQYTWY7Utzj ZdnZh30vuOP0BOUUfpX4BuFDUOgmZ1SrmC6QBd3Q064HBclYOz+1kOh9qx4VBXg+ u4+DGtIYtED7WMRkPUK3X4y44o= 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 :subject:to:references:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=1LSTKa7kSl4NCfk5AhL1xZw kSj0=; b=MS6/a1BX6mROke7o0j6Sri1CPBhftqLHZL0VyHciVY/BSOA5q69ewCc t8P+Jn8FnsJUrVXq7hRVcNN9p2GFCF76OW9pOuBMnJqAuUHaT8t8t6s8m9b+3Li2 l9NfCIGfy+FeN1VsGjnoXqgJdwwXeUCIOfOwNKsb1kMkl6sw9w+8= Received: (qmail 96817 invoked by alias); 13 Nov 2015 20:07:01 -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 96807 invoked by uid 89); 13 Nov 2015 20:07:01 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 required=5.0 tests=BAYES_00, FREEMAIL_FROM, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=no version=3.3.2 X-HELO: mail-yk0-f174.google.com Received: from mail-yk0-f174.google.com (HELO mail-yk0-f174.google.com) (209.85.160.174) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Fri, 13 Nov 2015 20:07:00 +0000 Received: by ykba77 with SMTP id a77so165440166ykb.2 for ; Fri, 13 Nov 2015 12:06:58 -0800 (PST) X-Received: by 10.129.146.7 with SMTP id j7mr24525445ywg.217.1447445217989; Fri, 13 Nov 2015 12:06:57 -0800 (PST) Received: from ?IPv6:2601:181:c000:c497:a2a8:cdff:fe3e:b48? ([2601:181:c000:c497:a2a8:cdff:fe3e:b48]) by smtp.googlemail.com with ESMTPSA id m65sm23276903ywf.56.2015.11.13.12.06.56 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Fri, 13 Nov 2015 12:06:57 -0800 (PST) Subject: Re: [ptx] partitioning optimization To: Bernd Schmidt , GCC Patches References: <564270D6.6090303@acm.org> <56432F50.6000208@redhat.com> <564349B5.5050202@acm.org> <56434E7D.4070803@redhat.com> From: Nathan Sidwell Message-ID: <564642E0.6090303@acm.org> Date: Fri, 13 Nov 2015 15:06:56 -0500 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 In-Reply-To: <56434E7D.4070803@redhat.com> On 11/11/15 09:19, Bernd Schmidt wrote: > On 11/11/2015 02:59 PM, Nathan Sidwell wrote: >> That's not the problem. How to conditionally enable the test is the >> difficulty. I suspect porting something concerning accel_compiler from >> the libgomp testsuite is needed? > > Maybe a check_effective_target_offload_nvptx which tries to see if > -foffload=nvptx gives an error (I would hope it does if it's unsupported). This patch seems to do the trick. tested on an offload-aware build (passes), and a regular x86_64-linux build (skipped as unsupported). ok? nathan 2015-11-13 Nathan Sidwell * lib/target-supports.exp (check_effective_target_offload_nvptx): New. * gcc.dg/goacc/nvptx-merged-loop.c: New. Index: testsuite/gcc.dg/goacc/nvptx-merged-loop.c =================================================================== --- testsuite/gcc.dg/goacc/nvptx-merged-loop.c (revision 0) +++ testsuite/gcc.dg/goacc/nvptx-merged-loop.c (working copy) @@ -0,0 +1,30 @@ +/* { dg-do link } */ +/* { dg-require-effective-target offload_nvptx } */ +/* { dg-options "-fopenacc -O2 -foffload=-fdump-rtl-mach\\ -dumpbase\\ nvptx-merged-loop.c\\ -Wa,--no-verify" } */ + +#define N (32*32*32+17) +void __attribute__ ((noinline)) Foo (int *ary) +{ + int ix; + +#pragma acc parallel num_workers(32) vector_length(32) copyout(ary[0:N]) + { + /* Loop partitioning should be merged. */ +#pragma acc loop worker vector + for (unsigned ix = 0; ix < N; ix++) + { + ary[ix] = ix; + } + } +} + +int main () +{ + int ary[N]; + + Foo (ary); + + return 0; +} + +/* { dg-final { scan-rtl-dump "Merging loop .* into " "mach" } } */ Index: testsuite/lib/target-supports.exp =================================================================== --- testsuite/lib/target-supports.exp (revision 230324) +++ testsuite/lib/target-supports.exp (working copy) @@ -6716,3 +6716,11 @@ proc check_effective_target_vect_max_red } return 0 } + +# Return 1 if there is an nvptx offload compiler. + +proc check_effective_target_offload_nvptx { } { + return [check_no_compiler_messages offload_nvptx object { + int main () {return 0;} + } "-foffload=nvptx-none" ] +}