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 | |
16 | extern "C" { |
17 | void *llvm_omp_target_lock_mem(void *ptr, size_t size, int device_num); |
18 | void llvm_omp_target_unlock_mem(void *ptr, int device_num); |
19 | } |
20 | |
21 | int 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 | |