offload/test/mapping/prelock.cpp

   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
  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 }