llvm-project

1
// RUN: %libomptarget-compilexx-generic
2
// RUN: %libomptarget-run-generic %fcheck-generic
3

4
// REQURIES: gpu
5
// UNSUPPORTED: nvidiagpu
6
// UNSUPPORTED: amdgpu
7

8
#include <cstdio>
9

10
#include <omp.h>
11

12
extern "C" {
13
void *llvm_omp_target_lock_mem(void *ptr, size_t size, int device_num);
14
void llvm_omp_target_unlock_mem(void *ptr, int device_num);
15
}
16

17
int 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("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("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("PASS\n");
63

64
  return err;
65
}
66