xref: /llvm-project/offload/test/unified_shared_memory/close_enter_exit.c (revision 330d8983d25d08580fc1642fea48b2473f47a9da)
1*330d8983SJohannes Doerfert // RUN: %libomptarget-compile-generic
2*330d8983SJohannes Doerfert // RUN: env HSA_XNACK=1 \
3*330d8983SJohannes Doerfert // RUN: %libomptarget-run-generic | %fcheck-generic
4*330d8983SJohannes Doerfert 
5*330d8983SJohannes Doerfert // REQUIRES: unified_shared_memory
6*330d8983SJohannes Doerfert // UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
7*330d8983SJohannes Doerfert 
8*330d8983SJohannes Doerfert // Fails on nvptx with error: an illegal memory access was encountered
9*330d8983SJohannes Doerfert // XFAIL: nvptx64-nvidia-cuda
10*330d8983SJohannes Doerfert // XFAIL: nvptx64-nvidia-cuda-LTO
11*330d8983SJohannes Doerfert 
12*330d8983SJohannes Doerfert #include <omp.h>
13*330d8983SJohannes Doerfert #include <stdio.h>
14*330d8983SJohannes Doerfert 
15*330d8983SJohannes Doerfert #pragma omp requires unified_shared_memory
16*330d8983SJohannes Doerfert 
17*330d8983SJohannes Doerfert #define N 1024
18*330d8983SJohannes Doerfert 
main(int argc,char * argv[])19*330d8983SJohannes Doerfert int main(int argc, char *argv[]) {
20*330d8983SJohannes Doerfert   int fails;
21*330d8983SJohannes Doerfert   void *host_alloc = 0, *device_alloc = 0;
22*330d8983SJohannes Doerfert   int *a = (int *)malloc(N * sizeof(int));
23*330d8983SJohannes Doerfert   int dev = omp_get_default_device();
24*330d8983SJohannes Doerfert 
25*330d8983SJohannes Doerfert   // Init
26*330d8983SJohannes Doerfert   for (int i = 0; i < N; ++i) {
27*330d8983SJohannes Doerfert     a[i] = 10;
28*330d8983SJohannes Doerfert   }
29*330d8983SJohannes Doerfert   host_alloc = &a[0];
30*330d8983SJohannes Doerfert 
31*330d8983SJohannes Doerfert   //
32*330d8983SJohannes Doerfert   // map + target no close
33*330d8983SJohannes Doerfert   //
34*330d8983SJohannes Doerfert #pragma omp target data map(tofrom : a[ : N]) map(tofrom : device_alloc)
35*330d8983SJohannes Doerfert   {
36*330d8983SJohannes Doerfert #pragma omp target map(tofrom : device_alloc)
37*330d8983SJohannes Doerfert     { device_alloc = &a[0]; }
38*330d8983SJohannes Doerfert   }
39*330d8983SJohannes Doerfert 
40*330d8983SJohannes Doerfert   // CHECK: a used from unified memory.
41*330d8983SJohannes Doerfert   if (device_alloc == host_alloc)
42*330d8983SJohannes Doerfert     printf("a used from unified memory.\n");
43*330d8983SJohannes Doerfert 
44*330d8983SJohannes Doerfert   //
45*330d8983SJohannes Doerfert   // map + target with close
46*330d8983SJohannes Doerfert   //
47*330d8983SJohannes Doerfert   device_alloc = 0;
48*330d8983SJohannes Doerfert #pragma omp target data map(close, tofrom : a[ : N]) map(tofrom : device_alloc)
49*330d8983SJohannes Doerfert   {
50*330d8983SJohannes Doerfert #pragma omp target map(tofrom : device_alloc)
51*330d8983SJohannes Doerfert     { device_alloc = &a[0]; }
52*330d8983SJohannes Doerfert   }
53*330d8983SJohannes Doerfert   // CHECK: a copied to device.
54*330d8983SJohannes Doerfert   if (device_alloc != host_alloc)
55*330d8983SJohannes Doerfert     printf("a copied to device.\n");
56*330d8983SJohannes Doerfert 
57*330d8983SJohannes Doerfert   //
58*330d8983SJohannes Doerfert   // map + use_device_ptr no close
59*330d8983SJohannes Doerfert   //
60*330d8983SJohannes Doerfert   device_alloc = 0;
61*330d8983SJohannes Doerfert #pragma omp target data map(tofrom : a[ : N]) use_device_ptr(a)
62*330d8983SJohannes Doerfert   { device_alloc = &a[0]; }
63*330d8983SJohannes Doerfert 
64*330d8983SJohannes Doerfert   // CHECK: a used from unified memory with use_device_ptr.
65*330d8983SJohannes Doerfert   if (device_alloc == host_alloc)
66*330d8983SJohannes Doerfert     printf("a used from unified memory with use_device_ptr.\n");
67*330d8983SJohannes Doerfert 
68*330d8983SJohannes Doerfert   //
69*330d8983SJohannes Doerfert   // map + use_device_ptr close
70*330d8983SJohannes Doerfert   //
71*330d8983SJohannes Doerfert   device_alloc = 0;
72*330d8983SJohannes Doerfert #pragma omp target data map(close, tofrom : a[ : N]) use_device_ptr(a)
73*330d8983SJohannes Doerfert   { device_alloc = &a[0]; }
74*330d8983SJohannes Doerfert 
75*330d8983SJohannes Doerfert   // CHECK: a used from device memory with use_device_ptr.
76*330d8983SJohannes Doerfert   if (device_alloc != host_alloc)
77*330d8983SJohannes Doerfert     printf("a used from device memory with use_device_ptr.\n");
78*330d8983SJohannes Doerfert 
79*330d8983SJohannes Doerfert   //
80*330d8983SJohannes Doerfert   // map enter/exit + close
81*330d8983SJohannes Doerfert   //
82*330d8983SJohannes Doerfert   device_alloc = 0;
83*330d8983SJohannes Doerfert #pragma omp target enter data map(close, to : a[ : N])
84*330d8983SJohannes Doerfert 
85*330d8983SJohannes Doerfert #pragma omp target map(from : device_alloc)
86*330d8983SJohannes Doerfert   {
87*330d8983SJohannes Doerfert     device_alloc = &a[0];
88*330d8983SJohannes Doerfert     a[0] = 99;
89*330d8983SJohannes Doerfert   }
90*330d8983SJohannes Doerfert 
91*330d8983SJohannes Doerfert   // 'close' is missing, so the runtime must check whether s is actually in
92*330d8983SJohannes Doerfert   // shared memory in order to determine whether to transfer data and delete the
93*330d8983SJohannes Doerfert   // allocation.
94*330d8983SJohannes Doerfert #pragma omp target exit data map(from : a[ : N])
95*330d8983SJohannes Doerfert 
96*330d8983SJohannes Doerfert   // CHECK: a has been mapped to the device.
97*330d8983SJohannes Doerfert   if (device_alloc != host_alloc)
98*330d8983SJohannes Doerfert     printf("a has been mapped to the device.\n");
99*330d8983SJohannes Doerfert 
100*330d8983SJohannes Doerfert   // CHECK: a[0]=99
101*330d8983SJohannes Doerfert   // CHECK: a is present: 0
102*330d8983SJohannes Doerfert   printf("a[0]=%d\n", a[0]);
103*330d8983SJohannes Doerfert   printf("a is present: %d\n", omp_target_is_present(a, dev));
104*330d8983SJohannes Doerfert 
105*330d8983SJohannes Doerfert   free(a);
106*330d8983SJohannes Doerfert 
107*330d8983SJohannes Doerfert   // CHECK: Done!
108*330d8983SJohannes Doerfert   printf("Done!\n");
109*330d8983SJohannes Doerfert 
110*330d8983SJohannes Doerfert   return 0;
111*330d8983SJohannes Doerfert }
112