1// RUN: %libomptarget-compilexx-generic
2// RUN: %libomptarget-run-generic %fcheck-generic
3
4// UNSUPPORTED: aarch64-unknown-linux-gnu
5// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
6// UNSUPPORTED: nvptx64-nvidia-cuda
7// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
8// UNSUPPORTED: x86_64-pc-linux-gnu
9// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
10// UNSUPPORTED: amdgcn-amd-amdhsa
11
12#include <cstdio>
13
14#include <omp.h>
15
16extern "C" {
17void *llvm_omp_target_lock_mem(void *ptr, size_t size, int device_num);
18void llvm_omp_target_unlock_mem(void *ptr, int device_num);
19}
20
21int main() {
22 int n = 100;
23 int *unlocked = new int[n];
24
25 for (int i = 0; i < n; i++)
26 unlocked[i] = i;
27
28 int *locked = (int *)llvm_omp_target_lock_mem(unlocked, n * sizeof(int),
29 omp_get_default_device());
30 if (!locked)
31 return 0;
32
33#pragma omp target teams distribute parallel for map(tofrom : unlocked[ : n])
34 for (int i = 0; i < n; i++)
35 unlocked[i] += 1;
36
37#pragma omp target teams distribute parallel for map(tofrom : unlocked[10 : 10])
38 for (int i = 10; i < 20; i++)
39 unlocked[i] += 1;
40
41#pragma omp target teams distribute parallel for map(tofrom : locked[ : n])
42 for (int i = 0; i < n; i++)
43 locked[i] += 1;
44
45#pragma omp target teams distribute parallel for map(tofrom : locked[10 : 10])
46 for (int i = 10; i < 20; i++)
47 locked[i] += 1;
48
49 llvm_omp_target_unlock_mem(unlocked, omp_get_default_device());
50
51 int err = 0;
52 for (int i = 0; i < n; i++) {
53 if (i < 10 || i > 19) {
54 if (unlocked[i] != i + 2) {
55 printf(format: "Err at %d, got %d, expected %d\n", i, unlocked[i], i + 1);
56 err++;
57 }
58 } else if (unlocked[i] != i + 4) {
59 printf(format: "Err at %d, got %d, expected %d\n", i, unlocked[i], i + 2);
60 err++;
61 }
62 }
63
64 // CHECK: PASS
65 if (err == 0)
66 printf(format: "PASS\n");
67
68 return err;
69}
70

source code of offload/test/mapping/prelock.cpp