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

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