From patchwork Tue Sep 17 17:21:52 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1163526 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-509130-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="XWrLvzJG"; 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 46Xqg06rGFz9s00 for ; Wed, 18 Sep 2019 03:22:20 +1000 (AEST) 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-transfer-encoding:content-type; q=dns; s=default; b=yUg FzL+NinSXCQStRgOuXei5oMvq1x/aPuSi7Ire3A52lXMuCIorSHNWy9vW3P4IThN XTzIzPKOeH5IFdS8E240CGHPqqzY0m1su91KwmTcpTO3UnV+2C2rcjoA5gqi7H8S 8kWTu/qxzX1HErlMKx3GYVMXkA1StXPerkI/y1fE= 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-transfer-encoding:content-type; s=default; bh=W5Jyqvway NHquPKaxpnFTvGl6Zg=; b=XWrLvzJGdFWu6CJl26FiSTTJ1lbbfA/t5f2cq7cWx cO4O1ZnojwRZ0OUBz0R8lp33jg7EdSb+UhFfkUUeRC36rmnN3BFdNvj5ZbobXCGj Ad22QJo2ljsTzcbEMR5UGeNpXstrM2Wwqh0zwZyWrD+rxe1b42KdhFxXTbyUZA+1 7Q= Received: (qmail 117989 invoked by alias); 17 Sep 2019 17:22:13 -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 117981 invoked by uid 89); 17 Sep 2019 17:22:12 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-21.9 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=ham version=3.3.1 spammy=queues, sk:changel, sk:ChangeL, lto1 X-HELO: esa4.mentor.iphmx.com Received: from esa4.mentor.iphmx.com (HELO esa4.mentor.iphmx.com) (68.232.137.252) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 17 Sep 2019 17:22:11 +0000 IronPort-SDR: oThGKjpwVQ/K45ahcaKuDwrMbqIc2xPqW2TDSasaAWFTwbFJtahGbmvQUUOaC0he9zudlpvvqf Avj8aKdqdQRqj8KK5X8CPkXx990AojVPHNRoRn3NuWaNa1Qa16+oiGcklD1lgjtFLW7jFMlSht jRmP6PEUrXLLnq3dn5ku7MYiKprdcnE4NtkefhL3TqkNIUsSGy3yLdULQHk2hrKGj0P7TGdB+/ YJwIYdPmD8qVtY0OTv5T0NoQyqDGU9rjgenuZyrKT/opssHGQUleQU57qnyuNehyEO1oowofPo aR8= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 17 Sep 2019 09:22:09 -0800 IronPort-SDR: e5j8O5joHY3Ky+R/+poYx/OBNXBouoSg3J5H2h0GMYYcqvL9egWbaOh0Qk2V3XfojabVh5Fc1M ewnkVwgsFl0DOfRVat5Q1bnAvIkJ95rg3PZaEZcFO/ZIQK5x/h0pXe6/RcLqjBHEu3jbCdStN6 VP9GPu+C5oz7aro0ATpEff6Q8U8VrGcCxDL5LPBeGXyUI9OQ4Mp4Q1kvH+bcDaBCw76bD1nb7F Hjku/A2FaoPxNvmQZrmQDpWYgpzWG5sSwATMtr/lvprsmVQStkb0UHhj3nGMUaeeijP/E4coLD u88= From: Julian Brown To: CC: Andrew Stubbs , Thomas Schwinge Subject: [PATCH] [og9] A couple of GCN-specific test fixes Date: Tue, 17 Sep 2019 10:21:52 -0700 Message-ID: <20190917172156.111727-1-julian@codesourcery.com> MIME-Version: 1.0 X-IsSubscribed: yes This patch provides a couple of AMD GCN-specific fixes for a few OpenACC offloading tests in libgomp. I will apply to the openacc-gcc-9-branch shortly. Julian ChangeLog libgomp/ * testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Only run NVidia-specific test on NVidia hardware. * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c (main): Initialise for acc_device_gcn if testing on AMD GCN. * testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c: Support AMD GCN. * testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c (check): Skip vector dimension test for AMD GCN. --- libgomp/ChangeLog.openacc | 11 +++++++++++ .../libgomp.oacc-c-c++-common/async_queue-1.c | 2 ++ .../libgomp.oacc-c-c++-common/asyncwait-nop-1.c | 2 ++ .../function-not-offloaded.c | 4 ++-- .../libgomp.oacc-c-c++-common/loop-dim-default.c | 11 ++++++++--- 5 files changed, 25 insertions(+), 5 deletions(-) diff --git a/libgomp/ChangeLog.openacc b/libgomp/ChangeLog.openacc index 14ed4e0ec2c..1a624af1ff9 100644 --- a/libgomp/ChangeLog.openacc +++ b/libgomp/ChangeLog.openacc @@ -1,3 +1,14 @@ +2019-09-17 Julian Brown + + * testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Only run + NVidia-specific test on NVidia hardware. + * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c (main): + Initialise for acc_device_gcn if testing on AMD GCN. + * testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c: Support + AMD GCN. + * testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c (check): Skip + vector dimension test for AMD GCN. + 2019-09-13 Tobias Burnus * plugin/plugin-gcn.c (hsa_warn, hsa_fatal, hsa_error): Ensure diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/async_queue-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/async_queue-1.c index 544b19fe663..4f9e53da85d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/async_queue-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/async_queue-1.c @@ -1,3 +1,5 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + /* Test mapping of async values to specific underlying queues. */ #undef NDEBUG diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c index 4ab67363ba6..840052fec12 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c @@ -26,6 +26,8 @@ main () acc_device_t d; #if defined ACC_DEVICE_TYPE_nvidia d = acc_device_nvidia; +#elif defined ACC_DEVICE_TYPE_gcn + d = acc_device_gcn; #elif defined ACC_DEVICE_TYPE_host d = acc_device_host; #else diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c index fdf4eb08f8a..517004a562d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c @@ -1,11 +1,11 @@ /* { dg-do link } */ -/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target openacc_nvidia_accel_selected } } */ +/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } */ int var; #pragma acc declare create (var) void __attribute__((noinline, noclone)) -foo () /* { dg-error "function 'foo' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target openacc_nvidia_accel_selected } } */ +foo () /* { dg-error "function 'foo' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } */ { var++; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c index dd8107c1acc..5cd0e3122be 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c @@ -80,13 +80,18 @@ int check (const int *ary, int size, int gp, int wp, int vp) exit = 1; } +#ifndef ACC_DEVICE_TYPE_gcn + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ for (ix = 0; ix < vp; ix++) if (vectors[ix] != vectors[0]) { - printf ("vector %d not used %d times\n", ix, vectors[0]); - exit = 1; + printf ("vector %d not used %d times\n", ix, vectors[0]); exit + = 1; } - +#endif + return exit; }