diff mbox series

[og9] A couple of GCN-specific test fixes

Message ID 20190917172156.111727-1-julian@codesourcery.com
State New
Headers show
Series [og9] A couple of GCN-specific test fixes | expand

Commit Message

Julian Brown Sept. 17, 2019, 5:21 p.m. UTC 
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 mbox series

Patch

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  <julian@codesourcery.com>
+
+	* 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  <tobias@codesourcery.com>
 
 	* 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;
 }