From d99111fd8e12deffdd9a965ce17e8a760d531ec3 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@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 ++++++++++++++++++++++
 .../testsuite/libgomp.c-c++-common/task-detach-6.c |  7 +++++
 libgomp/testsuite/libgomp.c/pr99555-1.c            | 19 ++++++++++++++
 .../testsuite/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 0000000..1c0753c
--- /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 e5c2291..4a3e4a2 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 0000000..9ba3309
--- /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 b2c476f..eda20e7 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.7.4