openmp/libomptarget/test/unified_shared_memory/close_member.c

   1 // RUN: %libomptarget-compile-run-and-check-generic
   2
   3 // REQUIRES: unified_shared_memory
   4 // UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
   5
   6 #include <omp.h>
   7 #include <stdio.h>
   8
   9 #pragma omp requires unified_shared_memory
  10
  11 struct S {
  12   int x;
  13   int y;
  14 };
  15
  16 int main(int argc, char *argv[]) {
  17   int dev = omp_get_default_device();
  18   struct S s = {10, 20};
  19
  20 #pragma omp target enter data map(close, to : s)
  21 #pragma omp target map(alloc : s)
  22   {
  23     s.x = 11;
  24     s.y = 21;
  25   }
  26 // To determine whether x needs to be transfered or deleted, the runtime
  27 // cannot simply check whether unified shared memory is enabled and the
  28 // 'close' modifier is specified.  It must check whether x was previously
  29 // placed in device memory by, for example, a 'close' modifier that isn't
  30 // specified here.  The following struct member case checks a special code
  31 // path in the runtime implementation where members are transferred before
  32 // deletion of the struct.
  33 #pragma omp target exit data map(from : s.x, s.y)
  34
  35   // CHECK: s.x=11, s.y=21
  36   printf("s.x=%d, s.y=%d\n", s.x, s.y);
  37   // CHECK: present: 0
  38   printf("present: %d\n", omp_target_is_present(&s, dev));
  39
  40   return 0;
  41 }