xref: /llvm-project/offload/test/unified_shared_memory/close_member.c (revision 330d8983d25d08580fc1642fea48b2473f47a9da)
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 
main(int argc,char * argv[])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 }
42