[PATCH] [og9] Use more appropriate var in localize_reductions call

[PATCH] [og9] Use more appropriate var in localize_reductions call

[Julian Brown]

This patch uses a more appropriate local variable in the call to
localize_reductions in gimplify_omp_for, for better self-documentation.

Tested with offloading to AMD GCN. I will apply to the
openacc-gcc-9-branch shortly.

Julian

ChangeLog

	gcc/
	* gimplify.c (gimplify_omp_for): Use for_stmt in call to
	localize_reductions.
---
 gcc/ChangeLog.openacc | 5 +++++
 gcc/gimplify.c        | 3 ++-
 2 files changed, 7 insertions(+), 1 deletion(-)

diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc
index ffe19bc5809..02667a2aed4 100644
--- a/gcc/ChangeLog.openacc
+++ b/gcc/ChangeLog.openacc
@@ -1,3 +1,8 @@
+2019-09-06  Julian Brown  <julian@codesourcery.com>
+
+	* gimplify.c (gimplify_omp_for): Use for_stmt in call to
+	localize_reductions.
+
 2019-09-06  Julian Brown  <julian@codesourcery.com>
 
 	* config/nvptx/nvptx.c (omp-sese.h): Include.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 3d869447d70..60761504a5e 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -11082,7 +11082,8 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	 here, because the code to remove reductions in kernels regions cannot
 	 handle that.  */
       if (outer && outer->region_type == ORT_ACC_PARALLEL)
-	localize_reductions (OMP_FOR_CLAUSES (*expr_p), OMP_FOR_BODY (*expr_p));
+	localize_reductions (OMP_FOR_CLAUSES (for_stmt),
+			     OMP_FOR_BODY (for_stmt));
     }
 
   /* Set OMP_CLAUSE_LINEAR_NO_COPYIN flag on explicit linear
-- 
2.22.0

[PATCH] [og9] OpenACC profiling support for AMD GCN

[Julian Brown]

This patch adds profiling support to the AMD GCN libgomp plugin, modeled
after the equivalent support in the NVPTX plugin. This gives a positive
test delta in AMD GCN offload testing.

I will apply to the openacc-gcc-9-branch shortly.

Julian

2019-09-06  Julian Brown  <julian@codesourcery.com>

	libgomp/
	* plugin/plugin-gcn.c (GOMP_OFFLOAD_alloc_by_agent,
	GOMP_OFFLOAD_free, gcn_exec): Add profiling support.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c: Add GCN
	support.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c: Likewise.
---
 libgomp/ChangeLog.openacc                     |  9 ++
 libgomp/plugin/plugin-gcn.c                   | 96 +++++++++++++++++++
 .../acc_prof-init-1.c                         |  2 +
 .../acc_prof-kernels-1.c                      |  4 +
 .../acc_prof-parallel-1.c                     | 12 +++
 5 files changed, 123 insertions(+)

diff --git a/libgomp/ChangeLog.openacc b/libgomp/ChangeLog.openacc
index d7a4c7a5f8a..8ed0a10a589 100644
--- a/libgomp/ChangeLog.openacc
+++ b/libgomp/ChangeLog.openacc
@@ -1,3 +1,12 @@
+2019-09-06  Julian Brown  <julian@codesourcery.com>
+
+	* plugin/plugin-gcn.c (GOMP_OFFLOAD_alloc_by_agent,
+	GOMP_OFFLOAD_free, gcn_exec): Add profiling support.
+	* testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c: Add GCN
+	support.
+	* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Likewise.
+	* testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c: Likewise.
+
 2019-09-06  Julian Brown  <julian@codesourcery.com>
 
 	* config/gcn/target.c (omp_pause_resource, omp_pause_resource_all): New
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index f0b22ebc3d7..2f273967bad 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -3024,6 +3024,35 @@ GOMP_OFFLOAD_alloc_by_agent (struct agent_info *agent, size_t size)
       return NULL;
     }
 
+  struct goacc_thread *thr = GOMP_PLUGIN_goacc_thread ();
+  bool profiling_dispatch_p
+    = __builtin_expect (thr != NULL && thr->prof_info != NULL, false);
+  if (profiling_dispatch_p)
+    {
+      acc_prof_info *prof_info = thr->prof_info;
+      acc_event_info data_event_info;
+      acc_api_info *api_info = thr->api_info;
+
+      prof_info->event_type = acc_ev_alloc;
+
+      data_event_info.data_event.event_type = prof_info->event_type;
+      data_event_info.data_event.valid_bytes
+	= _ACC_DATA_EVENT_INFO_VALID_BYTES;
+      data_event_info.data_event.parent_construct
+	= acc_construct_parallel;
+      data_event_info.data_event.implicit = 1;
+      data_event_info.data_event.tool_info = NULL;
+      data_event_info.data_event.var_name = NULL;
+      data_event_info.data_event.bytes = size;
+      data_event_info.data_event.host_ptr = NULL;
+      data_event_info.data_event.device_ptr = (void *) ptr;
+
+      api_info->device_api = acc_device_api_other;
+
+      GOMP_PLUGIN_goacc_profiling_dispatch (prof_info, &data_event_info,
+					    api_info);
+    }
+
   return ptr;
 }
 
@@ -3046,6 +3075,35 @@ GOMP_OFFLOAD_free (int device, void *ptr)
       return false;
     }
 
+  struct goacc_thread *thr = GOMP_PLUGIN_goacc_thread ();
+  bool profiling_dispatch_p
+    = __builtin_expect (thr != NULL && thr->prof_info != NULL, false);
+  if (profiling_dispatch_p)
+    {
+      acc_prof_info *prof_info = thr->prof_info;
+      acc_event_info data_event_info;
+      acc_api_info *api_info = thr->api_info;
+
+      prof_info->event_type = acc_ev_free;
+
+      data_event_info.data_event.event_type = prof_info->event_type;
+      data_event_info.data_event.valid_bytes
+	= _ACC_DATA_EVENT_INFO_VALID_BYTES;
+      data_event_info.data_event.parent_construct
+	= acc_construct_parallel;
+      data_event_info.data_event.implicit = 1;
+      data_event_info.data_event.tool_info = NULL;
+      data_event_info.data_event.var_name = NULL;
+      data_event_info.data_event.bytes = 0;
+      data_event_info.data_event.host_ptr = NULL;
+      data_event_info.data_event.device_ptr = (void *) ptr;
+
+      api_info->device_api = acc_device_api_other;
+
+      GOMP_PLUGIN_goacc_profiling_dispatch (prof_info, &data_event_info,
+					    api_info);
+    }
+
   return true;
 }
 
@@ -3276,6 +3334,35 @@ gcn_exec (struct kernel_info *kernel, size_t mapnum, void **hostaddrs,
      {1,       64, 16}
     };
 
+  struct goacc_thread *thr = GOMP_PLUGIN_goacc_thread ();
+  acc_prof_info *prof_info = thr->prof_info;
+  acc_event_info enqueue_launch_event_info;
+  acc_api_info *api_info = thr->api_info;
+  bool profiling_dispatch_p = __builtin_expect (prof_info != NULL, false);
+  if (profiling_dispatch_p)
+    {
+      prof_info->event_type = acc_ev_enqueue_launch_start;
+
+      enqueue_launch_event_info.launch_event.event_type
+	= prof_info->event_type;
+      enqueue_launch_event_info.launch_event.valid_bytes
+	= _ACC_LAUNCH_EVENT_INFO_VALID_BYTES;
+      enqueue_launch_event_info.launch_event.parent_construct
+	= acc_construct_parallel;
+      enqueue_launch_event_info.launch_event.implicit = 1;
+      enqueue_launch_event_info.launch_event.tool_info = NULL;
+      enqueue_launch_event_info.launch_event.kernel_name
+	= (char *) kernel->name;
+      enqueue_launch_event_info.launch_event.num_gangs = kla.gdims[0];
+      enqueue_launch_event_info.launch_event.num_workers = kla.gdims[2];
+      enqueue_launch_event_info.launch_event.vector_length = kla.gdims[1];
+
+      api_info->device_api = acc_device_api_other;
+
+      GOMP_PLUGIN_goacc_profiling_dispatch (prof_info,
+	&enqueue_launch_event_info, api_info);
+    }
+
   if (!async)
     {
       run_kernel (kernel, ind_da, &kla, NULL, false);
@@ -3289,6 +3376,15 @@ gcn_exec (struct kernel_info *kernel, size_t mapnum, void **hostaddrs,
 		   aq->agent->device_id, aq->id, ind_da);
       queue_push_callback (aq, gomp_offload_free, ind_da);
     }
+
+  if (profiling_dispatch_p)
+    {
+      prof_info->event_type = acc_ev_enqueue_launch_end;
+      enqueue_launch_event_info.launch_event.event_type = prof_info->event_type;
+      GOMP_PLUGIN_goacc_profiling_dispatch (prof_info,
+					    &enqueue_launch_event_info,
+					    api_info);
+    }
 }
 
 void
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 6a44e8ffb6a..cf980f1baec 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
@@ -224,6 +224,8 @@ static void cb_compute_construct_end (acc_prof_info *prof_info, acc_event_info *
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
+  else if (acc_device_type == acc_device_gcn)
+    assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
   assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
index 269b4398478..9c1cfbe292c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
@@ -107,6 +107,8 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e
     assert (event_info->launch_event.vector_length >= 1);
   else if (acc_device_type == acc_device_nvidia) /* ... is special.  */
     assert (event_info->launch_event.vector_length == 32);
+  else if (acc_device_type == acc_device_gcn) /* ...and so is this.  */
+    assert (event_info->launch_event.vector_length == 64);
   else
     {
 #ifdef __OPTIMIZE__
@@ -119,6 +121,8 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
+  else if (acc_device_type == acc_device_gcn)
+    assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
   assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
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 116b9b538a6..5d392511592 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
@@ -265,6 +265,8 @@ static void cb_enter_data_end (acc_prof_info *prof_info, acc_event_info *event_i
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
+  else if (acc_device_type == acc_device_gcn)
+    assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
   assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
@@ -319,6 +321,8 @@ static void cb_exit_data_start (acc_prof_info *prof_info, acc_event_info *event_
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
+  else if (acc_device_type == acc_device_gcn)
+    assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
   assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
@@ -371,6 +375,8 @@ static void cb_exit_data_end (acc_prof_info *prof_info, acc_event_info *event_in
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
+  else if (acc_device_type == acc_device_gcn)
+    assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
   assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
@@ -510,6 +516,8 @@ static void cb_compute_construct_end (acc_prof_info *prof_info, acc_event_info *
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
+  else if (acc_device_type == acc_device_gcn)
+    assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
   assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
@@ -573,6 +581,8 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
+  else if (acc_device_type == acc_device_gcn)
+    assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
   assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
@@ -637,6 +647,8 @@ static void cb_enqueue_launch_end (acc_prof_info *prof_info, acc_event_info *eve
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
+  else if (acc_device_type == acc_device_gcn)
+    assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
   assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
-- 
2.22.0

[PATCH] [og9] Add omp_pause_resource{,_all} for AMD GCN

[Julian Brown]

This patch adds some missing functions to the AMD GCN-specific target.c
file. This fixes several link errors in the testsuite.

Tested with offloading to AMD GCN. I will apply to the
openacc-gcc-9-branch shortly.

Julian

ChangeLog

	libgomp/
	* config/gcn/target.c (omp_pause_resource, omp_pause_resource_all): New
	functions, plus ialiases.
---
 libgomp/ChangeLog.openacc   |  5 +++++
 libgomp/config/gcn/target.c | 18 ++++++++++++++++++
 2 files changed, 23 insertions(+)

diff --git a/libgomp/ChangeLog.openacc b/libgomp/ChangeLog.openacc
index 438bd59b47b..d7a4c7a5f8a 100644
--- a/libgomp/ChangeLog.openacc
+++ b/libgomp/ChangeLog.openacc
@@ -1,3 +1,8 @@
+2019-09-06  Julian Brown  <julian@codesourcery.com>
+
+	* config/gcn/target.c (omp_pause_resource, omp_pause_resource_all): New
+	functions, plus ialiases.
+
 2019-09-05  Julian Brown  <julian@codesourcery.com>
 
 	* plugin/plugin-gcn.c (gcn_exec): Change default number of workers to
diff --git a/libgomp/config/gcn/target.c b/libgomp/config/gcn/target.c
index 5ec57778772..db00551e695 100644
--- a/libgomp/config/gcn/target.c
+++ b/libgomp/config/gcn/target.c
@@ -47,3 +47,21 @@ GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
     }
   gomp_num_teams_var = num_teams - 1;
 }
+
+int
+omp_pause_resource (omp_pause_resource_t kind, int device_num)
+{
+  (void) kind;
+  (void) device_num;
+  return -1;
+}
+
+int
+omp_pause_resource_all (omp_pause_resource_t kind)
+{
+  (void) kind;
+  return -1;
+}
+
+ialias (omp_pause_resource)
+ialias (omp_pause_resource_all)
-- 
2.22.0

Fix OpenACC/GCN 'acc_ev_enqueue_launch_end' position (was: [PATCH] [og9] OpenACC profiling support for AMD GCN)

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 3604 bytes --]

Hi!

On 2019-09-06T09:02:13-0700, Julian Brown <julian@codesourcery.com> wrote:
> This patch adds profiling support to the AMD GCN libgomp plugin, modeled
> after the equivalent support in the NVPTX plugin. This gives a positive
> test delta in AMD GCN offload testing.

Yay!  \o/

> I will apply to the openacc-gcc-9-branch shortly.

..., and later these changes got into master branch, via integration into
"[PATCH 7/7 libgomp,amdgcn] GCN Libgomp Plugin".

> --- a/libgomp/plugin/plugin-gcn.c
> +++ b/libgomp/plugin/plugin-gcn.c

|  static void
|  gomp_offload_free (void *ptr)
|  {
|    GCN_DEBUG ("Async thread ?:?: Freeing %p\n", ptr);
|    GOMP_OFFLOAD_free (0, ptr);
|  }

> @@ -3046,6 +3075,35 @@ GOMP_OFFLOAD_free (int device, void *ptr)
>        return false;
>      }
>
> +  struct goacc_thread *thr = GOMP_PLUGIN_goacc_thread ();
> +  bool profiling_dispatch_p
> +    = __builtin_expect (thr != NULL && thr->prof_info != NULL, false);
> +  if (profiling_dispatch_p)
> +    {
> +      [...]
> +      prof_info->event_type = acc_ev_free;
> +
> +      [...]
> +      GOMP_PLUGIN_goacc_profiling_dispatch (prof_info, &data_event_info,
> +                                         api_info);
> +    }
> +
>    return true;
>  }
>

> @@ -3276,6 +3334,35 @@ gcn_exec (struct kernel_info *kernel, size_t mapnum, void **hostaddrs,
>       {1,       64, 16}
>      };
>
> +  struct goacc_thread *thr = GOMP_PLUGIN_goacc_thread ();
> +  acc_prof_info *prof_info = thr->prof_info;
> +  acc_event_info enqueue_launch_event_info;
> +  acc_api_info *api_info = thr->api_info;
> +  bool profiling_dispatch_p = __builtin_expect (prof_info != NULL, false);
> +  if (profiling_dispatch_p)
> +    {
> +      prof_info->event_type = acc_ev_enqueue_launch_start;
> +
> +      [...]
> +      GOMP_PLUGIN_goacc_profiling_dispatch (prof_info,
> +     &enqueue_launch_event_info, api_info);
> +    }
> +
>    if (!async)
>      {
>        run_kernel (kernel, ind_da, &kla, NULL, false);
|        gomp_offload_free (ind_da);
|      }
|    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);
>      }
> +
> +  if (profiling_dispatch_p)
> +    {
> +      prof_info->event_type = acc_ev_enqueue_launch_end;
> +      enqueue_launch_event_info.launch_event.event_type = prof_info->event_type;
> +      GOMP_PLUGIN_goacc_profiling_dispatch (prof_info,
> +                                         &enqueue_launch_event_info,
> +                                         api_info);
> +    }
>  }

Per that, 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

Pushed to master branch commit 649f1939baf11f45fd3579b8b9601c7840a097b3
"Fix OpenACC/GCN 'acc_ev_enqueue_launch_end' position", see attached.


Grüße
 Thomas


-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Fix-OpenACC-GCN-acc_ev_enqueue_launch_end-position.patch --]
[-- Type: text/x-diff, Size: 12617 bytes --]

From 649f1939baf11f45fd3579b8b9601c7840a097b3 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Thu, 2 Mar 2023 10:39:09 +0100
Subject: [PATCH] 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'.
---
 libgomp/plugin/plugin-gcn.c                   |  23 +-
 .../acc_prof-parallel-1.c                     | 202 ++++++++++++++++--
 2 files changed, 195 insertions(+), 30 deletions(-)

diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index 11ce6b0fa8d..96920a48d2e 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -3192,18 +3192,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)
     {
@@ -3213,6 +3204,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 1f503861cb6..cbf23d7d83b 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);
@@ -703,6 +865,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);
@@ -725,9 +889,9 @@ int main()
 
       state_init = state;
     }
-    assert (state_init == 4);
+    assert (state_init == 5);
   }
-  assert (state == 10);
+  assert (state == 14);
 
   STATE_OP (state, = 100);
 
@@ -742,9 +906,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;
 }
-- 
2.25.1