xref: /llvm-project/offload/test/unified_shared_memory/close_member.c (revision 330d8983d25d08580fc1642fea48b2473f47a9da)

// 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;
}