[PATCH] xfail and improve some failing libgomp tests

[PATCH] xfail and improve some failing libgomp tests

[Harwath, Frederik]

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

Hi,
the libgomp testsuite contains some test cases (all in /libgomp/testsuite/libgomp.c/)
which fail with nvptx offloading because of some long standing issues:

* {target-32.c, thread-limit-2.c}:
no "usleep" implemented for nvptx. Cf. https://gcc.gnu.org/PR81690

* target-{33,34}.c:
no "GOMP_OFFLOAD_async_run" implemented in plugin-nvptx.c. Cf. https://gcc.gnu.org/PR81688

* target-link-1.c:
omp "target link" not implemented for nvptx. Cf. https://gcc.gnu.org/PR81689


All these issues have been known, at least, since 2016:

https://gcc.gnu.org/ml/gcc-patches/2016-11/msg00972.html

As suggested in this mail:
 "Short term, it should be possible to implement something like -foffload=^nvptx
to skip PTX (and only PTX) offloading on those tests."

Well, we can now skip/xfail tests for nvptx offloading using the effective target
"offload_target_nvptx" and the present patch uses this to xfail the tests for which
no short-term solution is in sight, i.e. the GOMP_OFFLOAD_async_run and the "target link"
related failures.

Regarding the "usleep" issue, I have decided to follow Jakub's suggestion
(cf. https://gcc.gnu.org/ml/gcc-patches/2016-11/msg01026.html) to
replace usleep by busy waiting. As noted by Tobias
(https://gcc.gnu.org/bugzilla/show_bug.cgi?id=81690#c4), this involves creating separate
test files for the cases with and without usleep. This solution is a bit cumbersome but I
think we can live with it, in particular, since the actual test case implementations do not
get duplicated (they have been moved into auxiliary header files which are shared by both
variants of the corresponding tests).

Since the "usleep" issue also concerns amdgcn, I have introduced an effective target
"offload_target_amdgcn" to add xfails for this offloading target, too. This behaves like
"offload_target_nvptx" but for amdgcn. Note that the existing amdgcn effective targets
cannot be used for our purpose since they are OpenACC-specific.

The new thread-limit-2-nosleep.c should now pass for both nvptx and amdgcn offloading
whereas thread-limit-2.c should xfail. The new target-32-nosleep.c passes with amdgcn
offloading, but xfails with nvptx offloading, because it also needs the unimplemented
GOMP_OFFLOAD_async_run.

With the patch, the detailed test summary now looks as follows for me:

nvptx offloading:

// Expected execution failures due to missing usleep
PASS: libgomp.c/target-32-nosleep.c (test for excess errors)
XFAIL: libgomp.c/target-32-nosleep.c execution test    // missing GOMP_OFFLOAD_async_run
XFAIL: libgomp.c/target-32.c (test for excess errors)
UNRESOLVED: libgomp.c/target-32.c compilation failed to produce executable

PASS: libgomp.c/thread-limit-2-nosleep.c (test for excess errors)
PASS: libgomp.c/thread-limit-2-nosleep.c execution test
XFAIL: libgomp.c/thread-limit-2.c (test for excess errors)
UNRESOLVED: libgomp.c/thread-limit-2.c compilation failed to produce executable

// Expected execution failures due to missing GOMP_OFFLOAD_async_run
PASS: libgomp.c/target-33.c (test for excess errors)
XFAIL: libgomp.c/target-33.c execution test
PASS: libgomp.c/target-34.c (test for excess errors)
XFAIL: libgomp.c/target-34.c execution test

// Expected compilation failures due to missing target link
XFAIL: libgomp.c/target-link-1.c (test for excess errors)
UNRESOLVED: libgomp.c/target-link-1.c compilation failed to produce executable


amdgcn offloading:

// Tests using usleep
PASS: libgomp.c/target-32-nosleep.c (test for excess errors)
PASS: libgomp.c/target-32-nosleep.c execution test
XFAIL: libgomp.c/target-32.c 7 blank line(s) in output
XFAIL: libgomp.c/target-32.c (test for excess errors)
UNRESOLVED: libgomp.c/target-32.c compilation failed to produce executable

PASS: libgomp.c/thread-limit-2-nosleep.c (test for excess errors)
PASS: libgomp.c/thread-limit-2-nosleep.c execution test
XFAIL: libgomp.c/thread-limit-2.c 1 blank line(s) in output
XFAIL: libgomp.c/thread-limit-2.c (test for excess errors)

// No failures since GOMP_OFFLOAD_async_run works on amdgcn
PASS: libgomp.c/target-33.c (test for excess errors)
PASS: libgomp.c/target-33.c execution test
PASS: libgomp.c/target-34.c (test for excess errors)
PASS: libgomp.c/target-34.c execution test

// No xfail here
PASS: libgomp.c/target-link-1.c (test for excess errors)
FAIL: libgomp.c/target-link-1.c execution test

Note that target-link-1.c execution does also fail on amdgcn.
Since - in contrast to nvptx - it seems that the cause of this failure
has not yet been investigated and discussed, I have not added an xfail
for amdgcn to this test.

All testing has been done with a x86_64-linux-gnu host and target.

Ok to commit this patch?

Best regards,
Frederik






[-- Attachment #2: 0001-xfail-and-improve-some-failing-libgomp-tests.patch --]
[-- Type: text/x-patch, Size: 11617 bytes --]

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

Re: [PATCH] xfail and improve some failing libgomp tests

[Jakub Jelinek]

On Fri, Feb 07, 2020 at 09:56:38AM +0100, Harwath, Frederik wrote:
> * {target-32.c, thread-limit-2.c}:
> no "usleep" implemented for nvptx. Cf. https://gcc.gnu.org/PR81690

Please don't, I want to deal with that using declare variant, just didn't
get yet around to finishing the last patch needed for that.  Will try next week.

> * target-{33,34}.c:
> no "GOMP_OFFLOAD_async_run" implemented in plugin-nvptx.c. Cf. https://gcc.gnu.org/PR81688
> 
> * target-link-1.c:
> omp "target link" not implemented for nvptx. Cf. https://gcc.gnu.org/PR81689

I guess this is ok, though of course the right thing would be to implement
both.  There has been even in some PR a suggestion that instead of failing
in nvptx async_run we should just ignore the nowait clause if the plugin
doesn't implement it properly.

	Jakub

Re: [PATCH] xfail and improve some failing libgomp tests

[Harwath, Frederik]

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

Hi Jakub,

On 07.02.20 16:29, Jakub Jelinek wrote:
> On Fri, Feb 07, 2020 at 09:56:38AM +0100, Harwath, Frederik wrote:
>> * {target-32.c, thread-limit-2.c}:
>> no "usleep" implemented for nvptx. Cf. https://gcc.gnu.org/PR81690
> 
> Please don't, I want to deal with that using declare variant, just didn't
> get yet around to finishing the last patch needed for that.  Will try next week.

Ok, great! looking forward to see a better solution.

>> * target-{33,34}.c:
>> no "GOMP_OFFLOAD_async_run" implemented in plugin-nvptx.c. Cf. https://gcc.gnu.org/PR81688
>>
>> * target-link-1.c:
>> omp "target link" not implemented for nvptx. Cf. https://gcc.gnu.org/PR81689
> 
> I guess this is ok, though of course the right thing would be to implement
> both
Ok, this means that I can commit the attached patch which contains only the changes to
target-{33,43}.c and target-link-1.c? Of course, I agree that those features should be
implemented.

> There has been even in some PR a suggestion that instead of failing
> in nvptx async_run we should just ignore the nowait clause if the plugin
> doesn't implement it properly.

This must be https://gcc.gnu.org/PR93481.

Best regards,
Frederik



[-- Attachment #2: 0001-Add-xfails-to-libgomp-tests-target-33-34-.c-target-l.patch --]
[-- Type: text/x-patch, Size: 2258 bytes --]

From e5165ccb143022614920dbd208f6f368b84b4382 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <frederik@codesourcery.com>
Date: Mon, 10 Feb 2020 08:08:00 +0100
Subject: [PATCH] Add xfails to libgomp tests target-{33,34}.c, target-link-1.c

Add xfails for nvptx offloading because
"no GOMP_OFFLOAD_async_run implemented in plugin-nvptx.c"
(https://gcc.gnu.org/PR81688) and because
"omp target link not implemented for nvptx"
(https://gcc.gnu.org/PR81689).

libgomp/
	* testsuite/libgomp.c/target-33.c: Add xfail for execution on
	offload_target_nvptx, cf. https://gcc.gnu.org/PR81688.
	* testsuite/libgomp.c/target-34.c: Likewise.
	* testsuite/libgomp.c/target-link-1.c: Add xfail for
	offload_target_nvptx, cf. https://gcc.gnu.org/PR81689.
---
 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 +++
 3 files changed, 9 insertions(+)

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

Re: [PATCH] xfail and improve some failing libgomp tests

[Jakub Jelinek]

On Mon, Feb 10, 2020 at 08:49:47AM +0100, Harwath, Frederik wrote:
> Add xfails for nvptx offloading because
> "no GOMP_OFFLOAD_async_run implemented in plugin-nvptx.c"
> (https://gcc.gnu.org/PR81688) and because
> "omp target link not implemented for nvptx"
> (https://gcc.gnu.org/PR81689).
> 
> libgomp/
> 	* testsuite/libgomp.c/target-33.c: Add xfail for execution on
> 	offload_target_nvptx, cf. https://gcc.gnu.org/PR81688.
> 	* testsuite/libgomp.c/target-34.c: Likewise.
> 	* testsuite/libgomp.c/target-link-1.c: Add xfail for
> 	offload_target_nvptx, cf. https://gcc.gnu.org/PR81689.

Ok, thanks.

	Jakub

[PATCH] openmp: ignore nowait if async execution is unsupported [PR93481]

[Harwath, Frederik]

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

Hi Jakub,

On 10.02.20 08:49, Harwath, Frederik wrote:

>> There has been even in some PR a suggestion that instead of failing
>> in nvptx async_run we should just ignore the nowait clause if the plugin
>> doesn't implement it properly.
> 
> This must be https://gcc.gnu.org/PR93481.

The attached patch implements the behavior that has been suggested in the PR.
It makes GOMP_OFFLOAD_async_run optional, removes the stub which produces
the error described in the PR from the nvptx plugin, and changes the nowait-handling
to ignore the clause if GOMP_OFFLOAD_async_run is not available for the executing
device's plugin. I have tested the patch by running the full libgomp testsuite with
nvptx-none offloading on x86_64-linux-gnu. I have observed no regressions.

Ok to push the commit to master?

For the record: Someone should implement GOMP_OFFLOAD_async_run properly
in the nvtpx plugin.

Best regards,
Frederik


[-- Attachment #2: 0001-openmp-ignore-nowait-if-async-execution-is-unsupport.patch --]
[-- Type: text/x-patch, Size: 5881 bytes --]

From 1258f713be317870e9171281e3f7c3a174773aa1 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <frederik@codesourcery.com>
Date: Thu, 13 Feb 2020 07:30:16 +0100
Subject: [PATCH] openmp: ignore nowait if async execution is unsupported
 [PR93481]

An OpenMP "nowait" clause on a target construct currently leads to
a call to GOMP_OFFLOAD_async_run in the plugin that is used for
offloading at execution time. The nvptx plugin contains only a stub
of this function that always produces a fatal error if called.

This commit changes the "nowait" implementation to ignore the clause
if the executing device's plugin does not implement GOMP_OFFLOAD_async_run.
The stub in the nvptx plugin is removed which effectively means that
programs containing "nowait" can now be executed with nvptx offloading
as if the clause had not been used.
This behavior is consistent with the OpenMP specification which says that
"[...] execution of the target task *may* be deferred" (emphasis added),
cf. OpenMP 5.0, page 172.

libgomp/

	* plugin/plugin-nvptx.c: Remove GOMP_OFFLOAD_async_run stub.
	* target.c (gomp_load_plugin_for_device): Make "async_run" loading
	optional.
	(gomp_target_task_fn): Assert "devicep->async_run_func".
	(clear_unsupported_flags): New function to remove unsupported flags
	(right now only GOMP_TARGET_FLAG_NOWAIT) that can be be ignored.
	(GOMP_target_ext): Apply clear_unsupported_flags to flags.
	(GOMP_target_update_ext): Likewise.
	(GOMP_target_enter_exit_data): Likewise.
	* testsuite/libgomp.c/target-33.c:
	Remove xfail for offload_target_nvptx.
	* testsuite/libgomp.c/target-34.c: Likewise.
---
 libgomp/plugin/plugin-nvptx.c           |  7 +------
 libgomp/target.c                        | 19 ++++++++++++++++++-
 libgomp/testsuite/libgomp.c/target-33.c |  3 ---
 libgomp/testsuite/libgomp.c/target-34.c |  3 ---
 4 files changed, 19 insertions(+), 13 deletions(-)

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 6033c71a9db..ec103a2f40b 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1931,9 +1931,4 @@ GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args)
   nvptx_stacks_free (stacks, teams * threads);
 }
 
-void
-GOMP_OFFLOAD_async_run (int ord, void *tgt_fn, void *tgt_vars, void **args,
-			void *async_data)
-{
-  GOMP_PLUGIN_fatal ("GOMP_OFFLOAD_async_run unimplemented");
-}
+/* TODO: Implement GOMP_OFFLOAD_async_run. */
diff --git a/libgomp/target.c b/libgomp/target.c
index 3df007283f4..4fbf963f305 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2022,6 +2022,16 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
   gomp_unmap_vars (tgt_vars, true);
 }
 
+static unsigned int
+clear_unsupported_flags (struct gomp_device_descr *devicep, unsigned int flags)
+{
+  /* If we cannot run asynchronously, simply ignore nowait.  */
+  if (devicep != NULL && devicep->async_run_func == NULL)
+    flags &= ~GOMP_TARGET_FLAG_NOWAIT;
+
+  return flags;
+}
+
 /* Like GOMP_target, but KINDS is 16-bit, UNUSED is no longer present,
    and several arguments have been added:
    FLAGS is a bitmask, see GOMP_TARGET_FLAG_* in gomp-constants.h.
@@ -2054,6 +2064,8 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
   size_t tgt_align = 0, tgt_size = 0;
   bool fpc_done = false;
 
+  flags = clear_unsupported_flags (devicep, flags);
+
   if (flags & GOMP_TARGET_FLAG_NOWAIT)
     {
       struct gomp_thread *thr = gomp_thread ();
@@ -2257,6 +2269,8 @@ GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs,
 {
   struct gomp_device_descr *devicep = resolve_device (device);
 
+  flags = clear_unsupported_flags (devicep, flags);
+
   /* If there are depend clauses, but nowait is not present,
      block the parent task until the dependencies are resolved
      and then just continue with the rest of the function as if it
@@ -2398,6 +2412,8 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
 {
   struct gomp_device_descr *devicep = resolve_device (device);
 
+  flags = clear_unsupported_flags (devicep, flags);
+
   /* If there are depend clauses, but nowait is not present,
      block the parent task until the dependencies are resolved
      and then just continue with the rest of the function as if it
@@ -2524,6 +2540,7 @@ gomp_target_task_fn (void *data)
 	}
       ttask->state = GOMP_TARGET_TASK_READY_TO_RUN;
 
+      assert (devicep->async_run_func);
       devicep->async_run_func (devicep->target_id, fn_addr, actual_arguments,
 			       ttask->args, (void *) ttask);
       return true;
@@ -3040,7 +3057,7 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device,
   if (device->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
     {
       DLSYM (run);
-      DLSYM (async_run);
+      DLSYM_OPT (async_run, async_run);
       DLSYM_OPT (can_run, can_run);
       DLSYM (dev2dev);
     }
diff --git a/libgomp/testsuite/libgomp.c/target-33.c b/libgomp/testsuite/libgomp.c/target-33.c
index 15d2d7e38ab..1bed4b6bc67 100644
--- a/libgomp/testsuite/libgomp.c/target-33.c
+++ b/libgomp/testsuite/libgomp.c/target-33.c
@@ -1,6 +1,3 @@
-/* { 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 5a3596424d8..66d9f54202b 100644
--- a/libgomp/testsuite/libgomp.c/target-34.c
+++ b/libgomp/testsuite/libgomp.c/target-34.c
@@ -1,6 +1,3 @@
-/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } }
-   Cf. https://gcc.gnu.org/PR81688.  */
-
 extern void abort (void);
 
 int
-- 
2.17.1

Re: [PATCH] openmp: ignore nowait if async execution is unsupported [PR93481]

[Jakub Jelinek]

On Thu, Feb 13, 2020 at 09:04:36AM +0100, Harwath, Frederik wrote:
> --- a/libgomp/target.c
> +++ b/libgomp/target.c
> @@ -2022,6 +2022,16 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
>    gomp_unmap_vars (tgt_vars, true);
>  }
>  
> +static unsigned int

Add inline?

> +clear_unsupported_flags (struct gomp_device_descr *devicep, unsigned int flags)
> +{
> +  /* If we cannot run asynchronously, simply ignore nowait.  */
> +  if (devicep != NULL && devicep->async_run_func == NULL)
> +    flags &= ~GOMP_TARGET_FLAG_NOWAIT;
> +
> +  return flags;
> +}
> +
>  /* Like GOMP_target, but KINDS is 16-bit, UNUSED is no longer present,
>     and several arguments have been added:
>     FLAGS is a bitmask, see GOMP_TARGET_FLAG_* in gomp-constants.h.
> @@ -2054,6 +2064,8 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
>    size_t tgt_align = 0, tgt_size = 0;
>    bool fpc_done = false;
>  
> +  flags = clear_unsupported_flags (devicep, flags);
> +
>    if (flags & GOMP_TARGET_FLAG_NOWAIT)
>      {
>        struct gomp_thread *thr = gomp_thread ();

LGTM.

> @@ -2257,6 +2269,8 @@ GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs,
>  {
>    struct gomp_device_descr *devicep = resolve_device (device);
>  
> +  flags = clear_unsupported_flags (devicep, flags);
> +
>    /* If there are depend clauses, but nowait is not present,
>       block the parent task until the dependencies are resolved
>       and then just continue with the rest of the function as if it
> @@ -2398,6 +2412,8 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
>  {
>    struct gomp_device_descr *devicep = resolve_device (device);
>  
> +  flags = clear_unsupported_flags (devicep, flags);
> +
>    /* If there are depend clauses, but nowait is not present,
>       block the parent task until the dependencies are resolved
>       and then just continue with the rest of the function as if it

I don't see why you need to do the above two.  GOMP_TARGET_TASK_DATA
is done on the host side, async_run callback isn't called in that case
and while we create a task, all we do is wait for the (host) dependencies
in there and then perform the data transfer we need.
I think it is perfectly fine to ignore nowait on target but honor it
on target update or target {enter,exit} data.

> @@ -2524,6 +2540,7 @@ gomp_target_task_fn (void *data)
>  	}
>        ttask->state = GOMP_TARGET_TASK_READY_TO_RUN;
>  
> +      assert (devicep->async_run_func);
>        devicep->async_run_func (devicep->target_id, fn_addr, actual_arguments,
>  			       ttask->args, (void *) ttask);
>        return true;
> @@ -3040,7 +3057,7 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device,
>    if (device->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
>      {
>        DLSYM (run);
> -      DLSYM (async_run);
> +      DLSYM_OPT (async_run, async_run);
>        DLSYM_OPT (can_run, can_run);
>        DLSYM (dev2dev);
>      }
> diff --git a/libgomp/testsuite/libgomp.c/target-33.c b/libgomp/testsuite/libgomp.c/target-33.c
> index 15d2d7e38ab..1bed4b6bc67 100644
> --- a/libgomp/testsuite/libgomp.c/target-33.c
> +++ b/libgomp/testsuite/libgomp.c/target-33.c
> @@ -1,6 +1,3 @@
> -/* { 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 5a3596424d8..66d9f54202b 100644
> --- a/libgomp/testsuite/libgomp.c/target-34.c
> +++ b/libgomp/testsuite/libgomp.c/target-34.c
> @@ -1,6 +1,3 @@
> -/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } }
> -   Cf. https://gcc.gnu.org/PR81688.  */
> -
>  extern void abort (void);
>  
>  int
> -- 
> 2.17.1
> 
> 

Otherwise LGTM.

	Jakub

Re: [PATCH] openmp: ignore nowait if async execution is unsupported [PR93481]

[Harwath, Frederik]

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

Hi Jakub,

On 13.02.20 09:30, Jakub Jelinek wrote:
> On Thu, Feb 13, 2020 at 09:04:36AM +0100, Harwath, Frederik wrote:
>> --- a/libgomp/target.c
>> +++ b/libgomp/target.c
>> @@ -2022,6 +2022,16 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
>>    gomp_unmap_vars (tgt_vars, true);
>>  }
>>  
>> +static unsigned int
> 
> Add inline?
> 

Added.

>> @@ -2257,6 +2269,8 @@ GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs,
>>  {
>>    struct gomp_device_descr *devicep = resolve_device (device);
>>  
>> +  flags = clear_unsupported_flags (devicep, flags);

>> @@ -2398,6 +2412,8 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
>>  {
>>    struct gomp_device_descr *devicep = resolve_device (device);
>>  
>> +  flags = clear_unsupported_flags (devicep, flags);

> I don't see why you need to do the above two.  GOMP_TARGET_TASK_DATA
> is done on the host side, async_run callback isn't called in that case
> and while we create a task, all we do is wait for the (host) dependencies
> in there and then perform the data transfer we need.
> I think it is perfectly fine to ignore nowait on target but honor it
> on target update or target {enter,exit} data.

I see. Removed.


> Otherwise LGTM.

Thanks for the review! I have committed the patch with those changes. I forgot to include the ChangeLog
entry which I had to add in a separate commit. Sorry for that! It seems that I have to adapt my workflow -
perhaps some pre-push hook ;-).

Best regards,
Frederik


[-- Attachment #2: 0001-openmp-ignore-nowait-if-async-execution-is-unsupport.patch --]
[-- Type: text/x-patch, Size: 4964 bytes --]

From 001ab12e620c6f117b2e93c77d188bd62fe7ba03 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <frederik@codesourcery.com>
Date: Thu, 13 Feb 2020 07:30:16 +0100
Subject: [PATCH 1/2] openmp: ignore nowait if async execution is unsupported
 [PR93481]

An OpenMP "nowait" clause on a target construct currently leads to
a call to GOMP_OFFLOAD_async_run in the plugin that is used for
offloading at execution time. The nvptx plugin contains only a stub
of this function that always produces a fatal error if called.

This commit changes the "nowait" implementation to ignore the clause
if the executing device's plugin does not implement GOMP_OFFLOAD_async_run.
The stub in the nvptx plugin is removed which effectively means that
programs containing "nowait" can now be executed with nvptx offloading
as if the clause had not been used.
This behavior is consistent with the OpenMP specification which says that
"[...] execution of the target task *may* be deferred" (emphasis added),
cf. OpenMP 5.0, page 172.

libgomp/

	* plugin/plugin-nvptx.c: Remove GOMP_OFFLOAD_async_run stub.
	* target.c (gomp_load_plugin_for_device): Make "async_run" loading
	optional.
	(gomp_target_task_fn): Assert "devicep->async_run_func".
	(clear_unsupported_flags): New function to remove unsupported flags
	(right now only GOMP_TARGET_FLAG_NOWAIT) that can be be ignored.
	(GOMP_target_ext): Apply clear_unsupported_flags to flags.
	* testsuite/libgomp.c/target-33.c:
	Remove xfail for offload_target_nvptx.
	* testsuite/libgomp.c/target-34.c: Likewise.
---
 libgomp/plugin/plugin-nvptx.c           |  7 +------
 libgomp/target.c                        | 15 ++++++++++++++-
 libgomp/testsuite/libgomp.c/target-33.c |  3 ---
 libgomp/testsuite/libgomp.c/target-34.c |  3 ---
 4 files changed, 15 insertions(+), 13 deletions(-)

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 6033c71a9db..ec103a2f40b 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1931,9 +1931,4 @@ GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args)
   nvptx_stacks_free (stacks, teams * threads);
 }
 
-void
-GOMP_OFFLOAD_async_run (int ord, void *tgt_fn, void *tgt_vars, void **args,
-			void *async_data)
-{
-  GOMP_PLUGIN_fatal ("GOMP_OFFLOAD_async_run unimplemented");
-}
+/* TODO: Implement GOMP_OFFLOAD_async_run. */
diff --git a/libgomp/target.c b/libgomp/target.c
index 3df007283f4..0ff727de47d 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2022,6 +2022,16 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
   gomp_unmap_vars (tgt_vars, true);
 }
 
+static inline unsigned int
+clear_unsupported_flags (struct gomp_device_descr *devicep, unsigned int flags)
+{
+  /* If we cannot run asynchronously, simply ignore nowait.  */
+  if (devicep != NULL && devicep->async_run_func == NULL)
+    flags &= ~GOMP_TARGET_FLAG_NOWAIT;
+
+  return flags;
+}
+
 /* Like GOMP_target, but KINDS is 16-bit, UNUSED is no longer present,
    and several arguments have been added:
    FLAGS is a bitmask, see GOMP_TARGET_FLAG_* in gomp-constants.h.
@@ -2054,6 +2064,8 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
   size_t tgt_align = 0, tgt_size = 0;
   bool fpc_done = false;
 
+  flags = clear_unsupported_flags (devicep, flags);
+
   if (flags & GOMP_TARGET_FLAG_NOWAIT)
     {
       struct gomp_thread *thr = gomp_thread ();
@@ -2524,6 +2536,7 @@ gomp_target_task_fn (void *data)
 	}
       ttask->state = GOMP_TARGET_TASK_READY_TO_RUN;
 
+      assert (devicep->async_run_func);
       devicep->async_run_func (devicep->target_id, fn_addr, actual_arguments,
 			       ttask->args, (void *) ttask);
       return true;
@@ -3040,7 +3053,7 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device,
   if (device->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
     {
       DLSYM (run);
-      DLSYM (async_run);
+      DLSYM_OPT (async_run, async_run);
       DLSYM_OPT (can_run, can_run);
       DLSYM (dev2dev);
     }
diff --git a/libgomp/testsuite/libgomp.c/target-33.c b/libgomp/testsuite/libgomp.c/target-33.c
index 15d2d7e38ab..1bed4b6bc67 100644
--- a/libgomp/testsuite/libgomp.c/target-33.c
+++ b/libgomp/testsuite/libgomp.c/target-33.c
@@ -1,6 +1,3 @@
-/* { 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 5a3596424d8..66d9f54202b 100644
--- a/libgomp/testsuite/libgomp.c/target-34.c
+++ b/libgomp/testsuite/libgomp.c/target-34.c
@@ -1,6 +1,3 @@
-/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } }
-   Cf. https://gcc.gnu.org/PR81688.  */
-
 extern void abort (void);
 
 int
-- 
2.17.1


[-- Attachment #3: 0002-Add-ChangeLog-entry-for-my-last-commit.patch --]
[-- Type: text/x-patch, Size: 1166 bytes --]

From 2d9eb4e4ca7be5cbe289eef390afb53f36d313b1 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <frederik@codesourcery.com>
Date: Thu, 13 Feb 2020 10:24:39 +0100
Subject: [PATCH 2/2] Add ChangeLog entry for my last commit

---
 libgomp/ChangeLog | 13 +++++++++++++
 1 file changed, 13 insertions(+)

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index f77ce955ac7..ef1c5dae66c 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,16 @@
+2020-02-13  Frederik Harwath  <frederik@codesourcery.com>
+
+	PR libgomp/93481
+	* plugin/plugin-nvptx.c: Remove GOMP_OFFLOAD_async_run stub.
+	* target.c (gomp_load_plugin_for_device): Make "async_run" loading
+	optional.
+	(gomp_target_task_fn): Assert "devicep->async_run_func".
+	(clear_unsupported_flags): New function to remove unsupported flags
+	(right now only GOMP_TARGET_FLAG_NOWAIT) that can be be ignored.
+	(GOMP_target_ext): Apply clear_unsupported_flags to flags.
+	* testsuite/libgomp.c/target-33.c:
+	Remove xfail for offload_target_nvptx.
+	* testsuite/libgomp.c/target-34.c: Likewise.
 
 2020-02-10  Frederik Harwath  <frederik@codesourcery.com>
 
-- 
2.17.1

Re: [PATCH] xfail and improve some failing libgomp tests

[Tom de Vries]

On 2/7/20 4:29 PM, Jakub Jelinek wrote:
> On Fri, Feb 07, 2020 at 09:56:38AM +0100, Harwath, Frederik wrote:
>> * {target-32.c, thread-limit-2.c}:
>> no "usleep" implemented for nvptx. Cf. https://gcc.gnu.org/PR81690
> 
> Please don't, I want to deal with that using declare variant, just didn't
> get yet around to finishing the last patch needed for that.  Will try next week.
> 

Hi Jakub,

Ping, any update on this?

Thanks,
- Tom

Re: [PATCH] xfail and improve some failing libgomp tests

[Tom de Vries]

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

On 10/5/20 3:15 PM, Tom de Vries wrote:
> On 2/7/20 4:29 PM, Jakub Jelinek wrote:
>> On Fri, Feb 07, 2020 at 09:56:38AM +0100, Harwath, Frederik wrote:
>>> * {target-32.c, thread-limit-2.c}:
>>> no "usleep" implemented for nvptx. Cf. https://gcc.gnu.org/PR81690
>>
>> Please don't, I want to deal with that using declare variant, just didn't
>> get yet around to finishing the last patch needed for that.  Will try next week.
>>
> 
> Hi Jakub,
> 
> Ping, any update on this?

FWIW, I've tried as in patch attached below, but I didn't get it
compiling, I still got:
...
FAIL: libgomp.c/target-32.c (test for excess errors)
Excess errors:
unresolved symbol usleep
...

Jakub, is this already supposed to work?

Thanks,
- Tom

[-- Attachment #2: tmp.patch --]
[-- Type: text/x-patch, Size: 1673 bytes --]

diff --git a/libgomp/testsuite/libgomp.c/target-32.c b/libgomp/testsuite/libgomp.c/target-32.c
index 233877b702b..7ddf8721ed3 100644
--- a/libgomp/testsuite/libgomp.c/target-32.c
+++ b/libgomp/testsuite/libgomp.c/target-32.c
@@ -1,6 +1,26 @@
 #include <stdlib.h>
 #include <unistd.h>
 
+extern void base_delay(int);
+extern void nvptx_delay(int);
+
+#pragma omp declare variant( nvptx_delay ) match( construct={target}, implementation={vendor(nvidia)} )
+void base_delay(int d)
+{
+  usleep (d);
+}
+
+void nvptx_delay(int 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");
+}
+
 int main ()
 {
   int a = 0, b = 0, c = 0, d[7];
@@ -18,28 +38,28 @@ int main ()
 
     #pragma omp target nowait map(alloc: b) depend(in: d[2]) depend(out: d[3])
     {
-      usleep (1000);
+      base_delay (1000);
       #pragma omp atomic update
       b |= 4;
     }
 
     #pragma omp target nowait map(alloc: b) depend(in: d[2]) depend(out: d[4])
     {
-      usleep (5000);
+      base_delay (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);
+      base_delay (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);
+      base_delay (1000);
       #pragma omp atomic update
       c |= 2;
     }

Re: [PATCH] xfail and improve some failing libgomp tests

[Jakub Jelinek]

On Tue, Oct 06, 2020 at 04:48:40PM +0200, Tom de Vries wrote:
> On 10/5/20 3:15 PM, Tom de Vries wrote:
> > On 2/7/20 4:29 PM, Jakub Jelinek wrote:
> >> On Fri, Feb 07, 2020 at 09:56:38AM +0100, Harwath, Frederik wrote:
> >>> * {target-32.c, thread-limit-2.c}:
> >>> no "usleep" implemented for nvptx. Cf. https://gcc.gnu.org/PR81690
> >>
> >> Please don't, I want to deal with that using declare variant, just didn't
> >> get yet around to finishing the last patch needed for that.  Will try next week.
> >>
> > 
> > Hi Jakub,
> > 
> > Ping, any update on this?

Not finished the last step, I run into LTO issues.  Will need to return to
that soon.
Last progress in "[RFH] LTO cgraph support for late declare variant resolution"
mail from May on gcc-patches.

> --- a/libgomp/testsuite/libgomp.c/target-32.c
> +++ b/libgomp/testsuite/libgomp.c/target-32.c
> @@ -1,6 +1,26 @@
>  #include <stdlib.h>
>  #include <unistd.h>
>  
> +extern void base_delay(int);

No need to declare this one early.

> +extern void nvptx_delay(int);

Space before (, and the definition could go here instead of
the declaration.

> +#pragma omp declare variant( nvptx_delay ) match( construct={target}, implementation={vendor(nvidia)} )

This isn't the right declare variant for what we want though,
we only provide gnu as accepted vendor, it is implementation's vendor,
not vendor of one of the hw components.
So, it ought to be instead
#pragma omp declare variant (nvptx_delay) match(construct={target},device={arch(nvptx)})

> +void base_delay(int d)
> +{
> +  usleep (d);
> +}

	Jakub

Re: [PATCH] xfail and improve some failing libgomp tests

[Tom de Vries]

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

On 10/6/20 5:02 PM, Jakub Jelinek wrote:
> On Tue, Oct 06, 2020 at 04:48:40PM +0200, Tom de Vries wrote:
>> On 10/5/20 3:15 PM, Tom de Vries wrote:
>>> On 2/7/20 4:29 PM, Jakub Jelinek wrote:
>>>> On Fri, Feb 07, 2020 at 09:56:38AM +0100, Harwath, Frederik wrote:
>>>>> * {target-32.c, thread-limit-2.c}:
>>>>> no "usleep" implemented for nvptx. Cf. https://gcc.gnu.org/PR81690
>>>>
>>>> Please don't, I want to deal with that using declare variant, just didn't
>>>> get yet around to finishing the last patch needed for that.  Will try next week.
>>>>
>>>
>>> Hi Jakub,
>>>
>>> Ping, any update on this?
> 
> Not finished the last step, I run into LTO issues.  Will need to return to
> that soon.
> Last progress in "[RFH] LTO cgraph support for late declare variant resolution"
> mail from May on gcc-patches.
> 

Ack, thanks for the update.

>> --- a/libgomp/testsuite/libgomp.c/target-32.c
>> +++ b/libgomp/testsuite/libgomp.c/target-32.c
>> @@ -1,6 +1,26 @@
>>  #include <stdlib.h>
>>  #include <unistd.h>
>>  
>> +extern void base_delay(int);
> 
> No need to declare this one early.
> 
>> +extern void nvptx_delay(int);
> 
> Space before (, and the definition could go here instead of
> the declaration.
> 
>> +#pragma omp declare variant( nvptx_delay ) match( construct={target}, implementation={vendor(nvidia)} )
> 
> This isn't the right declare variant for what we want though,
> we only provide gnu as accepted vendor, it is implementation's vendor,
> not vendor of one of the hw components.
> So, it ought to be instead
> #pragma omp declare variant (nvptx_delay) match(construct={target},device={arch(nvptx)})
> 
>> +void base_delay(int d)
>> +{
>> +  usleep (d);
>> +}

I've updated the patch accordingly.

FWIW, I now run into an ICE which looks like PR96680:
...
lto1: internal compiler error: in lto_fixup_prevailing_decls, at
lto/lto-common.c:2595^M
0x93afcd lto_fixup_prevailing_decls^M
        /home/vries/oacc/trunk/source-gcc/gcc/lto/lto-common.c:2595^M
0x93b1d6 lto_fixup_decls^M
        /home/vries/oacc/trunk/source-gcc/gcc/lto/lto-common.c:2645^M
0x93bcc4 read_cgraph_and_symbols(unsigned int, char const**)^M
        /home/vries/oacc/trunk/source-gcc/gcc/lto/lto-common.c:2897^M
0x910358 lto_main()^M
        /home/vries/oacc/trunk/source-gcc/gcc/lto/lto.c:625^M
...

Thanks,
- Tom

[-- Attachment #2: tmp.patch --]
[-- Type: text/x-patch, Size: 1599 bytes --]

diff --git a/libgomp/testsuite/libgomp.c/target-32.c b/libgomp/testsuite/libgomp.c/target-32.c
index 233877b702b..b8deae72b08 100644
--- a/libgomp/testsuite/libgomp.c/target-32.c
+++ b/libgomp/testsuite/libgomp.c/target-32.c
@@ -1,6 +1,25 @@
 #include <stdlib.h>
 #include <unistd.h>
 
+void
+nvptx_delay (int 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");
+}
+
+#pragma omp declare variant (nvptx_delay) match(construct={target},device={arch(nvptx)})
+void
+base_delay(int d)
+{
+  usleep (d);
+}
+
 int main ()
 {
   int a = 0, b = 0, c = 0, d[7];
@@ -18,28 +37,28 @@ int main ()
 
     #pragma omp target nowait map(alloc: b) depend(in: d[2]) depend(out: d[3])
     {
-      usleep (1000);
+      base_delay (1000);
       #pragma omp atomic update
       b |= 4;
     }
 
     #pragma omp target nowait map(alloc: b) depend(in: d[2]) depend(out: d[4])
     {
-      usleep (5000);
+      base_delay (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);
+      base_delay (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);
+      base_delay (1000);
       #pragma omp atomic update
       c |= 2;
     }

Re: [PATCH] xfail and improve some failing libgomp tests

[Jakub Jelinek]

On Tue, Oct 06, 2020 at 05:45:31PM +0200, Tom de Vries wrote:
> I've updated the patch accordingly.
> 
> FWIW, I now run into an ICE which looks like PR96680:

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.

The following should work instead (I've tested it without offloading and
Tobias with offloading):

2020-10-22  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.

--- gcc/libgomp/testsuite/libgomp.c/usleep.h.jj	2020-10-22 14:45:14.034196695 +0200
+++ gcc/libgomp/testsuite/libgomp.c/usleep.h	2020-10-22 14:48:05.186719495 +0200
@@ -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)
--- gcc/libgomp/testsuite/libgomp.c/target-32.c.jj	2020-01-12 11:54:39.037373820 +0100
+++ gcc/libgomp/testsuite/libgomp.c/target-32.c	2020-10-22 14:46:23.211195456 +0200
@@ -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;
     }
--- gcc/libgomp/testsuite/libgomp.c/thread-limit-2.c.jj	2020-01-12 11:54:39.037373820 +0100
+++ gcc/libgomp/testsuite/libgomp.c/thread-limit-2.c	2020-10-22 14:57:31.957516284 +0200
@@ -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;
 	}


	Jakub

Re: [PATCH] xfail and improve some failing libgomp tests

[Tom de Vries]

On 10/22/20 3:19 PM, Jakub Jelinek wrote:
> On Tue, Oct 06, 2020 at 05:45:31PM +0200, Tom de Vries wrote:
>> I've updated the patch accordingly.
>>
>> FWIW, I now run into an ICE which looks like PR96680:
> 
> 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.
> 
> The following should work instead (I've tested it without offloading and
> Tobias with offloading):
> 

I've tested this patch in combination with:
- "[PATCH] lto: LTO cgraph support for late declare variant resolution"
  https://gcc.gnu.org/pipermail/gcc-patches/2020-October/556793.html
- "[omp, simt] Handle alternative IV"
  https://gcc.gnu.org/pipermail/gcc-patches/2020-October/555352.html
on top of commit c26d7df1031 "OpenMP: Fortran - support omp flush's
memorder clauses".

The only FAILs I see are for PR97532 (
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=97532 ), 10 in total.

So, LGTM.

Thanks,
- Tom

> 2020-10-22  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.
> 
> --- gcc/libgomp/testsuite/libgomp.c/usleep.h.jj	2020-10-22 14:45:14.034196695 +0200
> +++ gcc/libgomp/testsuite/libgomp.c/usleep.h	2020-10-22 14:48:05.186719495 +0200
> @@ -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)
> --- gcc/libgomp/testsuite/libgomp.c/target-32.c.jj	2020-01-12 11:54:39.037373820 +0100
> +++ gcc/libgomp/testsuite/libgomp.c/target-32.c	2020-10-22 14:46:23.211195456 +0200
> @@ -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;
>      }
> --- gcc/libgomp/testsuite/libgomp.c/thread-limit-2.c.jj	2020-01-12 11:54:39.037373820 +0100
> +++ gcc/libgomp/testsuite/libgomp.c/thread-limit-2.c	2020-10-22 14:57:31.957516284 +0200
> @@ -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;
>  	}
> 
> 
> 	Jakub
>