Hi! On 2021-04-09T13:00:39+0200, I wrote: > 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.
Pushed "XFAIL OpenMP/nvptx execution-time hangs for simple nested OpenMP 'target'/'parallel'/'task' constructs [PR99555]" to master branch in commit 4dd9e1c541e0eb921d62c8652c854b1259e56aac, see attached. 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 4dd9e1c541e0eb921d62c8652c854b1259e56aac 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/lib/libgomp.exp (check_effective_target_offload_device_nvptx): New. * 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.30.2