diff mbox series

[openacc] Adjust tests for amdgcn offloading

Message ID fdd1affb-9a4c-d326-8edd-e2bdcb6a0b8b@codesourcery.com
State New
Headers show
Series [openacc] Adjust tests for amdgcn offloading | expand

Commit Message

Andrew Stubbs Nov. 19, 2019, 12:21 p.m. UTC 
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

Comments

Andrew Stubbs Dec. 13, 2019, 5:43 p.m. UTC | #1 
On 19/11/2019 12:21, Andrew Stubbs wrote:
> 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.

The patches aren't all that controversial, should only change the 
results on amdgcn, and Tobias already went and obsoleted half the patch, 
so I've gone ahead and committed the attached.

I've clarified the reason for skipping tile-1.c, and removed the 
obsolete bits, but otherwise it is as it was.

You can remove this patch from your review list. :-)

Andrew
Thomas Schwinge June 8, 2021, 9:28 a.m. UTC | #2 
Hi!

On 2019-12-13T17:43:57+0000, Andrew Stubbs <ams@codesourcery.com> wrote:
> On 19/11/2019 12:21, Andrew Stubbs wrote:
>> 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.
>
> The patches aren't all that controversial, should only change the
> results on amdgcn

Almost.  ;-)

> Update OpenACC tests for amdgcn

>       * testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Disable on GCN.

> --- 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 } } */

Actually that also disables it for 'acc_device_host'.

It's however trivial to make it work for all; pushed "Don't require
'openacc_nvidia_accel_selected' in
'libgomp.oacc-c-c++-common/async_queue-1.c'" to master branch in commit
89c1a427a1cfdb38e4b2354eeb1e28e0042af54c, see attached.


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
Thomas Schwinge June 8, 2021, 9:41 a.m. UTC | #3 
Hi!

On 2019-12-13T17:43:57+0000, Andrew Stubbs <ams@codesourcery.com> wrote:
> On 19/11/2019 12:21, Andrew Stubbs wrote:
>> This patch adds GCN special casing for most of the OpenACC libgomp tests
>> that require it.
>
> [...] I've gone ahead and committed the attached.

>       * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: [Handle gcn.]

> --- 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);

(Someone please scold me for making 'acc_device_api_cuda' the default
'else' case here, without 'if (acc_device_type == acc_device_nvidia)'...)

To make this testcase work in the current GCN setting, I've pushed "Fix
'libgomp.oacc-c-c++-common/acc_prof-kernels-1.c' for 'acc_device_radeon'"
to master branch in commit 984df1e1630f262d782c00cefad2643b8e8469f8, see
attached.  (... to be adjusted again, later...)


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
diff mbox series

Patch

Update OpenACC tests for amdgcn

2019-11-19  Andrew Stubbs  <ams@codesourcery.com>

	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 <stdio.h>
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 "" }