Hi!
Reported by Cesar for a test case similar to the one below, where we
observe:
acc_prof-cuda-1.exe: [...]/libgomp/oacc-profiling.c:592:
goacc_profiling_dispatch_p: Assertion `thr->prof_info == NULL' failed.
This is because of:
On Tue, 25 Jul 2017 20:51:05 +0800, Chung-Lin Tang <chunglin_t...@mentor.com>
wrote:
> --- libgomp/oacc-cuda.c (revision 250497)
> +++ libgomp/oacc-cuda.c (working copy)
> @@ -99,17 +99,12 @@ acc_get_cuda_stream (int async)
> prof_info.async_queue = prof_info.async;
> }
>
> - void *ret = NULL;
> if (thr && thr->dev && thr->dev->openacc.cuda.get_stream_func)
> - ret = thr->dev->openacc.cuda.get_stream_func (async);
> -
> - if (profiling_setup_p)
> {
> - thr->prof_info = NULL;
> - thr->api_info = NULL;
> + goacc_aq aq = lookup_goacc_asyncqueue (thr, false, async);
> + return aq ? thr->dev->openacc.cuda.get_stream_func (aq) : NULL;
> }
> -
> - return ret;
> + return NULL;
> }
Pushed to openacc-gcc-7-branch:
commit db149741171147fa86a9bfe708a9082f508115ac
Author: Thomas Schwinge <tho...@codesourcery.com>
Date: Tue Oct 10 19:25:19 2017 +0200
acc_get_cuda_stream: Clean up data of the OpenACC Profiling Interface
libgomp/
* oacc-cuda.c (acc_get_cuda_stream): Clean up data of the OpenACC
Profiling Interface.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-cuda-1.c: New file.
---
libgomp/oacc-cuda.c | 13 +++++++++++--
.../libgomp.oacc-c-c++-common/acc_prof-cuda-1.c | 16 ++++++++++++++++
2 files changed, 27 insertions(+), 2 deletions(-)
diff --git libgomp/oacc-cuda.c libgomp/oacc-cuda.c
index 1fbe77d..0ac93e9 100644
--- libgomp/oacc-cuda.c
+++ libgomp/oacc-cuda.c
@@ -99,12 +99,21 @@ acc_get_cuda_stream (int async)
prof_info.async_queue = prof_info.async;
}
+ void *ret = NULL;
if (thr && thr->dev && thr->dev->openacc.cuda.get_stream_func)
{
goacc_aq aq = lookup_goacc_asyncqueue (thr, false, async);
- return aq ? thr->dev->openacc.cuda.get_stream_func (aq) : NULL;
+ if (aq)
+ ret = thr->dev->openacc.cuda.get_stream_func (aq);
}
- return NULL;
+
+ if (profiling_setup_p)
+ {
+ thr->prof_info = NULL;
+ thr->api_info = NULL;
+ }
+
+ return ret;
}
int
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-cuda-1.c
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-cuda-1.c
new file mode 100644
index 0000000..63f5e49
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-cuda-1.c
@@ -0,0 +1,16 @@
+/* TODO: This is to eventually test dispatch of events to callbacks. */
+
+#include <openacc.h>
+
+int main()
+{
+ acc_init(acc_device_default);
+
+ (void) acc_get_cuda_stream(acc_async_default);
+ /* The following used to crash the runtime due to acc_get_cuda_stream not
+ cleaning up data of the OpenACC Profiling Interface. */
+#pragma acc data
+ ;
+
+ return 0;
+}
Grüße
Thomas
