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 // amdgpu runtime crash
7 // Fails on nvptx with error: an illegal memory access was encountered
8 // UNSUPPORTED: amdgcn-amd-amdhsa
9 // UNSUPPORTED: nvptx64-nvidia-cuda
10 // UNSUPPORTED: 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, *device_alloc;
22 void *host_data, *device_data;
23 int *alloc = (int *)malloc(N * sizeof(int));
24 int data[N];
25
26 for (int i = 0; i < N; ++i) {
27 alloc[i] = 10;
28 data[i] = 1;
29 }
30
31 host_data = &data[0];
32 host_alloc = &alloc[0];
33
34 //
35 // Test that updates on the device are not visible to host
36 // when only a TO mapping is used.
37 //
38 #pragma omp target map(tofrom : device_data, device_alloc) \
39 map(close, to : alloc[ : N], data[ : N])
40 {
41 device_data = &data[0];
42 device_alloc = &alloc[0];
43
44 for (int i = 0; i < N; i++) {
45 alloc[i] += 1;
46 data[i] += 1;
47 }
48 }
49
50 // CHECK: Address of alloc on device different from host address.
51 if (device_alloc != host_alloc)
52 printf("Address of alloc on device different from host address.\n");
53
54 // CHECK: Address of data on device different from host address.
55 if (device_data != host_data)
56 printf("Address of data on device different from host address.\n");
57
58 // On the host, check that the arrays have been updated.
59 // CHECK: Alloc host values not updated: Succeeded
60 fails = 0;
61 for (int i = 0; i < N; i++) {
62 if (alloc[i] != 10)
63 fails++;
64 }
65 printf("Alloc host values not updated: %s\n",
66 (fails == 0) ? "Succeeded" : "Failed");
67
68 // CHECK: Data host values not updated: Succeeded
69 fails = 0;
70 for (int i = 0; i < N; i++) {
71 if (data[i] != 1)
72 fails++;
73 }
74 printf("Data host values not updated: %s\n",
75 (fails == 0) ? "Succeeded" : "Failed");
76
77 //
78 // Test that updates on the device are visible on host
79 // when a from is used.
80 //
81
82 for (int i = 0; i < N; i++) {
83 alloc[i] += 1;
84 data[i] += 1;
85 }
86
87 #pragma omp target map(close, tofrom : alloc[ : N], data[ : N])
88 {
89 // CHECK: Alloc device values are correct: Succeeded
90 fails = 0;
91 for (int i = 0; i < N; i++) {
92 if (alloc[i] != 11)
93 fails++;
94 }
95 printf("Alloc device values are correct: %s\n",
96 (fails == 0) ? "Succeeded" : "Failed");
97 // CHECK: Data device values are correct: Succeeded
98 fails = 0;
99 for (int i = 0; i < N; i++) {
100 if (data[i] != 2)
101 fails++;
102 }
103 printf("Data device values are correct: %s\n",
104 (fails == 0) ? "Succeeded" : "Failed");
105
106 // Update values on the device
107 for (int i = 0; i < N; i++) {
108 alloc[i] += 1;
109 data[i] += 1;
110 }
111 }
112
113 // CHECK: Alloc host values updated: Succeeded
114 fails = 0;
115 for (int i = 0; i < N; i++) {
116 if (alloc[i] != 12)
117 fails++;
118 }
119 printf("Alloc host values updated: %s\n",
120 (fails == 0) ? "Succeeded" : "Failed");
121
122 // CHECK: Data host values updated: Succeeded
123 fails = 0;
124 for (int i = 0; i < N; i++) {
125 if (data[i] != 3)
126 fails++;
127 }
128 printf("Data host values updated: %s\n",
129 (fails == 0) ? "Succeeded" : "Failed");
130
131 free(alloc);
132
133 // CHECK: Done!
134 printf("Done!\n");
135
136 return 0;
137 }
138