[gcc/devel/omp/gcc-12] Revert "Revert changes to acc_prof-init-1.c and acc_prof-parallel-1.c"

[Thomas Schwinge]

https://gcc.gnu.org/g:1818bab2ce9f11d8dde5b378f580971b87a5c4ff

commit 1818bab2ce9f11d8dde5b378f580971b87a5c4ff
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Thu Mar 2 11:24:28 2023 +0100

    Revert "Revert changes to acc_prof-init-1.c and acc_prof-parallel-1.c"
    
    ... as a prerequisite for reverting
    "OpenACC profiling-interface fixes for asynchronous operations".
    
    This reverts og12 commit b845d2f62e7da1c4cfdfee99690de94b648d076d.
    
            libgomp/
            * testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c: Revert
            "Revert changes to acc_prof-init-1.c and acc_prof-parallel-1.c"
            changes.
            * testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c:
            Likewise.

Diff:
---
 libgomp/ChangeLog.omp                                |  8 ++++++++
 .../libgomp.oacc-c-c++-common/acc_prof-init-1.c      | 17 +++++++++++++++++
 .../libgomp.oacc-c-c++-common/acc_prof-parallel-1.c  | 20 ++++++++++++++++++++
 3 files changed, 45 insertions(+)

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 3ed90bb38f2..d55b0503920 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,11 @@
+2023-03-10  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c: Revert
+	"Revert changes to acc_prof-init-1.c and acc_prof-parallel-1.c"
+	changes.
+	* testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c:
+	Likewise.
+
 2023-03-01  Tobias Burnus  <tobias@codesourcery.com>
 
 	Backported from master:
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c
index 6bbe99df1ff..a33fac7556c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c
@@ -208,6 +208,21 @@ static void cb_compute_construct_end (acc_prof_info *prof_info, acc_event_info *
 
   assert (state == 11
 	  || state == 111);
+#if defined COPYIN
+  /* In an 'async' setting, this event may be triggered before actual 'async'
+     data copying has completed.  Given that 'state' appears in 'COPYIN', we
+     first have to synchronize (that is, let the 'async' 'COPYIN' read the
+     current 'state' value)...  */
+  if (acc_async != acc_async_sync)
+    {
+      /* "We're not yet accounting for the fact that _OpenACC events may occur
+	 during event processing_"; temporarily disable to avoid deadlock.  */
+      unreg (acc_ev_none, NULL, acc_toggle_per_thread);
+      acc_wait (acc_async);
+      reg (acc_ev_none, NULL, acc_toggle_per_thread);
+    }
+  /* ... before modifying it in the following.  */
+#endif
   STATE_OP (state, ++);
 
   assert (tool_info != NULL);
@@ -280,6 +295,7 @@ int main()
     {
       state_init = state;
     }
+    acc_async = acc_async_sync;
 #pragma acc wait
     assert (state_init == 11);
   }
@@ -306,6 +322,7 @@ int main()
     {
       state_init = state;
     }
+    acc_async = acc_async_sync;
 #pragma acc wait
     assert (state_init == 111);
   }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c
index 9a542b56fe5..663f7f724d5 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c
@@ -248,6 +248,25 @@ static void cb_enter_data_end (acc_prof_info *prof_info, acc_event_info *event_i
 
   assert (state == 4
 	  || state == 104);
+#if defined COPYIN
+  /* Conceptually, 'acc_ev_enter_data_end' marks the end of data copying,
+     before 'acc_ev_enqueue_launch_start' marks invoking the compute region.
+     That's the 'state_init = state;' intended to be captured in the compute
+     regions.  */
+  /* In an 'async' setting, this event may be triggered before actual 'async'
+     data copying has completed.  Given that 'state' appears in 'COPYIN', we
+     first have to synchronize (that is, let the 'async' 'COPYIN' read the
+     current 'state' value)...  */
+  if (acc_async != acc_async_sync)
+    {
+      /* "We're not yet accounting for the fact that _OpenACC events may occur
+	 during event processing_"; temporarily disable to avoid deadlock.  */
+      unreg (acc_ev_none, NULL, acc_toggle_per_thread);
+      acc_wait (acc_async);
+      reg (acc_ev_none, NULL, acc_toggle_per_thread);
+    }
+  /* ... before modifying it in the following.  */
+#endif
   STATE_OP (state, ++);
 
   assert (tool_info != NULL);
@@ -679,6 +698,7 @@ int main()
 
       state_init = state;
     }
+    acc_async = acc_async_sync;
 #pragma acc wait
     assert (state_init == 104);
   }