Hi!
On 2021-03-11T17:52:55+0100, I wrote:
> On 2021-02-23T22:52:38+0100, Jakub Jelinek via Gcc-patches
> <gcc-patches@gcc.gnu.org> wrote:
>> On Tue, Feb 23, 2021 at 09:43:51PM +0000, Kwok Cheung Yeung wrote:
>>> On 19/02/2021 7:12 pm, Kwok Cheung Yeung wrote:
>>> > I have included the current state of my patch. All task-detach-* tests
>>> > pass when executed without offloading or with offloading to GCN, but
>>> > with offloading to Nvidia, task-detach-6.* hangs consistently but
>>> > everything else passes (probably because of the missing
>>> > gomp_team_barrier_done?).
>>>
>>> It looks like the hang has nothing to do with the detach patch - this hangs
>>> consistently for me when offloaded to NVPTX:
>>>
>>> #include <omp.h>
>>>
>>> int main (void)
>>> {
>>> #pragma omp target
>>> #pragma omp parallel
>>> #pragma omp task
>>> ;
>>> }
>>>
>>> This doesn't hang when offloaded to GCN or the host device, or if
>>> num_threads(1) is specified on the omp parallel.
>
> So, I reproduced this the hard way;
> <https://gcc.gnu.org/bugzilla/show_bug.cgi?id=98738#c13> :-/
>
> Please always file issues when you run into such things. I've now filed
> PR99555 "[OpenMP/nvptx] Execution-time hang for simple nested OpenMP
> 'target'/'parallel'/'task' constructs".
>
>> Then it can be solved separately, I'll try to have a look if I see something
>> bad from the dumps, but I admit I don't have much experience with debugging
>> NVPTX offloaded code...
>
> Any luck?
>
>
> Until this gets resolved properly, OK to push something like the attached
> (currently testing) "Avoid OpenMP/nvptx execution-time hangs for simple
> nested OpenMP 'target'/'parallel'/'task' constructs [PR99555]"?
As posted, I've now pushed "Avoid OpenMP/nvptx execution-time hangs for
simple nested OpenMP 'target'/'parallel'/'task' constructs [PR99555]" to
master branch in commit d99111fd8e12deffdd9a965ce17e8a760d531ec3, see
attached. "... awaiting proper resolution, of course."
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
>From d99111fd8e12deffdd9a965ce17e8a760d531ec3 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Thu, 11 Mar 2021 17:01:22 +0100
Subject: [PATCH] Avoid OpenMP/nvptx execution-time hangs for simple nested
OpenMP 'target'/'parallel'/'task' constructs [PR99555]
... awaiting proper resolution, of course.
libgomp/
PR target/99555
* testsuite/lib/on_device_arch.c: New file.
* testsuite/libgomp.c/pr99555-1.c: Likewise.
* testsuite/libgomp.c-c++-common/task-detach-6.c: Until resolved,
skip for nvptx offloading, with error status.
* testsuite/libgomp.fortran/task-detach-6.f90: Likewise.
---
libgomp/testsuite/lib/on_device_arch.c | 30 +++++++++++++++++++
.../libgomp.c-c++-common/task-detach-6.c | 7 +++++
libgomp/testsuite/libgomp.c/pr99555-1.c | 19 ++++++++++++
.../libgomp.fortran/task-detach-6.f90 | 13 ++++++++
4 files changed, 69 insertions(+)
create mode 100644 libgomp/testsuite/lib/on_device_arch.c
create mode 100644 libgomp/testsuite/libgomp.c/pr99555-1.c
diff --git a/libgomp/testsuite/lib/on_device_arch.c b/libgomp/testsuite/lib/on_device_arch.c
new file mode 100644
index 000000000000..1c0753c31814
--- /dev/null
+++ b/libgomp/testsuite/lib/on_device_arch.c
@@ -0,0 +1,30 @@
+#include <gomp-constants.h>
+
+/* static */ int
+device_arch_nvptx (void)
+{
+ return GOMP_DEVICE_NVIDIA_PTX;
+}
+
+#pragma omp declare variant (device_arch_nvptx) match(construct={target},device={arch(nvptx)})
+/* static */ int
+device_arch (void)
+{
+ return GOMP_DEVICE_DEFAULT;
+}
+
+static int
+on_device_arch (int d)
+{
+ int d_cur;
+ #pragma omp target map(from:d_cur)
+ d_cur = device_arch ();
+
+ return d_cur == d;
+}
+
+int
+on_device_arch_nvptx ()
+{
+ return on_device_arch (GOMP_DEVICE_NVIDIA_PTX);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c
index e5c2291e6ff0..4a3e4a2a3d28 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c
@@ -1,5 +1,8 @@
/* { dg-do run } */
+/* { dg-additional-sources "../lib/on_device_arch.c" } */
+extern int on_device_arch_nvptx ();
+
#include <omp.h>
#include <assert.h>
@@ -9,6 +12,10 @@
int main (void)
{
+ //TODO See '../libgomp.c/pr99555-1.c'.
+ if (on_device_arch_nvptx ())
+ __builtin_abort (); //TODO Until resolved, skip, with error status.
+
int x = 0, y = 0, z = 0;
int thread_count;
omp_event_handle_t detach_event1, detach_event2;
diff --git a/libgomp/testsuite/libgomp.c/pr99555-1.c b/libgomp/testsuite/libgomp.c/pr99555-1.c
new file mode 100644
index 000000000000..9ba330959d80
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr99555-1.c
@@ -0,0 +1,19 @@
+// PR99555 "[OpenMP/nvptx] Execution-time hang for simple nested OpenMP 'target'/'parallel'/'task' constructs"
+
+// { dg-additional-options "-O0" }
+
+// { dg-additional-sources "../lib/on_device_arch.c" }
+extern int on_device_arch_nvptx ();
+
+int main (void)
+{
+ if (on_device_arch_nvptx ())
+ __builtin_abort (); //TODO Until resolved, skip, with error status.
+
+#pragma omp target
+#pragma omp parallel // num_threads(1)
+#pragma omp task
+ ;
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-6.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-6.f90
index b2c476fd6a6b..eda20e73bb84 100644
--- a/libgomp/testsuite/libgomp.fortran/task-detach-6.f90
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-6.f90
@@ -1,5 +1,8 @@
! { dg-do run }
+! { dg-additional-sources ../lib/on_device_arch.c }
+ ! { dg-prune-output "command-line option '-fintrinsic-modules-path=.*' is valid for Fortran but not for C" }
+
! Test tasks with detach clause on an offload device. Each device
! thread spawns off a chain of tasks, that can then be executed by
! any available thread.
@@ -11,6 +14,16 @@ program task_detach_6
integer :: x = 0, y = 0, z = 0
integer :: thread_count
+ interface
+ integer function on_device_arch_nvptx() bind(C)
+ end function on_device_arch_nvptx
+ end interface
+
+ !TODO See '../libgomp.c/pr99555-1.c'.
+ if (on_device_arch_nvptx () /= 0) then
+ error stop !TODO Until resolved, skip, with error status.
+ end if
+
!$omp target map (tofrom: x, y, z) map (from: thread_count)
!$omp parallel private (detach_event1, detach_event2)
!$omp single
--
2.30.2
