[flang-commits] [flang] [OpenMP][libomptarget] Add map checks when running under unified shared memory (PR #69005)

[via flang-commits]

================
@@ -0,0 +1,158 @@
+// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1
+// %libomptarget-run-generic 2>&1 | %fcheck-generic
+
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// REQUIRES: amdgcn-amd-amdhsa
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+#define N 1024
+
+int main(int argc, char *argv[]) {
+  int fails;
+  void *host_alloc, *device_alloc;
+  void *host_data, *device_data;
+  int *alloc = (int *)malloc(N * sizeof(int));
+  int data[N];
+
+  for (int i = 0; i < N; ++i) {
+    alloc[i] = 10;
+    data[i] = 1;
+  }
+
+  host_data = &data[0];
+  host_alloc = &alloc[0];
+
+// CHECK: Creating new map entry ONLY with
+// HstPtrBase=[[DEVICE_DATA_HST_PTR:0x.*]], HstPtrBegin=[[DEVICE_DATA_HST_PTR]],
+// TgtAllocBegin=[[DEVICE_DATA_HST_PTR]], TgtPtrBegin=[[DEVICE_DATA_HST_PTR]],
+// Size=8, DynRefCount=1, HoldRefCount=0 CHECK: Creating new map entry ONLY with
+// HstPtrBase=[[DATA_HST_PTR:0x.*]], HstPtrBegin=[[DATA_HST_PTR]],
+// TgtAllocBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096,
+// DynRefCount=1, HoldRefCount=0 CHECK: Creating new map entry ONLY with
+// HstPtrBase=[[DEVICE_ALLOC_HST_PTR:0x.*]],
+// HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtAllocBegin=[[DEVICE_ALLOC_HST_PTR]],
+// TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1, HoldRefCount=0
+
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_DATA_HST_PTR]],
+// TgtPtrBegin=[[DEVICE_DATA_HST_PTR]], Size=8, DynRefCount=1 (update
+// suppressed), HoldRefCount=0 CHECK: Mapping exists with
+// HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096,
+// DynRefCount=1 (update suppressed), HoldRefCount=0 CHECK: Mapping exists with
+// HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]],
+// Size=8, DynRefCount=1 (update suppressed), HoldRefCount=0
+
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks
+// and 256 threads in Generic mode
+
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]],
+// TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=0 (decremented,
+// delayed deletion), HoldRefCount=0 CHECK: Mapping exists with
+// HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096,
+// DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0 CHECK: Mapping
+// exists with HstPtrBegin=[[DEVICE_DATA_HST_PTR]],
+// TgtPtrBegin=[[DEVICE_DATA_HST_PTR]], Size=8, DynRefCount=0 (decremented,
+// delayed deletion), HoldRefCount=0
+
+// CHECK: Removing map entry with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]]{{.*}}
+// Size=8 CHECK: Removing map entry with HstPtrBegin=[[DATA_HST_PTR]]{{.*}}
+// Size=4096 CHECK: Removing map entry with
+// HstPtrBegin=[[DEVICE_DATA_HST_PTR]]{{.*}} Size=8
+
+// implicit mapping of data
+#pragma omp target map(tofrom : device_data, device_alloc)
+  {
+    device_data = &data[0];
+    device_alloc = &alloc[0];
+
+    for (int i = 0; i < N; i++) {
+      alloc[i] += 1;
+      data[i] += 1;
+    }
+  }
+
+  if (device_alloc == host_alloc)
+    printf("Address of alloc on device matches host address.\n");
+
+  if (device_data == host_data)
+    printf("Address of data on device matches host address.\n");
+
+  // On the host, check that the arrays have been updated.
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (alloc[i] != 11)
+      fails++;
+  }
+  printf("Alloc device values updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (data[i] != 2)
+      fails++;
+  }
+  printf("Data device values updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  //
+  // Test that updates on the host snd on the device are both visible.
----------------
carlobertolli wrote:

snd-->and

https://github.com/llvm/llvm-project/pull/69005