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