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