From 6e5e2d45f02235a0f72e6130dcd8d52f88f7b126 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <frederik@codesourcery.com>
Date: Fri, 7 Feb 2020 08:03:00 +0100
Subject: [PATCH] xfail and improve some failing libgomp tests

* libgomp.c/{target-32.c,thread-limit-2.c}

Regarding failures because "no usleep implemented for nvptx."
(cf. https://gcc.gnu.org/PR81690):

Create test copies using busy wait instead of usleep, add
xfails for nvptx and amdgcn (introduce new effective target
for the latter) to original tests.

* libgomp.c/target-{33,34}.c

Regarding "no GOMP_OFFLOAD_async_run implemented in
plugin-nvptx.c." (cf. https://gcc.gnu.org/PR81688):

Add xfails for nvptx.

* libgomp.c/target-link-1.c:

Regarding "omp target link not implemented for nvptx."
(cf. https://gcc.gnu.org/PR81689):

Add xfail for nvptx.

libgomp/
	* testsuite/lib/libgomp.exp (proc match_effective_offload_target):
	New proc extracted from check_effective_target_offload_target_nvptx.
	(proc check_effective_target_offload_target_nvptx): Change to use
	match_effective_offload_target.
	(proc check_effective_target_offload_target_amgcn): New proc.
	* testsuite/libgomp.c/target-32-aux.h: New file, extracted from
	target-32.c.
	* testsuite/libgomp.c/target-32-nosleep.c: New test, like target-32.c
	but with busy waiting instead of usleep.
	* testsuite/libgomp.c/target-32.c: Use target-32-aux.h.
	* testsuite/libgomp.c/target-33.c Add xfail for execution on
	offload_target_nvptx.
	* testsuite/libgomp.c/target-34.c: Add xfail for execution on
	offload_target_nvptx.
	* testsuite/libgomp.c/target-link-1.c: Add xfail for
	offload_target_nvptx.
	* testsuite/libgomp.c/thread-limit-2-aux.h: New file, extracted from
	thread-limit-2.c.
	* testsuite/libgomp.c/thread-limit-2-nosleep.c: New test, like
	thread-limit-2.c, but with busy waiting instead of usleep.
	* testsuite/libgomp.c/thread-limit-2.c: Use thread-limit-2-aux.h.
---
 libgomp/testsuite/lib/libgomp.exp             | 17 +++--
 .../testsuite/libgomp.c/target-32-nosleep.c   | 21 +++++++
 libgomp/testsuite/libgomp.c/target-32.c       | 61 ++++--------------
 libgomp/testsuite/libgomp.c/target-33.c       |  3 +
 libgomp/testsuite/libgomp.c/target-34.c       |  3 +
 libgomp/testsuite/libgomp.c/target-link-1.c   |  3 +
 .../libgomp.c/thread-limit-2-nosleep.c        | 22 +++++++
 libgomp/testsuite/libgomp.c/thread-limit-2.c  | 63 ++++---------------
 8 files changed, 89 insertions(+), 104 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.c/target-32-nosleep.c
 create mode 100644 libgomp/testsuite/libgomp.c/thread-limit-2-nosleep.c

diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index e7ce784314d..3f4ced6fe7a 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -338,9 +338,8 @@ proc offload_target_to_openacc_device_type { offload_target } {
 	}
     }
 }
-
-# Return 1 if compiling for offload target nvptx.
-proc check_effective_target_offload_target_nvptx { } {
+# Return 1 if compiling for an offload target matching the given target pattern.
+proc match_effective_offload_target { target } {
     # Consider all actual options, including the flags passed to
     # 'gcc-dg-runtest', or 'gfortran-dg-runtest' (see the 'libgomp.*/*.exp'
     # files; in particular, '-foffload', 'libgomp.oacc-*/*.exp'), which don't
@@ -353,13 +352,23 @@ proc check_effective_target_offload_target_nvptx { } {
     set gcc_output [libgomp_target_compile "" "" "none" $options]
     if [regexp "(?n)^OFFLOAD_TARGET_NAMES=(.*)" $gcc_output dummy offload_targets] {
 	verbose "compiling for offload targets: $offload_targets"
-	return [string match "*:nvptx*:*" ":$offload_targets:"]
+	return [string match "*:$target:*" ":$offload_targets:"]
     }
 
     verbose "not compiling for any offload targets"
     return 0
 }
 
+# Return 1 if compiling for offload target nvptx.
+proc check_effective_target_offload_target_nvptx { } {
+    return [match_effective_offload_target "nvptx*"]
+}
+
+# Return 1 if compiling for offload target amdgcn
+proc check_effective_target_offload_target_amdgcn { } {
+    return [match_effective_offload_target "amdgcn*"]
+}
+
 # Return 1 if offload device is available.
 proc check_effective_target_offload_device { } {
     return [check_runtime_nocache offload_device_available_ {
diff --git a/libgomp/testsuite/libgomp.c/target-32-nosleep.c b/libgomp/testsuite/libgomp.c/target-32-nosleep.c
new file mode 100644
index 00000000000..7534a2fdf44
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-32-nosleep.c
@@ -0,0 +1,21 @@
+/* This is a variation of test-32.c for offloading targets which do not support
+   usleep. */
+/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } }
+   Cf. https://gcc.gnu.org/PR81688.  */
+
+#pragma omp declare target
+
+/* This function serves as a replacement for usleep in
+   this test case. It does not even attempt to be functionally
+   equivalent  - we just want some sort of delay. */
+
+void delay (int d)
+{
+  int i;
+  int N = d * 2000;
+  for (i = 0; i < N; i++)
+    asm volatile ("" : : : "memory");
+}
+#pragma omp end declare target
+
+#include "target-32-aux.h"
diff --git a/libgomp/testsuite/libgomp.c/target-32.c b/libgomp/testsuite/libgomp.c/target-32.c
index 233877b702b..3366d06779d 100644
--- a/libgomp/testsuite/libgomp.c/target-32.c
+++ b/libgomp/testsuite/libgomp.c/target-32.c
@@ -1,54 +1,17 @@
-#include <stdlib.h>
-#include <unistd.h>
-
-int main ()
-{
-  int a = 0, b = 0, c = 0, d[7];
-
-  #pragma omp parallel
-  #pragma omp single
-  {
-    #pragma omp task depend(out: d[0])
-      a = 2;
-
-    #pragma omp target enter data nowait map(to: a,b,c) depend(in: d[0]) depend(out: d[1])
-
-    #pragma omp target nowait map(alloc: a) depend(in: d[1]) depend(out: d[2])
-      a++;
+/* { dg-xfail-if "usleep not implemented" { offload_target_amdgcn || offload_target_nvptx } }
+   Cf. https://gcc.gnu.org/PR81690.  */
+/* { dg-excess-errors "usleep not implemented" { xfail { offload_target_amdgcn || offload_target_nvptx } } } */
 
-    #pragma omp target nowait map(alloc: b) depend(in: d[2]) depend(out: d[3])
-    {
-      usleep (1000);
-      #pragma omp atomic update
-      b |= 4;
-    }
-
-    #pragma omp target nowait map(alloc: b) depend(in: d[2]) depend(out: d[4])
-    {
-      usleep (5000);
-      #pragma omp atomic update
-      b |= 1;
-    }
-
-    #pragma omp target nowait map(alloc: c) depend(in: d[3], d[4]) depend(out: d[5])
-    {
-      usleep (5000);
-      #pragma omp atomic update
-      c |= 8;
-    }
+#include <unistd.h>
 
-    #pragma omp target nowait map(alloc: c) depend(in: d[3], d[4]) depend(out: d[6])
-    {
-      usleep (1000);
-      #pragma omp atomic update
-      c |= 2;
-    }
+/* Use usleep for delays in the test case.
+   See also target-32-nosleep.c.  */
 
-    #pragma omp target exit data map(always,from: a,b,c) depend(in: d[5], d[6])
-  }
+void delay (int microseconds)
+{
+  usleep (microseconds);
+}
 
-  if (a != 3 || b != 5 || c != 10)
-    abort ();
+/* Include the actual test case definition.  */
 
-  return 0;
-}
+#include "target-32-aux.h"
diff --git a/libgomp/testsuite/libgomp.c/target-33.c b/libgomp/testsuite/libgomp.c/target-33.c
index 1bed4b6bc67..15d2d7e38ab 100644
--- a/libgomp/testsuite/libgomp.c/target-33.c
+++ b/libgomp/testsuite/libgomp.c/target-33.c
@@ -1,3 +1,6 @@
+/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } }
+   Cf. https://gcc.gnu.org/PR81688.  */
+
 extern void abort (void);
 
 int
diff --git a/libgomp/testsuite/libgomp.c/target-34.c b/libgomp/testsuite/libgomp.c/target-34.c
index 66d9f54202b..5a3596424d8 100644
--- a/libgomp/testsuite/libgomp.c/target-34.c
+++ b/libgomp/testsuite/libgomp.c/target-34.c
@@ -1,3 +1,6 @@
+/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } }
+   Cf. https://gcc.gnu.org/PR81688.  */
+
 extern void abort (void);
 
 int
diff --git a/libgomp/testsuite/libgomp.c/target-link-1.c b/libgomp/testsuite/libgomp.c/target-link-1.c
index 681677cc2aa..99ce33bc9b4 100644
--- a/libgomp/testsuite/libgomp.c/target-link-1.c
+++ b/libgomp/testsuite/libgomp.c/target-link-1.c
@@ -1,3 +1,6 @@
+/* { dg-xfail-if "#pragma omp target link not implemented" { offload_target_nvptx } }
+   Cf. https://gcc.gnu.org/PR81689.  */
+
 struct S { int s, t; };
 
 int a = 1, b = 1;
diff --git a/libgomp/testsuite/libgomp.c/thread-limit-2-nosleep.c b/libgomp/testsuite/libgomp.c/thread-limit-2-nosleep.c
new file mode 100644
index 00000000000..606db53d701
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/thread-limit-2-nosleep.c
@@ -0,0 +1,22 @@
+/* This is a variation of thread-limit-2.c for offloading targets which do not support
+   usleep.  */
+
+#pragma omp declare target
+
+/* This function serves as a replacement for usleep in
+   this test case. It does not even attempt to be functionally
+   equivalent  - we just want some sort of delay. */
+
+void delay (int d)
+{
+  int i;
+  int N = d * 2000;
+  for (i = 0; i < N; i++)
+    asm volatile ("" : : : "memory");
+}
+#pragma omp end declare target
+
+
+/* Include the actual test case definition.  */
+
+#include "thread-limit-2-aux.h"
diff --git a/libgomp/testsuite/libgomp.c/thread-limit-2.c b/libgomp/testsuite/libgomp.c/thread-limit-2.c
index 1a97fb62985..c64781dbbb4 100644
--- a/libgomp/testsuite/libgomp.c/thread-limit-2.c
+++ b/libgomp/testsuite/libgomp.c/thread-limit-2.c
@@ -1,58 +1,19 @@
 /* { dg-do run } */
 /* { dg-set-target-env-var OMP_THREAD_LIMIT "9" } */
+/* { dg-xfail-if "usleep not implemented" { offload_target_amdgcn || offload_target_nvptx } }
+   Cf. https://gcc.gnu.org/PR81690.  */
+/* { dg-excess-errors "usleep not implemented" { xfail { offload_target_amdgcn || offload_target_nvptx } } } */
 
-#include <stdlib.h>
 #include <unistd.h>
-#include <omp.h>
 
-int
-main ()
+/* Use usleep for delays in the test case.
+   See also thread-limit-2-nosleep.c.  */
+
+void delay (int microseconds)
 {
-  if (omp_get_thread_limit () != 9)
-    return 0;
-  omp_set_dynamic (0);
-  #pragma omp parallel num_threads (8)
-  if (omp_get_num_threads () != 8)
-    abort ();
-  #pragma omp parallel num_threads (16)
-  if (omp_get_num_threads () > 9)
-    abort ();
-  #pragma omp target if (0)
-  #pragma omp teams thread_limit (6)
-  {
-    if (omp_get_thread_limit () > 6)
-      abort ();
-    if (omp_get_thread_limit () == 6)
-      {
-	omp_set_dynamic (0);
-	omp_set_nested (1);
-	#pragma omp parallel num_threads (3)
-	if (omp_get_num_threads () != 3)
-	  abort ();
-	#pragma omp parallel num_threads (3)
-	if (omp_get_num_threads () != 3)
-	  abort ();
-	#pragma omp parallel num_threads (8)
-	if (omp_get_num_threads () > 6)
-	  abort ();
-	#pragma omp parallel num_threads (6)
-	if (omp_get_num_threads () != 6)
-	  abort ();
-	int cnt = 0;
-	#pragma omp parallel num_threads (5)
-	#pragma omp parallel num_threads (5)
-	#pragma omp parallel num_threads (2)
-	{
-	  int v;
-	  #pragma omp atomic capture
-	  v = ++cnt;
-	  if (v > 6)
-	    abort ();
-	  usleep (10000);
-	  #pragma omp atomic
-	  --cnt;
-	}
-      }
-  }
-  return 0;
+  usleep (microseconds);
 }
+
+/* Include the actual test case definition.  */
+
+#include "thread-limit-2-aux.h"
-- 
2.17.1

