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