From patchwork Tue Nov 19 12:21:59 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Andrew Stubbs X-Patchwork-Id: 1197422 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-514027-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="eB77AZaQ"; 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 47HQ1v21bLz9sPV for ; Tue, 19 Nov 2019 23:22:27 +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:to :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=j9MSoofx6imVcv0GYKexiw9P6qW5t1+cifG38HhWXiYClpgcsp +nNarXWYZva7mPCQ0xXDQXNGqj93YCpl+DcFNBlkjgQwgoXnx6cMtS/idsTySu4o ZvqxFGY/q31Sh7obzEpaGTUtiNr+AQGTgxd+Pmpf9b5ihES4sh6Jge9KM= 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:to :from:subject:message-id:date:mime-version:content-type; s= default; bh=S8QwAQ7pFnghIYBP3eKtLe6BHPo=; b=eB77AZaQMu8y4T19BVOq 9hYB7QKn7JwmmxOirULKdDPoq41vEHTmSG79DVk5Ybhty4PihQzM4trbWrm3qVNA nLbjTHs9XGonutjia0ZIG9qGoNzF8+OWZVlcYas5T89YXwj4o0BYmzwy1+HcXXvZ D74RhDfTM/qQWfFNEfTKa88= Received: (qmail 108671 invoked by alias); 19 Nov 2019 12:22:19 -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 108617 invoked by uid 89); 19 Nov 2019 12:22:19 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-18.4 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.1 spammy=STOP, offloaded, mkoffload X-HELO: esa1.mentor.iphmx.com Received: from esa1.mentor.iphmx.com (HELO esa1.mentor.iphmx.com) (68.232.129.153) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 19 Nov 2019 12:22:16 +0000 IronPort-SDR: toATZw/lOxCu3Rc9taA/qSvhsq66Wr/dFyQ64pBVY/nIXr/EchOQWDlg6kvK7dZvqM/ICRMmOl gTiR8D3IbasM1IyGcN1T8oNCwhH0V7xH6Di7y+rz5z9qdu1CLiLPO4mfYQISAmOaOGcQb947R5 yGQ34dsDatnbWm7IHpx2WSyjgWX509C/hzGeT0eDyrzgnogAoIbpvpOBzH6FIfSBEc4Gd7Ozp3 f7VA3IYdqdnmajOHfZdrQAlNPrY/XRwSQP0jTOgFhysXmhb3kvyskRXWuOlRunfkGqV9Zv/77M p6I= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 19 Nov 2019 04:22:05 -0800 IronPort-SDR: ziMR9jPPliiNJRLrKDYQQhF9fm2IDtMuuP/5cHzEBMatG034xp1veaVu3cUJOD5E80WK678W7G D5WJeNTSYZYXqak/IQPO2dRHuVgUvRVqRXnXM7BcScgeFyVbfQ6FP+JkA0TxROfyqWtC3uC1HH vMpz1Sms77MF6HOLdiNFWmpwqbJDtEaVb83X1YXG3BF2avizZNtuE/ncORrGZX7q4mi3/PTAgE 1+/9EPcQtWXx6Wtq41wY9MBwdYkVsXkk54AIDOxwPbx6uW2ur/iBr9zQsjLfEqVdM6M5/K619L avw= To: "gcc-patches@gcc.gnu.org" , Thomas Schwinge From: Andrew Stubbs Subject: [patch, openacc] Adjust tests for amdgcn offloading Message-ID: Date: Tue, 19 Nov 2019 12:21:59 +0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.1.2 MIME-Version: 1.0 This patch adds GCN special casing for most of the OpenACC libgomp tests that require it. It also disables one testcase that explicitly uses CUDA. OK to commit? Andrew Update OpenACC tests for amdgcn 2019-11-19 Andrew Stubbs libgomp/ * testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c: Handle gcn. * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/tile-1.c: Likewise. * testsuite/libgomp.oacc-fortran/error_stop-1.f: Likewise. * testsuite/libgomp.oacc-fortran/error_stop-2.f: Likewise. * testsuite/libgomp.oacc-fortran/error_stop-3.f: Likewise. * testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Disable on GCN. diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c index b356feb8108..e82a03e8f3c 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c @@ -224,6 +224,8 @@ static void cb_compute_construct_end (acc_prof_info *prof_info, acc_event_info * if (acc_device_type == acc_device_host) assert (api_info->device_api == acc_device_api_none); + else if (acc_device_type == acc_device_gcn) + assert (api_info->device_api == acc_device_api_other); else assert (api_info->device_api == acc_device_api_cuda); assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c index 7cfc364e411..ddf647cda9b 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c @@ -106,6 +106,8 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e assert (event_info->launch_event.vector_length >= 1); else if (acc_device_type == acc_device_nvidia) /* ... is special. */ assert (event_info->launch_event.vector_length == 32); + else if (acc_device_type == acc_device_gcn) /* ...and so is this. */ + assert (event_info->launch_event.vector_length == 64); else { #ifdef __OPTIMIZE__ @@ -118,6 +120,8 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e if (acc_device_type == acc_device_host) assert (api_info->device_api == acc_device_api_none); + else if (acc_device_type == acc_device_gcn) + assert (api_info->device_api == acc_device_api_other); else assert (api_info->device_api == acc_device_api_cuda); assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c index ac6eb48cbbe..dc7c7582ce2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c @@ -265,6 +265,8 @@ static void cb_enter_data_end (acc_prof_info *prof_info, acc_event_info *event_i if (acc_device_type == acc_device_host) assert (api_info->device_api == acc_device_api_none); + else if (acc_device_type == acc_device_gcn) + assert (api_info->device_api == acc_device_api_other); else assert (api_info->device_api == acc_device_api_cuda); assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES); @@ -319,6 +321,8 @@ static void cb_exit_data_start (acc_prof_info *prof_info, acc_event_info *event_ if (acc_device_type == acc_device_host) assert (api_info->device_api == acc_device_api_none); + else if (acc_device_type == acc_device_gcn) + assert (api_info->device_api == acc_device_api_other); else assert (api_info->device_api == acc_device_api_cuda); assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES); @@ -371,6 +375,8 @@ static void cb_exit_data_end (acc_prof_info *prof_info, acc_event_info *event_in if (acc_device_type == acc_device_host) assert (api_info->device_api == acc_device_api_none); + else if (acc_device_type == acc_device_gcn) + assert (api_info->device_api == acc_device_api_other); else assert (api_info->device_api == acc_device_api_cuda); assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES); @@ -510,6 +516,8 @@ static void cb_compute_construct_end (acc_prof_info *prof_info, acc_event_info * if (acc_device_type == acc_device_host) assert (api_info->device_api == acc_device_api_none); + else if (acc_device_type == acc_device_gcn) + assert (api_info->device_api == acc_device_api_other); else assert (api_info->device_api == acc_device_api_cuda); assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES); @@ -573,6 +581,8 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e if (acc_device_type == acc_device_host) assert (api_info->device_api == acc_device_api_none); + else if (acc_device_type == acc_device_gcn) + assert (api_info->device_api == acc_device_api_other); else assert (api_info->device_api == acc_device_api_cuda); assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES); @@ -637,6 +647,8 @@ static void cb_enqueue_launch_end (acc_prof_info *prof_info, acc_event_info *eve if (acc_device_type == acc_device_host) assert (api_info->device_api == acc_device_api_none); + else if (acc_device_type == acc_device_gcn) + assert (api_info->device_api == acc_device_api_other); else assert (api_info->device_api == acc_device_api_cuda); assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES); 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/tile-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c index 5130591dd81..076e3cd75fe 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c @@ -1,3 +1,6 @@ +/* AMD GCN does not use 32-lane vectors. + { dg-skip-if "unsuitable dimensions" { openacc_amdgcn_accel_selected } { "*" } { "" } } */ + /* { dg-additional-options "-fopenacc-dim=32" } */ #include diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f index 4965e674c27..95810a6ae93 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f @@ -15,6 +15,6 @@ ! { dg-output "ERROR STOP (\n|\r\n|\r)+" } ! PR85463. The "minimal" libgfortran implementation used with nvptx ! offloading is a little bit different. -! { dg-output "Error termination.*" { target { ! openacc_nvidia_accel_selected } } } +! { dg-output "Error termination.*" { target { { ! openacc_nvidia_accel_selected } && { ! openacc_amdgcn_accel_selected } } } } ! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } } ! { dg-shouldfail "" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f index 7103fdb5d8e..ce59bbda3c3 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f @@ -15,6 +15,6 @@ ! { dg-output "ERROR STOP 35(\n|\r\n|\r)+" } ! PR85463. The "minimal" libgfortran implementation used with nvptx ! offloading is a little bit different. -! { dg-output "Error termination.*" { target { ! openacc_nvidia_accel_selected } } } +! { dg-output "Error termination.*" { target { { ! openacc_nvidia_accel_selected } && { ! openacc_amdgcn_accel_selected } } } } ! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } } ! { dg-shouldfail "" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f index 9c217f14ea1..9b606c83ad9 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f @@ -15,6 +15,6 @@ ! { dg-output "ERROR STOP SiGN(\n|\r\n|\r)+" } ! PR85463. The "minimal" libgfortran implementation used with nvptx ! offloading is a little bit different. -! { dg-output "Error termination.*" { target { ! openacc_nvidia_accel_selected } } } +! { dg-output "Error termination.*" { target { { ! openacc_nvidia_accel_selected } && { ! openacc_amdgcn_accel_selected } } } } ! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } } ! { dg-shouldfail "" }