1 // RUN: %libomptarget-compile-run-and-check-generic
2
3 // REQUIRES: unified_shared_memory
4
5 // amdgpu runtime crash
6 // UNSUPPORTED: amdgcn-amd-amdhsa
7
8 #include <stdio.h>
9 #include <omp.h>
10
11 // ---------------------------------------------------------------------------
12 // Various definitions copied from OpenMP RTL
13
14 extern void __tgt_register_requires(int64_t);
15
16 // End of definitions copied from OpenMP RTL.
17 // ---------------------------------------------------------------------------
18
19 #pragma omp requires unified_shared_memory
20
21 #define N 1024
22
main(int argc,char * argv[])23 int main(int argc, char *argv[]) {
24 int fails;
25 void *host_alloc, *device_alloc;
26 void *host_data, *device_data;
27 int *alloc = (int *)malloc(N * sizeof(int));
28 int data[N];
29
30 // Manual registration of requires flags for Clang versions
31 // that do not support requires.
32 __tgt_register_requires(8);
33
34 for (int i = 0; i < N; ++i) {
35 alloc[i] = 10;
36 data[i] = 1;
37 }
38
39 host_data = &data[0];
40 host_alloc = &alloc[0];
41
42 // implicit mapping of data
43 #pragma omp target map(tofrom : device_data, device_alloc)
44 {
45 device_data = &data[0];
46 device_alloc = &alloc[0];
47
48 for (int i = 0; i < N; i++) {
49 alloc[i] += 1;
50 data[i] += 1;
51 }
52 }
53
54 // CHECK: Address of alloc on device matches host address.
55 if (device_alloc == host_alloc)
56 printf("Address of alloc on device matches host address.\n");
57
58 // CHECK: Address of data on device matches host address.
59 if (device_data == host_data)
60 printf("Address of data on device matches host address.\n");
61
62 // On the host, check that the arrays have been updated.
63 // CHECK: Alloc device values updated: Succeeded
64 fails = 0;
65 for (int i = 0; i < N; i++) {
66 if (alloc[i] != 11)
67 fails++;
68 }
69 printf("Alloc device values updated: %s\n",
70 (fails == 0) ? "Succeeded" : "Failed");
71
72 // CHECK: Data device values updated: Succeeded
73 fails = 0;
74 for (int i = 0; i < N; i++) {
75 if (data[i] != 2)
76 fails++;
77 }
78 printf("Data device values updated: %s\n",
79 (fails == 0) ? "Succeeded" : "Failed");
80
81 //
82 // Test that updates on the host snd on the device are both visible.
83 //
84
85 // Update on the host.
86 for (int i = 0; i < N; ++i) {
87 alloc[i] += 1;
88 data[i] += 1;
89 }
90
91 #pragma omp target
92 {
93 // CHECK: Alloc host values updated: Succeeded
94 fails = 0;
95 for (int i = 0; i < N; i++) {
96 if (alloc[i] != 12)
97 fails++;
98 }
99 printf("Alloc host values updated: %s\n",
100 (fails == 0) ? "Succeeded" : "Failed");
101 // CHECK: Data host values updated: Succeeded
102 fails = 0;
103 for (int i = 0; i < N; i++) {
104 if (data[i] != 3)
105 fails++;
106 }
107 printf("Data host values updated: %s\n",
108 (fails == 0) ? "Succeeded" : "Failed");
109 }
110
111 free(alloc);
112
113 printf("Done!\n");
114
115 return 0;
116 }
117