@@ -237,7 +237,31 @@ cuda_map_create (size_t size)
static void
cuda_map_destroy (struct cuda_map *map)
{
- CUDA_CALL_ASSERT (cuMemFree, map->d);
+ if (map->active)
+ /* Possible reasons for the map to be still active:
+ - the associated async kernel might still be running.
+ - the associated async kernel might have finished, but the
+ corresponding event that should trigger the pop_map has not been
+ processed by event_gc.
+ - the associated sync kernel might have aborted
+
+ The async cases could happen if the user specified an async region
+ without adding a corresponding wait that is guaranteed to be executed
+ (before returning from main, or in an atexit handler).
+ We do not want to deallocate a device pointer that is still being
+ used, so skip it.
+
+ In the sync case, the device pointer is no longer used, but deallocating
+ it using cuMemFree will not succeed, so skip it.
+
+ TODO: Handle this in a more constructive way, by f.i. waiting for streams
+ to finish before de-allocating them (PR88981), or by ensuring the CUDA
+ lib atexit handler is called before rather than after the libgomp plugin
+ atexit handler (PR83795). */
+ ;
+ else
+ CUDA_CALL_ASSERT (cuMemFree, map->d);
+
free (map);
}
@@ -268,7 +292,6 @@ static bool
map_fini (struct ptx_stream *s)
{
assert (s->map->next == NULL);
- assert (!s->map->active);
cuda_map_destroy (s->map);
new file mode 100644
@@ -0,0 +1,15 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+int
+main (void)
+{
+
+#pragma acc parallel async
+ ;
+
+ /* no #pragma acc wait */
+ return 0;
+}
+
Hi, There are currently two situations where this assert triggers: ... libgomp/plugin/plugin-nvptx.c: map_fini: Assertion `!s->map->active' failed. ... First, in abort-1.c, a parallel region triggering an abort: ... int main (void) { #pragma acc parallel abort (); return 0; } ... The abort is detected in nvptx_exec as the CUDA_ERROR_ILLEGAL_INSTRUCTION return status of the cuStreamSynchronize call after kernel launch, which is then handled by calling non-returning function GOMP_PLUGIN_fatal. Consequently, the map_pop in nvptx_exec that in case of cuStreamSynchronize success would remove or inactive the element added by the map_push earlier in nvptx_exec, does not trigger. With the element no longer active, but still marked active and a member of s->map, we run into the assert during GOMP_OFFLOAD_fini_device, which is triggered from atexit handler gomp_target_fini (which is triggered by the GOMP_PLUGIN_fatal mentioned above calling exit). Second, in pr88941.c, an async parallel region without wait: ... int main (void) { #pragma acc parallel async ; /* no #pragma acc wait */ return 0; } ... Because nvptx_exec is handling an async region, it does not call map_pop for the element added by map_push, but schedules an kernel execution completion event to call map_pop. Again, we run into the assert during GOMP_OFFLOAD_fini_device, which is triggered from atexit handler gomp_target_fini, but the exit in this case is triggered by returning from main. So either the kernel is still running, or the kernel has completed but the corresponding event that is supposed to call map_pop is stuck in the event queue, waiting for an event_gc. Fix this by removing the assert, and skipping the freeing of device memory if the map is still marked active (though in the async case, this is more a workaround than an fix). Committed to trunk. Thanks, - Tom [nvptx, libgomp] Fix assert (!s->map->active) in map_fini 2019-01-22 Tom de Vries <tdevries@suse.de> PR target/88941 PR target/88939 * plugin/plugin-nvptx.c (cuda_map_destroy): Handle map->active case. (map_fini): Remove "assert (!s->map->active)". * testsuite/libgomp.oacc-c-c++-common/pr88941.c: New test. --- libgomp/plugin/plugin-nvptx.c | 27 ++++++++++++++++++++-- .../testsuite/libgomp.oacc-c-c++-common/pr88941.c | 15 ++++++++++++ 2 files changed, 40 insertions(+), 2 deletions(-)