Hi!
On 2021-03-25T12:02:15+0100, I wrote:
> 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."
> + if (on_device_arch_nvptx ())
> + __builtin_abort (); //TODO Until resolved, skip, with error status.
Actually, we can do better: do try to execute this trivial OpenMP code
(expected to complete in no time), but for nvptx offloading "make sure
that we exit quickly, with error status", and XFAIL that. So that we'll
get XFAIL -> XPASS when this starts to work for nvptx offloading. Is
that attached "XFAIL OpenMP/nvptx execution-time hangs for simple nested
OpenMP 'target'/'parallel'/'task' constructs [PR99555]" OK to push?
There are other testcases that '#include <unistd.h>' -- do we have to
worry about 'alarm' not being available in some configurations where the
libgomp testsuite executes (and OpenMP 'target' doesn't already fail for
other reasons)?
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 ac247a5962955b20cbf5e4e1a5c4dad81591aeb7 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Wed, 7 Apr 2021 10:36:36 +0200
Subject: [PATCH] XFAIL OpenMP/nvptx execution-time hangs for simple nested
OpenMP 'target'/'parallel'/'task' constructs [PR99555]
... still awaiting proper resolution, of course.
libgomp/
PR target/99555
* testsuite/libgomp.c/pr99555-1.c <nvptx offload device>: Until
resolved, make sure that we exit quickly, with error status,
XFAILed.
* testsuite/libgomp.c-c++-common/task-detach-6.c: Likewise.
* testsuite/libgomp.fortran/task-detach-6.f90: Likewise.
---
libgomp/testsuite/lib/libgomp.exp | 12 ++++++++++++
.../testsuite/libgomp.c-c++-common/task-detach-6.c | 5 ++++-
libgomp/testsuite/libgomp.c/pr99555-1.c | 5 ++++-
libgomp/testsuite/libgomp.fortran/task-detach-6.f90 | 3 ++-
4 files changed, 22 insertions(+), 3 deletions(-)
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index 72d001186a5..14dcfdfd00a 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -401,6 +401,18 @@ proc check_effective_target_offload_device_shared_as { } {
} ]
}
+# Return 1 if using nvptx offload device.
+proc check_effective_target_offload_device_nvptx { } {
+ return [check_runtime_nocache offload_device_nvptx {
+ #include <omp.h>
+ #include "testsuite/libgomp.c-c++-common/on_device_arch.h"
+ int main ()
+ {
+ return !on_device_arch_nvptx ();
+ }
+ } ]
+}
+
# Return 1 if at least one Nvidia GPU is accessible.
proc check_effective_target_openacc_nvidia_accel_present { } {
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 119d7f52f8f..f18b57bf047 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c
@@ -2,6 +2,8 @@
#include <omp.h>
#include <assert.h>
+#include <unistd.h> // For 'alarm'.
+
#include "on_device_arch.h"
/* Test tasks with detach clause on an offload device. Each device
@@ -12,7 +14,8 @@ int main (void)
{
//TODO See '../libgomp.c/pr99555-1.c'.
if (on_device_arch_nvptx ())
- __builtin_abort (); //TODO Until resolved, skip, with error status.
+ alarm (4); /*TODO Until resolved, make sure that we exit quickly, with error status.
+ { dg-xfail-run-if "PR99555" { offload_device_nvptx } } */
int x = 0, y = 0, z = 0;
int thread_count;
diff --git a/libgomp/testsuite/libgomp.c/pr99555-1.c b/libgomp/testsuite/libgomp.c/pr99555-1.c
index 0dc17bfa337..bd33b93716b 100644
--- a/libgomp/testsuite/libgomp.c/pr99555-1.c
+++ b/libgomp/testsuite/libgomp.c/pr99555-1.c
@@ -2,12 +2,15 @@
// { dg-additional-options "-O0" }
+#include <unistd.h> // For 'alarm'.
+
#include "../libgomp.c-c++-common/on_device_arch.h"
int main (void)
{
if (on_device_arch_nvptx ())
- __builtin_abort (); //TODO Until resolved, skip, with error status.
+ alarm (4); /*TODO Until resolved, make sure that we exit quickly, with error status.
+ { dg-xfail-run-if "PR99555" { offload_device_nvptx } } */
#pragma omp target
#pragma omp parallel // num_threads(1)
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-6.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-6.f90
index bd0beb63179..e4373b4c6f1 100644
--- a/libgomp/testsuite/libgomp.fortran/task-detach-6.f90
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-6.f90
@@ -21,7 +21,8 @@ program task_detach_6
!TODO See '../libgomp.c/pr99555-1.c'.
if (on_device_arch_nvptx () /= 0) then
- error stop !TODO Until resolved, skip, with error status.
+ call alarm (4, 0); !TODO Until resolved, make sure that we exit quickly, with error status.
+ ! { dg-xfail-run-if "PR99555" { offload_device_nvptx } }
end if
!$omp target map (tofrom: x, y, z) map (from: thread_count)
--
2.17.1
