forked from OSchip/llvm-project
42 lines
1.1 KiB
C
42 lines
1.1 KiB
C
// 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;
|
|
}
|