diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 5d13262c398..30fb11d0290 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,11 @@
+2019-01-23  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.
+
 2019-01-23  Tom de Vries  <tdevries@suse.de>
 
 	PR target/87835
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index a220560b189..4a67191932e 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -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);
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr88941.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr88941.c
new file mode 100644
index 00000000000..e31bb527df3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr88941.c
@@ -0,0 +1,15 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+int
+main (void)
+{
+
+#pragma acc parallel async
+  ;
+
+  /* no #pragma acc wait */
+  return 0;
+}
+