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 --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 3ed90bb38f2aa48adc0c6c3280ccbd8d2c616c7a..d55b050392016db24640aaee7575e13533bccb87 100644 (file)
--- 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 6bbe99df1ff89a5c6f66997148513312b6037d24..a33fac7556cb02d6a33ee8038b4e00e9b2ad42b9 100644 (file)
--- 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 9a542b56fe5232fb165c07921b5b755b0cf2ce9c..663f7f724d5bd082c8342d4f8c807221dd1f5cb0 100644 (file)
--- 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);
   }