[patch, openacc] Adjust tests for amdgcn offloading

[patch, openacc] Adjust tests for amdgcn offloading

[Andrew Stubbs]

[-- Attachment #1: Type: text/plain, Size: 171 bytes --]

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

[-- Attachment #2: 191119-gcn-testsuite.patch --]
[-- Type: text/x-patch, Size: 10797 bytes --]

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 "" }

Re: [patch, openacc] Adjust tests for amdgcn offloading

[Andrew Stubbs]

[-- Attachment #1: Type: text/plain, Size: 558 bytes --]

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

[-- Attachment #2: 191213-update-openacc-tests.patch --]
[-- Type: text/x-patch, Size: 8278 bytes --]

Update OpenACC tests for amdgcn

2019-12-13  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/async_queue-1.c: Disable on GCN.
	* testsuite/libgomp.oacc-c-c++-common/tile-1.c: Likewise.

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..c019fe55c7a 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, so the expected use counts mismatch.
+   { dg-skip-if "unsuitable dimensions" { openacc_amdgcn_accel_selected } { "*" } { "" } } */
+
 /* { dg-additional-options "-fopenacc-dim=32" } */
 
 #include <stdio.h>

Re: [patch, openacc] Adjust tests for amdgcn offloading

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 1230 bytes --]

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

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Don-t-require-openacc_nvidia_accel_selected-in-libgo.patch --]
[-- Type: text/x-diff, Size: 1880 bytes --]

From 89c1a427a1cfdb38e4b2354eeb1e28e0042af54c Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 4 Jun 2021 15:27:55 +0200
Subject: [PATCH] Don't require 'openacc_nvidia_accel_selected' in
 'libgomp.oacc-c-c++-common/async_queue-1.c'

That is, re-enable it for host-fallback, and enable it for GCN offloading.

Fix-up for r279378 (commit 26b74ed0223d108d7d7818c3c860f20cfe81a4af)
"Update OpenACC tests for amdgcn".

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Don't
	require 'openacc_nvidia_accel_selected'.  Fix up for
	'ACC_DEVICE_TYPE_radeon'.
---
 .../testsuite/libgomp.oacc-c-c++-common/async_queue-1.c    | 7 +++++--
 1 file changed, 5 insertions(+), 2 deletions(-)

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 4f9e53da85d..533d498bcf7 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,5 +1,3 @@
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
 /* Test mapping of async values to specific underlying queues.  */
 
 #undef NDEBUG
@@ -29,6 +27,8 @@ int main(void)
   acc_device_t d;
 #if defined ACC_DEVICE_TYPE_nvidia
   d = acc_device_nvidia;
+#elif defined ACC_DEVICE_TYPE_radeon
+  d = acc_device_radeon;
 #elif defined ACC_DEVICE_TYPE_host
   d = acc_device_host;
 #else
@@ -88,6 +88,9 @@ int main(void)
 	assert (queues[i].cuda_stream == NULL);
       else
 	assert (queues[i].cuda_stream != NULL);
+#elif defined ACC_DEVICE_TYPE_radeon
+      /* For "acc_device_radeon" there are no CUDA streams.  */
+      assert (queues[i].cuda_stream == NULL);
 #elif defined ACC_DEVICE_TYPE_host
       /* For "acc_device_host" there are no CUDA streams.  */
       assert (queues[i].cuda_stream == NULL);
-- 
2.30.2

Re: [patch, openacc] Adjust tests for amdgcn offloading

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 2096 bytes --]

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

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Fix-libgomp.oacc-c-c-common-acc_prof-kernels-1.c-for.patch --]
[-- Type: text/x-diff, Size: 1401 bytes --]

From 984df1e1630f262d782c00cefad2643b8e8469f8 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Sun, 6 Jun 2021 10:41:18 +0200
Subject: [PATCH] Fix 'libgomp.oacc-c-c++-common/acc_prof-kernels-1.c' for
 'acc_device_radeon'

... on top of r279378 (commit 26b74ed0223d108d7d7818c3c860f20cfe81a4af)
"Update OpenACC tests for amdgcn".

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Fix
	for 'acc_device_radeon'.
---
 .../testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c | 5 +++++
 1 file changed, 5 insertions(+)

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 ad33f72e2fb..7f74ee922b7 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
@@ -93,6 +93,11 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e
     }
   if (num_workers < 1)
     assert (event_info->launch_event.num_workers >= 1);
+  /* GCN currently enforces 'num_workers (1)'.  */
+  else if (acc_device_type == acc_device_radeon
+	   /*TODO ... just not in the "Parallelism dimensions: variable" case.  */
+	   && /*TODO*/ num_gangs != 22)
+    assert (event_info->launch_event.num_workers == 1);
   else
     {
 #ifdef __OPTIMIZE__
-- 
2.30.2