1// RUN: %libomptarget-compile-generic
2// RUN: env HSA_XNACK=1 \
3// RUN: %libomptarget-run-generic | %fcheck-generic
4
5// REQUIRES: unified_shared_memory
6// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
7
8// Fails on nvptx with error: an illegal memory access was encountered
9// XFAIL: nvptx64-nvidia-cuda
10// XFAIL: 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
19int main(int argc, char *argv[]) {
20 int fails;
21 void *host_alloc = 0, *device_alloc = 0;
22 int *a = (int *)malloc(N * sizeof(int));
23 int dev = omp_get_default_device();
24
25 // Init
26 for (int i = 0; i < N; ++i) {
27 a[i] = 10;
28 }
29 host_alloc = &a[0];
30
31 //
32 // map + target no close
33 //
34#pragma omp target data map(tofrom : a[ : N]) map(tofrom : device_alloc)
35 {
36#pragma omp target map(tofrom : device_alloc)
37 { device_alloc = &a[0]; }
38 }
39
40 // CHECK: a used from unified memory.
41 if (device_alloc == host_alloc)
42 printf(format: "a used from unified memory.\n");
43
44 //
45 // map + target with close
46 //
47 device_alloc = 0;
48#pragma omp target data map(close, tofrom : a[ : N]) map(tofrom : device_alloc)
49 {
50#pragma omp target map(tofrom : device_alloc)
51 { device_alloc = &a[0]; }
52 }
53 // CHECK: a copied to device.
54 if (device_alloc != host_alloc)
55 printf(format: "a copied to device.\n");
56
57 //
58 // map + use_device_ptr no close
59 //
60 device_alloc = 0;
61#pragma omp target data map(tofrom : a[ : N]) use_device_ptr(a)
62 { device_alloc = &a[0]; }
63
64 // CHECK: a used from unified memory with use_device_ptr.
65 if (device_alloc == host_alloc)
66 printf(format: "a used from unified memory with use_device_ptr.\n");
67
68 //
69 // map + use_device_ptr close
70 //
71 device_alloc = 0;
72#pragma omp target data map(close, tofrom : a[ : N]) use_device_ptr(a)
73 { device_alloc = &a[0]; }
74
75 // CHECK: a used from device memory with use_device_ptr.
76 if (device_alloc != host_alloc)
77 printf(format: "a used from device memory with use_device_ptr.\n");
78
79 //
80 // map enter/exit + close
81 //
82 device_alloc = 0;
83#pragma omp target enter data map(close, to : a[ : N])
84
85#pragma omp target map(from : device_alloc)
86 {
87 device_alloc = &a[0];
88 a[0] = 99;
89 }
90
91 // 'close' is missing, so the runtime must check whether s is actually in
92 // shared memory in order to determine whether to transfer data and delete the
93 // allocation.
94#pragma omp target exit data map(from : a[ : N])
95
96 // CHECK: a has been mapped to the device.
97 if (device_alloc != host_alloc)
98 printf(format: "a has been mapped to the device.\n");
99
100 // CHECK: a[0]=99
101 // CHECK: a is present: 0
102 printf(format: "a[0]=%d\n", a[0]);
103 printf(format: "a is present: %d\n", omp_target_is_present(a, dev));
104
105 free(a);
106
107 // CHECK: Done!
108 printf(format: "Done!\n");
109
110 return 0;
111}
112

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