From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1643) id C811B3858D1E; Fri, 10 Mar 2023 11:32:06 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org C811B3858D1E DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1678447926; bh=XNhjXKiyGUjmXyk0a8R8yXnABJuAFe8VlTHP1OI9ovs=; h=From:To:Subject:Date:From; b=XSO/elgXXLXpNIEsaU2zjOaz/YGXrRjbTBmu1LHJWK/BJJmh4HPWonrrBwKK1Z2GI njoq3fuqMg+LovDMxZ1EQxPisKcEJQFymYR7FRWJVA+XmtcA6nATelgyi2abLdHeje DU1VOO9AtfbBA6WxgQstbVZWR/5gXLe/Vrx/p1Fg= Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Thomas Schwinge To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] Revert "Revert changes to acc_prof-init-1.c and acc_prof-parallel-1.c" X-Act-Checkin: gcc X-Git-Author: Thomas Schwinge X-Git-Refname: refs/heads/devel/omp/gcc-12 X-Git-Oldrev: 745775eb2c47dad4f2ce30f5559ae39a9b73f184 X-Git-Newrev: 1818bab2ce9f11d8dde5b378f580971b87a5c4ff Message-Id: <20230310113206.C811B3858D1E@sourceware.org> Date: Fri, 10 Mar 2023 11:32:06 +0000 (GMT) List-Id: https://gcc.gnu.org/g:1818bab2ce9f11d8dde5b378f580971b87a5c4ff commit 1818bab2ce9f11d8dde5b378f580971b87a5c4ff Author: Thomas Schwinge 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 + + * 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 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); }