openmp/libomptarget/test/mapping/prelock.cpp

   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("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("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("PASS\n");
  67
  68   return err;
  69 }