1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
|
// RUN: %libomptarget-compile-run-and-check-generic
// REQUIRES: unified_shared_memory
// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
#include <omp.h>
#include <stdio.h>
#pragma omp requires unified_shared_memory
struct S {
int x;
int y;
};
int main(int argc, char *argv[]) {
int dev = omp_get_default_device();
struct S s = {10, 20};
#pragma omp target enter data map(close, to: s)
#pragma omp target map(alloc: s)
{
s.x = 11;
s.y = 21;
}
// To determine whether x needs to be transfered or deleted, the runtime
// cannot simply check whether unified shared memory is enabled and the
// 'close' modifier is specified. It must check whether x was previously
// placed in device memory by, for example, a 'close' modifier that isn't
// specified here. The following struct member case checks a special code
// path in the runtime implementation where members are transferred before
// deletion of the struct.
#pragma omp target exit data map(from: s.x, s.y)
// CHECK: s.x=11, s.y=21
printf("s.x=%d, s.y=%d\n", s.x, s.y);
// CHECK: present: 0
printf("present: %d\n", omp_target_is_present(&s, dev));
return 0;
}
|
