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

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