public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] Fix OpenACC/GCN 'acc_ev_enqueue_launch_end' position
@ 2023-03-10 14:10 Thomas Schwinge
0 siblings, 0 replies; only message in thread
From: Thomas Schwinge @ 2023-03-10 14:10 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:c942b72532238e970ede66ffbc73e0a47edae0e3
commit c942b72532238e970ede66ffbc73e0a47edae0e3
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Thu Mar 2 10:39:09 2023 +0100
Fix OpenACC/GCN 'acc_ev_enqueue_launch_end' position
For an OpenACC compute construct, we've currently got:
- [...]
- acc_ev_enqueue_launch_start
- launch kernel
- free memory
- acc_ev_free
- acc_ev_enqueue_launch_end
This confused another thing that I'm working on, so I adjusted that to:
- [...]
- acc_ev_enqueue_launch_start
- launch kernel
- acc_ev_enqueue_launch_end
- free memory
- acc_ev_free
Correspondingly, verify 'acc_ev_alloc', 'acc_ev_free' in
'libgomp.oacc-c-c++-common/acc_prof-parallel-1.c'.
libgomp/
* plugin/plugin-gcn.c (gcn_exec): Fix 'acc_ev_enqueue_launch_end'
position.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c:
Verify 'acc_ev_alloc', 'acc_ev_free'.
(cherry picked from commit 649f1939baf11f45fd3579b8b9601c7840a097b3)
Diff:
---
libgomp/ChangeLog.omp | 8 +
libgomp/plugin/plugin-gcn.c | 23 +--
.../acc_prof-parallel-1.c | 202 +++++++++++++++++++--
3 files changed, 203 insertions(+), 30 deletions(-)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index be21ec39428..633811af5a4 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,5 +1,13 @@
2023-03-10 Thomas Schwinge <thomas@codesourcery.com>
+ Backported from master:
+ 2023-03-10 Thomas Schwinge <thomas@codesourcery.com>
+
+ * plugin/plugin-gcn.c (gcn_exec): Fix 'acc_ev_enqueue_launch_end'
+ position.
+ * testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c:
+ Verify 'acc_ev_alloc', 'acc_ev_free'.
+
PR other/76739
* libgomp.h (goacc_map_vars): Add 'struct goacc_ncarray_info *'
formal parameter.
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index ca89ba658fd..028fd1c1b3b 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -3347,18 +3347,9 @@ gcn_exec (struct kernel_info *kernel, size_t mapnum, void **hostaddrs,
}
if (!async)
- {
- run_kernel (kernel, ind_da, &kla, NULL, false);
- gomp_offload_free (ind_da);
- }
+ run_kernel (kernel, ind_da, &kla, NULL, false);
else
- {
- queue_push_launch (aq, kernel, ind_da, &kla);
- if (DEBUG_QUEUES)
- GCN_DEBUG ("queue_push_callback %d:%d gomp_offload_free, %p\n",
- aq->agent->device_id, aq->id, ind_da);
- queue_push_callback (aq, gomp_offload_free, ind_da);
- }
+ queue_push_launch (aq, kernel, ind_da, &kla);
if (profiling_dispatch_p)
{
@@ -3368,6 +3359,16 @@ gcn_exec (struct kernel_info *kernel, size_t mapnum, void **hostaddrs,
&enqueue_launch_event_info,
api_info);
}
+
+ if (!async)
+ gomp_offload_free (ind_da);
+ else
+ {
+ if (DEBUG_QUEUES)
+ GCN_DEBUG ("queue_push_callback %d:%d gomp_offload_free, %p\n",
+ aq->agent->device_id, aq->id, ind_da);
+ queue_push_callback (aq, gomp_offload_free, ind_da);
+ }
}
/* }}} */
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 28a47ccc27d..711ead588c4 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
@@ -195,6 +195,139 @@ static void cb_device_init_end (acc_prof_info *prof_info, acc_event_info *event_
#endif
}
+static void cb_alloc (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
+{
+ DEBUG_printf ("%s\n", __FUNCTION__);
+
+#if DEVICE_INIT_INSIDE_COMPUTE_CONSTRUCT
+# error TODO
+#else
+ assert (state == 4
+ || state == 6
+ || state == 104
+ || state == 106);
+ STATE_OP (state, ++);
+
+ if (state == 5
+ || state == 105)
+ {
+ assert (tool_info != NULL);
+ assert (tool_info->event_info.other_event.event_type == acc_ev_compute_construct_start);
+ assert (tool_info->nested != NULL);
+ assert (tool_info->nested->event_info.other_event.event_type == acc_ev_enter_data_start);
+ assert (tool_info->nested->nested == NULL);
+ }
+ else if (state == 7
+ || state == 107)
+ {
+ assert (tool_info != NULL);
+ assert (tool_info->event_info.other_event.event_type == acc_ev_compute_construct_start);
+ assert (tool_info->nested == NULL);
+ }
+ else
+ abort ();
+#endif
+
+ assert (prof_info->event_type == acc_ev_alloc);
+ assert (prof_info->valid_bytes == _ACC_PROF_INFO_VALID_BYTES);
+ assert (prof_info->version == _ACC_PROF_INFO_VERSION);
+ assert (prof_info->device_type == acc_device_type);
+ assert (prof_info->device_number == acc_device_num);
+ assert (prof_info->thread_id == -1);
+ assert (prof_info->async == acc_async);
+ assert (prof_info->async_queue == prof_info->async);
+ assert (prof_info->src_file == NULL);
+ assert (prof_info->func_name == NULL);
+ assert (prof_info->line_no == -1);
+ assert (prof_info->end_line_no == -1);
+ assert (prof_info->func_line_no == -1);
+ assert (prof_info->func_end_line_no == -1);
+
+ assert (event_info->data_event.event_type == prof_info->event_type);
+ assert (event_info->data_event.valid_bytes == _ACC_DATA_EVENT_INFO_VALID_BYTES);
+ assert (event_info->data_event.parent_construct == acc_construct_parallel);
+ assert (event_info->data_event.implicit == 1);
+ assert (event_info->data_event.tool_info == NULL);
+ assert (event_info->data_event.var_name == NULL);
+ assert (event_info->data_event.bytes != 0);
+ assert (event_info->data_event.host_ptr == NULL);
+ assert (event_info->data_event.device_ptr != NULL);
+
+ assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
+ assert (api_info->device_type == prof_info->device_type);
+ assert (api_info->vendor == -1);
+ assert (api_info->device_handle == NULL);
+ assert (api_info->context_handle == NULL);
+ assert (api_info->async_handle == NULL);
+}
+
+static void cb_free (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
+{
+ DEBUG_printf ("%s\n", __FUNCTION__);
+
+#if DEVICE_INIT_INSIDE_COMPUTE_CONSTRUCT
+# error TODO
+#else
+ assert (state == 9
+ || state == 11);
+ STATE_OP (state, ++);
+
+ if (state == 10)
+ {
+ assert (tool_info != NULL);
+ assert (tool_info->event_info.other_event.event_type == acc_ev_compute_construct_start);
+ assert (tool_info->nested == NULL);
+ }
+ else if (state == 12)
+ {
+ assert (tool_info != NULL);
+ assert (tool_info->event_info.other_event.event_type == acc_ev_compute_construct_start);
+ assert (tool_info->nested != NULL);
+ assert (tool_info->nested->event_info.other_event.event_type == acc_ev_exit_data_start);
+ assert (tool_info->nested->nested == NULL);
+ }
+ else
+ abort ();
+#endif
+
+ assert (prof_info->event_type == acc_ev_free);
+ assert (prof_info->valid_bytes == _ACC_PROF_INFO_VALID_BYTES);
+ assert (prof_info->version == _ACC_PROF_INFO_VERSION);
+ assert (prof_info->device_type == acc_device_type);
+ assert (prof_info->device_number == acc_device_num);
+ assert (prof_info->thread_id == -1);
+ assert (prof_info->async == acc_async);
+ assert (prof_info->async_queue == prof_info->async);
+ assert (prof_info->src_file == NULL);
+ assert (prof_info->func_name == NULL);
+ assert (prof_info->line_no == -1);
+ assert (prof_info->end_line_no == -1);
+ assert (prof_info->func_line_no == -1);
+ assert (prof_info->func_end_line_no == -1);
+
+ assert (event_info->data_event.event_type == prof_info->event_type);
+ assert (event_info->data_event.valid_bytes == _ACC_DATA_EVENT_INFO_VALID_BYTES);
+ assert (event_info->data_event.parent_construct == acc_construct_parallel);
+ assert (event_info->data_event.implicit == 1);
+ assert (event_info->data_event.tool_info == NULL);
+ assert (event_info->data_event.var_name == NULL);
+ if (acc_device_type == acc_device_nvidia)
+ assert (event_info->data_event.bytes == (size_t) -1);
+ else if (acc_device_type == acc_device_radeon)
+ assert (event_info->data_event.bytes == 0);
+ else
+ abort ();
+ assert (event_info->data_event.host_ptr == NULL);
+ assert (event_info->data_event.device_ptr != NULL);
+
+ assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
+ assert (api_info->device_type == prof_info->device_type);
+ assert (api_info->vendor == -1);
+ assert (api_info->device_handle == NULL);
+ assert (api_info->context_handle == NULL);
+ assert (api_info->async_handle == NULL);
+}
+
static void cb_enter_data_start (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
{
DEBUG_printf ("%s\n", __FUNCTION__);
@@ -246,8 +379,8 @@ static void cb_enter_data_end (acc_prof_info *prof_info, acc_event_info *event_i
{
DEBUG_printf ("%s\n", __FUNCTION__);
- assert (state == 4
- || state == 104);
+ assert (state == 5
+ || state == 105);
#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.
@@ -316,9 +449,19 @@ static void cb_exit_data_start (acc_prof_info *prof_info, acc_event_info *event_
{
DEBUG_printf ("%s\n", __FUNCTION__);
- assert (state == 7
#if ASYNC_EXIT_DATA
- || state == 107
+ if (acc_async != acc_async_sync)
+ {
+ /* Compensate for the deferred 'acc_ev_free'. */
+ state += 1;
+ }
+#else
+# error TODO
+#endif
+
+ assert (state == 10
+#if ASYNC_EXIT_DATA
+ || state == 110
#endif
);
STATE_OP (state, ++);
@@ -366,15 +509,25 @@ static void cb_exit_data_start (acc_prof_info *prof_info, acc_event_info *event_
tool_info->nested->event_info.other_event.event_type = event_info->other_event.event_type;
event_info->other_event.tool_info = tool_info->nested;
+
+#if ASYNC_EXIT_DATA
+ if (acc_async != acc_async_sync)
+ {
+ /* Compensate for the deferred 'acc_ev_free'. */
+ state += 1;
+ }
+#else
+# error TODO
+#endif
}
static void cb_exit_data_end (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
{
DEBUG_printf ("%s\n", __FUNCTION__);
- assert (state == 8
+ assert (state == 12
#if ASYNC_EXIT_DATA
- || state == 108
+ || state == 112
#endif
);
STATE_OP (state, ++);
@@ -488,6 +641,8 @@ static void cb_compute_construct_start (acc_prof_info *prof_info, acc_event_info
{
/* Compensate for the missing 'acc_ev_enter_data_start'. */
state += 1;
+ /* Compensate for the missing 'acc_ev_alloc'. */
+ state += 1;
}
}
@@ -499,12 +654,19 @@ static void cb_compute_construct_end (acc_prof_info *prof_info, acc_event_info *
{
/* Compensate for the missing 'acc_ev_enter_data_end'. */
state += 1;
+ /* Compensate for the missing 'acc_ev_alloc'. */
+ state += 1;
/* Compensate for the missing 'acc_ev_enqueue_launch_start' and
'acc_ev_enqueue_launch_end'. */
state += 2;
- /* Compensate for the missing 'acc_ev_exit_data_start' and
- 'acc_ev_exit_data_end'. */
- state += 2;
+ /* Compensate for the missing 'acc_ev_free'. */
+ state += 1;
+ /* Compensate for the missing 'acc_ev_exit_data_start'. */
+ state += 1;
+ /* Compensate for the missing 'acc_ev_free'. */
+ state += 1;
+ /* Compensate for the missing 'acc_ev_exit_data_end'. */
+ state += 1;
}
#if !ASYNC_EXIT_DATA
else if (acc_async != acc_async_sync)
@@ -514,8 +676,8 @@ static void cb_compute_construct_end (acc_prof_info *prof_info, acc_event_info *
state += 2;
}
#endif
- assert (state == 9
- || state == 109);
+ assert (state == 13
+ || state == 113);
STATE_OP (state, ++);
assert (tool_info != NULL);
@@ -569,8 +731,8 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e
assert (acc_device_type != acc_device_host);
- assert (state == 5
- || state == 105);
+ assert (state == 7
+ || state == 107);
STATE_OP (state, ++);
assert (tool_info != NULL);
@@ -638,8 +800,8 @@ static void cb_enqueue_launch_end (acc_prof_info *prof_info, acc_event_info *eve
assert (acc_device_type != acc_device_host);
- assert (state == 6
- || state == 106);
+ assert (state == 8
+ || state == 108);
STATE_OP (state, ++);
assert (tool_info != NULL);
@@ -701,6 +863,8 @@ int main()
STATE_OP (state, = 0);
reg (acc_ev_device_init_start, cb_device_init_start, acc_reg);
reg (acc_ev_device_init_end, cb_device_init_end, acc_reg);
+ reg (acc_ev_alloc, cb_alloc, acc_reg);
+ reg (acc_ev_free, cb_free, acc_reg);
reg (acc_ev_enter_data_start, cb_enter_data_start, acc_reg);
reg (acc_ev_enter_data_end, cb_enter_data_end, acc_reg);
reg (acc_ev_exit_data_start, cb_exit_data_start, acc_reg);
@@ -723,9 +887,9 @@ int main()
state_init = state;
}
- assert (state_init == 4);
+ assert (state_init == 5);
}
- assert (state == 10);
+ assert (state == 14);
STATE_OP (state, = 100);
@@ -740,9 +904,9 @@ int main()
}
acc_async = acc_async_sync;
#pragma acc wait
- assert (state_init == 104);
+ assert (state_init == 105);
}
- assert (state == 110);
+ assert (state == 114);
return 0;
}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2023-03-10 14:10 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-03-10 14:10 [gcc/devel/omp/gcc-12] Fix OpenACC/GCN 'acc_ev_enqueue_launch_end' position Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).