With USM, the testcase libgomp.c++/target-std__multimap-concurrent-usm.C
segfaults on the host in the destructor of the multimap _map.
Freeing device allocated memory on the device fixes this.
I intent to commit it later today, unless there are comments
implying otherwise.
Thanks,
Tobias
libgomp.c++/target-std__multimap-concurrent.C: Fix USM memory freeing
Fix the unified-shared memory test,
libgomp.c++/target-std__multimap-concurrent-usm.C
added in commit r16-1010-g83ca283853f195
libgomp: Add testcases for concurrent access to standard C++ containers
on offload targets, a number of USM variants
This tests includes the actual code of target-std__multimap-concurrent.C.
The issue is that multimap.insert allocates memory – which is freed by
the destructor. However, if the memory is allocated on a device
('insert'), it also needs to be freed there ('clear') as in general
freeing device-allocated memory is not possible on the host.
libgomp/ChangeLog:
* testsuite/libgomp.c++/target-std__multimap-concurrent.C: Fix memory
freeing of device allocated memory with USM.
.../testsuite/libgomp.c++/target-std__multimap-concurrent.C | 11 +++++++++++
1 file changed, 11 insertions(+)
diff --git a/libgomp/testsuite/libgomp.c++/target-std__multimap-concurrent.C b/libgomp/testsuite/libgomp.c++/target-std__multimap-concurrent.C
index 6a4a4e80d34..8dbc912868a 100644
--- a/libgomp/testsuite/libgomp.c++/target-std__multimap-concurrent.C
+++ b/libgomp/testsuite/libgomp.c++/target-std__multimap-concurrent.C
@@ -4,6 +4,7 @@
#include <stdlib.h>
#include <time.h>
#include <map>
+#include <omp.h>
// Make sure that KEY_MAX is less than N to ensure some duplicate keys.
#define N 3000
@@ -53,6 +54,16 @@ int main (void)
for (auto it = range.first; it != range.second; ++it)
sum += (long long) it->first * it->second;
}
+#ifdef MEM_SHARED
+ /* Even with USM, memory allocated on the device (with _map.insert)
+ must be freed on the device. */
+ if (omp_get_default_device () != omp_initial_device
+ && omp_get_default_device () != omp_get_num_devices ())
+ {
+ #pragma omp target
+ _map.clear ();
+ }
+#endif
#ifndef MEM_SHARED
#pragma omp target
