[libomptarget] PR38704: Fix erase of ShadowPtrMap

erase() invalidates the iterator and returns a new one pointing
to the following element. The code now follows the example at
https://en.cppreference.com/w/cpp/container/map/erase.
(The added testcase crashes without this patch.)

Reported by David Binderman (https://llvm.org/PR38704)!

Differential Revision: https://reviews.llvm.org/D51623

2 files changed, 45 insertions, 2 deletions

diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 14f45312030..447820deaaf 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -419,7 +419,7 @@ int target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base,
       uintptr_t ub = (uintptr_t) HstPtrBegin + data_size;
       Device.ShadowMtx.lock();
       for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin();
-          it != Device.ShadowPtrMap.end(); ++it) {
+           it != Device.ShadowPtrMap.end();) {
         void **ShadowHstPtrAddr = (void**) it->first;
 
         // An STL map is sorted on its keys; use this property
@@ -439,7 +439,9 @@ int target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base,
         // If the struct is to be deallocated, remove the shadow entry.
         if (DelEntry) {
           DP("Removing shadow pointer " DPxMOD "\n", DPxPTR(ShadowHstPtrAddr));
-          Device.ShadowPtrMap.erase(it);
+          it = Device.ShadowPtrMap.erase(it);
+        } else {
+          ++it;
         }
       }
       Device.ShadowMtx.unlock();
diff --git a/openmp/libomptarget/test/mapping/pr38704.c b/openmp/libomptarget/test/mapping/pr38704.c
new file mode 100644
index 00000000000..d1b3cc42e2f
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/pr38704.c
@@ -0,0 +1,41 @@
+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+// Clang 6.0 doesn't use the new map interface, undefined behavior when
+// the compiler emits "old" interface code for structures.
+// UNSUPPORTED: clang-6
+
+#include <stdio.h>
+#include <stdlib.h>
+
+typedef struct {
+  int *ptr1;
+  int *ptr2;
+} StructWithPtrs;
+
+int main(int argc, char *argv[]) {
+  StructWithPtrs s;
+  s.ptr1 = malloc(sizeof(int));
+  s.ptr2 = malloc(2 * sizeof(int));
+
+#pragma omp target map(s.ptr1[0:1], s.ptr2[0:2])
+  {
+    s.ptr1[0] = 1;
+    s.ptr2[0] = 2;
+    s.ptr2[1] = 3;
+  }
+
+  // CHECK: s.ptr1[0] = 1
+  // CHECK: s.ptr2[0] = 2
+  // CHECK: s.ptr2[1] = 3
+  printf("s.ptr1[0] = %d\n", s.ptr1[0]);
+  printf("s.ptr2[0] = %d\n", s.ptr2[0]);
+  printf("s.ptr2[1] = %d\n", s.ptr2[1]);
+
+  free(s.ptr1);
+  free(s.ptr2);
+
+  return 0;
+}