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
21int 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

source code of offload/test/unified_shared_memory/shared_update.c