xfail and improve some failing libgomp tests [PR81690]

With the patch I've posted today to fix up declare variant LTO handling,
Tobias reported the patch still doesn't work, and there are two
reasons for that.
One is that when the base function is marked implicitly as declare target,
we don't mark also implicitly the variants.  I'll need to ask on omp-lang
about details for that, but generally the compiler should do it some way.
The other one is that the way base_delay is written, it will always
call the usleep function, which is undesirable for nvptx.  While the
compiler will replace all direct calls to base_delay to nvptx_delay,
the base_delay definition which calls usleep stays.

2020-10-28  Jakub Jelinek  <jakub@redhat.com>
	    Tom de Vries  <tdevries@suse.de>

	PR testsuite/81690
	* testsuite/libgomp.c/usleep.h: New file.
	* testsuite/libgomp.c/target-32.c: Include usleep.h.
	(main): Use tgt_usleep instead of usleep.
	* testsuite/libgomp.c/thread-limit-2.c: Include usleep.h.
	(main): Use tgt_usleep instead of usleep.

3 files changed, 31 insertions, 5 deletions

diff --git a/libgomp/testsuite/libgomp.c/target-32.c b/libgomp/testsuite/libgomp.c/target-32.c
index 233877b..1444523 100644
--- a/libgomp/testsuite/libgomp.c/target-32.c
+++ b/libgomp/testsuite/libgomp.c/target-32.c
@@ -1,5 +1,6 @@
 #include <stdlib.h>
 #include <unistd.h>
+#include "usleep.h"
 
 int main ()
 {
@@ -18,28 +19,28 @@ int main ()
 
     #pragma omp target nowait map(alloc: b) depend(in: d[2]) depend(out: d[3])
     {
-      usleep (1000);
+      tgt_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);
+      tgt_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);
+      tgt_usleep (5000);
       #pragma omp atomic update
       c |= 8;
     }
 
     #pragma omp target nowait map(alloc: c) depend(in: d[3], d[4]) depend(out: d[6])
     {
-      usleep (1000);
+      tgt_usleep (1000);
       #pragma omp atomic update
       c |= 2;
     }
diff --git a/libgomp/testsuite/libgomp.c/thread-limit-2.c b/libgomp/testsuite/libgomp.c/thread-limit-2.c
index 1a97fb6..dc247a7 100644
--- a/libgomp/testsuite/libgomp.c/thread-limit-2.c
+++ b/libgomp/testsuite/libgomp.c/thread-limit-2.c
@@ -4,6 +4,7 @@
 #include <stdlib.h>
 #include <unistd.h>
 #include <omp.h>
+#include "usleep.h"
 
 int
 main ()
@@ -48,7 +49,7 @@ main ()
 	  v = ++cnt;
 	  if (v > 6)
 	    abort ();
-	  usleep (10000);
+	  tgt_usleep (10000);
 	  #pragma omp atomic
 	  --cnt;
 	}
diff --git a/libgomp/testsuite/libgomp.c/usleep.h b/libgomp/testsuite/libgomp.c/usleep.h
new file mode 100644
index 0000000..c01aaa0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/usleep.h
@@ -0,0 +1,24 @@
+#include <unistd.h>
+
+int
+nvptx_usleep (useconds_t d)
+{
+  /* 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. */
+  int i;
+  int N = d * 2000;
+  for (i = 0; i < N; i++)
+    asm volatile ("" : : : "memory");
+  return 0;
+}
+
+#pragma omp declare variant (nvptx_usleep) match(construct={target},device={arch(nvptx)})
+#pragma omp declare variant (usleep) match(user={condition(1)})
+int
+tgt_usleep (useconds_t d)
+{
+  return 0;
+}
+
+#pragma omp declare target to (nvptx_usleep, tgt_usleep)