// RUN: %libomptarget-compile-run-and-check-generic

// REQUIRES: unified_shared_memory
// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9

#include <omp.h>
#include <stdio.h>

#pragma omp requires unified_shared_memory

struct S {
  int x;
  int y;
};

int main(int argc, char *argv[]) {
  int dev = omp_get_default_device();
  struct S s = {10, 20};

  #pragma omp target enter data map(close, to: s)
  #pragma omp target map(alloc: s)
  {
    s.x = 11;
    s.y = 21;
  }
  // To determine whether x needs to be transfered or deleted, the runtime
  // cannot simply check whether unified shared memory is enabled and the
  // 'close' modifier is specified.  It must check whether x was previously
  // placed in device memory by, for example, a 'close' modifier that isn't
  // specified here.  The following struct member case checks a special code
  // path in the runtime implementation where members are transferred before
  // deletion of the struct.
  #pragma omp target exit data map(from: s.x, s.y)

  // CHECK: s.x=11, s.y=21
  printf("s.x=%d, s.y=%d\n", s.x, s.y);
  // CHECK: present: 0
  printf("present: %d\n", omp_target_is_present(&s, dev));

  return 0;
}